From 4f3f21f72602ab4ce49caa6e7264d369afc98977 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Mon, 18 May 2026 18:55:12 +0200 Subject: Particle --- lib/core/c++/iterator.hpp | 28 ++++++++++++++++++++++++++++ lib/core/c++/particle/particle.hpp | 7 +++++++ lib/core/tests/iterator.cpp | 17 +++++++++++++++++ 3 files changed, 52 insertions(+) diff --git a/lib/core/c++/iterator.hpp b/lib/core/c++/iterator.hpp index 866543a..696c521 100644 --- a/lib/core/c++/iterator.hpp +++ b/lib/core/c++/iterator.hpp @@ -50,6 +50,34 @@ public: } }; +template +struct ct_tuple_iterator final { +private: + template + static constexpr saw::error_or apply_i( + Func& func, + saw::data>& tup + ){ + if constexpr ( i < sizeof...(Tup) ){ + auto eov = func(tup.template get()); + if(eov.is_error()){ + return eov; + } + return apply_i(func,tup); + } + + return saw::make_void(); + } +public: + template + static constexpr saw::error_or apply( + Func&& func, + saw::data>& tup + ){ + return apply_i(func,tup); + } +}; + /* Ambiguous template void iterate_over(Func&& func, const saw::data>& start, const saw::data>& end, const saw::data>& dist = {{{0u,0u,0u}}}){ diff --git a/lib/core/c++/particle/particle.hpp b/lib/core/c++/particle/particle.hpp index 7af43dc..01429b2 100644 --- a/lib/core/c++/particle/particle.hpp +++ b/lib/core/c++/particle/particle.hpp @@ -238,5 +238,12 @@ constexpr auto broadphase_collision_check = []( ) -> bool{ return broadphase_collision_distance_squared(left,coll_l,right,coll_r).first; }; + + + +namespace impl { + +} + } } diff --git a/lib/core/tests/iterator.cpp b/lib/core/tests/iterator.cpp index 261765a..919c12c 100644 --- a/lib/core/tests/iterator.cpp +++ b/lib/core/tests/iterator.cpp @@ -20,6 +20,23 @@ SAW_TEST("Old Iterate"){ }, start, end); } +SAW_TEST("CT Tuple Iterator"){ + using namespace kel; + + saw::data< + sch::Tuple< + sch::Float32, + sch::Int16, + sch::Float64, + sch::UInt32 + > + > tup; + + lbm::ct_tuple_iterator::Schema>::apply([](auto& val) -> saw::error_or { + return saw::make_void(); + },tup); +} + SAW_TEST("Old Iterate with Distance 1"){ using namespace kel; -- cgit v1.2.3 From 81b7d6d2c09bbff6d9c1738fbb52f7c9a90b10de Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Tue, 19 May 2026 15:04:59 +0200 Subject: IBM reimplement --- examples/poiseulle_particles_2d_ibm_gpu/sim.cpp | 87 ++++++++++++++++++------- examples/settling_cubes_2d_ibm_gpu/sim.cpp | 6 +- lib/core/c++/iterator.hpp | 5 +- lib/core/c++/term_renderer.hpp | 6 -- 4 files changed, 72 insertions(+), 32 deletions(-) delete mode 100644 lib/core/c++/term_renderer.hpp diff --git a/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp b/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp index 9905906..ea49d51 100644 --- a/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp @@ -54,12 +54,21 @@ using MacroStruct = Struct< //using ParticleArray = Array< // Particle //>; + +template +using ParticleSpheroidGroup = ParticleGroup< + T, + Desc::D, + sch::ParticleCollisionSpheroid +>; + } template saw::error_or setup_initial_conditions( saw::data>& fields, - saw::data>& macros + saw::data>& macros, + saw::data>& particles ){ auto& info_f = fields.template get<"info">(); // Set everything as walls @@ -136,25 +145,41 @@ saw::error_or setup_initial_conditions( df_f.get_dims(), {{1u,1u}} ); - - iterator::apply( - [&](auto& index){ - saw::data> middle, ind_vec; - middle.at({{0u}}) = dim_x * 0.25; - middle.at({{1u}}) = dim_y * 0.5; - ind_vec.at({{0u}}) = index.at({{0u}}).template cast_to(); - ind_vec.at({{1u}}) = index.at({{1u}}).template cast_to(); + // Particles + { + saw::data> dense_p; + dense_p.at({}).set(1); - auto dist = middle - ind_vec; - auto dist_2 = saw::math::dot(dist,dist); - if(dist_2.at({}).get() < dim_y*dim_y*0.01){ - info_f.at(index) = 1u; - } - }, - {},// 0-index - df_f.get_dims() - ); + auto& spheroid_grp = particles; + + spheroid_grp = create_spheroid_particle_group( + dense_p, + {32u} + ); + + { + auto& p = spheroid_grp.template get<"particles">(); + p = {{{1u}}}; + + iterator<1u>::apply( + [&](auto& index){ + auto& p_ind = p.at(index); + auto& p_rb = p_ind.template get<"rigid_body">(); + auto& p_pos = p_rb.template get<"position">(); + auto& p_pos_old = p_rb.template get<"position_old">(); + + // Set position + p_pos.at({{0u}}) = dim_x * 0.25; + p_pos.at({{1u}}) = dim_y * 0.5; + + p_pos_old = p_pos; + }, + {}, + p.meta() + ); + } + } return saw::make_void(); } @@ -163,6 +188,7 @@ 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 ){ @@ -172,9 +198,10 @@ saw::error_or step( // auto coll_ev = q.submit([&](acpp::sycl::handler& h){ // Need nicer things to handle the flow. I see improvement here - saw::data> f; - f.at({{0u}}) = 0.0; - f.at({{1u}}) = -1.0; + saw::data> f; + f.at({{0u}}) = 0.0; + f.at({{1u}}) = -1.0; + component> collision{0.65,f}; component> bb; component,encode::Sycl> abb; @@ -237,6 +264,11 @@ saw::error_or step( // h.depends_on(collision_ev); }).wait(); */ + q.submit([&](acpp::sycl::handler& h){ + h.parallel_for(acpp::sycl::range<1u>{1u}, [=](acpp::sycl::id<1u> idx){ + + }); + }).wait(); return saw::make_void(); } @@ -285,6 +317,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_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); + auto eov = setup_initial_conditions(*lbm_data_ptr,*lbm_macro_data_ptr,*lbm_particle_ptr); if(eov.is_error()){ return eov; } @@ -324,6 +357,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_particle{sycl_q}; sycl_q.wait(); { @@ -338,9 +372,16 @@ saw::error_or lbm_main(int argc, char** argv){ return eov; } } + { + auto eov = dev.copy_to_device(*lbm_particle_ptr,lbm_sycl_particle); + 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_particle); saw::data time_steps{16u*4096ul}; auto& info_f = lsd_view.template get<"info">(); @@ -348,7 +389,7 @@ saw::error_or lbm_main(int argc, char** argv){ 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; } diff --git a/examples/settling_cubes_2d_ibm_gpu/sim.cpp b/examples/settling_cubes_2d_ibm_gpu/sim.cpp index d7b402a..9fdea8c 100644 --- a/examples/settling_cubes_2d_ibm_gpu/sim.cpp +++ b/examples/settling_cubes_2d_ibm_gpu/sim.cpp @@ -145,8 +145,7 @@ saw::error_or setup_initial_conditions( } } // Particle in hacky flavour - { - } + {} return saw::make_void(); } @@ -197,6 +196,8 @@ saw::error_or step( auto& p_pos = p_rb.template get<"position">(); auto& p_rot = p_rb.template get<"rotation">(); + auto& p_acc = p_rb.template get<"acceleration">(); + iterator::apply( [&](auto& m_ind){ saw::data> index_shift; @@ -238,6 +239,7 @@ saw::error_or step( forces.at(p_cell_pos) = forces.at(p_cell_pos) + force; // TODO APPLY FORCE TO PARTICLE + p_acc = p_acc + force; // TODO divide by mass }, {}, p_mask.meta() diff --git a/lib/core/c++/iterator.hpp b/lib/core/c++/iterator.hpp index 696c521..7fd6f58 100644 --- a/lib/core/c++/iterator.hpp +++ b/lib/core/c++/iterator.hpp @@ -50,8 +50,11 @@ public: } }; +template +struct ct_tuple_iterator; + template -struct ct_tuple_iterator final { +struct ct_tuple_iterator> final { private: template static constexpr saw::error_or apply_i( diff --git a/lib/core/c++/term_renderer.hpp b/lib/core/c++/term_renderer.hpp deleted file mode 100644 index 5cbb551..0000000 --- a/lib/core/c++/term_renderer.hpp +++ /dev/null @@ -1,6 +0,0 @@ -#pragma once - -namespace kel { -namespace lbm { -} -} -- cgit v1.2.3 From 8e22658bf3791fb5b1f9a11b5a5408d5e4e3f732 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Tue, 19 May 2026 15:05:59 +0200 Subject: Adding particle to hlbm --- lib/core/c++/hlbm.hpp | 25 +++++++++++++++++++++++++ 1 file changed, 25 insertions(+) diff --git a/lib/core/c++/hlbm.hpp b/lib/core/c++/hlbm.hpp index 41e2e27..7af4f7e 100644 --- a/lib/core/c++/hlbm.hpp +++ b/lib/core/c++/hlbm.hpp @@ -8,6 +8,7 @@ namespace kel { namespace lbm { namespace cmpt { struct HLBM {}; +struct HlbmParticle {}; } /** @@ -65,5 +66,29 @@ public: } }; +template +class component final { +private: + typename saw::native_data_type::type relaxation_; + saw::data frequency_; +public: + + template + void apply(const saw::data& field, const saw::data& macros, const saw::data& particles, saw::data> index, saw::data time_step) const { + /// Figure out how to access the particle list + auto& p = particles.at(i); + + /// Iterate over the grid bounds + auto& grid = p.template get<"grid">(); + + iterator::apply([&](const auto& index){ + auto& g_i = grid.at(index); + + /// Pick up forces and calculate to rigid body TODO + + }); + + } +}; } } -- cgit v1.2.3 From f07f88088cfacaad3f94dc667f08ee2c2f38a093 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Wed, 20 May 2026 14:06:41 +0200 Subject: Particle issues in hlbm --- lib/core/c++/hlbm.hpp | 17 +++++------------ 1 file changed, 5 insertions(+), 12 deletions(-) diff --git a/lib/core/c++/hlbm.hpp b/lib/core/c++/hlbm.hpp index 7af4f7e..196de73 100644 --- a/lib/core/c++/hlbm.hpp +++ b/lib/core/c++/hlbm.hpp @@ -7,7 +7,7 @@ namespace kel { namespace lbm { namespace cmpt { -struct HLBM {}; +struct Hlbm {}; struct HlbmParticle {}; } @@ -15,7 +15,7 @@ struct HlbmParticle {}; * HLBM collision operator for LBM */ template -class component final { +class component final { private: typename saw::native_data_type::type relaxation_; saw::data frequency_; @@ -73,20 +73,13 @@ private: saw::data frequency_; public: - template + template void apply(const saw::data& field, const saw::data& macros, const saw::data& particles, saw::data> index, saw::data time_step) const { /// Figure out how to access the particle list - auto& p = particles.at(i); + // auto& p = particles.at(i); /// Iterate over the grid bounds - auto& grid = p.template get<"grid">(); - - iterator::apply([&](const auto& index){ - auto& g_i = grid.at(index); - - /// Pick up forces and calculate to rigid body TODO - - }); + // auto& grid = p.template get<"grid">(); } }; -- cgit v1.2.3 From b5d8593b9a2f0f58cb228444dcd09a2c5002e039 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Thu, 21 May 2026 15:00:09 +0200 Subject: Prep for proper com in circle group --- examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp | 2 +- examples/poiseulle_particles_2d_ibm_gpu/sim.cpp | 35 +++++++++++++++++++++++- lib/core/c++/math/n_linear.hpp | 3 ++ lib/core/c++/particle/particle.hpp | 14 +++++++++- 4 files changed, 51 insertions(+), 3 deletions(-) diff --git a/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp b/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp index e12b0d8..9375078 100644 --- a/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp @@ -176,7 +176,7 @@ saw::error_or step( // auto coll_ev = q.submit([&](acpp::sycl::handler& h){ - component> collision{0.65}; + component> collision{0.65}; component> bb; component,encode::Sycl> abb; diff --git a/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp b/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp index ea49d51..e68d7da 100644 --- a/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp @@ -1,6 +1,7 @@ #include #include #include +#include #include #include @@ -266,7 +267,39 @@ saw::error_or step( */ q.submit([&](acpp::sycl::handler& h){ h.parallel_for(acpp::sycl::range<1u>{1u}, [=](acpp::sycl::id<1u> idx){ - + auto& vel = macros.template get<"velocity">(); + + auto& ps = particles; + auto& mask = ps.template get<"mask">(); + auto& dense = ps.template get<"density">().at({}); + auto& com = ps.template get<"center_of_mass">(); + + auto& parts = ps.template get<"particles">(); + + auto& p_i = parts.at({{0u}}); + + auto& p_i_rb = p_i.template get<"rigid_body">(); + /// 0. Iterate over mask and calculate position in LBM grid + /// In this case it's simple since I'm too lazy to do scaling and rotation + /// Technically scale => rotate => translate + /// Here it's only translate + auto& p_i_rb_pos = p_i_rb.template get<"position">(); + + iterator::apply([&](const auto& index){ + /// Calculate the shift from the mask + saw::data> index_shift; + for(uint64_t i = 0u; i < Desc::D; ++i){ + index_shift.at({{i}}) = index.at({i}).template cast_to() - com.at({{i}}); + } + + + /// TODO 1. Calculate force pickup from neigbouring u_vel cells + auto inter_vel = n_linear_interpolate(vel,index_shift); + + /// TODO 3. Distribute force to fluid + + + }, {}, mask.meta()); }); }).wait(); diff --git a/lib/core/c++/math/n_linear.hpp b/lib/core/c++/math/n_linear.hpp index 8fb0600..b378440 100644 --- a/lib/core/c++/math/n_linear.hpp +++ b/lib/core/c++/math/n_linear.hpp @@ -122,6 +122,9 @@ auto n_linear_interpolate( } }, {}, ones_ind); + + /// TODO I need to actually calc stuff + return field.at({}); } template diff --git a/lib/core/c++/particle/particle.hpp b/lib/core/c++/particle/particle.hpp index 01429b2..5110893 100644 --- a/lib/core/c++/particle/particle.hpp +++ b/lib/core/c++/particle/particle.hpp @@ -55,6 +55,8 @@ template, "mask">, Member,1u>, "density">, + Member, "center_of_mass">, + Member, "total_mass">, Member>, "particles"> >; } @@ -76,10 +78,20 @@ saw::data>> cre for(uint64_t i = 0u; i < D; ++i){ mask_dims.at({i}) = mask_resolution; } + saw::data> mask_step; + saw::data dia_d{radius*2}; + mask_step.at({}) = dia_d / mask_resolution.template cast_to(); mask = {mask_dims}; + + auto& com = part.template get<"center_of_mass">(); + for(uint64_t i = 0u; i < D; ++i){ + com.at({{i}}) = {}; + } + saw::data ele_ctr{0u}; iterator::apply([&](const auto& index){ - + ++ele_ctr; + },{},mask_dims); return part; -- cgit v1.2.3