From eeb01452fe46fcb5efdc9c34b660305262097ca4 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Wed, 10 Dec 2025 11:26:13 +0100 Subject: Dangling changes --- examples/poiseulle_2d_gpu/poiseulle_2d_gpu.cpp | 364 +++++-------------------- 1 file changed, 66 insertions(+), 298 deletions(-) (limited to 'examples/poiseulle_2d_gpu/poiseulle_2d_gpu.cpp') diff --git a/examples/poiseulle_2d_gpu/poiseulle_2d_gpu.cpp b/examples/poiseulle_2d_gpu/poiseulle_2d_gpu.cpp index a319122..d52af6a 100644 --- a/examples/poiseulle_2d_gpu/poiseulle_2d_gpu.cpp +++ b/examples/poiseulle_2d_gpu/poiseulle_2d_gpu.cpp @@ -1,10 +1,14 @@ -#include "../c++/lbm.hpp" -#include "../c++/collision.hpp" -#include "../c++/boundary.hpp" -#include "../c++/iterator.hpp" +#include +#include #include +template +using SyclHostAlloc = acpp::sycl::usm_allocator, acpp::sycl::usm::alloc::host>; + +template +using SyclDeviceAlloc = acpp::sycl::usm_allocator, acpp::sycl::usm::alloc::device>; + namespace kel { namespace lbm { namespace sch { @@ -19,24 +23,19 @@ using namespace saw::schema; * D factor * Q factor */ -using T = Float32; -using D2Q5 = Descriptor<2u,5u>; +using T = Float64; using D2Q9 = Descriptor<2u,9u>; -template -using DfCell = Cell; - -template +using DfCell = Cell; using CellInfo = Cell; /** * Basic type for simulation */ -template using CellStruct = Struct< - Member, "dfs">, - Member, "dfs_old">, - Member, "info"> + Member, + Member, + Member >; template @@ -45,12 +44,6 @@ using MacroStruct = Struct< Member >; -template -using GeometryStruct = Struct< - Member ->; - -using CavityFieldD2Q9 = CellField>; } namespace cmpt { @@ -70,169 +63,40 @@ struct PressureBoundaryRestrictedVelocityTo {}; * 0 - 2 - 2 * */ -template -struct component> { -private: - saw::data pressure_setting_; - saw::data rho_setting_; -public: - component(const saw::data& pressure_setting__): - pressure_setting_{pressure_setting__}, - rho_setting_{pressure_setting__ * df_info::inv_cs2} - {} - - template - void apply(saw::data& field, saw::data> index, uint64_t time_step){ - using dfi = df_info; - - bool is_even = ((time_step % 2) == 0); - auto& cell = field(index); - - auto& info = cell.template get<"info">(); - if(info({0u}).get() == 0u){ - return; - } - auto& dfs_old = (is_even) ? cell.template get<"dfs_old">() : cell.template get<"dfs">(); - // auto& dfs = (not is_even) ? cell.template get<"dfs_old">() : cell.template get<"dfs">(); - - /** - * Sum all known DFs - */ - saw::data sum_df{0}; - for(saw::data k{0u}; k < saw::data{Descriptor::Q}; ++k){ - auto c_k = dfi::directions[k.get()]; - auto& cell_n = field({{index.at({0u})+c_k[0u], index.at({1u})+c_k[1u]}}); - auto& info_n = cell_n.template get<"info">(); - auto info_n_val = info_n({0u}); - auto k_opp = dfi::opposite_index[k.get()]; - - if(info_n_val.get() > 0u){ - sum_df += dfs_old({k_opp}); - } - } - /** - * Get the sum of the unknown dfs and precalculate the direction - */ - constexpr int known_dir = East ? 1 : -1; - auto sum_unknown_dfs = (rho_setting_ - sum_df) * saw::data{known_dir}; - - for(saw::data k{0u}; k < saw::data{Descriptor::Q}; ++k){ - auto c_k = dfi::directions[k.get()]; - auto& cell_n = field({{index.at({0u})+c_k[0u], index.at({1u})+c_k[1u]}}); - auto& info_n = cell_n.template get<"info">(); - auto info_n_val = info_n({0u}); - // auto k_opp = dfi::opposite_index[k.get()]; - - if(info_n_val.get() > 0u){ - sum_unknown_dfs += dfs_old({k}) * c_k[0u]; - } - } - - auto vel_x = sum_unknown_dfs / rho_setting_; - - if constexpr (East) { - dfs_old({2u}) = dfs_old({1u}) + saw::data{2.0 / 3.0} * rho_setting_ * vel_x; - dfs_old({6u}) = dfs_old({5u}) + saw::data{1.0 / 6.0} * rho_setting_ * vel_x + saw::data{0.5} * (dfs_old({4u}) - dfs_old({3u})); - dfs_old({8u}) = dfs_old({7u}) + saw::data{1.0 / 6.0} * rho_setting_ * vel_x + saw::data{0.5} * (dfs_old({3u}) - dfs_old({4u})); - }else if constexpr (not East){ - dfs_old({1u}) = dfs_old({2u}) - saw::data{2.0 / 3.0} * rho_setting_ * vel_x; - dfs_old({5u}) = dfs_old({6u}) - saw::data{1.0 / 6.0} * rho_setting_ * vel_x + saw::data{0.5} * (dfs_old({3u}) - dfs_old({4u})); - dfs_old({7u}) = dfs_old({8u}) - saw::data{1.0 / 6.0} * rho_setting_ * vel_x + saw::data{0.5} * (dfs_old({4u}) - dfs_old({3u})); - } - } -}; -} -} -void set_geometry(saw::data& latt){ +void set_geometry( + std::vector, SyclDeviceAlloc>& dfs, + std::vector, SyclDeviceAlloc>& dfs_old, + std::vector, SyclDeviceAlloc>& info, + const saw::data>& meta, + acpp::sycl::queue& sycl_q +){ using namespace kel::lbm; - auto meta = latt.meta(); - - /** - * Set ghost - */ - iterate_over([&](const saw::data>& index){ - auto& cell = latt(index); - auto& info = cell.template get<"info">(); - - info({0u}).set(0u); - - }, {{0u,0u}}, meta); - - /** - * Set wall - */ - iterate_over([&](const saw::data>& index){ - auto& cell = latt(index); - auto& info = cell.template get<"info">(); - - info({0u}).set(2u); - - }, {{0u,0u}}, meta, {{1u,1u}}); - - /** - * Set fluid - */ - iterate_over([&](const saw::data>& index){ - auto& cell = latt(index); - auto& info = cell.template get<"info">(); - - info({0u}).set(1u); - - }, {{0u,0u}}, meta, {{2u,2u}}); - - /** - * Set inflow - */ - iterate_over([&](const saw::data>& index){ - auto& cell = latt(index); - auto& info = cell.template get<"info">(); - - info({0u}).set(3u); - - }, {{1u,0u}}, {{2u,meta.at({1u})}}, {{0u,2u}}); - - /** - * Set outflow - */ - iterate_over([&](const saw::data>& index){ - auto& cell = latt(index); - auto& info = cell.template get<"info">(); - - info({0u}).set(4u); - - }, {{meta.at({0u})-2u,0u}}, {{meta.at({0u})-1u, meta.at({1u})}}, {{0u,2u}}); } -void set_initial_conditions(saw::data& latt){ +void set_initial_conditions( + std::vector, SyclDeviceAlloc>& dfs, + std::vector, SyclDeviceAlloc>& dfs_old, + std::vector, SyclDeviceAlloc>& info, + const saw::data>& meta, + acpp::sycl::queue& sycl_q +){ using namespace kel::lbm; saw::data rho{1.0}; saw::data> vel{{0.0,0.0}}; auto eq = equilibrium(rho, vel); - auto meta = latt.meta(); - - /** - * Set distribution - */ - iterate_over([&](const saw::data>& index){ - auto& cell = latt(index); - auto& dfs = cell.template get<"dfs">(); - auto& dfs_old = cell.template get<"dfs_old">(); - - for(saw::data k = 0; k < saw::data{sch::D2Q9::Q}; ++k){ - dfs(k) = eq.at(k); - dfs_old(k) = eq.at(k); - } - - }, {{0u,0u}}, meta); } void lbm_step( - saw::data& latt, - uint64_t time_step + std::vector, SyclDeviceAlloc>& dfs, + std::vector, SyclDeviceAlloc>& dfs_old, + std::vector, SyclDeviceAlloc>& info, + const saw::data>& meta, + uint64_t time_step, + acpp::sycl::queue& sycl_q ){ using namespace kel::lbm; using dfi = df_info; @@ -246,55 +110,8 @@ void lbm_step( component> inlet{1.01 * dfi::cs2}; component> outlet{1.0 * dfi::cs2}; - auto meta = latt.meta(); - - /** - * Collision - */ - iterator::apply([&](const saw::data>& index){ - auto& cell = latt(index); - auto& info = cell.template get<"info">(); - - switch(info({0u}).get()){ - case 1u: { - coll.apply(latt, index, time_step); - break; - } - case 2u: { - bb.apply(latt, index, time_step); - break; - } - case 3u: { - inlet.apply(latt, index, time_step); - break; - } - case 4u: { - outlet.apply(latt, index, time_step); - break; - } - default: - break; - } - - }, {{0u,0u}}, meta); - - // Stream - iterator::apply([&](const saw::data>& index){ - auto& cell = latt(index); - auto& df_new = even_step ? cell.template get<"dfs">() : cell.template get<"dfs_old">(); - auto& info_new = cell.template get<"info">(); - - if(info_new({0u}).get() > 0u){ - for(uint64_t k = 0u; k < sch::D2Q9::Q; ++k){ - auto dir = dfi::directions[dfi::opposite_index[k]]; - auto& cell_dir_old = latt({{index.at({0u})+dir[0],index.at({1u})+dir[1]}}); - - auto& df_old = even_step ? cell_dir_old.template get<"dfs_old">() : cell_dir_old.template get<"dfs">(); - - df_new({k}) = df_old({k}); - } - } - }, {{0u,0u}}, meta); +} +} } template @@ -307,100 +124,51 @@ saw::error_or kel_main(int argc, char** argv){ if(eo_lbm_dir.is_error()){ return std::move(eo_lbm_dir.get_error()); } - - return saw::make_void(); -} - -int main(int argc, char** argv){ - using namespace kel::lbm; - using dfi = df_info; - auto eo_lbm_dir = output_directory(); - if(eo_lbm_dir.is_error()){ - return -1; - } auto& lbm_dir = eo_lbm_dir.get_value(); - auto out_dir = lbm_dir / "poiseulle_particles_channel_2d"; - - std::string_view cfg_file_name = "config.json"; - if(argc > 1){ - cfg_file_name = argv[1]; - } - - auto eo_conf = load_lbm_config>(cfg_file_name); - if(eo_conf.is_error()){ - auto& err = eo_conf.get_error(); - std::cerr<<"[Error]: "< conv { - {conf.template get<"delta_x">()}, - {conf.template get<"delta_t">()} + // + lbm::converter conv { + // delta_x + {{1.0}}, + // delta_t + {{1.0}} }; - print_lbm_meta>(conv, {conf.template get<"kinematic_viscosity">()}); + uint64_t x_d = 1024u; + uint64_t y_d = 128u; - saw::data> dim{{1024u, 128u}}; - saw::data lattice{dim}; - auto meta = lattice.meta(); + acpp::sycl::queue sycl_q; + SyclHostAlloc sycl_host_alloc{sycl_q}; + SyclDeviceAlloc sycl_dev_alloc{sycl_q}; - /** - * Setup geometry - */ - set_geometry(lattice); + std::vector, SyclHostAlloc> cells{x_d * y_d,sycl_alloc}; + std::vector, SyclDeviceAlloc> cells{x_d * y_d,sycl_alloc}; { - - saw::data, sch::D2Q9::D>> geo{dim}; - - iterate_over([&](const saw::data>& index){ - auto& cell = lattice(index); - auto& info = cell.template get<"info">(); - - geo(index).template get<"info">().set(info({0u}).get()); - }, {{0u,0u}}, dim); - - std::string vtk_f_name{"tmp/geometry.vtk"}; - write_vtk_file(vtk_f_name, geo); + auto eov = set_geometry(cells); + if(eov.is_error()){ + return eov; + } } - /** - * Setup DFs - */ - set_initial_conditions(lattice); - - saw::data,sch::D2Q9::D>> macros{dim}; - for(uint64_t i = 0u; i < 4096u*16u; ++i){ - bool even_step = ((i % 2u) == 0u); - - { - // Stream - iterate_over([&](const saw::data>& index){ - auto& cell = lattice(index); - auto& dfs = even_step ? cell.template get<"dfs">() : cell.template get<"dfs_old">(); - - auto& rho = macros.at(index).template get<"pressure">(); - auto& vel = macros.at(index).template get<"velocity">(); - compute_rho_u(dfs,rho,vel); - rho = rho * saw::data{dfi::cs2}; + return saw::make_void(); +} - }, {{0u,0u}}, meta); - std::string vtk_f_name{"poiseulle_2d_"}; - vtk_f_name += std::to_string(i) + ".vtk"; - write_vtk_file(out_dir / vtk_f_name, macros); +int main(int argc, char** argv){ + auto eov = kel_main(argc, argv); + if(eov.is_error()){ + auto& err = eov.get_error(); + std::cerr<<"[Error] "< 0u){ + std::cerr<<" - "<