From 0a11f4fbdbdf62d153904c9ed7de6c82a8916b7d Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Thu, 26 Mar 2026 12:08:19 +0100 Subject: Dangling changes. Restructuring particle setup --- .../.nix/derivation.nix | 2 +- examples/poiseulle_particles_2d_bgk_gpu/sim.cpp | 24 ++---------- examples/poiseulle_particles_2d_psm_gpu/sim.cpp | 43 ++-------------------- examples/settling_cubes_2d_ibm_gpu/sim.cpp | 35 +++++++++++------- lib/core/c++/boundary.hpp | 8 ++-- lib/core/c++/particle/particle.hpp | 29 +++++++++------ 6 files changed, 51 insertions(+), 90 deletions(-) diff --git a/examples/poiseulle_particles_2d_bgk_gpu/.nix/derivation.nix b/examples/poiseulle_particles_2d_bgk_gpu/.nix/derivation.nix index 517b6b8..8aca15e 100644 --- a/examples/poiseulle_particles_2d_bgk_gpu/.nix/derivation.nix +++ b/examples/poiseulle_particles_2d_bgk_gpu/.nix/derivation.nix @@ -11,7 +11,7 @@ }: stdenv.mkDerivation { - pname = pname + "-examples-" + "poiseulle_2d_gpu"; + pname = pname + "-examples-" + "poiseulle_particles_2d_bgk_gpu"; inherit version; src = ./..; diff --git a/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp b/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp index 6bf4f84..7d1a624 100644 --- a/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp @@ -57,8 +57,7 @@ using MacroStruct = Struct< template saw::error_or setup_initial_conditions( saw::data>& fields, - saw::data>& macros, - saw::data, particle_amount>>& particles + saw::data>& macros ){ auto& info_f = fields.template get<"info">(); // Set everything as walls @@ -162,21 +161,12 @@ template saw::error_or step( saw::data>,encode::Sycl>& fields, saw::data>,encode::Sycl>& macros, - saw::data, particle_amount>>& particles, saw::data t_i, device& dev ){ auto& q = dev.get_handle(); auto& info_f = fields.template get<"info">(); - { - auto& p = particles.at({{0u}}); - - auto& p_coll = p.template get<"collision">(); - auto& p_rad = p_coll.template get<"radius">(); - } - - // auto coll_ev = q.submit([&](acpp::sycl::handler& h){ // Need nicer things to handle the flow. I see improvement here @@ -283,7 +273,6 @@ 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, particle_amount>>>(); 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,*lbm_particle_data_ptr); + auto eov = setup_initial_conditions(*lbm_data_ptr,*lbm_macro_data_ptr); if(eov.is_error()){ return eov; } @@ -323,7 +312,6 @@ 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, particle_amount>, encode::Sycl> lbm_sycl_particle_data{sycl_q}; sycl_q.wait(); auto lsd_view = make_chunk_struct_view(lbm_sycl_data); @@ -340,12 +328,6 @@ saw::error_or lbm_main(int argc, char** argv){ return eov; } } - { - auto eov = dev.copy_to_device(*lbm_particle_data_ptr,lbm_sycl_particle_data); - if(eov.is_error()){ - return eov; - } - } sycl_q.wait(); saw::data time_steps{16u*4096ul}; @@ -354,7 +336,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,*lbm_particle_data_ptr,i,dev); + auto eov = step(lsd_view,lsdm_view,i,dev); if(eov.is_error()){ return eov; } diff --git a/examples/poiseulle_particles_2d_psm_gpu/sim.cpp b/examples/poiseulle_particles_2d_psm_gpu/sim.cpp index 7702808..7e76da6 100644 --- a/examples/poiseulle_particles_2d_psm_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_psm_gpu/sim.cpp @@ -60,8 +60,7 @@ using MacroStruct = Struct< template saw::error_or setup_initial_conditions( saw::data>& fields, - saw::data>& macros, - saw::data, particle_amount>>& particles + saw::data>& macros ){ auto& info_f = fields.template get<"info">(); auto& porous_f = macros.template get<"porosity">(); @@ -178,23 +177,6 @@ saw::error_or setup_initial_conditions( df_f.get_dims() ); - for(saw::data i{0u}; i < saw::data{particle_amount}; ++i){ - auto& part = particles.at(i); - - saw::data> pos; - pos.at({{0u}}) = dim_x * 0.25; - pos.at({{1u}}) = dim_y * 0.5; - saw::data> rad, dense, dt; - rad.at({}) = dim_y * 0.1; - dense.at({}) = 1.0; - dt.at({}) = 1.0; - part = create_spheroid_particle( - pos,{},{}, - {},{},{}, - rad, dense,dt - ); - } - return saw::make_void(); } @@ -202,7 +184,6 @@ template saw::error_or step( saw::data>,encode::Sycl>& fields, saw::data>,encode::Sycl>& macros, - saw::data, particle_amount>>& particles, saw::data t_i, device& dev ){ @@ -210,14 +191,6 @@ saw::error_or step( auto& info_f = fields.template get<"info">(); auto& porous_f = macros.template get<"porosity">(); - { - auto& p = particles.at({{0u}}); - - auto& p_coll = p.template get<"collision">(); - auto& p_rad = p_coll.template get<"radius">(); - } - - q.submit([&](acpp::sycl::handler& h){ component> collision{0.8}; component> bb; @@ -326,7 +299,6 @@ 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, particle_amount>>>(); 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,*lbm_particle_data_ptr); + auto eov = setup_initial_conditions(*lbm_data_ptr,*lbm_macro_data_ptr); if(eov.is_error()){ return eov; } @@ -366,7 +338,6 @@ 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, particle_amount>, encode::Sycl> lbm_sycl_particle_data{sycl_q}; sycl_q.wait(); auto lsd_view = make_chunk_struct_view(lbm_sycl_data); @@ -383,12 +354,6 @@ saw::error_or lbm_main(int argc, char** argv){ return eov; } } - { - auto eov = dev.copy_to_device(*lbm_particle_data_ptr,lbm_sycl_particle_data); - if(eov.is_error()){ - return eov; - } - } sycl_q.wait(); saw::data time_steps{16u*4096ul}; @@ -397,13 +362,13 @@ 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,*lbm_particle_data_ptr,i,dev); + auto eov = step(lsd_view,lsdm_view,i,dev); if(eov.is_error()){ return eov; } } sycl_q.wait(); - if(i.get() % 1 == 0u){ + if(i.get() % 32u == 0u){ { auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr); if(eov.is_error()){ diff --git a/examples/settling_cubes_2d_ibm_gpu/sim.cpp b/examples/settling_cubes_2d_ibm_gpu/sim.cpp index bb6ab59..aec8bb9 100644 --- a/examples/settling_cubes_2d_ibm_gpu/sim.cpp +++ b/examples/settling_cubes_2d_ibm_gpu/sim.cpp @@ -49,15 +49,10 @@ using MacroStruct = Struct< Member, "force"> >; -//template -//using ParticleArray = Array< -// Particle -//>; - template using ParticleGroups = Tuple< ParticleGroup< - T,Desc::D + T,Desc::D,sch::ParticleCollisionSpheroid > >; } @@ -65,7 +60,8 @@ using ParticleGroups = Tuple< 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 @@ -105,12 +101,20 @@ saw::error_or setup_initial_conditions( df_f.get_dims() ); - return saw::make_void(); -} + // Particles + { + saw::data> rad_p, dense_p; + rad_p.at({}).set(2.0); + dense_p.at({}).set(1); + auto& spheroid_group = particles.template get<0u>(); + spheroid_group = create_spheroid_particle_group( + rad_p, + dense_p, + {64u} + ); + } -template -void add_particles(saw::data>>& particles){ - ////// TODO + return saw::make_void(); } template @@ -129,6 +133,8 @@ saw::error_or step( // 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; @@ -205,7 +211,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_particles_data = saw::data>(); + 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_particles_data); if(eov.is_error()){ return eov; } diff --git a/lib/core/c++/boundary.hpp b/lib/core/c++/boundary.hpp index 4dbbdf8..01ae7b5 100644 --- a/lib/core/c++/boundary.hpp +++ b/lib/core/c++/boundary.hpp @@ -211,12 +211,12 @@ public: if constexpr (Dir) { dfs_old.at({2u}) = dfs_old.at({1u}) + saw::data{2.0 / 3.0} * rho_vel_x; - dfs_old.at({6u}) = dfs_old.at({5u}) + saw::data{1.0 / 6.0} * rho_vel_x + saw::data{0.5} * (dfs_old.at({4u}) - dfs_old.at({3u})); - dfs_old.at({8u}) = dfs_old.at({7u}) + saw::data{1.0 / 6.0} * rho_vel_x + saw::data{0.5} * (dfs_old.at({3u}) - dfs_old.at({4u})); + dfs_old.at({6u}) = dfs_old.at({5u}) + saw::data{1.0 / 6.0} * rho_vel_x + saw::data{0.5} * (dfs_old.at({3u}) - dfs_old.at({4u})); + dfs_old.at({8u}) = dfs_old.at({7u}) + saw::data{1.0 / 6.0} * rho_vel_x + saw::data{0.5} * (dfs_old.at({4u}) - dfs_old.at({3u})); }else if constexpr (not Dir){ dfs_old.at({1u}) = dfs_old.at({2u}) - saw::data{2.0 / 3.0} * rho_vel_x; - dfs_old.at({5u}) = dfs_old.at({6u}) - saw::data{1.0 / 6.0} * rho_vel_x + saw::data{0.5} * (dfs_old.at({3u}) - dfs_old.at({4u})); - dfs_old.at({7u}) = dfs_old.at({8u}) - saw::data{1.0 / 6.0} * rho_vel_x + saw::data{0.5} * (dfs_old.at({4u}) - dfs_old.at({3u})); + dfs_old.at({5u}) = dfs_old.at({6u}) - saw::data{1.0 / 6.0} * rho_vel_x + saw::data{0.5} * (dfs_old.at({4u}) - dfs_old.at({3u})); + dfs_old.at({7u}) = dfs_old.at({8u}) - saw::data{1.0 / 6.0} * rho_vel_x + saw::data{0.5} * (dfs_old.at({3u}) - dfs_old.at({4u})); } } }; diff --git a/lib/core/c++/particle/particle.hpp b/lib/core/c++/particle/particle.hpp index a95a173..c0d115f 100644 --- a/lib/core/c++/particle/particle.hpp +++ b/lib/core/c++/particle/particle.hpp @@ -1,8 +1,10 @@ #pragma once -#include -#include #include +#include +#include + +#include "../iterator.hpp" namespace kel { namespace lbm { @@ -52,25 +54,20 @@ using Particle = Struct< template> using ParticleGroup = Struct< - Member, "mask">, + Member, "mask">, Member, Member, "mass">, Member>, "particles"> >; - -template -using ParticleGroupTuple = Tuple< - PartGroups... ->; - } template saw::data>> create_spheroid_particle_group( saw::data> rad_p, - saw::data> density_p + saw::data> density_p, + const saw::data& mask_resolution ){ - saw::data,sch::ParticleCollisionSpheroid> part; + saw::data>> part; auto& mask = part.template get<"mask">(); auto& collision = part.template get<"collision">(); @@ -92,6 +89,16 @@ saw::data>> create_sph static_assert(D == 0u or D > 3u, "Dimensions only supported for Dim 1,2 & 3."); } + saw::data> mask_dims; + for(uint64_t i = 0u; i < D; ++i){ + mask_dims.at({i}) = mask_resolution; + } + + mask = {mask_dims}; + iterator::apply([&](const auto& index){ + + },{},mask_dims); + return part; } -- cgit v1.2.3