From 1d81ab6f66b5efd6b4d291762935866930731152 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Wed, 28 Jan 2026 11:24:30 +0100 Subject: Reworked fetching due to unstable fetchTarball --- lib/sycl/c++/data.hpp | 163 +++++++++++++++++++++++++++++++++++++++++++++++++- 1 file changed, 161 insertions(+), 2 deletions(-) (limited to 'lib') diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp index d5adb95..9b15678 100644 --- a/lib/sycl/c++/data.hpp +++ b/lib/sycl/c++/data.hpp @@ -27,7 +27,8 @@ private: acpp::sycl::queue* q_; data* values_; - //SAW_FORBID_COPY(data); + SAW_FORBID_COPY(data); + SAW_DEFAULT_MOVE(data); public: data(acpp::sycl::queue& q__): q_{&q__}, @@ -64,6 +65,41 @@ public: } }; +template +class data>, kel::lbm::encode::Sycl> final { +public: + using Schema = schema::Ptr>; +private: + data* values_; + + SAW_DEFAULT_COPY(data); + SAW_DEFAULT_MOVE(data); +public: + data(): + values_{nullptr} + {} + + data(data* values__): + values_{values__} + {} + + static constexpr data> get_dims() { + return saw::data>{{Dims...}}; + } + + constexpr data& at(data>& index){ + return values_[kel::lbm::flatten_index::apply(index,get_dims()).get()]; + } + + constexpr data& at(data>& index) const{ + return values_[kel::lbm::flatten_index::apply(index,get_dims()).get()]; + } + + constexpr data* flat_data() { + return values_; + } +}; + template class data,kel::lbm::encode::Sycl> final { public: @@ -119,13 +155,74 @@ public: } }; +template +class data>,kel::lbm::encode::Sycl> final { +public: + using Schema = schema::Ptr>; +private: + using InnerSchema = typename kel::lbm::sch::Chunk::InnerSchema; + using ValueSchema = typename InnerSchema::ValueType; + + data,kel::lbm::encode::Sycl> values_; + + SAW_DEFAULT_MOVE(data); + SAW_DEFAULT_COPY(data); +public: + data(): + values_{nullptr} + {} + + data(data* values__): + values_{values__} + {} + + data>& ghost_at(const data>& index){ + return values_.at(index); + } + + data>& ghost_at(const data>& index) const { + return values_.at(index); + } + + static constexpr auto get_ghost_dims() { + return data>::get_dims(); + } + + data& at(const data>& index){ + std::decay_t ind; + for(uint64_t i = 0u; i < sizeof...(Sides); ++i){ + ind.at({i}) = index.at({i}) + Ghost; + } + return values_.at(ind); + } + + data& at(const data>& index) const { + std::decay_t ind; + for(uint64_t i = 0u; i < sizeof...(Sides); ++i){ + ind.at({i}) = index.at({i}) + Ghost; + } + return values_.at(ind); + } + + static constexpr auto get_dims(){ + return data,kel::lbm::encode::Sycl>{{Sides...}}; + } + + auto flat_data(){ + return values_.flat_data(); + } + + static constexpr auto flat_size() { + return data>::flat_size(); + } +}; + template class data, kel::lbm::encode::Sycl > final { public: using StorageT = std::tuple>...>; using Schema = schema::Struct; private: - /** * @todo Check by static assert that the members all have the same dimensions. Alternatively * Do it here by specializing. @@ -153,6 +250,32 @@ public: return std::get::value>(members_); } }; + +template +class data...>>, kel::lbm::encode::Sycl > final { +public: + using StorageT = std::tuple,kel::lbm::encode::Sycl>...>; + using Schema = schema::Struct,Keys>...>; +private: + /** + * @todo Check by static assert that the members all have the same dimensions. Alternatively + * Do it here by specializing. + */ + StorageT members_; + +public: + data() = default; + + template + auto& get(){ + return std::get::value>(members_); + } + + template + const auto& get() const { + return std::get::value>(members_); + } +}; } namespace kel { @@ -217,7 +340,43 @@ struct sycl_copy_helper, Encode> final { return copy_to_host_member<0u>(sycl_data, host_data, q); } }; + +template +struct make_chunk_struct_view_helper; + +template +struct make_chunk_struct_view_helper...>, encode::Sycl> { +private: + template + static void apply_i( + saw::data...>, encode::Sycl>& dat, + saw::data...>>, encode::Sycl>& dat_view) + { + if constexpr (i < sizeof...(Sch)){ + using M = typename saw::parameter_pack_type...>::type; + + dat_view.template get() = {dat.template get().flat_data()}; + + apply_i(dat,dat_view); + } + } +public: + static auto apply(saw::data...>, encode::Sycl>& dat){ + saw::data...>>, encode::Sycl> str_view; + + apply_i<0u>(dat,str_view); + + return str_view; + } +}; } + +// Ptr => Ptr>> => Ptr +template +auto make_chunk_struct_view(saw::data& dat){ + return impl::make_chunk_struct_view_helper::apply(dat); +} + class device final { private: sycl::queue q_; -- cgit v1.2.3