From 2b153a0e0263e5a76550b15382e2609f7b5a5d41 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Mon, 24 Aug 2026 17:15:52 +0200 Subject: Stuff --- default.nix | 46 ++++----- examples/one_particle_sedimentation_2d/common.hpp | 14 +-- examples/one_particle_sedimentation_2d/init.hpp | 4 +- examples/one_particle_sedimentation_2d/sim.cpp | 2 +- examples/one_particle_sedimentation_2d/step.hpp | 1 - .../schaefer_turek_durst_krause_rannbacher/sim.cpp | 9 +- modules/core/c++/abstract/schema.hpp | 6 -- modules/core/c++/abstract/stringify_schema.hpp | 20 ++-- modules/core/c++/fplbm.hpp | 105 +-------------------- modules/core/tests/schema.cpp | 2 + 10 files changed, 55 insertions(+), 154 deletions(-) diff --git a/default.nix b/default.nix index d3692b6..b945aab 100644 --- a/default.nix +++ b/default.nix @@ -86,7 +86,7 @@ in rec { }; }; - examples = { + ex = { cavity_2d_gpu = pkgs.callPackage ./examples/cavity_2d_gpu/.nix/derivation.nix { inherit pname version stdenv forstio adaptive-cpp; inherit kel; @@ -165,11 +165,11 @@ in rec { poiseulle_particles_2d_all_gpu = pkgs.symlinkJoin { name = "poiseulle_particles_2d_all_gpu"; paths = [ - examples.poiseulle_particles_2d_bgk_gpu - examples.poiseulle_particles_2d_psm_gpu - examples.poiseulle_particles_2d_hlbm_gpu - examples.poiseulle_particles_2d_fplbm_gpu - examples.poiseulle_particles_2d_ibm_gpu + ex.poiseulle_particles_2d_bgk_gpu + ex.poiseulle_particles_2d_psm_gpu + ex.poiseulle_particles_2d_hlbm_gpu + ex.poiseulle_particles_2d_fplbm_gpu + ex.poiseulle_particles_2d_ibm_gpu ]; }; @@ -241,7 +241,7 @@ in rec { settling_particles = pkgs.symlinkJoin { name = "kel-lbm-settling_particles-${version}"; paths = [ - examples.settling_cubes_2d_ibm_gpu + ex.settling_cubes_2d_ibm_gpu ]; }; }; @@ -255,36 +255,36 @@ in rec { inherit kel; }; - examples = pkgs.symlinkJoin { + ex = pkgs.symlinkJoin { name = "kel-lbm-sims-${version}"; paths = [ kel.lbm.core kel.lbm.sycl - examples.meta - examples.poiseulle_particles_2d_bgk_gpu - examples.poiseulle_particles_2d_psm_gpu - examples.poiseulle_particles_2d_hlbm_gpu - examples.poiseulle_particles_2d_fplbm_gpu - examples.poiseulle_particles_2d_ibm_gpu + ex.meta + ex.poiseulle_particles_2d_bgk_gpu + ex.poiseulle_particles_2d_psm_gpu + ex.poiseulle_particles_2d_hlbm_gpu + ex.poiseulle_particles_2d_fplbm_gpu + ex.poiseulle_particles_2d_ibm_gpu ]; }; poiseulle_particles_2d = pkgs.symlinkJoin { name = "kel-lbm-sims-poiseulle_particles_2d-${version}"; paths = [ - examples.poiseulle_particles_2d_bgk_gpu - examples.poiseulle_particles_2d_psm_gpu - examples.poiseulle_particles_2d_hlbm_gpu - examples.poiseulle_particles_2d_fplbm_gpu + ex.poiseulle_particles_2d_bgk_gpu + ex.poiseulle_particles_2d_psm_gpu + ex.poiseulle_particles_2d_hlbm_gpu + ex.poiseulle_particles_2d_fplbm_gpu ]; }; stokes_drag_2d = pkgs.symlinkJoin { name = "kel-lbm-sims-stokes_drag_2d-${version}"; paths = [ - examples.stokes_drag_particle_2d_bgk_gpu - examples.stokes_drag_particle_2d_psm_gpu - examples.stokes_drag_particle_2d_hlbm_gpu - examples.stokes_drag_particle_2d_fplbm_gpu + ex.stokes_drag_particle_2d_bgk_gpu + ex.stokes_drag_particle_2d_psm_gpu + ex.stokes_drag_particle_2d_hlbm_gpu + ex.stokes_drag_particle_2d_fplbm_gpu ]; }; @@ -297,5 +297,5 @@ in rec { }; }; - default = release.examples; + default = release.ex; } diff --git a/examples/one_particle_sedimentation_2d/common.hpp b/examples/one_particle_sedimentation_2d/common.hpp index 37e0402..954557d 100644 --- a/examples/one_particle_sedimentation_2d/common.hpp +++ b/examples/one_particle_sedimentation_2d/common.hpp @@ -14,15 +14,17 @@ namespace kel { namespace lbm { /** - * length 2.2m - * height 0.41m + * length 20mm + * height 80mm * viscosity 10^-3 - * particle diameter 0.1m - * position at (0.16m,0.2m) + * particle radius 1.25mm + * position at (10mm,40mm) + * particle density 0.00125g/mm³ + * fluid density 0.001g/mm³ */ -constexpr uint64_t dim_y = 512ul; -constexpr uint64_t dim_x = 256ul; +constexpr uint64_t dim_y = 800ul; +constexpr uint64_t dim_x = 200ul; constexpr uint64_t particle_amount = 1ul; diff --git a/examples/one_particle_sedimentation_2d/init.hpp b/examples/one_particle_sedimentation_2d/init.hpp index a35c6e8..d9b8a37 100644 --- a/examples/one_particle_sedimentation_2d/init.hpp +++ b/examples/one_particle_sedimentation_2d/init.hpp @@ -87,8 +87,8 @@ saw::error_or init( auto& prb = p.template get<"rigid_body">(); auto& p_pos = prb.template get<"position">(); auto& p_posold = prb.template get<"position_old">(); - p_pos.at({{0u}}) = conv.meter_si_to_lbm({{0.16}}).handle(); - p_pos.at({{1u}}) = conv.meter_si_to_lbm({{0.2}}).handle(); + p_pos.at({{0u}}) = conv.meter_si_to_lbm({{0.01}}).handle(); + p_pos.at({{1u}}) = conv.meter_si_to_lbm({{0.04}}).handle(); p_posold = p_pos; } diff --git a/examples/one_particle_sedimentation_2d/sim.cpp b/examples/one_particle_sedimentation_2d/sim.cpp index 7c9a56c..b0b374c 100644 --- a/examples/one_particle_sedimentation_2d/sim.cpp +++ b/examples/one_particle_sedimentation_2d/sim.cpp @@ -46,7 +46,7 @@ saw::error_or lbm_main(const saw::data& args){ converter conv { // delta_x - {{0.41 / dim_y}}, + {{0.02 / dim_x}}, // delta_t {{1.067e-4}} }; diff --git a/examples/one_particle_sedimentation_2d/step.hpp b/examples/one_particle_sedimentation_2d/step.hpp index 07f51d1..20b7f4b 100644 --- a/examples/one_particle_sedimentation_2d/step.hpp +++ b/examples/one_particle_sedimentation_2d/step.hpp @@ -30,7 +30,6 @@ saw::error_or step( for(uint64_t i = 0u; i < Desc::D; ++i){ index.at({{i}}).set(idx[i]); } - hlbm_reset.apply(fields,macros,index,t_i); }); }).wait(); diff --git a/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp b/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp index b9b3778..4d015c9 100644 --- a/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp +++ b/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp @@ -144,7 +144,7 @@ saw::error_or lbm_main(const saw::data& args){ } } { - auto eov = write_csv_file(out_dir,"stdkr_2d",i.get(), *lbm_macro_data_ptr); + auto eov = write_vtk_file(out_dir,std::string{"stdkr_2d_"}+an.template get<"coupling">().stl_string(),i.get(), *lbm_macro_data_ptr); if(eov.is_error()){ return eov; } @@ -189,7 +189,12 @@ saw::error_or lbm_main(const saw::data& args){ return eov; } } - + { + auto eov = write_csv_file(out_dir,"stdkr_2d",time_steps.get(), *lbm_macro_data_ptr); + if(eov.is_error()){ + return eov; + } + } sycl_q.wait(); return saw::make_void(); } diff --git a/modules/core/c++/abstract/schema.hpp b/modules/core/c++/abstract/schema.hpp index 8bc23b4..89d9c2e 100644 --- a/modules/core/c++/abstract/schema.hpp +++ b/modules/core/c++/abstract/schema.hpp @@ -42,12 +42,6 @@ using Float16 = Primitive; using Float32 = Primitive; using Float64 = Primitive; -template -struct MixedPrecision { - using Meta = Void; - using Inner = Sch; -}; - template struct FixedArray { using Meta = Void; diff --git a/modules/core/c++/abstract/stringify_schema.hpp b/modules/core/c++/abstract/stringify_schema.hpp index 3d6b1d1..dd218d0 100644 --- a/modules/core/c++/abstract/stringify_schema.hpp +++ b/modules/core/c++/abstract/stringify_schema.hpp @@ -10,36 +10,36 @@ template struct schema_to_string_helper { static_assert(always_false, "schema_to_string_helper specialization not handled."); - static constexpr string_literal apply(){ - return {}; + static constexpr auto apply(){ + return string_literal{""}; } }; template<> struct schema_to_string_helper final { - static constexpr string_literal apply(){ - return "SignedInteger"; + static constexpr auto apply(){ + return string_literal{"SignedInteger"}; } }; template<> struct schema_to_string_helper final { - static constexpr string_literal apply(){ - return "UnsignedInteger"; + static constexpr auto apply(){ + return string_literal{"UnsignedInteger"}; } }; template<> struct schema_to_string_helper final { - static constexpr string_literal apply(){ - return "FloatingPoint"; + static constexpr auto apply(){ + return string_literal{"FloatingPoint"}; } }; template struct schema_to_string_helper> final { - static constexpr string_literal apply(){ - return "Primitive<"+schema_to_string_helper::apply()+","+N+">"; + static constexpr auto apply(){ + return string_literal{"Primitive<"}+schema_to_string_helper::apply()+string_literal{","}+N+string_literal{">"}; } }; } diff --git a/modules/core/c++/fplbm.hpp b/modules/core/c++/fplbm.hpp index 382d067..c9ed6e9 100644 --- a/modules/core/c++/fplbm.hpp +++ b/modules/core/c++/fplbm.hpp @@ -204,119 +204,18 @@ public: // vel_s is technically time the density of the particle? - force = ( vel_s * rho - vel * rho ) /* two */ * flip_por; + force = ( vel_s * rho - vel * rho ) * two * flip_por; force_p = force_p - force; }, aabb.template get<"a">(), aabb.template get<"b">()); auto& pirb_acc = pirb.template get<"acceleration">(); - pirb_acc = force_p / pg.template get<"total_mass">().at({}); + pirb_acc = force_p / (two * pg.template get<"total_mass">().at({})); for(saw::data i{0u}; i < sub_steps; ++i){ verlet_step_lambda(pi,sss); } } }; - -template -class component final { -private: - saw::data relaxation_; - saw::data frequency_; -public: - component(typename saw::native_data_type::type relaxation__): - relaxation_{{relaxation__}}, - frequency_{saw::data{1} / relaxation_} - {} - - component(const saw::data& relaxation__): - relaxation_{relaxation__}, - frequency_{saw::data{1} / relaxation_} - {} - - using Component = cmpt::FpLbmOneParticleNoVelocity; - - template - void apply(const saw::data& field, const saw::data& macros, saw::data> index, saw::data time_step) const { - // void apply(saw::data& field, saw::data> index, saw::data time_step){ - bool is_even = ((time_step.get() % 2) == 0); - - auto& dfs_old_f = (is_even) ? field.template get<"dfs_old">() : field.template get<"dfs">(); - auto& dfs = dfs_old_f.at(index); - - auto& rho_f = macros.template get<"density">(); - auto& vel_f = macros.template get<"momentum">(); - auto& por_f = macros.template get<"porosity">(); - - // Temporary find a better way - auto& force_f = macros.template get<"force">(); - - - /** - * The other methods all work with a flipped porosity. - * For compat reasons this is also flipped - */ - auto porosity = por_f.at(index); - saw::data> one; - one.at({}) = 1.0; - auto flip_porosity = one - porosity; - - saw::data>& rho = rho_f.at(index); - - saw::data> half; - half.at({}).set(0.5); - saw::data>& vel = vel_f.at(index);// + total_force * ( half / rho ); - - compute_rho_u(dfs_old_f.at(index),rho,vel); - auto eq = equilibrium(rho,vel); - - using dfi = df_info; - - saw::data> min_two; - min_two.at({}).set(-2); - - auto& force = force_f.at(index); - force = vel * rho * min_two * flip_porosity; - - saw::data> dfi_inv_cs2; - dfi_inv_cs2.at({}).set(dfi::inv_cs2); - - for(uint64_t i = 0u; i < Descriptor::Q; ++i){ - // saw::data ci_min_u{0}; - saw::data> ci; - for(uint64_t d = 0u; d < Descriptor::D; ++d){ - ci.at({{d}}).set(static_cast::type>(dfi::directions[i][d])); - } - auto ci_dot_u = saw::math::dot(ci,vel); - - // saw::data> F_i; - // F_i = f * ((c_i - u) * ics2 + * c_i * ics2 * ics2) * w_i; - saw::data> w; - w.at({}).set(dfi::weights[i]); - - /* - saw::data> F_i_sum; - for(uint64_t d = 0u; d < Descriptor::D; ++d){ - saw::data> F_i_d; - F_i_d.at({}) = F_i.at({{d}}); - F_i_sum = F_i_sum + F_i_d; - } - */ - auto term1 = (ci-vel) * dfi_inv_cs2; - auto term2 = ci * (ci_dot_u * dfi_inv_cs2 * dfi_inv_cs2); - - auto force_projection = saw::math::dot(term1 + term2, force); - - auto F_i = w * force_projection; - - dfs.at({i}) = dfs.at({i}) - + frequency_ - * (eq.at(i) - dfs.at({i}) ) - + F_i.at({}) - * (saw::data{1} - saw::data{0.5f} * frequency_); - } - } - -}; } } diff --git a/modules/core/tests/schema.cpp b/modules/core/tests/schema.cpp index db82b97..6649bad 100644 --- a/modules/core/tests/schema.cpp +++ b/modules/core/tests/schema.cpp @@ -18,6 +18,7 @@ SAW_TEST("Abstract/Schema/Ptr"){ } */ +/* SAW_TEST("Abstract/SchemaStringify/PrimitiveKind"){ using namespace kel::lbm; @@ -25,4 +26,5 @@ SAW_TEST("Abstract/SchemaStringify/PrimitiveKind"){ static constexpr auto str = schema_to_string::value; } +*/ } -- cgit v1.2.3 From f5d88c1ae8948724f11d30c1c2af385cc093eed2 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Mon, 24 Aug 2026 19:02:51 +0200 Subject: Fixing FpLBM, because I messed up guo forcing --- .../schaefer_turek_durst_krause_rannbacher/sim.cpp | 2 +- .../step.hpp | 23 ++++++++++++++------ modules/core/c++/fplbm.hpp | 2 +- modules/core/c++/macroscopic.hpp | 25 ++++++++++++++++++++++ 4 files changed, 44 insertions(+), 8 deletions(-) diff --git a/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp b/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp index 4d015c9..ff6a43a 100644 --- a/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp +++ b/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp @@ -48,7 +48,7 @@ saw::error_or lbm_main(const saw::data& args){ // delta_x {{0.41 / dim_y}}, // delta_t - {{1.067e-4}} + {{1.067e-4 * 0.5}} }; print_lbm_meta(conv,{1e-3},{1e-4},{0.5 * dim_y}); diff --git a/examples/schaefer_turek_durst_krause_rannbacher/step.hpp b/examples/schaefer_turek_durst_krause_rannbacher/step.hpp index cbdfd10..9748ec8 100644 --- a/examples/schaefer_turek_durst_krause_rannbacher/step.hpp +++ b/examples/schaefer_turek_durst_krause_rannbacher/step.hpp @@ -8,6 +8,7 @@ namespace lbm { template saw::error_or step( const converter& conv, + const saw::data>& tau, saw::data>,encode::Sycl>& fields, saw::data>,encode::Sycl>& macros, saw::data>,encode::Sycl>& particles, @@ -48,7 +49,7 @@ saw::error_or step( }).wait(); q.submit([&](acpp::sycl::handler& h){ - component> collision{1.0}; + component> collision{0.75}; component> bb; component,encode::Sycl> flow_out{1.0}; @@ -74,9 +75,19 @@ saw::error_or step( component,encode::Sycl> flow_in{ [&]() -> saw::data> { saw::data> vel; + saw::data scale; + { + scale.set([&](){ + if(t_i.get() > 2048u){ + return 1.0f; + } + + return t_i.get() / 2048.0f; + }()); + } { auto y_si = conv.meter_lbm_to_si(index.at({{1u}}).template cast_to()).handle(); - vel.at({{0u}}) = conv.velocity_si_to_lbm({{static_cast::type>((1.2 / 0.41) * y_si.get() - (1.2 / (0.41*0.41)) * y_si.get() * y_si.get())}}).handle(); + vel.at({{0u}}) = conv.velocity_si_to_lbm({{static_cast::type>((1.2 / 0.41) * y_si.get() - (1.2 / (0.41*0.41)) * y_si.get() * y_si.get())}}).handle() * scale; vel.at({{1u}}) = 0.0f; } return vel; @@ -123,8 +134,8 @@ saw::error_or step( // auto coll_ev = q.submit([&](acpp::sycl::handler& h){ - component> bgk{1.0}; - component> collision{1.0}; + component> bgk{0.75}; + component> collision{0.75}; component> bb; component,encode::Sycl> flow_out{1.0}; @@ -199,8 +210,8 @@ saw::error_or step( // auto coll_ev = q.submit([&](acpp::sycl::handler& h){ - component> bgk{1.0}; - component> collision{1.0}; + component> bgk{0.75}; + component> collision{0.75}; component> bb; component,encode::Sycl> flow_out{1.0}; diff --git a/modules/core/c++/fplbm.hpp b/modules/core/c++/fplbm.hpp index c9ed6e9..d8c075b 100644 --- a/modules/core/c++/fplbm.hpp +++ b/modules/core/c++/fplbm.hpp @@ -92,7 +92,7 @@ public: auto& force_f = macros.template get<"force">(); auto& force = force_f.at(index); - vel = vel + force * half; + vel = vel + force * half / rho; // compute_rho_u(dfs,rho,vel); auto eq = equilibrium(rho,vel); diff --git a/modules/core/c++/macroscopic.hpp b/modules/core/c++/macroscopic.hpp index 768fef0..9db1538 100644 --- a/modules/core/c++/macroscopic.hpp +++ b/modules/core/c++/macroscopic.hpp @@ -33,5 +33,30 @@ void compute_rho_u ( vel().at({{i}}) = vel().at({{i}}) / rho().at({}); } } + +/** + * Calculate the macroscopic variables rho and u in Lattice Units. + */ +template +void compute_rho_momentum ( + const saw::data>& dfs, + saw::ref>> rho, + saw::ref>> mom + ) +{ + using dfi = df_info; + + rho().at({}).set(0); + for(uint64_t i = 0; i < Desc::D; ++i){ + mom().at({{i}}).set(0); + } + + for(size_t j = 0; j < Desc::Q; ++j){ + rho().at({}) = rho().at({}) + dfs.at({j}); + for(size_t i = 0; i < Desc::D; ++i){ + mom().at({{i}}) = mom().at({{i}}) + saw::data{static_cast::type>(dfi::directions[j][i])} * dfs.at({j}); + } + } +} } } -- cgit v1.2.3 From 76dea4c1035bf063802ca0cc650d69aef8047b9f Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Mon, 24 Aug 2026 19:19:45 +0200 Subject: Fixing examples --- examples/schaefer_turek_durst_krause_rannbacher/sim.cpp | 4 +++- examples/schaefer_turek_durst_krause_rannbacher/step.hpp | 7 ++----- 2 files changed, 5 insertions(+), 6 deletions(-) diff --git a/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp b/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp index ff6a43a..59f61d9 100644 --- a/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp +++ b/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp @@ -122,7 +122,7 @@ saw::error_or lbm_main(const saw::data& args){ auto lsdm_view = make_view(lbm_sycl_macro_data); auto lsdp_view = make_view(lbm_sycl_particle_data); - saw::data time_steps{64u*1024ul}; + saw::data time_steps{128u*1024ul}; auto& info_f = lsd_view.template get<"info">(); @@ -136,6 +136,7 @@ saw::error_or lbm_main(const saw::data& args){ } sycl_q.wait(); + /* if(i.get() % 64u == 0u){ { auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr); @@ -150,6 +151,7 @@ saw::error_or lbm_main(const saw::data& args){ } } } + */ // Stream sycl_q.submit([&](acpp::sycl::handler& h){ component> stream; diff --git a/examples/schaefer_turek_durst_krause_rannbacher/step.hpp b/examples/schaefer_turek_durst_krause_rannbacher/step.hpp index 9748ec8..beb123b 100644 --- a/examples/schaefer_turek_durst_krause_rannbacher/step.hpp +++ b/examples/schaefer_turek_durst_krause_rannbacher/step.hpp @@ -8,7 +8,6 @@ namespace lbm { template saw::error_or step( const converter& conv, - const saw::data>& tau, saw::data>,encode::Sycl>& fields, saw::data>,encode::Sycl>& macros, saw::data>,encode::Sycl>& particles, @@ -20,8 +19,6 @@ saw::error_or step( auto& q = dev.get_handle(); auto& info_f = fields.template get<"info">(); auto& porous_f = macros.template get<"porosity">(); - - if constexpr ( std::is_same_v ){ q.submit([&](acpp::sycl::handler& h){ @@ -78,11 +75,11 @@ saw::error_or step( saw::data scale; { scale.set([&](){ - if(t_i.get() > 2048u){ + if(t_i.get() > 1024u){ return 1.0f; } - return t_i.get() / 2048.0f; + return t_i.get() / 1024.0f; }()); } { -- cgit v1.2.3 From 692b413fec56fbb5c6198813df369106721c33f5 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Mon, 24 Aug 2026 19:47:14 +0200 Subject: Test --- default.nix | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/default.nix b/default.nix index b945aab..04b2616 100644 --- a/default.nix +++ b/default.nix @@ -36,7 +36,7 @@ let inherit stdenv llvmPackages lld; config = { - rocm = false; + rocm = true; }; }; -- cgit v1.2.3