From f3f2139d76648cca0e58688b8dbbc386efbebc06 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Mon, 29 Jun 2026 03:58:28 +0200 Subject: Writing in the writeout helper --- lib/core/c++/particle/particle_opa.hpp | 4 +++- lib/core/c++/particle/porosity.hpp | 35 +++++++++++++++++++++++++--------- 2 files changed, 29 insertions(+), 10 deletions(-) diff --git a/lib/core/c++/particle/particle_opa.hpp b/lib/core/c++/particle/particle_opa.hpp index 4588a55..e6396d6 100644 --- a/lib/core/c++/particle/particle_opa.hpp +++ b/lib/core/c++/particle/particle_opa.hpp @@ -39,7 +39,9 @@ public: auto pos_ind = saw::math::vectorize_data(index); auto diff = pos_ind - pos_; - auto diff_dot = saw::math::dot(diff,diff); + + // Write out value + porous = particle_porosity<>::calculate(diff, rad_, eps_); } }; } diff --git a/lib/core/c++/particle/porosity.hpp b/lib/core/c++/particle/porosity.hpp index f555cae..f5966c4 100644 --- a/lib/core/c++/particle/porosity.hpp +++ b/lib/core/c++/particle/porosity.hpp @@ -9,7 +9,6 @@ template class particle_porosity { public: - saw::data> calculate(const saw::data>& part_group, uint64_t p_i, const saw::data>& lbm_pos){ auto& mask = part_group.template get<"mask">(); @@ -32,10 +31,32 @@ public: template::type radius, typename saw::native_data_type::type eps> class particle_porosity> final { public: - saw::data> calculate(const saw::data>& lbm_pos, saw::data> rad) const { - saw::data> pos; - + saw::data> calculate(const saw::data>& lbm_rel_dist, saw::data> rad, saw::data> eps) const { + saw::data> por; + auto s_dist_2 = saw::math::dot(lbm_rel_dist,lbm_rel_dist); + auto s_dist = saw::math::sqrt(s_dist_2); + + auto rad_low = rad - eps; + auto rad_high = rad + eps; + + if(s_dist < rad_low){ + por.at({}).set(1); + return por; + } + if(s_dist > rad_high){ + por.at({}).set(0); + return por; + } + + { + typename saw::native_data_type::type inner = (std::pi / ( 2.0 * eps.at({}).get() )) * (s_dist - rad_low).at({}).get(); + auto cos_inner = std::cos(inner); + auto cos_i_2 = cos_inner * cos_inner; + por.at({}).set(cos_i_2); + } + + return por; } saw::data> calculate(const saw::data > >& part_group, uint64_t i, const saw::data>& lbm_pos) const { @@ -51,11 +72,7 @@ public: // Basically the queried position minus the center of the particle auto dist = lbm_pos - pi_rb_pos; - saw::data> dist_len; - for(uint64_t i{0u}; i < D; ++i){ - auto& d_i = dist.at({{i}}); - dist_len.at({}) = d_i * d_i; - } + saw::data> dist_len = saw::math::dot(dist,dist); dist_len.at({}).set(std::sqrt(dist_len.at({}).get()); saw::data> rad_d{radius}; -- cgit v1.2.3 From bec95825e78dc1171c337f2c40790e1ad5676f54 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Mon, 29 Jun 2026 19:26:19 +0200 Subject: Dangling --- default.nix | 11 +++++++ .../poiseulle_moving_particle_2d_psm_gpu/sim.cpp | 38 ++++++++++------------ examples/poiseulle_particles_2d_bgk_gpu/sim.cpp | 14 ++++++++ examples/poiseulle_particles_2d_psm_gpu/sim.cpp | 37 +++++++++++++++++++++ lib/core/c++/particle/particle_opa.hpp | 3 +- scripts/python/csv_l2_norm.py | 26 +++++++++++++++ 6 files changed, 108 insertions(+), 21 deletions(-) create mode 100755 scripts/python/csv_l2_norm.py diff --git a/default.nix b/default.nix index dc75b57..02f9213 100644 --- a/default.nix +++ b/default.nix @@ -151,6 +151,17 @@ in rec { inherit pname version stdenv forstio adaptive-cpp; inherit kel; }; + + 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 + ]; + }; poiseulle_moving_particle_2d_psm_gpu = pkgs.callPackage ./examples/poiseulle_moving_particle_2d_psm_gpu/.nix/derivation.nix { inherit pname version stdenv forstio adaptive-cpp; diff --git a/examples/poiseulle_moving_particle_2d_psm_gpu/sim.cpp b/examples/poiseulle_moving_particle_2d_psm_gpu/sim.cpp index 0c10d38..cae8a4d 100644 --- a/examples/poiseulle_moving_particle_2d_psm_gpu/sim.cpp +++ b/examples/poiseulle_moving_particle_2d_psm_gpu/sim.cpp @@ -1,6 +1,7 @@ #include #include #include +#include #include #include @@ -157,25 +158,6 @@ 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.5; - 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(); - - auto dist = middle - ind_vec; - auto dist_2 = saw::math::dot(dist,dist); - if(dist_2.at({}).get() < dim_y*dim_y*0.01){ - porous_f.at(index).at({}) = 0.0; - } - }, - {},// 0-index - df_f.get_dims() - ); return saw::make_void(); } @@ -190,6 +172,8 @@ saw::error_or step( auto& q = dev.get_handle(); auto& info_f = fields.template get<"info">(); auto& porous_f = macros.template get<"porosity">(); + sch::data> eps; + eps.at({}).set(1.5f); q.submit([&](acpp::sycl::handler& h){ component> collision{0.8}; @@ -199,7 +183,7 @@ saw::error_or step( component,encode::Sycl> flow_in{1.0}; component,encode::Sycl> flow_out{1.0}; - component> opa{{},{}}; + component> opa{{},{},eps}; h.parallel_for(acpp::sycl::range{dim_x,dim_y}, [=](acpp::sycl::id idx){ saw::data> index; @@ -371,6 +355,7 @@ saw::error_or lbm_main(int argc, char** argv){ } } */ + /* if(i.get() % 32u == 0u){ { auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr); @@ -385,6 +370,7 @@ saw::error_or lbm_main(int argc, char** argv){ } } } + */ // Stream sycl_q.submit([&](acpp::sycl::handler& h){ component> stream; @@ -424,6 +410,18 @@ saw::error_or lbm_main(int argc, char** argv){ return eov; } } + { + auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr); + if(eov.is_error()){ + return eov; + } + } + { + auto eov = write_csv_file(out_dir,"m",time_steps.get(), *lbm_macro_data_ptr); + if(eov.is_error()){ + return eov; + } + } sycl_q.wait(); return saw::make_void(); diff --git a/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp b/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp index 9cc77ae..73964b7 100644 --- a/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp @@ -366,6 +366,7 @@ saw::error_or lbm_main(int argc, char** argv){ } } */ + /* if(i.get() % 32u == 0u){ { auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr); @@ -380,6 +381,7 @@ saw::error_or lbm_main(int argc, char** argv){ } } } + */ // Stream sycl_q.submit([&](acpp::sycl::handler& h){ component> stream; @@ -419,6 +421,18 @@ saw::error_or lbm_main(int argc, char** argv){ return eov; } } + { + auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr); + if(eov.is_error()){ + return eov; + } + } + { + auto eov = write_csv_file(out_dir,"m",time_steps.get(), *lbm_macro_data_ptr); + if(eov.is_error()){ + return eov; + } + } sycl_q.wait(); return saw::make_void(); diff --git a/examples/poiseulle_particles_2d_psm_gpu/sim.cpp b/examples/poiseulle_particles_2d_psm_gpu/sim.cpp index 0337158..e91686c 100644 --- a/examples/poiseulle_particles_2d_psm_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_psm_gpu/sim.cpp @@ -1,6 +1,7 @@ #include #include #include +#include #include #include @@ -157,7 +158,27 @@ saw::error_or setup_initial_conditions( df_f.get_dims(), {{1u,1u}} ); + + saw::data> eps; + eps.at({}).set(1.5f); + saw::data> rad; + rad.at({}).set(dim_y*0.1); + saw::data> p_pos; + { + p_pos.at({{0u}}) = dim_x * 0.25; + p_pos.at({{1u}}) = dim_y * 0.5; + } + + component opa{p_pos,rad,eps}; + iterator::apply( + [&](auto& index){ + opa.apply(macros,index,{}); + }, + {},// 0-index + df_f.get_dims() + ); + /* iterator::apply( [&](auto& index){ saw::data> middle, ind_vec; @@ -176,6 +197,7 @@ saw::error_or setup_initial_conditions( {},// 0-index df_f.get_dims() ); + */ return saw::make_void(); } @@ -203,6 +225,18 @@ saw::error_or step( component> equi{rho_b,vel_b}; + + saw::data> eps; + eps.at({}).set(1.5f); + saw::data> rad; + rad.at({}).set(dim_y*0.1); + saw::data> p_pos; + { + p_pos.at({{0u}}) = dim_x * 0.25; + p_pos.at({{1u}}) = dim_y * 0.5; + } + component> opa{p_pos,rad,eps}; + component,encode::Sycl> flow_in{ [&](){ uint64_t target_t_i = 64u; @@ -230,6 +264,7 @@ saw::error_or step( bb.apply(fields,index,t_i); break; case 2u: + opa.apply(macros,index,t_i); collision.apply(fields,macros,index,t_i); break; case 3u: @@ -385,6 +420,7 @@ saw::error_or lbm_main(int argc, char** argv){ } } */ + /* if(i.get() % 32u == 0u){ { auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr); @@ -399,6 +435,7 @@ saw::error_or lbm_main(int argc, char** argv){ } } } + */ // Stream sycl_q.submit([&](acpp::sycl::handler& h){ component> stream; diff --git a/lib/core/c++/particle/particle_opa.hpp b/lib/core/c++/particle/particle_opa.hpp index e6396d6..bfd2e3c 100644 --- a/lib/core/c++/particle/particle_opa.hpp +++ b/lib/core/c++/particle/particle_opa.hpp @@ -35,12 +35,13 @@ public: auto& porous = porous_f.at(index); - auto pos_ind = saw::math::vectorize_data(index); auto diff = pos_ind - pos_; // Write out value + auto diff = pos_ind.template cast_to() - pos_; + auto diff_dot = saw::math::dot(diff,diff); porous = particle_porosity<>::calculate(diff, rad_, eps_); } }; diff --git a/scripts/python/csv_l2_norm.py b/scripts/python/csv_l2_norm.py new file mode 100755 index 0000000..88f4817 --- /dev/null +++ b/scripts/python/csv_l2_norm.py @@ -0,0 +1,26 @@ +#!/usr/bin/env python3 + +import argparse +import numpy as np + + +def main(): + parser = argparse.ArgumentParser(description="Compute the L2 Norm"); + parser.add_argument("one"); + parser.add_argument("two"); + args = parser.parse_args(); + + a = np.loadtxt(args.one, delimiter=","); + b = np.loadtxt(args.two, delimiter=","); + + if a.shape != b.shape: + raise ValueError(f"Shape mismatch: {a.shape} vs {b.shape}"); + #endif + + print(np.linalg.norm(a-b)); + +#enddef + +if __name__ == "__main__": + main(); +#enddef -- cgit v1.2.3 From b88d59477a7973bdee102aaf0e26c13c9059048b Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Mon, 29 Jun 2026 22:44:03 +0200 Subject: Fixing psm resolution into fiedl --- examples/poiseulle_particles_2d_psm_gpu/sim.cpp | 2 + lib/core/c++/particle/particle_opa.hpp | 8 ++-- lib/core/c++/particle/porosity.hpp | 62 ++++++++++++++++--------- lib/core/c++/particle/schema.hpp | 1 + 4 files changed, 46 insertions(+), 27 deletions(-) diff --git a/examples/poiseulle_particles_2d_psm_gpu/sim.cpp b/examples/poiseulle_particles_2d_psm_gpu/sim.cpp index e91686c..b3ea7ae 100644 --- a/examples/poiseulle_particles_2d_psm_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_psm_gpu/sim.cpp @@ -169,6 +169,7 @@ saw::error_or setup_initial_conditions( p_pos.at({{1u}}) = dim_y * 0.5; } + /* component opa{p_pos,rad,eps}; iterator::apply( [&](auto& index){ @@ -177,6 +178,7 @@ saw::error_or setup_initial_conditions( {},// 0-index df_f.get_dims() ); + */ /* iterator::apply( diff --git a/lib/core/c++/particle/particle_opa.hpp b/lib/core/c++/particle/particle_opa.hpp index bfd2e3c..a5e6c28 100644 --- a/lib/core/c++/particle/particle_opa.hpp +++ b/lib/core/c++/particle/particle_opa.hpp @@ -1,7 +1,9 @@ #pragma once -#include "common.hpp" #include "../component.hpp" +#include "common.hpp" +#include "schema.hpp" +#include "porosity.hpp" namespace kel { namespace lbm { @@ -37,12 +39,10 @@ public: auto pos_ind = saw::math::vectorize_data(index); - auto diff = pos_ind - pos_; - // Write out value auto diff = pos_ind.template cast_to() - pos_; auto diff_dot = saw::math::dot(diff,diff); - porous = particle_porosity<>::calculate(diff, rad_, eps_); + porous = particle_porosity>::calculate(diff, rad_, eps_); } }; } diff --git a/lib/core/c++/particle/porosity.hpp b/lib/core/c++/particle/porosity.hpp index f5966c4..11185c5 100644 --- a/lib/core/c++/particle/porosity.hpp +++ b/lib/core/c++/particle/porosity.hpp @@ -1,15 +1,21 @@ #pragma once -#include "particle.hpp" +#include "common.hpp" +#include "schema.hpp" #include "../math/n_closest.hpp" namespace kel { namespace lbm { +namespace por { +template +struct ParticleSpheroid {}; +} + template class particle_porosity { public: - saw::data> calculate(const saw::data>& part_group, uint64_t p_i, const saw::data>& lbm_pos){ + static saw::data> calculate(const saw::data>& part_group, uint64_t p_i, const saw::data>& lbm_pos){ auto& mask = part_group.template get<"mask">(); auto& particles = part_group.template get<"particles">(); @@ -26,31 +32,32 @@ public: } }; - - -template::type radius, typename saw::native_data_type::type eps> -class particle_porosity> final { +template +class particle_porosity> final { public: - saw::data> calculate(const saw::data>& lbm_rel_dist, saw::data> rad, saw::data> eps) const { + static saw::data> calculate(const saw::data>& lbm_rel_dist, saw::data> rad, saw::data> eps){ saw::data> por; auto s_dist_2 = saw::math::dot(lbm_rel_dist,lbm_rel_dist); auto s_dist = saw::math::sqrt(s_dist_2); - auto rad_low = rad - eps; - auto rad_high = rad + eps; + saw::data> eps_h; + eps_h.at({}) = eps.at({}).get() / 2; - if(s_dist < rad_low){ + auto rad_low = rad - eps_h; + auto rad_high = rad + eps_h; + + if(s_dist.at({}) <= rad_low.at({})){ por.at({}).set(1); return por; } - if(s_dist > rad_high){ + if(s_dist.at({}) >= rad_high.at({})){ por.at({}).set(0); return por; } { - typename saw::native_data_type::type inner = (std::pi / ( 2.0 * eps.at({}).get() )) * (s_dist - rad_low).at({}).get(); + typename saw::native_data_type::type inner = (std::numbers::pi / ( eps.at({}).get() )) * (s_dist - rad_low).at({}).get(); auto cos_inner = std::cos(inner); auto cos_i_2 = cos_inner * cos_inner; por.at({}).set(cos_i_2); @@ -58,10 +65,14 @@ public: return por; } +}; + - saw::data> calculate(const saw::data > >& part_group, uint64_t i, const saw::data>& lbm_pos) const { +template +class particle_porosity> final { +public: + static saw::data> calculate(const saw::data > >& part_group, uint64_t i, const saw::data>& lbm_pos) { saw::data> por; - por.at({}); auto& parts = part_group.template get<"particles">(); auto& pi = parts.at({i}); @@ -73,13 +84,18 @@ public: auto dist = lbm_pos - pi_rb_pos; saw::data> dist_len = saw::math::dot(dist,dist); - dist_len.at({}).set(std::sqrt(dist_len.at({}).get()); + dist_len.at({}).set(std::sqrt(dist_len.at({}).get())); - saw::data> rad_d{radius}; - saw::data> eps_half{eps *0.5}; + auto& coll = part_group.template get<"collision">(); + const saw::data>& rad_d = coll.template get<"radius">(); + const auto& eps = part_group.template get<"epsilon">(); + // Move this somewhere - saw::data> rad_d_eps_p = rad_d + eps_half; - saw::data> rad_d_eps_n = rad_d - eps_half; + saw::data> eps_h; + eps_h.at({}) = eps.at({}).get() / 2; + + saw::data> rad_d_eps_p = rad_d + eps_h; + saw::data> rad_d_eps_n = rad_d - eps_h; // saw::data> rad_d_eps_p_2 = rad_d_eps_p * rad_d_eps_p; // saw::data> rad_d_eps_n_2 = rad_d_eps_n * rad_d_eps_n; @@ -88,19 +104,19 @@ public: saw::data> inside; if(dist_len.at({}) <= rad_d_eps_n.at({})){ - inside.at({}).set(1); + inside.at({}).set(0); return inside; } if(dist_len.at({}) > rad_d_eps_p.at({})){ - inside.at({}).set(0); - return inside + inside.at({}).set(1); + return inside; } /* * cos^2 ( ||x-X(t)||_2 - (R-eps/2) ) */ { - typename saw::native_data_type::type inner = (std::pi / ( 2.0 * eps )) * (dist_len - rad_d_eps_n).at({}).get(); + typename saw::native_data_type::type inner = (std::numbers::pi / (eps)) * (dist_len - rad_d_eps_n).at({}).get(); auto cos_inner = std::cos(inner); auto cos_i_2 = cos_inner * cos_inner; inside.at({}).set(cos_i_2); diff --git a/lib/core/c++/particle/schema.hpp b/lib/core/c++/particle/schema.hpp index 18a697a..714a16f 100644 --- a/lib/core/c++/particle/schema.hpp +++ b/lib/core/c++/particle/schema.hpp @@ -56,6 +56,7 @@ template> using ParticleGroup = Struct< Member, "mask">, Member, "collision">, + Member,1u>, "epsilon">, Member,1u>, "mask_step">, Member,1u>, "density">, Member,1u>, "center_of_mass">, -- cgit v1.2.3