From 43d795c0afa3c613fb3aee10de7ddf579cf78645 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Tue, 7 Jul 2026 20:33:25 +0200 Subject: Fixed Stokes example (yikes) --- examples/stokes_drag_particle_2d_fplbm_gpu/sim.cpp | 33 +++++++++++++----- examples/stokes_drag_particle_2d_hlbm_gpu/sim.cpp | 33 +++++++++++++----- examples/stokes_drag_particle_2d_psm_gpu/sim.cpp | 31 ++++++++++++----- modules/core/c++/boundary.hpp | 39 ++++++++++++++++++---- 4 files changed, 103 insertions(+), 33 deletions(-) diff --git a/examples/stokes_drag_particle_2d_fplbm_gpu/sim.cpp b/examples/stokes_drag_particle_2d_fplbm_gpu/sim.cpp index fb31f47..5de963d 100644 --- a/examples/stokes_drag_particle_2d_fplbm_gpu/sim.cpp +++ b/examples/stokes_drag_particle_2d_fplbm_gpu/sim.cpp @@ -150,7 +150,7 @@ saw::error_or setup_initial_conditions( rho.at({}) = {1}; auto& vel = vel_f.at(index); if(info_f.at(index).get() == 2u){ - vel.at({{0u}}) = 0.0; + vel.at({{0u}}) = 0.0002f; } auto eq = equilibrium(rho,vel); @@ -214,17 +214,30 @@ saw::error_or step( } component> opa{p_pos,rad,eps}; - component,encode::Sycl> flow_in{ - [&](){ - uint64_t target_t_i = 16u; + component,encode::Sycl> flow_in{ + [&]() -> saw::data>{ + saw::data> vel; + uint64_t target_t_i = 8u; if(t_i.get() < target_t_i){ - return 1.0 + (0.00002 / target_t_i) * t_i.get(); + vel.at({{0u}}) = (0.0002 / target_t_i) * t_i.get(); + }else{ + vel.at({{0u}}) = 0.0002; } - return 1.00002; + return vel; + }() + }; + component,encode::Sycl> flow_out{ + [&]() -> saw::data>{ + saw::data> vel; + uint64_t target_t_i = 8u; + if(t_i.get() < target_t_i){ + vel.at({{0u}}) = (0.0002 / target_t_i) * t_i.get(); + }else{ + vel.at({{0u}}) = 0.0002; + } + return vel; }() }; - component,encode::Sycl> flow_out{1.0}; - h.parallel_for(acpp::sycl::range{dim_x,dim_y}, [=](acpp::sycl::id idx){ saw::data> index; @@ -365,7 +378,9 @@ saw::error_or lbm_main(int argc, char** argv){ sycl_q.wait(); auto lsd_view = make_view(lbm_sycl_data); auto lsdm_view = make_view(lbm_sycl_macro_data); - saw::data time_steps{16u*4096ul}; + + saw::data time_steps{32u*4096ul}; + auto& info_f = lsd_view.template get<"info">(); for(saw::data i{0u}; i < time_steps and krun; ++i){ diff --git a/examples/stokes_drag_particle_2d_hlbm_gpu/sim.cpp b/examples/stokes_drag_particle_2d_hlbm_gpu/sim.cpp index 9a04211..68c0b35 100644 --- a/examples/stokes_drag_particle_2d_hlbm_gpu/sim.cpp +++ b/examples/stokes_drag_particle_2d_hlbm_gpu/sim.cpp @@ -147,7 +147,7 @@ saw::error_or setup_initial_conditions( rho.at({}) = {1}; auto& vel = vel_f.at(index); if(info_f.at(index).get() == 2u){ - vel.at({{0u}}) = 0.0; + vel.at({{0u}}) = 0.0002f; } auto eq = equilibrium(rho,vel); @@ -214,17 +214,30 @@ saw::error_or step( component> fg; - component,encode::Sycl> flow_in{ - [&](){ - uint64_t target_t_i = 16u; + component,encode::Sycl> flow_in{ + [&]() -> saw::data>{ + saw::data> vel; + uint64_t target_t_i = 8u; if(t_i.get() < target_t_i){ - return 1.0 + (0.00002 / target_t_i) * t_i.get(); + vel.at({{0u}}) = (0.0002 / target_t_i) * t_i.get(); + }else{ + vel.at({{0u}}) = 0.0002; } - return 1.00002; + return vel; + }() + }; + component,encode::Sycl> flow_out{ + [&]() -> saw::data>{ + saw::data> vel; + uint64_t target_t_i = 8u; + if(t_i.get() < target_t_i){ + vel.at({{0u}}) = (0.0002 / target_t_i) * t_i.get(); + }else{ + vel.at({{0u}}) = 0.0002; + } + return vel; }() }; - component,encode::Sycl> flow_out{1.0}; - h.parallel_for(acpp::sycl::range{dim_x,dim_y}, [=](acpp::sycl::id idx){ saw::data> index; @@ -367,7 +380,9 @@ saw::error_or lbm_main(int argc, char** argv){ sycl_q.wait(); auto lsd_view = make_view(lbm_sycl_data); auto lsdm_view = make_view(lbm_sycl_macro_data); - saw::data time_steps{16u*4096ul}; + + saw::data time_steps{32u*4096ul}; + auto& info_f = lsd_view.template get<"info">(); for(saw::data i{0u}; i < time_steps and krun; ++i){ diff --git a/examples/stokes_drag_particle_2d_psm_gpu/sim.cpp b/examples/stokes_drag_particle_2d_psm_gpu/sim.cpp index 9781691..3ab3dda 100644 --- a/examples/stokes_drag_particle_2d_psm_gpu/sim.cpp +++ b/examples/stokes_drag_particle_2d_psm_gpu/sim.cpp @@ -134,6 +134,7 @@ saw::error_or setup_initial_conditions( por_f.at(index).at({}) = {1}; rho.at({}) = {1}; auto& vel = vel_f.at(index); + auto eq = equilibrium(rho,vel); df = eq; @@ -149,7 +150,7 @@ saw::error_or setup_initial_conditions( rho.at({}) = {1}; auto& vel = vel_f.at(index); if(info_f.at(index).get() == 2u){ - vel.at({{0u}}) = 0.0; + vel.at({{0u}}) = 0.0002f; } auto eq = equilibrium(rho,vel); @@ -233,16 +234,30 @@ saw::error_or step( } component> opa{p_pos,rad,eps}; - component,encode::Sycl> flow_in{ - [&](){ - uint64_t target_t_i = 16u; + component,encode::Sycl> flow_in{ + [&]() -> saw::data>{ + saw::data> vel; + uint64_t target_t_i = 8u; + if(t_i.get() < target_t_i){ + vel.at({{0u}}) = (0.0002 / target_t_i) * t_i.get(); + }else{ + vel.at({{0u}}) = 0.0002; + } + return vel; + }() + }; + component,encode::Sycl> flow_out{ + [&]() -> saw::data>{ + saw::data> vel; + uint64_t target_t_i = 8u; if(t_i.get() < target_t_i){ - return 1.0 + (0.00002 / target_t_i) * t_i.get(); + vel.at({{0u}}) = (0.0002 / target_t_i) * t_i.get(); + }else{ + vel.at({{0u}}) = 0.0002; } - return 1.00002; + return vel; }() }; - component,encode::Sycl> flow_out{1.0}; component> fg; @@ -389,7 +404,7 @@ saw::error_or lbm_main(int argc, char** argv){ auto lsd_view = make_view(lbm_sycl_data); auto lsdm_view = make_view(lbm_sycl_macro_data); - saw::data time_steps{16u*4096ul}; + saw::data time_steps{32u*4096ul}; auto& info_f = lsd_view.template get<"info">(); diff --git a/modules/core/c++/boundary.hpp b/modules/core/c++/boundary.hpp index 01ae7b5..22a6f34 100644 --- a/modules/core/c++/boundary.hpp +++ b/modules/core/c++/boundary.hpp @@ -221,14 +221,14 @@ public: } }; -/* -template -class component, Encode> final { + +template +class component, Encode> final { private: - saw::data> velocity_; + saw::data> velocity_; public: component( - saw::data> velocity__ + saw::data> velocity__ ): velocity_{velocity__} {} @@ -237,13 +237,38 @@ public: void apply(const saw::data& field, const saw::data>& index, saw::data time_step) const { bool is_even = ((time_step.get() % 2u) == 0u); - using dfi = df_info; + using dfi = df_info; auto& dfs_old_f = (is_even) ? field.template get<"dfs_old">() : field.template get<"dfs">(); + + auto& dfs = dfs_old_f.at(index); + + saw::data dir_sum; + if constexpr (East) { + dir_sum = dfs.at({1u}) + dfs.at({5u}) + dfs.at({7u}); + } else if constexpr(not East){ + dir_sum = dfs.at({2u}) + dfs.at({6u}) + dfs.at({8u}); + } + + saw::data> rho; + rho.at({}) = (dfs.at({0u}) + dfs.at({4u}) + dfs.at({3u}) + dir_sum * 2) / (velocity_.at({{0u}}) + saw::data{static_cast::type>(1.0)}); + + if constexpr (East) { + dfs.at({2u}) = dfs.at({1u}) + saw::data{static_cast::type>(2.0 / 3.0)} * rho.at({}) * velocity_.at({{0u}}); + dfs.at({6u}) = dfs.at({5u}) + saw::data{static_cast::type>(1.0 / 6.0)} * rho.at({}) * velocity_.at({{0u}}) + + (dfs.at({3u}) - dfs.at({4u}) + rho.at({}) * velocity_.at({{1u}})) * 0.5f; + dfs.at({8u}) = dfs.at({7u}) + saw::data{static_cast::type>(1.0 / 6.0)} * rho.at({}) * velocity_.at({{0u}}) + + (dfs.at({4u}) - dfs.at({3u}) + rho.at({}) * velocity_.at({{1u}})) * 0.5f; + }else if constexpr ( not East ){ + dfs.at({1u}) = dfs.at({2u}) + saw::data{static_cast::type>(2.0 / 3.0)} * rho.at({}) * velocity_.at({{0u}}); + dfs.at({5u}) = dfs.at({6u}) + saw::data{static_cast::type>(1.0 / 6.0)} * rho.at({}) * velocity_.at({{0u}}) + + (dfs.at({4u}) - dfs.at({3u}) + rho.at({}) * velocity_.at({{1u}}))*0.5f; + dfs.at({7u}) = dfs.at({8u}) + saw::data{static_cast::type>(1.0 / 6.0)} * rho.at({}) * velocity_.at({{0u}}) + + (dfs.at({3u}) - dfs.at({4u}) + rho.at({}) * velocity_.at({{1u}}))*0.5f; + } } }; -*/ } } -- cgit v1.2.3