From 624f5d8181a449f36950503859271e1a052e8a65 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Tue, 27 Jan 2026 09:17:33 +0100 Subject: Double free, but constexpr things compile now --- examples/poiseulle_particles_2d_gpu/sim.cpp | 27 ++++++++++------ lib/core/c++/collision.hpp | 34 +++----------------- lib/core/c++/macroscopic.hpp | 25 +++++++++++++++ lib/sycl/c++/data.hpp | 48 ++++++++++++++--------------- 4 files changed, 71 insertions(+), 63 deletions(-) diff --git a/examples/poiseulle_particles_2d_gpu/sim.cpp b/examples/poiseulle_particles_2d_gpu/sim.cpp index bb1fca5..4a28541 100644 --- a/examples/poiseulle_particles_2d_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_gpu/sim.cpp @@ -8,7 +8,7 @@ namespace kel { namespace lbm { constexpr uint64_t dim_x = 256u; -constexpr uint64_t dim_y = 16u; +constexpr uint64_t dim_y = 32u; namespace sch { using namespace saw::schema; @@ -20,8 +20,9 @@ using DfChunk = Chunk, 1u, dim_x, dim_y>; template using ChunkStruct = Struct< - Member - //,Member, "dfs"> + Member, + Member, "dfs">, + Member, "dfs_old"> >; } @@ -66,7 +67,6 @@ saw::error_or setup_initial_conditions(saw::data> {{dim_x, dim_y}}, {{0u,1u}} ); -/* // auto& df_f = fields.template get<"dfs">(); saw::data rho{1}; @@ -82,7 +82,6 @@ saw::error_or setup_initial_conditions(saw::data> {},// 0-index df_f.get_dims() ); -*/ return saw::make_void(); } @@ -90,6 +89,19 @@ saw::error_or setup_initial_conditions(saw::data> template saw::error_or step(saw::data,encode::Sycl>& fields, saw::data t_i, device& dev){ auto& q = dev.get_handle(); + auto& info_f = fields.template get<"info">(); + + q.submit([&](acpp::sycl::handler& h){ + component> collision{0.6}; + h.parallel_for(acpp::sycl::range{dim_x,dim_y}, [=](acpp::sycl::id idx){ + saw::data> index; + for(uint64_t i = 0u; i < Desc::D; ++i){ + index.at({{i}}).set(idx[i]); + } + + collision.apply(fields,index,t_i); + }); + }).wait(); return saw::make_void(); } @@ -132,11 +144,8 @@ saw::error_or lbm_main(int argc, char** argv){ } { - std::cout<<"Hey"<, encode::Sycl> lbm_sycl_data{sycl_q}; sycl_q.wait(); - std::cout<<"Hey2"< lbm_main(int argc, char** argv){ } } sycl_q.wait(); - std::cout<<"Hey3"< i{0u}; i < saw::data{32ul}; ++i){ auto eov = step(lbm_sycl_data,i,dev); @@ -155,7 +163,6 @@ saw::error_or lbm_main(int argc, char** argv){ } sycl_q.wait(); - std::cout<<"Hey4"<::apply( [&](auto& index){ diff --git a/lib/core/c++/collision.hpp b/lib/core/c++/collision.hpp index 2f20601..85c3af9 100644 --- a/lib/core/c++/collision.hpp +++ b/lib/core/c++/collision.hpp @@ -45,43 +45,19 @@ public: * Raw setup */ template - void apply(saw::data& field, saw::data> index, saw::data time_step){ + void apply(const saw::data& field, saw::data> index, saw::data time_step) const { bool is_even = ((time_step.get() % 2) == 0); - auto& cell = field.at(index); - 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">(); + // auto& dfs_f = (is_even) ? field.template get<"dfs">() : field.template get<"dfs_old">(); + auto& dfs_old_f = (is_even) ? field.template get<"dfs_old">() : field.template get<"dfs">(); saw::data rho; saw::data> vel; - compute_rho_u(dfs_old,rho,vel); - auto eq = equilibrium(rho,vel); - - 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())); - } - } - - template - void apply( - saw::data* field, - const saw::data>& meta, - const saw::data>& index, - saw::data time_step - ){ - bool is_even = ((time_step.get() % 2) == 0); - auto& cell = field[0u]; - - auto& dfs_old = (is_even) ? cell.template get<"dfs_old">() : cell.template get<"dfs">(); - - saw::data rho; - saw::data> vel; - compute_rho_u(dfs_old,rho,vel); + compute_rho_u(dfs_old_f.at(index),rho,vel); auto eq = equilibrium(rho,vel); for(uint64_t i = 0u; i < Descriptor::Q; ++i){ - dfs_old({i}) = dfs_old({i}) + frequency_ * (eq.at(i) - dfs_old({i})); + dfs_old_f.at(index).at({i}) = dfs_old_f.at(index).at({i}) + frequency_ * (eq.at(i) - dfs_old_f.at(index).at({i})); } } }; diff --git a/lib/core/c++/macroscopic.hpp b/lib/core/c++/macroscopic.hpp index 51368e9..1532873 100644 --- a/lib/core/c++/macroscopic.hpp +++ b/lib/core/c++/macroscopic.hpp @@ -4,6 +4,31 @@ namespace kel { namespace lbm { +template +void compute_rho_u ( + const saw::data>& dfs, + saw::data& rho, + saw::data>& vel + ) +{ + using dfi = df_info; + + rho = 0; + for(size_t j = 0; j < Desc::D; ++j){ + vel.at({{j}}).set(0); + } + + for(size_t j = 0; j < Desc::Q; ++j){ + rho = rho + dfs.at({j}); + for(size_t i = 0; i < Desc::D; ++i){ + vel.at({{i}}) = vel.at({{i}}) + dfi::directions[j][i] * dfs.at({j}).get(); + } + } + + for(size_t i = 0; i < Desc::D; ++i){ + vel.at({i}) = vel.at({i}) / rho; + } +} /** * Calculate the macroscopic variables rho and u in Lattice Units. */ diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp index cffeb38..d5adb95 100644 --- a/lib/sycl/c++/data.hpp +++ b/lib/sycl/c++/data.hpp @@ -25,15 +25,14 @@ public: using Schema = schema::FixedArray; private: acpp::sycl::queue* q_; - data* values_; + data* values_; - SAW_FORBID_COPY(data); + //SAW_FORBID_COPY(data); public: data(acpp::sycl::queue& q__): q_{&q__}, values_{nullptr} { - std::cout<<"Hey: "<::value<>(ct_multiply::value,*q_); SAW_ASSERT(values_ and q_); } @@ -49,14 +48,18 @@ public: } static constexpr data> get_dims() { - return {std::array{Dims...}}; + return saw::data>{{Dims...}}; } - data& at(data>& index){ + constexpr data& at(data>& index){ return values_[kel::lbm::flatten_index::apply(index,get_dims()).get()]; } - constexpr data* flat_data() { + constexpr data& at(data>& index) const{ + return values_[kel::lbm::flatten_index::apply(index,get_dims()).get()]; + } + + constexpr data* flat_data() { return values_; } }; @@ -79,7 +82,7 @@ public: return values_.at(index); } - const data>& ghost_at(const data>& index) const { + data>& ghost_at(const data>& index) const { return values_.at(index); } @@ -87,7 +90,7 @@ public: return data>::get_dims(); } - data>& at(const data>& index){ + data& at(const data>& index){ std::decay_t ind; for(uint64_t i = 0u; i < sizeof...(Sides); ++i){ ind.at({i}) = index.at({i}) + Ghost; @@ -95,7 +98,7 @@ public: return values_.at(ind); } - const data>& at(const data>& index) const { + data& at(const data>& index) const { std::decay_t ind; for(uint64_t i = 0u; i < sizeof...(Sides); ++i){ ind.at({i}) = index.at({i}) + Ghost; @@ -116,12 +119,11 @@ public: } }; -template -class data, Keys>...>, kel::lbm::encode::Sycl > final { +template +class data, kel::lbm::encode::Sycl > final { public: - static constexpr data> meta = {{Meta...}}; - using StorageT = std::tuple,kel::lbm::encode::Sycl>...>; - using Schema = schema::Struct, Keys>...>; + using StorageT = std::tuple>...>; + using Schema = schema::Struct; private: /** @@ -133,24 +135,24 @@ private: template constexpr data(acpp::sycl::queue& q, std::index_sequence): members_{(static_cast(Is), q)...} - { - - } + {} public: data(acpp::sycl::queue& q__): - data{q__, std::make_index_sequence{}} -// members_(q__)//data,kel::lbm::encode::Sycl>(q__)...) + data{q__, std::make_index_sequence{}} { q__.wait(); } template auto& get(){ - return std::get::value>(members_); + return std::get::value>(members_); + } + + template + const auto& get() const { + return std::get::value>(members_); } }; - - } namespace kel { @@ -175,8 +177,6 @@ struct sycl_copy_helper, Encode> final { static_assert(sizeof(std::decay_t) == sizeof(std::decay_t), "Unequal size"); - std::cout<