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 +- 2 files changed, 68 insertions(+), 25 deletions(-) (limited to 'examples') 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() -- 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 +++++++++++++++++++++++- 2 files changed, 35 insertions(+), 2 deletions(-) (limited to 'examples') 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(); -- cgit v1.2.3