diff options
| -rw-r--r-- | default.nix | 31 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_gpu/sim.cpp | 7 | ||||
| -rw-r--r-- | 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<void> setup_initial_conditions(saw::data<sch::ChunkStruct<T,Desc>> } template<typename T, typename Desc> -saw::error_or<void> step(saw::data<sch::ChunkStruct<T,Desc>,encode::Sycl<saw::encode::Native>>& fields, saw::data<sch::UInt64> t_i, device& dev){ +saw::error_or<void> step(saw::data<sch::Ptr<sch::ChunkStruct<T,Desc>>,encode::Sycl<saw::encode::Native>>& fields, saw::data<sch::UInt64> t_i, device& dev){ auto& q = dev.get_handle(); auto& info_f = fields.template get<"info">(); @@ -99,7 +99,7 @@ saw::error_or<void> step(saw::data<sch::ChunkStruct<T,Desc>,encode::Sycl<saw::en index.at({{i}}).set(idx[i]); } - collision.apply(fields,index,t_i); + // collision.apply(fields,index,t_i); }); }).wait(); @@ -146,6 +146,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){ saw::data<sch::ChunkStruct<T,Desc>, encode::Sycl<saw::encode::Native>> 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<void> lbm_main(int argc, char** argv){ sycl_q.wait(); for(saw::data<sch::UInt64> i{0u}; i < saw::data<sch::UInt64>{32ul}; ++i){ - auto eov = step<T,Desc>(lbm_sycl_data,i,dev); + auto eov = step<T,Desc>(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<Sch,Encode>* 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<typename Sch, uint64_t... Dims, typename Encode> +class data<schema::Ptr<schema::FixedArray<Sch,Dims...>>, kel::lbm::encode::Sycl<Encode>> final { +public: + using Schema = schema::Ptr<schema::FixedArray<Sch,Dims...>>; +private: + data<Sch,Encode>* values_; + + SAW_DEFAULT_COPY(data); + SAW_DEFAULT_MOVE(data); +public: + data(): + values_{nullptr} + {} + + data(data<Sch,Encode>* values__): + values_{values__} + {} + + static constexpr data<schema::FixedArray<schema::UInt64, sizeof...(Dims)>> get_dims() { + return saw::data<schema::FixedArray<schema::UInt64, sizeof...(Dims)>>{{Dims...}}; + } + + constexpr data<Sch,Encode>& at(data<schema::FixedArray<schema::UInt64,sizeof...(Dims)>>& index){ + return values_[kel::lbm::flatten_index<schema::UInt64,sizeof...(Dims)>::apply(index,get_dims()).get()]; + } + + constexpr data<Sch,Encode>& at(data<schema::FixedArray<schema::UInt64,sizeof...(Dims)>>& index) const{ + return values_[kel::lbm::flatten_index<schema::UInt64,sizeof...(Dims)>::apply(index,get_dims()).get()]; + } + + constexpr data<Sch,Encode>* flat_data() { + return values_; + } +}; + template<typename Sch, uint64_t Ghost, uint64_t... Sides, typename Encode> class data<kel::lbm::sch::Chunk<Sch,Ghost,Sides...>,kel::lbm::encode::Sycl<Encode>> final { public: @@ -119,13 +155,74 @@ public: } }; +template<typename Sch, uint64_t Ghost, uint64_t... Sides, typename Encode> +class data<schema::Ptr<kel::lbm::sch::Chunk<Sch,Ghost,Sides...>>,kel::lbm::encode::Sycl<Encode>> final { +public: + using Schema = schema::Ptr<kel::lbm::sch::Chunk<Sch,Ghost,Sides...>>; +private: + using InnerSchema = typename kel::lbm::sch::Chunk<Sch,Ghost,Sides...>::InnerSchema; + using ValueSchema = typename InnerSchema::ValueType; + + data<schema::Ptr<InnerSchema>,kel::lbm::encode::Sycl<Encode>> values_; + + SAW_DEFAULT_MOVE(data); + SAW_DEFAULT_COPY(data); +public: + data(): + values_{nullptr} + {} + + data(data<Sch,Encode>* values__): + values_{values__} + {} + + data<ValueSchema, kel::lbm::encode::Sycl<Encode>>& ghost_at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index){ + return values_.at(index); + } + + data<ValueSchema, kel::lbm::encode::Sycl<Encode>>& ghost_at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index) const { + return values_.at(index); + } + + static constexpr auto get_ghost_dims() { + return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::get_dims(); + } + + data<ValueSchema, Encode>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index){ + std::decay_t<decltype(index)> ind; + for(uint64_t i = 0u; i < sizeof...(Sides); ++i){ + ind.at({i}) = index.at({i}) + Ghost; + } + return values_.at(ind); + } + + data<ValueSchema, Encode>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index) const { + std::decay_t<decltype(index)> 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<schema::FixedArray<schema::UInt64, sizeof...(Sides)>,kel::lbm::encode::Sycl<Encode>>{{Sides...}}; + } + + auto flat_data(){ + return values_.flat_data(); + } + + static constexpr auto flat_size() { + return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::flat_size(); + } +}; + template<typename... Members, typename Encode> class data<schema::Struct<Members...>, kel::lbm::encode::Sycl<Encode> > final { public: using StorageT = std::tuple<data<typename Members::ValueType,kel::lbm::encode::Sycl<Encode>>...>; using Schema = schema::Struct<Members...>; 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<parameter_key_pack_index<K, Members::KeyLiteral...>::value>(members_); } }; + +template<typename... Sch, saw::string_literal... Keys, typename Encode> +class data<schema::Ptr<schema::Struct<schema::Member<Sch,Keys>...>>, kel::lbm::encode::Sycl<Encode> > final { +public: + using StorageT = std::tuple<data<schema::Ptr<Sch>,kel::lbm::encode::Sycl<Encode>>...>; + using Schema = schema::Struct<schema::Member<schema::Ptr<Sch>,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<saw::string_literal K> + auto& get(){ + return std::get<parameter_key_pack_index<K, Keys...>::value>(members_); + } + + template<saw::string_literal K> + const auto& get() const { + return std::get<parameter_key_pack_index<K, Keys...>::value>(members_); + } +}; } namespace kel { @@ -217,7 +340,43 @@ struct sycl_copy_helper<sch::Struct<Members...>, Encode> final { return copy_to_host_member<0u>(sycl_data, host_data, q); } }; + +template<typename Schema, typename Encode> +struct make_chunk_struct_view_helper; + +template<typename... Sch, saw::string_literal... Keys, typename Encode> +struct make_chunk_struct_view_helper<sch::Struct<sch::Member<Sch,Keys>...>, encode::Sycl<Encode>> { +private: + template<uint64_t i> + static void apply_i( + saw::data<sch::Struct<sch::Member<Sch,Keys>...>, encode::Sycl<Encode>>& dat, + saw::data<sch::Ptr<sch::Struct<sch::Member<Sch,Keys>...>>, encode::Sycl<Encode>>& dat_view) + { + if constexpr (i < sizeof...(Sch)){ + using M = typename saw::parameter_pack_type<i,sch::Member<Sch,Keys>...>::type; + + dat_view.template get<M::KeyLiteral>() = {dat.template get<M::KeyLiteral>().flat_data()}; + + apply_i<i+1u>(dat,dat_view); + } + } +public: + static auto apply(saw::data<sch::Struct<sch::Member<Sch,Keys>...>, encode::Sycl<Encode>>& dat){ + saw::data<sch::Ptr<sch::Struct<sch::Member<Sch,Keys>...>>, encode::Sycl<Encode>> str_view; + + apply_i<0u>(dat,str_view); + + return str_view; + } +}; } + +// Ptr => Ptr<Chunk<T,Ghost,Dims...>>> => Ptr<FixedArray<T,Dims+Ghost...> +template<typename Schema,typename Encode> +auto make_chunk_struct_view(saw::data<Schema,Encode>& dat){ + return impl::make_chunk_struct_view_helper<Schema,Encode>::apply(dat); +} + class device final { private: sycl::queue q_; |