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 ++++++-------------------- 1 file changed, 6 insertions(+), 20 deletions(-) (limited to 'lib/core/c++/particle') 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){ -- cgit v1.3.1 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 --- default.nix | 2 +- examples/settling_cubes_2d_ibm_gpu/sim.cpp | 16 +- lib/core/c++/particle/particle.hpp | 13 +- lib/sycl/c++/data.hpp | 232 ++++++++++++++++++++--------- 4 files changed, 177 insertions(+), 86 deletions(-) (limited to 'lib/core/c++/particle') diff --git a/default.nix b/default.nix index 05ca001..2562f1d 100644 --- a/default.nix +++ b/default.nix @@ -71,7 +71,7 @@ let src = builtins.fetchurl { url = "https://git.keldu.de/forstio-forstio/snapshot/master.tar.gz"; - sha256 = "sha256:026scvd9015blphbj49d0cv5wfrxh1szzdv3p212wzk4iw643akj"; + sha256 = "sha256:1bpglvf0h90dqdm25kp3h1rvk3f7fyasp43xsgki09n1yp55iw9w"; }; phases = [ "unpackPhase" "installPhase" ]; diff --git a/examples/settling_cubes_2d_ibm_gpu/sim.cpp b/examples/settling_cubes_2d_ibm_gpu/sim.cpp index 3c698bb..b4f204e 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 > >; } @@ -103,12 +103,10 @@ saw::error_or setup_initial_conditions( // Particles { - saw::data> rad_p, dense_p; - rad_p.at({}).set(2.0); + saw::data> dense_p; dense_p.at({}).set(1); auto& spheroid_group = particles.template get<0u>(); - spheroid_group = create_spheroid_particle_group( - rad_p, + spheroid_group = create_spheroid_particle_group( dense_p, {64u} ); @@ -165,7 +163,7 @@ saw::error_or step( }).wait(); q.submit([&](acpp::sycl::handler& h){ - h.parallel_for(acpp::sycl::range{dim_x}, [=](acpp::sycl::id idx){ + h.parallel_for(acpp::sycl::range<1u>{dim_x}, [=](acpp::sycl::id<1u> idx){ }); }).wait(); // Step @@ -257,9 +255,9 @@ saw::error_or lbm_main(int argc, char** argv){ saw::data, encode::Sycl> lbm_sycl_particle_data{sycl_q}; sycl_q.wait(); - auto lsd_view = make_chunk_struct_view(lbm_sycl_data); - auto lsdm_view = make_chunk_struct_view(lbm_sycl_macro_data); - auto lsdp_view = make_chunk_struct_view(lbm_sycl_particle_data); + auto lsd_view = make_view(lbm_sycl_data); + auto lsdm_view = make_view(lbm_sycl_macro_data); + auto lsdp_view = make_view(lbm_sycl_particle_data); { auto eov = dev.copy_to_device(*lbm_data_ptr,lbm_sycl_data); if(eov.is_error()){ 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.3.1