From 6cf8735733efa535fd2c18fdbaac1deaff345041 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Tue, 21 Oct 2025 14:40:47 +0200 Subject: Going forward with testing sycl setup --- default.nix | 12 +-- examples/cavity_2d_gpu/.nix/derivation.nix | 16 ++-- examples/cavity_2d_gpu/cavity_2d_gpu.cpp | 118 ++++++++++++----------------- 3 files changed, 62 insertions(+), 84 deletions(-) diff --git a/default.nix b/default.nix index b6f8f5b..bd3182a 100644 --- a/default.nix +++ b/default.nix @@ -12,10 +12,10 @@ let inherit clang-tools; }).forstio; - adaptive-cpp = pkgs.callPackage ./.nix/adaptive-cpp.nix { - inherit stdenv; - llvmPackages = pkgs.llvmPackages_19; - lld = pkgs.lld_19; + adaptive-cpp-dev = pkgs.callPackage .nix/adaptive-cpp.nix { + inherit stdenv; + llvmPackages = pkgs.llvmPackages_19; + lld = pkgs.lld_19; }; pname = "kel-lbm"; @@ -29,7 +29,7 @@ in rec { examples = { cavity_2d_gpu = pkgs.callPackage ./examples/cavity_2d_gpu/.nix/derivation.nix { - inherit pname version stdenv forstio adaptive-cpp; + inherit pname version stdenv forstio adaptive-cpp-dev; kel-lbm = lbm; }; @@ -45,6 +45,6 @@ in rec { }; debug = { - inherit adaptive-cpp; + # inherit foo; }; } diff --git a/examples/cavity_2d_gpu/.nix/derivation.nix b/examples/cavity_2d_gpu/.nix/derivation.nix index 73b81b7..2748a18 100644 --- a/examples/cavity_2d_gpu/.nix/derivation.nix +++ b/examples/cavity_2d_gpu/.nix/derivation.nix @@ -3,9 +3,10 @@ , scons , clang-tools , forstio -, adaptive-cpp +, keldu , pname , version +, adaptive-cpp-dev , kel-lbm }: @@ -20,14 +21,13 @@ stdenv.mkDerivation { ]; buildInputs = [ - forstio.core - forstio.async - forstio.codec - forstio.codec-unit - forstio.remote - forstio.remote-sycl + forstio.core + forstio.async + forstio.codec + forstio.codec-unit + forstio.remote forstio.codec-json - adaptive-cpp + adaptive-cpp-dev kel-lbm ]; diff --git a/examples/cavity_2d_gpu/cavity_2d_gpu.cpp b/examples/cavity_2d_gpu/cavity_2d_gpu.cpp index c4836d0..ed7042f 100644 --- a/examples/cavity_2d_gpu/cavity_2d_gpu.cpp +++ b/examples/cavity_2d_gpu/cavity_2d_gpu.cpp @@ -1,7 +1,6 @@ #include #include - #include #include @@ -17,7 +16,7 @@ struct Sycl { template class data> { private: - cl::sycl::buffer> data_; + acpp::sycl::buffer> data_; data size_; public: data(const data& data__): @@ -37,13 +36,13 @@ public: return size_; } - template - auto access(cl::sycl::handler& h){ + template + auto access(acpp::sycl::handler& h){ return data_.template get_access(h); } - template - auto access(cl::sycl::handler& h) const { + template + auto access(acpp::sycl::handler& h) const { return data_.template get_access(h); } }; @@ -53,7 +52,7 @@ class data, encode::Sycl> { public: using Schema = schema::Array; private: - cl::sycl::buffer> data_; + acpp::sycl::buffer> data_; data size_; public: data(const data& host_data__): @@ -73,13 +72,13 @@ public: return size_; } - template - auto access(cl::sycl::handler& h){ + template + auto access(acpp::sycl::handler& h){ return data_.template get_access(h); } - template - auto access(cl::sycl::handler& h) const { + template + auto access(acpp::sycl::handler& h) const { return data_.template get_access(h); } }; @@ -184,7 +183,6 @@ public: } }; -constexpr size_t dim_size = 2; constexpr size_t dim_x = 128; constexpr size_t dim_y = 128; @@ -227,7 +225,6 @@ void set_geometry(saw::data& lattice){ */ iterate_over([&](const saw::data>& index){ auto& info = info_field.at(index); - info({0u}).set(3u); }, {{0u,1u}}, {{meta.at({0u}), 2u}}, {{2u,0u}}); @@ -265,26 +262,29 @@ struct sycl_component_context { acpp::sycl::handler* handler = nullptr; }; +template void lbm_step( - saw::data>, saw::encode::Sycl>& info, - saw::data>, saw::encode::Sycl>& dfs, - saw::data>, saw::encode::Sycl>& dfs_old, - bool even_step, + acpp::sycl::buffer>>& info, + acpp::sycl::buffer>>& dfs, + acpp::sycl::buffer>>& dfs_old, + bool is_even, uint64_t time_step, acpp::sycl::queue& sycl_q ){ using namespace kel::lbm; using namespace acpp; - using dfi = df_info; + using dfi = df_info; - //component coll{0.59}; - //component bb; - component bb_lid; + //component coll{0.59}; + //component bb; + component bb_lid; bb_lid.lid_vel = {0.1, 0.0}; + /* const size_t Nx = latt.template get_dim_size<0>().get(); const size_t Ny = latt.template get_dim_size<1>().get(); + */ // Submit collision kernel sycl_q.submit([&](sycl::handler& cgh) { @@ -294,67 +294,45 @@ void lbm_step( auto dfs_old_acc = dfs_old.template get_access(cgh); cgh.parallel_for( - sycl::range<2>{Nx, Ny}, [=](sycl::id<2> idx) { + sycl::range<2>{dim_x, dim_y}, [=](sycl::id<2> idx) { size_t i = idx[0]; size_t j = idx[1]; + size_t acc_id = i * dim_x + j; + // Read cell info - auto& info = info_acc[idx]; + auto& info = info_acc[acc_id]; switch (info({0u}).get()) { case 1u: { // bb.apply(latt_acc, {i, j}, time_step); - // bool is_even = ((time_step % 2) == 0); - auto& dfs_old = (is_even) ? dfs_old_acc[idx] : dfs_acc[idx]; - // auto& dfs = (not is_even) ? cell.template get<"dfs_old">() : cell.template get<"dfs">(); + auto& dfs_old = (is_even) ? dfs_old_acc[acc_id] : dfs_acc[acc_id]; + // auto& dfs = (not is_even) ? cell.template get<"dfs_old">() : cell.template get<"dfs">(); + auto df_cpy = dfs_old.copy(); - saw::data rho; - saw::data> vel; - compute_rho_u(dfs_old,rho,vel); - auto eq = equilibrium(rho,vel); + for(uint64_t i = 1u; i < Desc::Q; ++i){ + dfs_old({i}) = df_cpy({dfi::opposite_index.at(i)}); + } - for(uint64_t i = 0u; i < Descriptor::Q; ++i){ - dfs_old({i}) = dfs_old({i}) + frequency_ * (eq.at(i) - dfs_old({i})); - // dfs_old({i}).set(dfs_old({i}).get() + (1.0 / relaxation_) * (eq.at(i).get() - dfs_old({i}).get())); - } break; } case 2u: { // coll.apply(latt_acc, {i, j}, time_step); - // bool is_even = ((time_step % 2) == 0); - auto& dfs_old = (is_even) ? dfs_old_acc[idx] : dfs_acc[idx]; - // auto& dfs = (not is_even) ? cell.template get<"dfs_old">() : cell.template get<"dfs">(); + auto& dfs_old = (is_even) ? dfs_old_acc[acc_id] : dfs_acc[acc_id]; - saw::data rho; - saw::data> vel; - compute_rho_u(dfs_old,rho,vel); - auto eq = equilibrium(rho,vel); + saw::data rho; + saw::data> vel; + compute_rho_u(dfs_old,rho,vel); + auto eq = equilibrium(rho,vel); - using dfi = df_info; - - auto& force = cell.template get<"force">(); - - for(uint64_t i = 0u; i < Descriptor::Q; ++i){ - // saw::data ci_min_u{0}; - saw::data ci_dot_u{0}; - for(uint64_t d = 0u; d < Descriptor::D; ++d){ - ci_dot_u = ci_dot_u + vel.at({d}) * saw::data{static_cast::type>(dfi::directions[i][d])}; + for(uint64_t i = 0u; i < Desc::Q; ++i){ + dfs_old({i}) = dfs_old({i}) + frequency_ * (eq.at(i) - dfs_old({i})); + // dfs_old({i}).set(dfs_old({i}).get() + (1.0 / relaxation_) * (eq.at(i).get() - dfs_old({i}).get())); + } + break; } - saw::data F_i{0};// = saw::data{dfi::weights[i]} * (ci_dot_F * dfi::inv_cs2 + ci_dot_u * ci_dot_F * (dfi::inv_cs2 * dfi::inv_cs2)); - - for(uint64_t d = 0u; d < Descriptor::D; ++d){ - F_i = F_i + - saw::data{dfi::weights[i]} * ((saw::data{static_cast::type>(dfi::directions[i][d])} - - vel.at({d}) ) * dfi::inv_cs2 + ci_dot_u * saw::data{static_cast::type>(dfi::directions[i][d])} * dfi::inv_cs2 * dfi::inv_cs2 ) * force({d}); - } - - - dfs_old({i}) = dfs_old({i}) + frequency_ * (eq.at(i) - dfs_old({i}) ) + F_i * (saw::data{1} - saw::data{0.5f} * frequency_); - } - break; - } default: // Do nothing break; @@ -375,7 +353,7 @@ void lbm_step( auto& info_new = cell.template get<"info">(); if (info_new({0u}).get() > 0u && info_new({0u}).get() != 3u) { - for (uint64_t k = 0u; k < sch::D2Q9::Q; ++k) { + for (uint64_t k = 0u; k < Desc::Q; ++k) { auto dir = dfi::directions[dfi::opposite_index[k]]; auto& cell_dir_old = latt_acc({i + dir[0], j + dir[1]}); @@ -401,8 +379,8 @@ void lbm_step( int main(){ using namespace kel::lbm; - saw::data> dim{{dim_x,info_field.at dim_y}}; - const dim_size = dim_x * dim_y; + saw::data> dim{{dim_x, dim_y}}; + constexpr size_t dim_size = dim_x * dim_y; saw::data> lattice{dim}; @@ -435,9 +413,9 @@ int main(){ auto& dfs_old_field = lattice.template get<"dfs_old">(); acpp::sycl::queue sycl_q{acpp::sycl::default_selector_v, acpp::sycl::property::queue::in_order{}}; - saw::data>, saw::encode::Sycl> sycl_info_field{info_field}; - saw::data>, saw::encode::Sycl> sycl_dfs_field{dfs_field}; - saw::data>, saw::encode::Sycl> sycl_dfs_old_field{dfs_old_field}; + acpp::sycl::buffer>> info_sycl{info_field.internal_data(), info_field.internal_size()}; + acpp::sycl::buffer>> dfs_sycl{dfs_field.internal_data(), dfs_field.internal_size()}; + acpp::sycl::buffer>> dfs_old_sycl{dfs_old_field.internal_data(), dfs_old_field.internal_size()}; /** * Timeloop @@ -458,7 +436,7 @@ int main(){ write_vtk_file(vtk_f_name, macros); } - lbm_step(sycl_info_field, sycl_dfs_field, sycl_dfs_old_field, (i%2u == 0u), i, sycl_q); + lbm_step(info_sycl, dfs_sycl, dfs_old_sycl, (i%2u == 0u), i, sycl_q); } return 0; } -- cgit v1.2.3