From 20e46825069b3974d5cc883163d7e37b4836b2af Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Mon, 13 Apr 2026 19:26:57 +0200 Subject: Progress? --- lib/core/c++/particle/particle.hpp | 26 +--- lib/sycl/c++/data.hpp | 300 ++++++++++++++++++++++++++++++++++++- 2 files changed, 300 insertions(+), 26 deletions(-) (limited to 'lib') diff --git a/lib/core/c++/particle/particle.hpp b/lib/core/c++/particle/particle.hpp index c0d115f..e48cf08 100644 --- a/lib/core/c++/particle/particle.hpp +++ b/lib/core/c++/particle/particle.hpp @@ -55,14 +55,13 @@ using Particle = Struct< template> using ParticleGroup = Struct< Member, "mask">, - Member, - Member, "mass">, + Member,1u>, "density">, Member>, "particles"> >; } -template -saw::data>> create_spheroid_particle_group( +template +saw::data>> create_spheroid_particle_group( saw::data> rad_p, saw::data> density_p, const saw::data& mask_resolution @@ -71,23 +70,10 @@ saw::data>> create_sph auto& mask = part.template get<"mask">(); auto& collision = part.template get<"collision">(); - auto& mass = part.template get<"mass">(); + auto& density = part.template get<"density">().at({{0u}}); - if constexpr ( D == 1u ){ - saw::data> c; - c.at({}).set(2.0); - mass = rad_p * c * density_p; - } else if constexpr ( D == 2u){ - saw::data> pi; - pi.at({}).set(3.14159); - mass = rad_p * rad_p * pi * density_p; - } else if constexpr ( D == 3u ){ - saw::data> c; - c.at({}).set(3.14159 * 4.0 / 3.0); - mass = rad_p * rad_p * rad_p * c * density_p; - } else { - static_assert(D == 0u or D > 3u, "Dimensions only supported for Dim 1,2 & 3."); - } + static_assert(D == 0u or D > 3u, "Dimensions only supported for Dim 1,2 & 3."); + density = density_p; saw::data> mask_dims; for(uint64_t i = 0u; i < D; ++i){ diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp index 50adb4f..661c313 100644 --- a/lib/sycl/c++/data.hpp +++ b/lib/sycl/c++/data.hpp @@ -29,13 +29,15 @@ private: SAW_FORBID_COPY(data); SAW_FORBID_MOVE(data); + public: data(acpp::sycl::queue& q__): q_{&q__}, values_{nullptr} { + SAW_ASSERT(q_); values_ = acpp::sycl::malloc_device>(ct_multiply::value,*q_); - SAW_ASSERT(values_ and q_); + SAW_ASSERT(values_); } ~data(){ @@ -63,6 +65,7 @@ public: constexpr data* flat_data() { return values_; } + }; template @@ -81,7 +84,7 @@ public: {} data(data, kel::lbm::encode::Sycl>& values__): - values_{&values__.at({})} + values_{values__.flat_data()} {} data(data* values__): @@ -105,6 +108,113 @@ public: } }; +template +class data, kel::lbm::encode::Sycl> final { +public: + using Schema = schema::Array; +private: + static_assert(Dims > 0u, "Zero Dim Arrays make no sense here. If you meant to use this for math style approaches then use Tensor instead of Array"); + + data* values_; + data,Encode> meta_; + acpp::sycl::queue* q_; + + SAW_FORBID_COPY(data); + SAW_FORBID_MOVE(data); +public: + data(const data, Encode>& meta__,acpp::sycl::queue& q__): + values_{nullptr}, + meta_{meta__}, + q_{&q__} + { + SAW_ASSERT(q_); + values_ = acpp::sycl::malloc_device>(flat_size().get(),*q_); + SAW_ASSERT(values_); + } + + ~data(){ + if(not values_){ + return; + } + SAW_ASSERT(q_); + + acpp::sycl::free(values_,*q_); + values_ = nullptr; + } + + constexpr data, Encode> meta() const { + return meta_; + } + + constexpr data& at(const data, Encode>& index){ + return values_[kel::lbm::flatten_index::apply(index,meta()).get()]; + } + + constexpr data& at(const data, Encode>& index) const{ + return values_[kel::lbm::flatten_index::apply(index,meta()).get()]; + } + + constexpr data* flat_data() { + return values_; + } + + constexpr data flat_size() const { + data mult{1u}; + + for(uint64_t i{0u}; i < Dims; ++i){ + mult = mult * meta_.at({i}); + } + + return mult; + } +}; + +template +class data>, kel::lbm::encode::Sycl> final { +public: + using Schema = schema::Ptr>; +private: + static_assert(Dims > 0u, "Zero Dim Arrays make no sense here. If you meant to use this for math style approaches then use Tensor instead of Array"); + + data* values_; + data,Encode> meta_; + + SAW_FORBID_COPY(data); + SAW_FORBID_MOVE(data); +public: + data(const data, kel::lbm::encode::Sycl>& values__): + values_{values__.flat_data()}, + meta_{values__.meta()} + { + } + + constexpr data, Encode> meta() const { + return meta_; + } + + constexpr data& at(const data, Encode>& index){ + return values_[kel::lbm::flatten_index::apply(index,meta()).get()]; + } + + constexpr data& at(const data, Encode>& index) const{ + return values_[kel::lbm::flatten_index::apply(index,meta()).get()]; + } + + constexpr data* flat_data() { + return values_; + } + + constexpr data flat_size() const { + data mult{1u}; + + for(uint64_t i{0u}; i < Dims; ++i){ + mult = mult * meta_.at({i}); + } + + return mult; + } +}; + template class data,kel::lbm::encode::Sycl> final { public: @@ -223,6 +333,64 @@ public: } }; +template +struct data, kel::lbm::encode::Sycl> final { +public: + using StorageT = std::tuple>...>; + using Schema = schema::Tuple; +private: + StorageT members_; + + /** + * A helper constructor to forward the sycl queue to the inner "default" constructors + */ + template + constexpr data(acpp::sycl::queue& q, std::index_sequence): + members_{(static_cast(Is), q)...} + {} +public: + data(acpp::sycl::queue& q__): + data{q__, std::make_index_sequence{}} + { + q__.wait(); + } + + template + auto& get(){ + return std::get(members_); + } + + template + const auto& get() const { + return std::get(members_); + } +}; + +template +struct data>, kel::lbm::encode::Sycl> final { +public: + using StorageT = std::tuple,kel::lbm::encode::Sycl>...>; + using Schema = schema::Tuple; +private: + StorageT members_; + + /** + * A helper constructor to forward the sycl queue to the inner "default" constructors + */ +public: + data() = default; + + template + auto& get(){ + return std::get(members_); + } + + template + const auto& get() const { + return std::get(members_); + } +}; + template class data, kel::lbm::encode::Sycl > final { public: @@ -235,6 +403,9 @@ private: */ StorageT members_; + /** + * A helper constructor to forward the sycl queue to the inner "default" constructors + */ template constexpr data(acpp::sycl::queue& q, std::index_sequence): members_{(static_cast(Is), q)...} @@ -246,6 +417,16 @@ public: q__.wait(); } + template + auto& get(){ + return std::get(members_); + } + + template + const auto& get() const { + return std::get(members_); + } + template auto& get(){ return std::get::value>(members_); @@ -347,6 +528,60 @@ struct sycl_copy_helper, Encode> final { } }; +template +struct sycl_copy_helper, Encode> final { + using Schema = sch::Tuple; + + template + static saw::error_or copy_to_device_member(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ + if constexpr (i < sizeof...(Members)){ + auto& host_member_data = host_data.template get(); + auto& sycl_member_data = sycl_data.template get(); + + auto host_ptr = host_member_data.flat_data(); + auto sycl_ptr = sycl_member_data.flat_data(); + + static_assert(sizeof(std::decay_t) == sizeof(std::decay_t), "Unequal size"); + + q.submit([&](acpp::sycl::handler& h){ + h.copy(host_ptr,sycl_ptr, host_member_data.flat_size().get()); + }).wait(); + + return copy_to_device_member(host_data,sycl_data,q); + } + + return saw::make_void(); + } + + static saw::error_or copy_to_device(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ + return copy_to_device_member<0u>(host_data, sycl_data, q); + } + + template + static saw::error_or copy_to_host_member(saw::data>& sycl_data, saw::data& host_data, sycl::queue& q){ + if constexpr (i < sizeof...(Members)){ + auto& host_member_data = host_data.template get(); + auto& sycl_member_data = sycl_data.template get(); + + auto host_ptr = host_member_data.flat_data(); + auto sycl_ptr = sycl_member_data.flat_data(); + + q.submit([&](acpp::sycl::handler& h){ + h.copy(sycl_ptr,host_ptr, host_member_data.flat_size().get()); + }).wait(); + + return copy_to_host_member(sycl_data,host_data,q); + } + + return saw::make_void(); + } + + + static saw::error_or copy_to_host(saw::data>& sycl_data, saw::data& host_data, sycl::queue& q){ + return copy_to_host_member<0u>(sycl_data, host_data, q); + } +}; + template struct sycl_copy_helper, Encode> final { using Schema = sch::FixedArray; @@ -376,12 +611,41 @@ struct sycl_copy_helper, Encode> final { } }; +template +struct sycl_copy_helper, Encode> final { + using Schema = sch::Array; + + static saw::error_or copy_to_host(saw::data>& sycl_data, saw::data& host_data, sycl::queue& q){ + auto host_ptr = host_data.flat_data(); + auto sycl_ptr = sycl_data.flat_data(); + + static_assert(sizeof(std::decay_t) == sizeof(std::decay_t), "Unequal size"); + + q.submit([&](acpp::sycl::handler& h){ + h.copy(sycl_ptr,host_ptr, host_data.flat_size()); + }).wait(); + return saw::make_void(); + } + + static saw::error_or copy_to_device(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ + auto host_ptr = host_data.flat_data(); + auto sycl_ptr = sycl_data.flat_data(); + + static_assert(sizeof(std::decay_t) == sizeof(std::decay_t), "Unequal size"); + + q.submit([&](acpp::sycl::handler& h){ + h.copy(host_ptr,sycl_ptr, host_data.flat_size()); + }).wait(); + return saw::make_void(); + } +}; + template -struct make_chunk_struct_view_helper; +struct make_view_helper; template -struct make_chunk_struct_view_helper...>, encode::Sycl> { +struct make_view_helper...>, encode::Sycl> { private: template static void apply_i( @@ -405,12 +669,36 @@ public: return str_view; } }; + +template +struct make_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)){ + 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); +auto make_view(saw::data& dat){ + return impl::make_view_helper::apply(dat); } class device final { -- cgit v1.2.3 From ebc2e26a1b3498363bb7522c241de2925bb7f627 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Tue, 14 Apr 2026 15:33:11 +0200 Subject: Fixed compilation issues with make_view for sycl --- lib/core/c++/particle/particle.hpp | 13 +-- lib/sycl/c++/data.hpp | 232 ++++++++++++++++++++++++++----------- 2 files changed, 169 insertions(+), 76 deletions(-) (limited to 'lib') diff --git a/lib/core/c++/particle/particle.hpp b/lib/core/c++/particle/particle.hpp index e48cf08..7af43dc 100644 --- a/lib/core/c++/particle/particle.hpp +++ b/lib/core/c++/particle/particle.hpp @@ -40,9 +40,8 @@ using ParticleRigidBody = Struct< Member::Schema, "angular_acceleration"> >; -template +template::type radius = 1.0f> using ParticleCollisionSpheroid = Struct< - Member, "radius"> >; template @@ -52,7 +51,7 @@ using Particle = Struct< // Member, "mask">, >; -template> +template> using ParticleGroup = Struct< Member, "mask">, Member,1u>, "density">, @@ -60,19 +59,17 @@ using ParticleGroup = Struct< >; } -template +template::type radius> saw::data>> create_spheroid_particle_group( - saw::data> rad_p, saw::data> density_p, const saw::data& mask_resolution ){ - saw::data>> part; + saw::data>> part; auto& mask = part.template get<"mask">(); - auto& collision = part.template get<"collision">(); auto& density = part.template get<"density">().at({{0u}}); - static_assert(D == 0u or D > 3u, "Dimensions only supported for Dim 1,2 & 3."); + static_assert(D >= 1u and D <= 3u, "Dimensions only supported for Dim 1,2 & 3."); density = density_p; saw::data> mask_dims; diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp index 661c313..4e5129b 100644 --- a/lib/sycl/c++/data.hpp +++ b/lib/sycl/c++/data.hpp @@ -62,7 +62,7 @@ public: return values_[kel::lbm::flatten_index::apply(index,get_dims()).get()]; } - constexpr data* flat_data() { + constexpr data* flat_data() const { return values_; } @@ -103,7 +103,7 @@ public: return values_[kel::lbm::flatten_index::apply(index,get_dims()).get()]; } - constexpr data* flat_data() { + constexpr data* flat_data() const { return values_; } }; @@ -122,6 +122,14 @@ private: SAW_FORBID_COPY(data); SAW_FORBID_MOVE(data); public: + data(acpp::sycl::queue& q__): + values_{nullptr}, + meta_{}, + q_{&q__} + { + SAW_ASSERT(q_); + } + data(const data, Encode>& meta__,acpp::sycl::queue& q__): values_{nullptr}, meta_{meta__}, @@ -154,7 +162,7 @@ public: return values_[kel::lbm::flatten_index::apply(index,meta()).get()]; } - constexpr data* flat_data() { + constexpr data* flat_data() const { return values_; } @@ -179,9 +187,14 @@ private: data* values_; data,Encode> meta_; - SAW_FORBID_COPY(data); - SAW_FORBID_MOVE(data); public: + SAW_DEFAULT_COPY(data); + SAW_DEFAULT_MOVE(data); + + data(): + values_{nullptr} + {} + data(const data, kel::lbm::encode::Sycl>& values__): values_{values__.flat_data()}, meta_{values__.meta()} @@ -200,7 +213,7 @@ public: return values_[kel::lbm::flatten_index::apply(index,meta()).get()]; } - constexpr data* flat_data() { + constexpr data* flat_data() const { return values_; } @@ -261,7 +274,7 @@ public: return data,Encode>{{Sides...}}; } - auto flat_data(){ + auto flat_data() const { return values_.flat_data(); } @@ -288,6 +301,10 @@ public: values_{nullptr} {} + data(const data, kel::lbm::encode::Sycl>& values__): + values_{values__.flat_data()} + {} + data(data* values__): values_{values__} {} @@ -361,7 +378,7 @@ public: } template - const auto& get() const { + auto& get() const { return std::get(members_); } }; @@ -423,7 +440,7 @@ public: } template - const auto& get() const { + auto& get() const { return std::get(members_); } @@ -433,7 +450,7 @@ public: } template - const auto& get() const { + auto& get() const { return std::get::value>(members_); } }; @@ -449,17 +466,26 @@ private: * Do it here by specializing. */ StorageT members_; - public: data() = default; + template + auto& get(){ + return std::get(members_); + } + + template + auto& get() const { + return std::get(members_); + } + template auto& get(){ return std::get::value>(members_); } template - const auto& get() const { + auto& get() const { return std::get::value>(members_); } }; @@ -482,14 +508,10 @@ struct sycl_copy_helper, Encode> final { auto& host_member_data = host_data.template get(); auto& sycl_member_data = sycl_data.template get(); - auto host_ptr = host_member_data.flat_data(); - auto sycl_ptr = sycl_member_data.flat_data(); - - static_assert(sizeof(std::decay_t) == sizeof(std::decay_t), "Unequal size"); - - q.submit([&](acpp::sycl::handler& h){ - h.copy(host_ptr,sycl_ptr, host_member_data.flat_size().get()); - }).wait(); + auto eov = sycl_copy_helper::copy_to_device(host_member_data,sycl_member_data,q); + if(eov.is_error()){ + return eov; + } return copy_to_device_member(host_data,sycl_data,q); } @@ -509,12 +531,10 @@ struct sycl_copy_helper, Encode> final { auto& host_member_data = host_data.template get(); auto& sycl_member_data = sycl_data.template get(); - auto host_ptr = host_member_data.flat_data(); - auto sycl_ptr = sycl_member_data.flat_data(); - - q.submit([&](acpp::sycl::handler& h){ - h.copy(sycl_ptr,host_ptr, host_member_data.flat_size().get()); - }).wait(); + auto eov = sycl_copy_helper::copy_to_host(sycl_member_data,host_member_data,q); + if(eov.is_error()){ + return eov; + } return copy_to_host_member(sycl_data,host_data,q); } @@ -535,17 +555,14 @@ struct sycl_copy_helper, Encode> final { template static saw::error_or copy_to_device_member(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ if constexpr (i < sizeof...(Members)){ + using M = typename saw::parameter_pack_type::type; auto& host_member_data = host_data.template get(); auto& sycl_member_data = sycl_data.template get(); - auto host_ptr = host_member_data.flat_data(); - auto sycl_ptr = sycl_member_data.flat_data(); - - static_assert(sizeof(std::decay_t) == sizeof(std::decay_t), "Unequal size"); - - q.submit([&](acpp::sycl::handler& h){ - h.copy(host_ptr,sycl_ptr, host_member_data.flat_size().get()); - }).wait(); + auto eov = sycl_copy_helper::copy_to_device(host_member_data,sycl_member_data,q); + if(eov.is_error()){ + return eov; + } return copy_to_device_member(host_data,sycl_data,q); } @@ -560,15 +577,15 @@ struct sycl_copy_helper, Encode> final { template static saw::error_or copy_to_host_member(saw::data>& sycl_data, saw::data& host_data, sycl::queue& q){ if constexpr (i < sizeof...(Members)){ + using M = typename saw::parameter_pack_type::type; + auto& host_member_data = host_data.template get(); auto& sycl_member_data = sycl_data.template get(); - auto host_ptr = host_member_data.flat_data(); - auto sycl_ptr = sycl_member_data.flat_data(); - - q.submit([&](acpp::sycl::handler& h){ - h.copy(sycl_ptr,host_ptr, host_member_data.flat_size().get()); - }).wait(); + auto eov = sycl_copy_helper::copy_to_host(sycl_member_data,host_member_data,q); + if(eov.is_error()){ + return eov; + } return copy_to_host_member(sycl_data,host_data,q); } @@ -611,6 +628,35 @@ struct sycl_copy_helper, Encode> final { } }; +template +struct sycl_copy_helper, Encode> final { + using Schema = sch::Chunk; + + static saw::error_or copy_to_host(saw::data>& sycl_data, saw::data& host_data, sycl::queue& q){ + auto host_ptr = host_data.flat_data(); + auto sycl_ptr = sycl_data.flat_data(); + + static_assert(sizeof(std::decay_t) == sizeof(std::decay_t), "Unequal size"); + + q.submit([&](acpp::sycl::handler& h){ + h.copy(sycl_ptr,host_ptr, saw::ct_multiply::value); + }).wait(); + return saw::make_void(); + } + + static saw::error_or copy_to_device(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ + auto host_ptr = host_data.flat_data(); + auto sycl_ptr = sycl_data.flat_data(); + + static_assert(sizeof(std::decay_t) == sizeof(std::decay_t), "Unequal size"); + + q.submit([&](acpp::sycl::handler& h){ + h.copy(host_ptr,sycl_ptr, saw::ct_multiply::value); + }).wait(); + return saw::make_void(); + } +}; + template struct sycl_copy_helper, Encode> final { using Schema = sch::Array; @@ -622,7 +668,7 @@ struct sycl_copy_helper, Encode> final { static_assert(sizeof(std::decay_t) == sizeof(std::decay_t), "Unequal size"); q.submit([&](acpp::sycl::handler& h){ - h.copy(sycl_ptr,host_ptr, host_data.flat_size()); + h.copy(sycl_ptr,host_ptr, host_data.flat_size().get()); }).wait(); return saw::make_void(); } @@ -634,7 +680,7 @@ struct sycl_copy_helper, Encode> final { static_assert(sizeof(std::decay_t) == sizeof(std::decay_t), "Unequal size"); q.submit([&](acpp::sycl::handler& h){ - h.copy(host_ptr,sycl_ptr, host_data.flat_size()); + h.copy(host_ptr,sycl_ptr, host_data.flat_size().get()); }).wait(); return saw::make_void(); } @@ -644,53 +690,99 @@ struct sycl_copy_helper, Encode> final { template struct make_view_helper; +template +struct make_view_helper, encode::Sycl> { +public: +static saw::error_or apply( + saw::data, encode::Sycl>& dat, + saw::data>, encode::Sycl>& dat_view +){ + dat_view = {dat}; + return saw::make_void(); +} +}; + +template +struct make_view_helper, encode::Sycl> { +public: + static saw::error_or apply( + saw::data, encode::Sycl>& dat, + saw::data>, encode::Sycl>& dat_view + ){ + dat_view = {dat}; + return saw::make_void(); + } +}; + +template +struct make_view_helper, encode::Sycl> { +public: + static saw::error_or apply( + saw::data, encode::Sycl>& dat, + saw::data>, encode::Sycl>& dat_view + ){ + dat_view = {dat}; + return saw::make_void(); + } +}; + template struct make_view_helper...>, encode::Sycl> { private: - template - static void apply_i( +template +static saw::error_or apply_i( saw::data...>, encode::Sycl>& dat, - saw::data...>>, encode::Sycl>& dat_view) - { + 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; + auto eov = make_view_helper>::apply(dat.template get(),dat_view.template get()); + if(eov.is_error()){ + return eov; + } - apply_i<0u>(dat,str_view); + return apply_i(dat,dat_view); + } - return str_view; + return saw::make_void(); } +public: +static saw::error_or apply( + saw::data...>, encode::Sycl>& dat, + saw::data...>>, encode::Sycl>& dat_view +){ + return apply_i<0u>(dat,dat_view); +} }; template struct make_view_helper, encode::Sycl> { private: template - static void apply_i( + static saw::error_or apply_i( saw::data, encode::Sycl>& dat, - saw::data>, encode::Sycl>& dat_view) - { + saw::data>, encode::Sycl>& dat_view + ){ if constexpr (i < sizeof...(Sch)){ - dat_view.template get() = {dat.template get().flat_data()}; + using M = typename saw::parameter_pack_type::type; + + auto eov = make_view_helper>::apply(dat.template get(), dat_view.template get()); + if(eov.is_error()){ + return eov; + } - apply_i(dat,dat_view); + return apply_i(dat,dat_view); } + + return saw::make_void(); } public: - static auto apply(saw::data, encode::Sycl>& dat){ - saw::data>, encode::Sycl> str_view; - - apply_i<0u>(dat,str_view); - - return str_view; + static saw::error_or apply( + saw::data, encode::Sycl>& dat, + saw::data>, encode::Sycl> dat_view + ){ + return apply_i<0u>(dat,dat_view); } }; } @@ -698,7 +790,11 @@ public: // Ptr => Ptr>> => Ptr template auto make_view(saw::data& dat){ - return impl::make_view_helper::apply(dat); + saw::data,Encode> dat_view; + auto eov = impl::make_view_helper::apply(dat,dat_view); + (void) eov; + + return dat_view; } class device final { -- cgit v1.2.3 From c61ba8f8eb86f66915a54551fcc39dfbeab1fad9 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Tue, 14 Apr 2026 21:09:42 +0200 Subject: Working on meta schema instantiation, because I'm stupid --- lib/core/c++/abstract/data.hpp | 26 ++++++++-- lib/sycl/c++/data.hpp | 112 +++++++++++++++++++++++++++++++++++++++-- lib/sycl/tests/data.cpp | 27 ++++++++++ 3 files changed, 157 insertions(+), 8 deletions(-) (limited to 'lib') diff --git a/lib/core/c++/abstract/data.hpp b/lib/core/c++/abstract/data.hpp index e8f1757..0075718 100644 --- a/lib/core/c++/abstract/data.hpp +++ b/lib/core/c++/abstract/data.hpp @@ -4,28 +4,48 @@ namespace kel { namespace sch { +struct Void {}; + struct UnsignedInteger {}; struct SignedInteger {}; struct FloatingPoint {}; template struct MixedPrecision { + using Meta = Void; using StorageType = StorageT; using InterfaceType = InterfaceT; }; template struct Primitive { - using PrimitiveType = PrimType; + using Meta = Void; + using Type = PrimType; static constexpr uint64_t Bytes = N; }; template -struct Array { - using InnerType = T; +struct FixedArray { + using Meta = Void; + using Inner = T; static constexpr std::array Dimensions{Dims...}; }; +template +struct Array { + using Meta = FixedArray; + using Inner = T; + static constexpr std::array Dimensions{Dims}; +}; + +template +struct Tuple { +}; } + +template +struct schema { + using Type = Sch; +}; } diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp index 4e5129b..3ac51e0 100644 --- a/lib/sycl/c++/data.hpp +++ b/lib/sycl/c++/data.hpp @@ -29,8 +29,17 @@ private: SAW_FORBID_COPY(data); SAW_FORBID_MOVE(data); - public: + data(const data::MetaSchema>& meta__, acpp::sycl::queue& q__): + q_{&q__}, + values_{nullptr} + { + (void) meta__; + SAW_ASSERT(q_); + values_ = acpp::sycl::malloc_device>(ct_multiply::value,*q_); + SAW_ASSERT(values_); + } + data(acpp::sycl::queue& q__): q_{&q__}, values_{nullptr} @@ -65,7 +74,6 @@ public: constexpr data* flat_data() const { return values_; } - }; template @@ -122,6 +130,20 @@ private: SAW_FORBID_COPY(data); SAW_FORBID_MOVE(data); public: + data(const data::MetaSchema>& meta__, acpp::sycl::queue& q__): + q_{&q__}, + values_{nullptr} + { + SAW_ASSERT(q_); + /// TODO use meta + data m{1u}; + for(uint64_t i = 0u; i < Dims; ++i){ + m = m * meta__.at({i}); + } + values_ = acpp::sycl::malloc_device>(m.get(),*q_); + SAW_ASSERT(values_); + } + data(acpp::sycl::queue& q__): values_{nullptr}, meta_{}, @@ -198,8 +220,7 @@ public: data(const data, kel::lbm::encode::Sycl>& values__): values_{values__.flat_data()}, meta_{values__.meta()} - { - } + {} constexpr data, Encode> meta() const { return meta_; @@ -238,6 +259,10 @@ private: data> values_; public: + data(const data::MetaSchema>& meta__, acpp::sycl::queue& q__): + values_{meta__,q__} + {} + data(acpp::sycl::queue& q__): values_{q__} {} @@ -366,6 +391,11 @@ private: members_{(static_cast(Is), q)...} {} public: + data(data::MetaSchema>& meta__, acpp::sycl::queue& q__): + data{q__, std::make_index_sequence{}} + { + } + data(acpp::sycl::queue& q__): data{q__, std::make_index_sequence{}} { @@ -546,6 +576,29 @@ struct sycl_copy_helper, Encode> final { static saw::error_or copy_to_host(saw::data>& sycl_data, saw::data& host_data, sycl::queue& q){ return copy_to_host_member<0u>(sycl_data, host_data, q); } + + template + static saw::error_or malloc_on_device_member(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ + + if constexpr (i < sizeof...(Members)){ + using M = typename saw::parameter_pack_type::type; + auto& host_member_data = host_data.template get(); + auto& sycl_member_data = sycl_data.template get(); + + auto eov = sycl_copy_helper::malloc_on_device(host_member_data,sycl_member_data,q); + if(eov.is_error()){ + return eov; + } + + return malloc_on_device_member(host_data,sycl_data,q); + } + + return saw::make_void(); + } + + static saw::error_or malloc_on_device(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ + return malloc_on_device_member<0u>(host_data,sycl_data,q); + } }; template @@ -594,9 +647,36 @@ struct sycl_copy_helper, Encode> final { } - static saw::error_or copy_to_host(saw::data>& sycl_data, saw::data& host_data, sycl::queue& q){ + static saw::error_or copy_to_host( + saw::data& host_data, + saw::data>& sycl_data, + sycl::queue& q + ){ return copy_to_host_member<0u>(sycl_data, host_data, q); } + + template + static saw::error_or malloc_on_device_member(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ + + if constexpr (i < sizeof...(Members)){ + using M = typename saw::parameter_pack_type::type; + auto& host_member_data = host_data.template get(); + auto& sycl_member_data = sycl_data.template get(); + + auto eov = sycl_copy_helper::malloc_on_device(host_member_data,sycl_member_data,q); + if(eov.is_error()){ + return eov; + } + + return malloc_on_device_member(host_data,sycl_data,q); + } + + return saw::make_void(); + } + + static saw::error_or malloc_on_device(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ + return malloc_on_device_member<0u>(host_data,sycl_data,q); + } }; template @@ -626,6 +706,13 @@ struct sycl_copy_helper, Encode> final { }).wait(); return saw::make_void(); } + + static saw::error_or malloc_on_device(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ + (void) host_data; + (void) sycl_data; + (void) q; + return saw::make_void(); + } }; template @@ -684,6 +771,11 @@ struct sycl_copy_helper, Encode> final { }).wait(); return saw::make_void(); } + + static saw::error_or malloc_on_device(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ + sycl_data = {host_data.meta(),q}; + return saw::make_void(); + } }; @@ -817,6 +909,16 @@ public: return impl::sycl_copy_helper::copy_to_host(sycl_data, host_data, q_); } + template + saw::error_or malloc_on_device( + saw::data& host_data, + saw::data>& sycl_data + ){ + auto eov = impl::sycl_copy_helper::malloc_on_device(host_data, sycl_data, q_); + q_.wait(); + return eov; + } + auto& get_handle(){ return q_; } diff --git a/lib/sycl/tests/data.cpp b/lib/sycl/tests/data.cpp index 3073a22..6b17622 100644 --- a/lib/sycl/tests/data.cpp +++ b/lib/sycl/tests/data.cpp @@ -2,6 +2,24 @@ #include "../c++/lbm.hpp" +namespace { + +namespace sch { +using namespace kel::lbm::sch; +using TestObjSchema = Tuple< + Member, "foo">, + Member, "bar">, + Member< + Array< + Struct< + Member,"pos"> + > + >, + "baz" + > +>; +} + SAW_TEST("Sycl Data Compilation"){ acpp::sycl::queue q; saw::data< @@ -19,3 +37,12 @@ SAW_TEST("Sycl Data Compilation"){ // test_f.at({}).set(1); // SAW_EXPECT(test_f.at({}).get() == 1, "Value check failed"); } + +SAW_TEST("Sycl Data Compilation for Particle Similacrum"){ + acpp::sycl::queue q; + + saw::data< + sch::TestObjSchema + > a; +} +} -- cgit v1.2.3 From 30ff1caf073b4341fd0614e0974c67a8588c8931 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Wed, 15 Apr 2026 19:11:21 +0200 Subject: Feierabend --- lib/sycl/c++/data.hpp | 56 ++++++++++++++++++++++++++++++------------------- lib/sycl/tests/data.cpp | 8 +++++++ 2 files changed, 43 insertions(+), 21 deletions(-) (limited to 'lib') diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp index 3ac51e0..0206833 100644 --- a/lib/sycl/c++/data.hpp +++ b/lib/sycl/c++/data.hpp @@ -131,16 +131,12 @@ private: SAW_FORBID_MOVE(data); public: data(const data::MetaSchema>& meta__, acpp::sycl::queue& q__): - q_{&q__}, - values_{nullptr} + values_{nullptr}, + meta_{meta__}, + q_{&q__} { SAW_ASSERT(q_); - /// TODO use meta - data m{1u}; - for(uint64_t i = 0u; i < Dims; ++i){ - m = m * meta__.at({i}); - } - values_ = acpp::sycl::malloc_device>(m.get(),*q_); + values_ = acpp::sycl::malloc_device>(flat_size().get(),*q_); SAW_ASSERT(values_); } @@ -152,16 +148,6 @@ public: SAW_ASSERT(q_); } - data(const data, Encode>& meta__,acpp::sycl::queue& q__): - values_{nullptr}, - meta_{meta__}, - q_{&q__} - { - SAW_ASSERT(q_); - values_ = acpp::sycl::malloc_device>(flat_size().get(),*q_); - SAW_ASSERT(values_); - } - ~data(){ if(not values_){ return; @@ -184,6 +170,19 @@ public: return values_[kel::lbm::flatten_index::apply(index,meta()).get()]; } + constexpr error_or reset_to(const data::MetaSchema>& meta_arg){ + SAW_ASSERT(q_); + meta_ = meta_arg; + + if(values_){ + acpp::sycl::free(values_,*q_); + } + values_ = acpp::sycl::malloc_device>(flat_size().get(),*q_); + SAW_ASSERT(q_); + + return make_void(); + } + constexpr data* flat_data() const { return values_; } @@ -391,10 +390,11 @@ private: members_{(static_cast(Is), q)...} {} public: + /* data(data::MetaSchema>& meta__, acpp::sycl::queue& q__): data{q__, std::make_index_sequence{}} - { - } + {} + */ data(acpp::sycl::queue& q__): data{q__, std::make_index_sequence{}} @@ -754,6 +754,7 @@ struct sycl_copy_helper, Encode> final { static_assert(sizeof(std::decay_t) == sizeof(std::decay_t), "Unequal size"); + SAW_ASSERT(host_data.flat_size() == sycl_data.flat_size()); q.submit([&](acpp::sycl::handler& h){ h.copy(sycl_ptr,host_ptr, host_data.flat_size().get()); }).wait(); @@ -761,14 +762,27 @@ struct sycl_copy_helper, Encode> final { } static saw::error_or copy_to_device(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ + + { + auto hm = host_data.meta(); + auto sm = sycl_data.meta(); + bool equ{true}; + for(uint64_t i{0u}; i < Dims; ++i){ + equ &= (hm.at({i}).get() == sm.at({i}).get()); + } + if(not equ){ + sycl_data.reset_to(hm); + } + } + auto host_ptr = host_data.flat_data(); auto sycl_ptr = sycl_data.flat_data(); - static_assert(sizeof(std::decay_t) == sizeof(std::decay_t), "Unequal size"); q.submit([&](acpp::sycl::handler& h){ h.copy(host_ptr,sycl_ptr, host_data.flat_size().get()); }).wait(); + return saw::make_void(); } diff --git a/lib/sycl/tests/data.cpp b/lib/sycl/tests/data.cpp index 6b17622..4321a0d 100644 --- a/lib/sycl/tests/data.cpp +++ b/lib/sycl/tests/data.cpp @@ -38,6 +38,13 @@ SAW_TEST("Sycl Data Compilation"){ // SAW_EXPECT(test_f.at({}).get() == 1, "Value check failed"); } +SAW_TEST("Sycl Data Array of Struct"){ + acpp::sycl::queue q; + + saw::data, kel::lbm::encode::Sycl> a{{{2u}},q}; +} + +/* SAW_TEST("Sycl Data Compilation for Particle Similacrum"){ acpp::sycl::queue q; @@ -45,4 +52,5 @@ SAW_TEST("Sycl Data Compilation for Particle Similacrum"){ sch::TestObjSchema > a; } +*/ } -- cgit v1.2.3 From d01f8faadcec62593e0af5304bcc2db3d22ed0c5 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Thu, 16 Apr 2026 16:20:23 +0200 Subject: Brain fried --- lib/core/c++/particle.hpp | 5 +++++ lib/sycl/c++/data.hpp | 16 ++++++++++++++++ 2 files changed, 21 insertions(+) (limited to 'lib') diff --git a/lib/core/c++/particle.hpp b/lib/core/c++/particle.hpp index b098ecc..691a74b 100644 --- a/lib/core/c++/particle.hpp +++ b/lib/core/c++/particle.hpp @@ -11,7 +11,12 @@ struct Particle {}; template class component { +private: + saw::data> dt_; public: + component(saw::data> dt__): + dt_{dt__} + {} template void apply(const saw::data& particles, const saw::data& macros, saw::data index, saw::data time_step) const { diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp index 0206833..71627de 100644 --- a/lib/sycl/c++/data.hpp +++ b/lib/sycl/c++/data.hpp @@ -278,6 +278,10 @@ public: return data>::get_dims(); } + static constexpr auto ghost_meta() { + return data>::meta(); + } + data& at(const data>& index){ std::decay_t ind; for(uint64_t i = 0u; i < sizeof...(Sides); ++i){ @@ -298,6 +302,10 @@ public: return data,Encode>{{Sides...}}; } + static constexpr auto meta(){ + return data,Encode>{{Sides...}}; + } + auto flat_data() const { return values_.flat_data(); } @@ -345,6 +353,10 @@ public: return data>::get_dims(); } + static constexpr auto ghost_meta() { + return data>::meta(); + } + data& at(const data>& index){ std::decay_t ind; for(uint64_t i = 0u; i < sizeof...(Sides); ++i){ @@ -361,6 +373,10 @@ public: return values_.at(ind); } + static constexpr auto meta(){ + return data,Encode>{{Sides...}}; + } + static constexpr auto get_dims(){ return data,Encode>{{Sides...}}; } -- cgit v1.2.3 From 45ebf7411d687ab5530431ab1bcc74edb0499c69 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Wed, 22 Apr 2026 16:48:50 +0200 Subject: Working on equalizing poiseulle_particle channel methods --- lib/core/c++/write_vtk.hpp | 4 ++-- lib/sycl/c++/data.hpp | 10 +++++++--- 2 files changed, 9 insertions(+), 5 deletions(-) (limited to 'lib') diff --git a/lib/core/c++/write_vtk.hpp b/lib/core/c++/write_vtk.hpp index f925361..e852172 100644 --- a/lib/core/c++/write_vtk.hpp +++ b/lib/core/c++/write_vtk.hpp @@ -136,7 +136,7 @@ struct lbm_vtk_writer> { template struct lbm_vtk_writer> { static saw::error_or apply_header(std::ostream& vtk_file, std::string_view name){ - vtk_file<<"SCALARS "<, if constexpr (sizeof...(MemberT) > 0u){ // POINT DATA { - vtk_file << "POINT_DATA " << pd_size.get() <<"\n"; + vtk_file << "\nPOINT_DATA " << pd_size.get() <<"\n"; } // HEADER TO BODY diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp index 71627de..9f43848 100644 --- a/lib/sycl/c++/data.hpp +++ b/lib/sycl/c++/data.hpp @@ -381,7 +381,7 @@ public: return data,Encode>{{Sides...}}; } - auto flat_data(){ + auto flat_data() const { return values_.flat_data(); } @@ -741,8 +741,10 @@ struct sycl_copy_helper, Encode> final { static_assert(sizeof(std::decay_t) == sizeof(std::decay_t), "Unequal size"); + auto flat_size = host_data.flat_size(); + q.submit([&](acpp::sycl::handler& h){ - h.copy(sycl_ptr,host_ptr, saw::ct_multiply::value); + h.copy(sycl_ptr,host_ptr, flat_size.get()); }).wait(); return saw::make_void(); } @@ -753,8 +755,10 @@ struct sycl_copy_helper, Encode> final { static_assert(sizeof(std::decay_t) == sizeof(std::decay_t), "Unequal size"); + auto flat_size = host_data.flat_size(); + q.submit([&](acpp::sycl::handler& h){ - h.copy(host_ptr,sycl_ptr, saw::ct_multiply::value); + h.copy(host_ptr,sycl_ptr, flat_size.get()); }).wait(); return saw::make_void(); } -- cgit v1.2.3