From 44cb41bbe6128f2c7ca92b06e2267d52e3121721 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Wed, 15 Jul 2026 15:39:41 +0200 Subject: Preping two way --- .../moving_poiseulle_particles_2d_hlbm_gpu/sim.cpp | 81 +++++++++++++--------- examples/settling_spheres_2d_hlbm_gpu/sim.cpp | 14 +--- modules/core/c++/hlbm.hpp | 2 +- modules/core/c++/particle/aabb.hpp | 14 ++-- modules/core/c++/particle/particle.hpp | 6 +- modules/core/c++/particle/particle_opa.hpp | 2 +- modules/core/c++/particle/porosity.hpp | 12 ++-- modules/core/c++/particle/schema.hpp | 4 +- modules/core/tests/particles.cpp | 10 +-- 9 files changed, 79 insertions(+), 66 deletions(-) diff --git a/examples/moving_poiseulle_particles_2d_hlbm_gpu/sim.cpp b/examples/moving_poiseulle_particles_2d_hlbm_gpu/sim.cpp index daa397a..75dd666 100644 --- a/examples/moving_poiseulle_particles_2d_hlbm_gpu/sim.cpp +++ b/examples/moving_poiseulle_particles_2d_hlbm_gpu/sim.cpp @@ -57,14 +57,14 @@ using MacroStruct = Struct< //>; template -using ParticleSpheroidGroup = ParticleGroup>; +using ParticleSpheroidGroup = ParticleGroup>; } template saw::error_or setup_initial_conditions( - saw::data>& fields, - saw::data>& macros, - saw::data>& particles + saw::data>& fields, + saw::data>& macros, + saw::data>& particles ){ auto& info_f = fields.template get<"info">(); auto& porous_f = macros.template get<"porosity">(); @@ -145,6 +145,12 @@ saw::error_or setup_initial_conditions( {{1u,1u}} ); + saw::data> rad; + rad.at({}) = 2.0f; + saw::data> dense; + dense.at({}) = 1.0f; + particles = create_spheroid_particle_group(rad,dense,{16u}); + return saw::make_void(); } @@ -152,7 +158,7 @@ template saw::error_or step( saw::data>,encode::Sycl>& fields, saw::data>,encode::Sycl>& macros, - saw::data>,encode::Sycl>& particles + saw::data>,encode::Sycl>& particles, saw::data t_i, device& dev ){ @@ -160,26 +166,35 @@ saw::error_or step( auto& info_f = fields.template get<"info">(); auto& porous_f = macros.template get<"porosity">(); + q.submit([&](acpp::sycl::handler& h){ + component> hlbm_one_part; + + h.parallel_for(acpp::sycl::range<1u>{1u}, [=](acpp::sycl::id<1u> idx){ + saw::data> index; + for(uint64_t i = 0u; i < 1u; ++i){ + index.at({{i}}).set(idx[i]); + } + + hlbm_one_part.apply(fields,macros,particles,index,t_i); + + // auto aabb = particle_aabb<>::cacl(); + + }); + }).wait(); + // auto coll_ev = q.submit([&](acpp::sycl::handler& h){ component> collision{0.65}; component> bb; component,encode::Sycl> abb; - saw::data> rho_b; - rho_b.at({}) = 1.0; - saw::data> vel_b; - vel_b.at({{0u}}) = 0.015; - - component> equi{rho_b,vel_b}; - component,encode::Sycl> flow_in{ [&](){ - uint64_t target_t_i = 64u; + uint64_t target_t_i = 8u; if(t_i.get() < target_t_i){ - return 1.0 + (0.0002 / target_t_i) * t_i.get(); + return 1.0 + (0.0005 / target_t_i) * t_i.get(); } - return 1.0002; + return 1.0005; }() }; component,encode::Sycl> flow_out{1.0}; @@ -265,6 +280,8 @@ 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_parts_data_ptr = saw::heap>>(); + //saw::data>,encode::Sycl>& particles 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); + auto eov = setup_initial_conditions(*lbm_data_ptr,*lbm_macro_data_ptr,*lbm_parts_data_ptr); if(eov.is_error()){ return eov; } @@ -304,6 +321,7 @@ saw::error_or lbm_main(int argc, char** argv){ saw::data, encode::Sycl> lbm_sycl_data{sycl_q}; saw::data, encode::Sycl> lbm_sycl_macro_data{sycl_q}; + saw::data, encode::Sycl> lbm_sycl_parts_data{sycl_q}; sycl_q.wait(); { @@ -318,35 +336,30 @@ saw::error_or lbm_main(int argc, char** argv){ return eov; } } + { + auto eov = dev.copy_to_device(*lbm_parts_data_ptr,lbm_sycl_parts_data); + if(eov.is_error()){ + return eov; + } + } sycl_q.wait(); auto lsd_view = make_view(lbm_sycl_data); auto lsdm_view = make_view(lbm_sycl_macro_data); + auto lsdp_view = make_view(lbm_sycl_parts_data); + saw::data time_steps{16u*4096ul}; auto& info_f = lsd_view.template get<"info">(); for(saw::data i{0u}; i < time_steps and krun; ++i){ // BC + Collision { - auto eov = step(lsd_view,lsdm_view,i,dev); + auto eov = step(lsd_view,lsdm_view,lsdp_view,i,dev); if(eov.is_error()){ return eov; } } sycl_q.wait(); - if(i.get() % 32u == 0u){ - { - auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr); - if(eov.is_error()){ - return eov; - } - } - { - auto eov = write_csv_file(out_dir,"m",i.get(), *lbm_macro_data_ptr); - if(eov.is_error()){ - return eov; - } - } - } + // Stream sycl_q.submit([&](acpp::sycl::handler& h){ component> stream; @@ -386,6 +399,12 @@ saw::error_or lbm_main(int argc, char** argv){ return eov; } } + { + auto eov = write_csv_file(out_dir,"m",time_steps.get(), *lbm_macro_data_ptr); + if(eov.is_error()){ + return eov; + } + } sycl_q.wait(); return saw::make_void(); diff --git a/examples/settling_spheres_2d_hlbm_gpu/sim.cpp b/examples/settling_spheres_2d_hlbm_gpu/sim.cpp index a004a92..f7e47b6 100644 --- a/examples/settling_spheres_2d_hlbm_gpu/sim.cpp +++ b/examples/settling_spheres_2d_hlbm_gpu/sim.cpp @@ -13,7 +13,6 @@ namespace lbm { constexpr uint64_t dim_x = 512ul; constexpr uint64_t dim_y = dim_x * 2ul; - constexpr uint64_t particle_amount = 16ul; namespace sch { @@ -59,15 +58,9 @@ template using ParticleSpheroidGroup = ParticleGroup< T, Desc::D, - sch::ParticleCollisionSpheroid ->; - -template -using ParticleGroups = Tuple< - ParticleSpheroidGroup + particle_amount, + coll::Spheroid >; - - } template @@ -186,9 +179,6 @@ saw::error_or step( for(uint64_t i = 0u; i < Desc::D; ++i){ index.at({{i}}).set(idx[i]); } - - // Reset the force to zero - forces.at(index) = {}; }); }).wait(); diff --git a/modules/core/c++/hlbm.hpp b/modules/core/c++/hlbm.hpp index 9a89697..9ddc6c6 100644 --- a/modules/core/c++/hlbm.hpp +++ b/modules/core/c++/hlbm.hpp @@ -125,7 +125,7 @@ public: auto& mvel = macros.template get<"velocity">(); { auto& parts = part_spheroid_group.template get<"particles">(); - auto parts_size = parts.meta().at({0u}); + auto parts_size = parts.dims().at({0u}); auto& pi = parts.at(index); auto& pirb = pi.template get<"rigid_body">(); diff --git a/modules/core/c++/particle/aabb.hpp b/modules/core/c++/particle/aabb.hpp index 18426d5..a838401 100644 --- a/modules/core/c++/particle/aabb.hpp +++ b/modules/core/c++/particle/aabb.hpp @@ -10,12 +10,12 @@ class particle_aabb final { static_assert(saw::always_false, "Not supported"); }; -template +template class particle_aabb< - sch::ParticleGroup> + sch::ParticleGroup> > final { public: - using Schema = sch::ParticleGroup>; + using Schema = sch::ParticleGroup>; using AABB = sch::Struct< sch::Member,"a">, @@ -23,7 +23,11 @@ public: >; public: template - static constexpr saw::data calculate(const saw::data& p_grp, const saw::data>& i, const saw::data>& meta){ + static constexpr saw::error_or> calculate(const saw::data& p_grp, const saw::data>& i, const saw::data>& meta){ + static_assert(PC > 0u, "Can't calculate from no particles"); + if(not (i.at({0u}).get() < PC) ){ + return saw::make_error("Too large i in particle_aabb"); + } saw::data aabb; auto& parts = p_grp.template get<"particles">(); @@ -35,7 +39,7 @@ public: auto& a = aabb.template get<"a">(); auto& b = aabb.template get<"b">(); - const saw::data>& rad_d = p_grp.template get<"collision">().at({0u}).template get<"radius">(); + const saw::data>& rad_d = p_grp.template get<"collision">().at(i).template get<"radius">(); saw::data> lower; saw::data> upper; diff --git a/modules/core/c++/particle/particle.hpp b/modules/core/c++/particle/particle.hpp index 8e75e5a..3c3630a 100644 --- a/modules/core/c++/particle/particle.hpp +++ b/modules/core/c++/particle/particle.hpp @@ -13,13 +13,13 @@ namespace kel { namespace lbm { -template -saw::data>> create_spheroid_particle_group( +template +saw::data>> create_spheroid_particle_group( saw::data> radius_p, saw::data> density_p, const saw::data& mask_resolution ){ - saw::data>> part; + saw::data>> part; auto& rad_s = part.template get<"collision">().at({0u}).template get<"radius">(); rad_s = radius_p; diff --git a/modules/core/c++/particle/particle_opa.hpp b/modules/core/c++/particle/particle_opa.hpp index 30a8ade..ce1ecb7 100644 --- a/modules/core/c++/particle/particle_opa.hpp +++ b/modules/core/c++/particle/particle_opa.hpp @@ -41,7 +41,7 @@ public: // Write out value auto diff = pos_ind.template cast_to() - pos_; auto diff_dot = saw::math::dot(diff,diff); - porous = particle_porosity>::calculate(diff, rad_, eps_); + porous = particle_porosity>::calculate(diff, rad_, eps_); } }; } diff --git a/modules/core/c++/particle/porosity.hpp b/modules/core/c++/particle/porosity.hpp index 2f39d2e..353cd58 100644 --- a/modules/core/c++/particle/porosity.hpp +++ b/modules/core/c++/particle/porosity.hpp @@ -11,11 +11,11 @@ template struct ParticleSpheroid {}; } -template +template class particle_porosity { public: - static saw::data> calculate(const saw::data>& part_group, uint64_t p_i, const saw::data>& lbm_pos){ + static saw::data> calculate(const saw::data>& part_group, uint64_t p_i, const saw::data>& lbm_pos){ auto& mask = part_group.template get<"mask">(); auto& particles = part_group.template get<"particles">(); @@ -33,7 +33,7 @@ public: }; template -class particle_porosity> final { +class particle_porosity> final { public: static saw::data> calculate(const saw::data>& lbm_rel_dist, saw::data> rad, saw::data> eps){ saw::data> por; @@ -68,10 +68,10 @@ public: }; -template -class particle_porosity> final { +template +class particle_porosity> final { public: - static saw::data> calculate(const saw::data > >& part_group, uint64_t i, const saw::data>& lbm_pos) { + static saw::data> calculate(const saw::data > >& part_group, uint64_t i, const saw::data>& lbm_pos) { saw::data> por; auto& parts = part_group.template get<"particles">(); diff --git a/modules/core/c++/particle/schema.hpp b/modules/core/c++/particle/schema.hpp index 714a16f..221137a 100644 --- a/modules/core/c++/particle/schema.hpp +++ b/modules/core/c++/particle/schema.hpp @@ -52,7 +52,7 @@ using Particle = Struct< // Member, "mask">, >; -template> +template> using ParticleGroup = Struct< Member, "mask">, Member, "collision">, @@ -61,7 +61,7 @@ using ParticleGroup = Struct< Member,1u>, "density">, Member,1u>, "center_of_mass">, Member,1u>, "total_mass">, - Member,1u>, "particles"> + Member,PartAmount>, "particles"> >; } } diff --git a/modules/core/tests/particles.cpp b/modules/core/tests/particles.cpp index d11fa4f..e060097 100644 --- a/modules/core/tests/particles.cpp +++ b/modules/core/tests/particles.cpp @@ -278,15 +278,14 @@ SAW_TEST("Spheroid Particle / AABB"){ using T = sch::Float64; constexpr uint64_t D = 2u; + constexpr uint64_t PC = 1u; saw::data> radi,dense; radi.at({}) = 2.2f; dense.at({}) = 1.0f; - auto spheroid_pgrp = lbm::create_spheroid_particle_group(radi,dense,{8u}); + auto spheroid_pgrp = lbm::create_spheroid_particle_group(radi,dense,{8u}); auto& rb_array = spheroid_pgrp.template get<"particles">(); - rb_array = {1u}; - auto& rb = rb_array.at({0u}).template get<"rigid_body">(); auto& rb_pos = rb.template get<"position">(); @@ -295,8 +294,9 @@ SAW_TEST("Spheroid Particle / AABB"){ rb_pos.at({{1u}}).set(2.9f); } - auto res = lbm::particle_aabb>>::calculate(spheroid_pgrp,{{0u}},{{64u,32u}}); - + auto eo_res = lbm::particle_aabb>>::calculate(spheroid_pgrp,{{0u}},{{64u,32u}}); + SAW_EXPECT(eo_res.is_value(), "particle_aabb has error value"); + auto& res = eo_res.get_value(); { auto res_check = [&](const auto& tbp){ lbm::iterator<1u>::apply([&](const auto& index){ -- cgit v1.2.3 From aedeb28b41ffbe0229d065c69f7d2e4570968ace Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Wed, 15 Jul 2026 16:55:56 +0200 Subject: Reworking particle_aabb --- modules/core/c++/hlbm.hpp | 4 +++- modules/core/c++/particle/aabb.hpp | 5 +++-- 2 files changed, 6 insertions(+), 3 deletions(-) diff --git a/modules/core/c++/hlbm.hpp b/modules/core/c++/hlbm.hpp index 9ddc6c6..eccc2cd 100644 --- a/modules/core/c++/hlbm.hpp +++ b/modules/core/c++/hlbm.hpp @@ -105,7 +105,7 @@ class component final { private: template void apply_i(const saw::data& field, const saw::data& macros, const saw::data& part_groups, saw::data> index, saw::data time_step) const { - + // if constexpr ( i < ) } public: /* @@ -136,6 +136,8 @@ public: auto aabb = particle_aabb::calculate(part_spheroid_group,index,mvel.meta()); /// Ok, I iterate over the space which covers our particle? So lower bounds to upper bounds + start = aabb.template get<"a">(); + stop = aabb.template get<"b">(); iterator::apply([&](const auto& index){ // ask for the d_k value here. diff --git a/modules/core/c++/particle/aabb.hpp b/modules/core/c++/particle/aabb.hpp index a838401..983fd97 100644 --- a/modules/core/c++/particle/aabb.hpp +++ b/modules/core/c++/particle/aabb.hpp @@ -22,8 +22,8 @@ public: sch::Member,"b"> >; public: - template - static constexpr saw::error_or> calculate(const saw::data& p_grp, const saw::data>& i, const saw::data>& meta){ + template + static constexpr saw::error_or> calculate(const saw::data& p_grp, const saw::data>& i, const saw::data>& meta){ static_assert(PC > 0u, "Can't calculate from no particles"); if(not (i.at({0u}).get() < PC) ){ return saw::make_error("Too large i in particle_aabb"); @@ -55,5 +55,6 @@ public: } }; + } } -- cgit v1.2.3