From 30ff1caf073b4341fd0614e0974c67a8588c8931 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Wed, 15 Apr 2026 19:11:21 +0200 Subject: Feierabend --- examples/settling_cubes_2d_ibm_gpu/sim.cpp | 158 ++++++++++++++++++++++++----- 1 file changed, 131 insertions(+), 27 deletions(-) (limited to 'examples/settling_cubes_2d_ibm_gpu') diff --git a/examples/settling_cubes_2d_ibm_gpu/sim.cpp b/examples/settling_cubes_2d_ibm_gpu/sim.cpp index e1ba012..de40c54 100644 --- a/examples/settling_cubes_2d_ibm_gpu/sim.cpp +++ b/examples/settling_cubes_2d_ibm_gpu/sim.cpp @@ -6,6 +6,7 @@ #include #include #include +#include namespace kel { namespace lbm { @@ -49,19 +50,26 @@ using MacroStruct = Struct< Member, "force"> >; +template +using ParticleSpheroidGroup = ParticleGroup< + T, + Desc::D, + sch::ParticleCollisionSpheroid +>; + template using ParticleGroups = Tuple< - ParticleGroup< - T,Desc::D,sch::ParticleCollisionSpheroid - > + ParticleSpheroidGroup >; + + } template saw::error_or setup_initial_conditions( saw::data>& fields, saw::data>& macros, - saw::data>& particles + saw::data>& particles ){ auto& info_f = fields.template get<"info">(); // Set everything as walls @@ -105,11 +113,39 @@ saw::error_or setup_initial_conditions( { saw::data> dense_p; dense_p.at({}).set(1); - auto& spheroid_group = particles.template get<0u>(); + // auto& spheroid_group = particles.template get<0u>(); + auto& spheroid_group = particles; + spheroid_group = create_spheroid_particle_group( dense_p, {64u} ); + + { + auto& p = spheroid_group.template get<"particles">(); + + p = {{{16u}}}; + + iterator<1u>::apply( + [&](auto& index){ + // Set Pos here? + auto& p_ind = p.at(index); + + auto& p_rb = p_ind.template get<"rigid_body">(); + auto& p_pos = p_rb.template get<"position">(); + + // TODO CONTINUE HERE NEED to init pos here !!!! + + auto& p_pos_old = p_rb.template get<"position_old">(); + p_pos_old = p_pos; + }, + {}, + p.meta() + ); + } + } + // Particle in hacky flavour + { } return saw::make_void(); @@ -119,20 +155,98 @@ template saw::error_or step( saw::data>,encode::Sycl>& fields, saw::data>,encode::Sycl>& macros, - saw::data>,encode::Sycl>& particles, + saw::data>,encode::Sycl>& p_group, saw::data t_i, device& dev ){ auto& q = dev.get_handle(); auto& info_f = fields.template get<"info">(); - { - } + auto& parts = p_group.template get<"particles">(); + auto& p_mask = p_group.template get<"mask">(); + auto& vels = macros.template get<"velocity">(); + auto& forces = macros.template get<"force">(); + + auto p_meta = parts.meta(); + q.submit([&](acpp::sycl::handler& h){ + + h.parallel_for(acpp::sycl::range{dim_x,dim_y}, [=](acpp::sycl::id idx){ + saw::data> index; + for(uint64_t i = 0u; i < Desc::D; ++i){ + index.at({{i}}).set(idx[i]); + } + + // Reset the force to zero + forces.at(index) = {}; + }); + }).wait(); + + q.submit([&](acpp::sycl::handler& h){ + h.parallel_for(acpp::sycl::range<1u>{p_meta.at({0u}).get()}, [=](acpp::sycl::id<1u> idx){ + + saw::data> index; + for(uint64_t i = 0u; i < 1u; ++i){ + index.at({{i}}).set(idx[i]); + } + + auto& p = parts.at(index); + auto& p_rb = p.template get<"rigid_body">(); + saw::data> delta_t; + delta_t.at({}).set(1.0f); + + auto& p_pos = p_rb.template get<"position">(); + auto& p_rot = p_rb.template get<"rotation">(); + + iterator::apply( + [&](auto& m_ind){ + saw::data> index_shift; + for(uint64_t i{0u}; i < Desc::D; ++i){ + index_shift.at({{i}}) = m_ind.at({i}).template cast_to() - (p_mask.meta().at({i})+1u).template cast_to() * 0.5; + } + + saw::data> transformed_pos; + for(uint64_t i{0u}; i < Desc::D; ++i){ + // TODO add rotation, scaling here. + transformed_pos.at({{i}}) = index_shift.at({{i}}); + } + + // Lagrange indicator position + auto p_pos_lag = p_pos + transformed_pos; + + // Pick the closest velocity + saw::data> p_cell_pos; + saw::data> p_cell_pos_vec; + for(uint64_t i{0u}; i < Desc::D; ++i){ + p_cell_pos.at({{i}}) = (p_pos_lag.at({{i}}) + 0.5).template cast_to(); + p_cell_pos.at({{i}}).set(std::max(1ul,std::min(p_cell_pos.at({{i}}).get(), p_meta.at({{i}}).get() - 2ul))); + p_cell_pos_vec.at({{i}}) = p_cell_pos.at({{i}}); + } + + auto& u_fluid = vels.at(p_cell_pos); + + // this is our relative position to the particle + auto rel_cell_to_part_pos = p_cell_pos_vec.template cast_to() - p_pos; + + auto p_vel = (p_pos - p_rb.template get<"position_old">()) * delta_t; + auto u_solid = p_vel + saw::math::cross(p_rot,rel_cell_to_part_pos); + + + // Force + auto force = (u_solid - u_fluid) / delta_t; + + // TODO HERE ATOMIC! !!!! + forces.at(p_cell_pos) = forces.at(p_cell_pos) + force; + }, + {}, + p_mask.meta() + ); + + verlet_step_lambda(p,delta_t); + }); + }).wait(); // auto coll_ev = q.submit([&](acpp::sycl::handler& h){ - saw::data> force; - force.at({{1}}).set(-1.0); // Need nicer things to handle the flow. I see improvement here component> collision{0.8}; component> bb; @@ -157,15 +271,10 @@ saw::error_or step( default: break; } - }); - - }).wait(); - - q.submit([&](acpp::sycl::handler& h){ - h.parallel_for(acpp::sycl::range<1u>{dim_x}, [=](acpp::sycl::id<1u> idx){ }); }).wait(); + // Step /* q.submit([&](acpp::sycl::handler& h){ @@ -212,9 +321,11 @@ 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_ptr = saw::heap>>(); + auto lbm_particle_data_ptr = saw::heap>>(); + // auto lbm_particles_ptr = saw::heap,part_count>>>(); + // saw::data> p_mask; - std::cout<<"Estimated Bytes: "<,sch::MacroStruct>().get()<,sch::MacroStruct>().get()< 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_data{sycl_q}; - { - auto eov = dev.malloc_on_device(*lbm_particle_data_ptr,lbm_sycl_particle_data); - if(eov.is_error()){ - return eov; - } - } - sycl_q.wait(); + saw::data, encode::Sycl> lbm_sycl_particle_data{sycl_q}; { auto eov = dev.copy_to_device(*lbm_data_ptr,lbm_sycl_data); @@ -285,7 +389,7 @@ saw::error_or lbm_main(int argc, char** argv){ auto lsdm_view = make_view(lbm_sycl_macro_data); auto lsdp_view = make_view(lbm_sycl_particle_data); - saw::data time_steps{16u*4096ul}; + saw::data time_steps{4u*4096ul}; auto& info_f = lsd_view.template get<"info">(); -- cgit v1.2.3