From 7fd9bfd5946472230a3b74c52f88e19c15741faf Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Mon, 1 Jun 2026 17:21:44 +0200 Subject: I seem to have no clue what I'm doing --- default.nix | 2 +- examples/poiseulle_particles_2d_ibm_gpu/sim.cpp | 152 ++++++++++++++---------- lib/core/c++/collision.hpp | 3 +- lib/core/c++/math/n_closest.hpp | 18 ++- 4 files changed, 107 insertions(+), 68 deletions(-) diff --git a/default.nix b/default.nix index 4303f93..d233d90 100644 --- a/default.nix +++ b/default.nix @@ -46,7 +46,7 @@ let src = builtins.fetchurl { url = "https://git.keldu.de/forstio-forstio/snapshot/master.tar.gz"; - sha256 = "sha256:0khnwmrhdric5i701wfbkbdrfzp9vrswg6raan4njmwlx146vbrf"; + sha256 = "sha256:15iqzmymza47jjx4wpc19mbg3zzwmkabpssf5y968f566n0fnb9a"; }; phases = [ "unpackPhase" "installPhase" ]; diff --git a/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp b/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp index d4bf053..e1bd3ba 100644 --- a/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp @@ -114,6 +114,7 @@ saw::error_or setup_initial_conditions( auto& df_f = fields.template get<"dfs_old">(); auto& rho_f = macros.template get<"density">(); auto& vel_f = macros.template get<"velocity">(); + auto& force_f = macros.template get<"force">(); iterator::apply( [&](auto& index){ @@ -135,6 +136,9 @@ saw::error_or setup_initial_conditions( auto& rho = rho_f.at(index); rho.at({}) = {1}; auto& vel = vel_f.at(index); + auto& force = force_f.at(index); + force = {}; + if(info_f.at(index).get() == 2u){ vel.at({{0u}}) = 0.0; } @@ -195,15 +199,79 @@ saw::error_or step( ){ auto& q = dev.get_handle(); auto& info_f = fields.template get<"info">(); + auto& force_f = macros.template get<"force">(); + + 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]); + } + + auto& force = force_f.at(index); + + for(uint64_t i{0u}; i < Desc::D; ++i){ + force.at({{i}}) = 0.0; + } + }); + }).wait(); + + q.submit([&](acpp::sycl::handler& h){ + h.parallel_for(acpp::sycl::range<1u>{1u}, [=](acpp::sycl::id<1u> idx){ + auto& vel_f = macros.template get<"velocity">(); + auto& dense_f = macros.template get<"density">(); + + auto& ps = particles; + auto& mask = ps.template get<"mask">(); + auto& mask_step = ps.template get<"mask_step">().at({}); + auto& p_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({{idx[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({}).at({{i}}); + // Scale to LBM Grid + index_shift.at({{i}}) = index_shift.at({{i}}) * mask_step.at({}); + } + + // Shift our pos into the index + auto p_i_rb_pos_ind = p_i_rb_pos + index_shift; + + /// Calculate force pickup from neigbouring u_vel cells + // auto inter_vel_fluid = n_linear_interpolate(vel_f,p_i_rb_pos_ind); + auto inter_vel_fluid = n_closest_read(vel_f,p_i_rb_pos_ind); + + // Technically TODO to use moment + auto inter_moment_fluid = inter_vel_fluid; + + // Technically Particles can have more timesteps than the fluid + auto force_response = -inter_moment_fluid; + + /// Distribute force to fluid + // n_linear_spread(force_f,p_i_rb_pos_ind, force_response); + n_closest_add(force_f,p_i_rb_pos_ind, force_response); + + }, {}, mask.meta()); + }); + }).wait(); // 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; - - component> collision{0.65,f}; + component> collision{0.65}; component> bb; component,encode::Sycl> abb; @@ -233,7 +301,7 @@ saw::error_or step( } auto info = info_f.at(index); - + switch(info.get()){ case 0u: break; @@ -256,6 +324,8 @@ saw::error_or step( break; } }); + + }).wait(); @@ -265,61 +335,6 @@ saw::error_or step( // h.depends_on(collision_ev); }).wait(); */ - - auto& force_f = macros.template get<"force">(); - - q.submit([&](acpp::sycl::handler& h){ - h.parallel_for(acpp::sycl::range<1u>{1u}, [=](acpp::sycl::id<1u> idx){ - auto& vel_f = macros.template get<"velocity">(); - auto& dense_f = macros.template get<"density">(); - - auto& ps = particles; - auto& mask = ps.template get<"mask">(); - auto& mask_step = ps.template get<"mask_step">().at({}); - 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({{idx[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({}).at({{i}}); - // Scale to LBM Grid - index_shift.at({{i}}) = index_shift.at({{i}}) * mask_step.at({}); - } - - // Shift our pos into the index - auto p_i_rb_pos_ind = p_i_rb_pos + index_shift; - - /// Calculate force pickup from neigbouring u_vel cells - // auto inter_vel_fluid = n_linear_interpolate(vel_f,p_i_rb_pos_ind); - auto inter_vel_fluid = n_closest_read(vel_f,p_i_rb_pos_ind); - - // Technically TODO to use moment - auto inter_moment_fluid = inter_vel_fluid; - - // Technically Particles can have more timesteps than the fluid - auto force_response = -inter_moment_fluid; - - /// Distribute force to fluid - // n_linear_spread(force_f,p_i_rb_pos_ind, force_response); - n_closest_add(force_f,p_i_rb_pos_ind, force_response); - - }, {}, mask.meta()); - }); - }).wait(); - return saw::make_void(); } } @@ -493,8 +508,21 @@ saw::error_or lbm_main(int argc, char** argv){ }); }).wait(); wait.poll(); + + // PRINT STATUS ON SIGUSR1 if(print_status){ - std::cout<<"Status: "<().get() * 100 / time_steps.get())<<"%"<> half; half.at({}).set(0.5); - saw::data> vel = vel_f.at(index) + total_force * ( half / rho ); + auto& vel = vel_f.at(index); + vel = vel + total_force * ( half / rho ); compute_rho_u(dfs_old_f.at(index),rho,vel); auto eq = equilibrium(rho,vel); diff --git a/lib/core/c++/math/n_closest.hpp b/lib/core/c++/math/n_closest.hpp index 13414e2..ac0fe2f 100644 --- a/lib/core/c++/math/n_closest.hpp +++ b/lib/core/c++/math/n_closest.hpp @@ -7,7 +7,7 @@ namespace kel { namespace lbm { template -saw::data n_closest_read(const saw::data,Encode>& f, const saw::data>& frac_ind){ +saw::data n_closest_read(const saw::data,Encode>& f, const saw::data>& frac_ind){ auto shift_frac_ind = frac_ind; for(uint64_t i{0u}; i < D; ++i){ @@ -18,13 +18,16 @@ saw::data n_closest_read(const saw::data } } - auto shift_ind = frac_ind.template cast_to(); + saw::data> shift_ind; + for(uint64_t i{0u}; i < D; ++i){ + shift_ind.at({i}) = frac_ind.at({{i}}).template cast_to(); + } return f.at(shift_ind); } template -void n_closest_add(saw::data,Encode>& f, const saw::data>& frac_ind, const saw::data& val){ +void n_closest_add(const saw::data,Encode>& f, const saw::data>& frac_ind, const saw::data& val){ auto shift_frac_ind = frac_ind; for(uint64_t i{0u}; i < D; ++i){ @@ -34,7 +37,14 @@ void n_closest_add(saw::data,Encode>& f, const saw::data(); + auto f_meta = f.meta(); + saw::data> shift_ind; + for(uint64_t i{0u}; i < D; ++i){ + shift_ind.at({i}) = frac_ind.at({{i}}).template cast_to(); + if(shift_ind.at({i}) < f_meta.at({i})){ + shift_ind.at({i}) = f_meta.at({i}) - 1u; + } + } auto& f_i = f.at(shift_ind); f_i = f_i + val; -- cgit v1.2.3