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 --- lib/core/c++/collision.hpp | 34 +++++-------------------------- lib/core/c++/macroscopic.hpp | 25 +++++++++++++++++++++++ lib/sycl/c++/data.hpp | 48 ++++++++++++++++++++++---------------------- 3 files changed, 54 insertions(+), 53 deletions(-) (limited to 'lib') 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<