From 1d81ab6f66b5efd6b4d291762935866930731152 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Wed, 28 Jan 2026 11:24:30 +0100 Subject: Reworked fetching due to unstable fetchTarball --- default.nix | 31 ++++-- examples/poiseulle_particles_2d_gpu/sim.cpp | 7 +- lib/sycl/c++/data.hpp | 163 +++++++++++++++++++++++++++- 3 files changed, 188 insertions(+), 13 deletions(-) diff --git a/default.nix b/default.nix index d46eed5..9068e76 100644 --- a/default.nix +++ b/default.nix @@ -36,14 +36,29 @@ let llvmPackages = pkgs.llvmPackages_19; lld = pkgs.lld_19; }; - - forstio = (import ((builtins.fetchTarball { - url = "https://git.keldu.de/forstio-forstio/snapshot/master.tar.gz"; - sha256 = "sha256:0078544k33h6ihhzz1150xs9zcywdmb6y8p16038v488kblvh3nc"; - }) + "/default.nix"){ - inherit stdenv; - inherit clang-tools; - inherit adaptive-cpp; + forstio = let + forstioSrc = stdenv.mkDerivation { + name = "forstio-src"; + + src = builtins.fetchurl { + url = "https://git.keldu.de/forstio-forstio/snapshot/master.tar.gz"; + sha256 = "e91c18fef798dd7b3afbd1615c2e320b90d74aa2d7ef726801a76e3f7f77ae81"; + }; + + phases = [ "unpackPhase" "installPhase" ]; + + unpackPhase = '' + mkdir source + tar -xzf "$src" -C source --strip-components=1 + ''; + + installPhase = '' + cp -r source $out + ''; + }; + in + (import "${forstioSrc}/default.nix" { + inherit stdenv clang-tools adaptive-cpp; }).forstio; pname = "kel-lbm"; diff --git a/examples/poiseulle_particles_2d_gpu/sim.cpp b/examples/poiseulle_particles_2d_gpu/sim.cpp index 4a28541..8c7f1ea 100644 --- a/examples/poiseulle_particles_2d_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_gpu/sim.cpp @@ -87,7 +87,7 @@ saw::error_or setup_initial_conditions(saw::data> } template -saw::error_or step(saw::data,encode::Sycl>& fields, saw::data t_i, device& dev){ +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">(); @@ -99,7 +99,7 @@ saw::error_or step(saw::data,encode::Sycl lbm_main(int argc, char** argv){ saw::data, encode::Sycl> lbm_sycl_data{sycl_q}; sycl_q.wait(); + auto lsd_view = make_chunk_struct_view(lbm_sycl_data); { auto eov = dev.copy_to_device(*lbm_data_ptr,lbm_sycl_data); if(eov.is_error()){ @@ -155,7 +156,7 @@ saw::error_or lbm_main(int argc, char** argv){ sycl_q.wait(); for(saw::data i{0u}; i < saw::data{32ul}; ++i){ - auto eov = step(lbm_sycl_data,i,dev); + auto eov = step(lsd_view,i,dev); if(eov.is_error()){ return eov; } diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp index d5adb95..9b15678 100644 --- a/lib/sycl/c++/data.hpp +++ b/lib/sycl/c++/data.hpp @@ -27,7 +27,8 @@ private: acpp::sycl::queue* q_; data* values_; - //SAW_FORBID_COPY(data); + SAW_FORBID_COPY(data); + SAW_DEFAULT_MOVE(data); public: data(acpp::sycl::queue& q__): q_{&q__}, @@ -64,6 +65,41 @@ public: } }; +template +class data>, kel::lbm::encode::Sycl> final { +public: + using Schema = schema::Ptr>; +private: + data* values_; + + SAW_DEFAULT_COPY(data); + SAW_DEFAULT_MOVE(data); +public: + data(): + values_{nullptr} + {} + + data(data* values__): + values_{values__} + {} + + static constexpr data> get_dims() { + return saw::data>{{Dims...}}; + } + + constexpr data& at(data>& index){ + return values_[kel::lbm::flatten_index::apply(index,get_dims()).get()]; + } + + constexpr data& at(data>& index) const{ + return values_[kel::lbm::flatten_index::apply(index,get_dims()).get()]; + } + + constexpr data* flat_data() { + return values_; + } +}; + template class data,kel::lbm::encode::Sycl> final { public: @@ -119,13 +155,74 @@ public: } }; +template +class data>,kel::lbm::encode::Sycl> final { +public: + using Schema = schema::Ptr>; +private: + using InnerSchema = typename kel::lbm::sch::Chunk::InnerSchema; + using ValueSchema = typename InnerSchema::ValueType; + + data,kel::lbm::encode::Sycl> values_; + + SAW_DEFAULT_MOVE(data); + SAW_DEFAULT_COPY(data); +public: + data(): + values_{nullptr} + {} + + data(data* values__): + values_{values__} + {} + + data>& ghost_at(const data>& index){ + return values_.at(index); + } + + data>& ghost_at(const data>& index) const { + return values_.at(index); + } + + static constexpr auto get_ghost_dims() { + return data>::get_dims(); + } + + 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; + } + return values_.at(ind); + } + + 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; + } + return values_.at(ind); + } + + static constexpr auto get_dims(){ + return data,kel::lbm::encode::Sycl>{{Sides...}}; + } + + auto flat_data(){ + return values_.flat_data(); + } + + static constexpr auto flat_size() { + return data>::flat_size(); + } +}; + template class data, kel::lbm::encode::Sycl > final { public: using StorageT = std::tuple>...>; using Schema = schema::Struct; private: - /** * @todo Check by static assert that the members all have the same dimensions. Alternatively * Do it here by specializing. @@ -153,6 +250,32 @@ public: return std::get::value>(members_); } }; + +template +class data...>>, kel::lbm::encode::Sycl > final { +public: + using StorageT = std::tuple,kel::lbm::encode::Sycl>...>; + using Schema = schema::Struct,Keys>...>; +private: + /** + * @todo Check by static assert that the members all have the same dimensions. Alternatively + * Do it here by specializing. + */ + StorageT members_; + +public: + data() = default; + + template + auto& get(){ + return std::get::value>(members_); + } + + template + const auto& get() const { + return std::get::value>(members_); + } +}; } namespace kel { @@ -217,7 +340,43 @@ struct sycl_copy_helper, Encode> final { return copy_to_host_member<0u>(sycl_data, host_data, q); } }; + +template +struct make_chunk_struct_view_helper; + +template +struct make_chunk_struct_view_helper...>, encode::Sycl> { +private: + template + static void apply_i( + saw::data...>, encode::Sycl>& dat, + saw::data...>>, encode::Sycl>& dat_view) + { + if constexpr (i < sizeof...(Sch)){ + using M = typename saw::parameter_pack_type...>::type; + + dat_view.template get() = {dat.template get().flat_data()}; + + apply_i(dat,dat_view); + } + } +public: + static auto apply(saw::data...>, encode::Sycl>& dat){ + saw::data...>>, encode::Sycl> str_view; + + apply_i<0u>(dat,str_view); + + return str_view; + } +}; } + +// Ptr => Ptr>> => Ptr +template +auto make_chunk_struct_view(saw::data& dat){ + return impl::make_chunk_struct_view_helper::apply(dat); +} + class device final { private: sycl::queue q_; -- cgit v1.2.3