From 76c505727034567a5e8fad647f73af3b2f2141cd Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Wed, 5 Nov 2025 15:41:15 +0100 Subject: Moving to find out how to nicely build a wrapper for data --- examples/poiseulle_3d/.nix/derivation.nix | 3 +- examples/poiseulle_3d/poiseulle_3d.cpp | 64 +++++++++++++++++++------------ 2 files changed, 41 insertions(+), 26 deletions(-) diff --git a/examples/poiseulle_3d/.nix/derivation.nix b/examples/poiseulle_3d/.nix/derivation.nix index fdf4ad8..d53b5fe 100644 --- a/examples/poiseulle_3d/.nix/derivation.nix +++ b/examples/poiseulle_3d/.nix/derivation.nix @@ -28,8 +28,7 @@ stdenv.mkDerivation { forstio.remote forstio.remote-sycl adaptive-cpp - kel-lbm.core - kel-lbm.sycl + kel-lbm.core ]; preferLocalBuild = true; diff --git a/examples/poiseulle_3d/poiseulle_3d.cpp b/examples/poiseulle_3d/poiseulle_3d.cpp index 2c268d4..707212d 100644 --- a/examples/poiseulle_3d/poiseulle_3d.cpp +++ b/examples/poiseulle_3d/poiseulle_3d.cpp @@ -3,6 +3,7 @@ #include namespace saw { +/* template class data, encode::Sycl> final { public: @@ -54,6 +55,7 @@ public: return inner_.internal_data(); } }; +*/ /* template class data>, encode::Sycl> final { @@ -159,18 +161,23 @@ using CellStruct = Struct< Member >; } +using SyclKelAllocator = acpp::sycl::usm_allocator, acpp::sycl::usm::alloc::shared>; -template -saw::error_or set_geometry(saw::data& latt){ - using namespace kel::lbm; +uint64_t to_flat_index(const saw::data>& index, const saw::data>& meta){ + auto flat = index.at({2u}) + meta.at({2u})*(index.at({1u}) + index.at({0u})*meta.at({1u})); + return flat.get(); +} - auto meta = latt.meta(); +template +saw::error_or set_geometry(std::vector,SyclKelAllocator>& latt, const saw::data>& meta){ + using namespace kel::lbm; /** * Set ghost */ iterator::apply([&](const saw::data>& index){ - auto& cell = latt.at(index); + auto findex = to_flat_index(index,meta); + auto& cell = latt.at(findex); auto& info = cell.template get<"info">(); info({0u}).set(0u); @@ -187,7 +194,8 @@ saw::error_or set_geometry(saw::data& latt){ /* */ iterator::apply([&](const saw::data>& index){ - auto& cell = latt.at(index); + auto findex = to_flat_index(index,meta); + auto& cell = latt.at(findex); auto& info = cell.template get<"info">(); info({0u}).set(2u); @@ -202,7 +210,8 @@ saw::error_or set_geometry(saw::data& latt){ /* */ iterator::apply([&](const saw::data>& index){ - auto& cell = latt.at(index); + auto findex = to_flat_index(index,meta); + auto& cell = latt.at(findex); auto& info = cell.template get<"info">(); info({0u}).set(1u); @@ -221,7 +230,8 @@ saw::error_or set_geometry(saw::data& latt){ /* */ iterator::apply([&](const saw::data>& index){ - auto& cell = latt.at(index); + auto findex = to_flat_index(index,meta); + auto& cell = latt.at(findex); auto& info = cell.template get<"info">(); info({0u}).set(3u); @@ -238,7 +248,8 @@ saw::error_or set_geometry(saw::data& latt){ /* */ iterator::apply([&](const saw::data>& index){ - auto& cell = latt.at(index); + auto findex = to_flat_index(index,meta); + auto& cell = latt.at(findex); auto& info = cell.template get<"info">(); info({0u}).set(4u); @@ -248,21 +259,20 @@ saw::error_or set_geometry(saw::data& latt){ return saw::make_void(); } -template -saw::error_or set_initial_conditions(saw::data& latt){ +template +saw::error_or set_initial_conditions(std::vector,SyclKelAllocator>& latt, const saw::data>& meta){ saw::data rho{1.0}; saw::data> vel{}; vel.at({0u}) = {0.1f}; auto eq = equilibrium(rho, vel); - auto meta = latt.meta(); - /** * Set distribution */ iterator::apply([&](const saw::data>& index){ - auto& cell = latt.at(index); + auto findex = to_flat_index(index,meta); + auto& cell = latt.at(findex); auto& dfs = cell.template get<"dfs">(); auto& dfs_old = cell.template get<"dfs_old">(); @@ -274,14 +284,14 @@ saw::error_or set_initial_conditions(saw::data& return saw::make_void(); } -template +template saw::error_or lbm_step( - saw::data& latt, + std::vector,SyclKelAllocator>& latt, + const saw::data>& meta, saw::data time_step, acpp::sycl::queue& sycl_q ){ using dfi = df_info; - auto meta = latt.meta(); bool even_step = ((time_step.get() % 2u) == 0u); component> coll{0.5384}; @@ -291,7 +301,8 @@ saw::error_or lbm_step( * Collision */ iterator::apply([&](const saw::data>& index){ - auto& cell = latt.at(index); + auto findex = to_flat_index(index,meta); + auto& cell = latt.at(findex); auto& info = cell.template get<"info">(); switch(info({0u}).get()){ @@ -319,7 +330,8 @@ saw::error_or lbm_step( // Stream iterator::apply([&](const saw::data>& index){ - auto& cell = latt(index); + auto findex = to_flat_index(index,meta); + auto& cell = latt.at(findex); auto& df_new = even_step ? cell.template get<"dfs">() : cell.template get<"dfs_old">(); auto& info_new = cell.template get<"info">(); @@ -330,7 +342,9 @@ saw::error_or lbm_step( for(uint64_t i = 0; i < sch::Desc::D; ++i){ index_dir.at({i}) = index_dir.at({i}) + dir[i]; } - auto& cell_dir_old = latt.at(index_dir); + + auto findex_dir = to_flat_index(index_dir,meta); + auto& cell_dir_old = latt.at(findex_dir); auto& df_old = even_step ? cell_dir_old.template get<"dfs_old">() : cell_dir_old.template get<"dfs">(); @@ -356,13 +370,15 @@ saw::error_or real_main(int argc, char** argv){ * Set the meta and initialize the lattice */ saw::data> meta{{2048u,128u,128u}}; - saw::data, saw::encode::Sycl> lattice{meta,sycl_q}; + + SyclKelAllocator sycl_shared_alloc{sycl_q}; + std::vector,SyclKelAllocator> lattice{(meta.at({0u})*meta.at({1u})*meta.at({2u})).get(), sycl_shared_alloc}; /** * Set the geometry */ { - auto eov = set_geometry(lattice); + auto eov = set_geometry(lattice, meta); if(eov.is_error()){ return eov; } @@ -372,14 +388,14 @@ saw::error_or real_main(int argc, char** argv){ * Set initial conditions in the simulations */ { - auto eov = set_initial_conditions(lattice); + auto eov = set_initial_conditions(lattice, meta); if(eov.is_error()){ return eov; } } for(saw::data ts{0u}; ts < saw::data{1024u*32u}; ++ts){ - auto eov = lbm_step(lattice,ts,sycl_q); + auto eov = lbm_step(lattice,meta,ts,sycl_q); if(eov.is_error()){ return eov; } -- cgit v1.2.3