From 20e46825069b3974d5cc883163d7e37b4836b2af Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Mon, 13 Apr 2026 19:26:57 +0200 Subject: Progress? --- examples/settling_cubes_2d_ibm_gpu/sim.cpp | 16 +- lib/core/c++/particle/particle.hpp | 26 +-- lib/sycl/c++/data.hpp | 300 ++++++++++++++++++++++++++++- 3 files changed, 308 insertions(+), 34 deletions(-) diff --git a/examples/settling_cubes_2d_ibm_gpu/sim.cpp b/examples/settling_cubes_2d_ibm_gpu/sim.cpp index 3c1e96d..3c698bb 100644 --- a/examples/settling_cubes_2d_ibm_gpu/sim.cpp +++ b/examples/settling_cubes_2d_ibm_gpu/sim.cpp @@ -52,7 +52,7 @@ using MacroStruct = Struct< template using ParticleGroups = Tuple< ParticleGroup< - T,Desc::D,sch::ParticleCollisionSpheroid + T,Desc::D,sch::ParticleCollisionSpheroid > >; } @@ -119,11 +119,11 @@ saw::error_or setup_initial_conditions( template saw::error_or step( - saw::data>,encode::Sycl>& fields, - saw::data>,encode::Sycl>& macros, - saw::data, encode::Sycl>& particles, - saw::data t_i, - device& dev + saw::data>,encode::Sycl>& fields, + saw::data>,encode::Sycl>& macros, + saw::data>,encode::Sycl>& particles, + saw::data t_i, + device& dev ){ auto& q = dev.get_handle(); auto& info_f = fields.template get<"info">(); @@ -214,7 +214,7 @@ saw::error_or lbm_main(int argc, char** argv){ // saw::data> meta{{dim_x,dim_y}}; auto lbm_data_ptr = saw::heap>>(); auto lbm_macro_data_ptr = saw::heap>>(); - auto lbm_particle_data = saw::heap>>(); + auto lbm_particle_data_ptr = saw::heap>>(); std::cout<<"Estimated Bytes: "<,sch::MacroStruct>().get()< lbm_main(int argc, char** argv){ sycl_q.wait(); { - auto eov = setup_initial_conditions(*lbm_data_ptr,*lbm_macro_data_ptr,lbm_particle_data); + auto eov = setup_initial_conditions(*lbm_data_ptr,*lbm_macro_data_ptr,*lbm_particle_data_ptr); if(eov.is_error()){ return eov; } 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