diff options
| -rw-r--r-- | default.nix | 2 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_gpu/sim.cpp | 63 | ||||
| -rw-r--r-- | lib/core/c++/chunk.hpp | 14 | ||||
| -rw-r--r-- | lib/core/c++/flatten.hpp | 2 | ||||
| -rw-r--r-- | lib/sycl/c++/data.hpp | 205 | ||||
| -rw-r--r-- | lib/sycl/tests/data.cpp | 17 |
6 files changed, 252 insertions, 51 deletions
diff --git a/default.nix b/default.nix index a8ec430..fd36323 100644 --- a/default.nix +++ b/default.nix @@ -39,7 +39,7 @@ let forstio = (import ((builtins.fetchTarball { url = "https://git.keldu.de/forstio-forstio/snapshot/master.tar.gz"; - sha256 = "sha256:09xxs9dlnd52k1082lb5p034ik4jfg9zd5zc1x5da2clqhdb6s4m"; + sha256 = "sha256:0078544k33h6ihhzz1150xs9zcywdmb6y8p16038v488kblvh3nc"; }) + "/default.nix"){ inherit stdenv; inherit clang-tools; diff --git a/examples/poiseulle_particles_2d_gpu/sim.cpp b/examples/poiseulle_particles_2d_gpu/sim.cpp index be6b4f0..9646b0a 100644 --- a/examples/poiseulle_particles_2d_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_gpu/sim.cpp @@ -7,8 +7,8 @@ namespace kel { namespace lbm { -constexpr uint64_t dim_x = 16u; -constexpr uint64_t dim_y = 4u; +constexpr uint64_t dim_x = 256u; +constexpr uint64_t dim_y = 16u; namespace sch { using namespace saw::schema; @@ -31,29 +31,29 @@ using DfChunk = Chunk<FixedArray<T,Desc::Q>, 1u, dim_x, dim_y>; template<typename T, typename Desc> using ChunkStruct = Struct< Member<InfoChunk, "info"> + //,Member<DfChunk<T,Desc>, "dfs"> >; - } template<typename T, typename Desc> saw::error_or<void> setup_initial_conditions(saw::data<sch::ChunkStruct<T,Desc>>& fields){ - auto& info_f = fields.get<"info">(); + auto& info_f = fields.template get<"info">(); // Set everything as walls iterator<Desc::D>::apply( [&](auto& index){ info_f.at(index).set(2u); }, - {{0u,0u}}, - {{dim_x, dim_y}}, - {{0u,0u}} + {}, + info_f.get_dims(), + {} ); // Fluid iterator<Desc::D>::apply( [&](auto& index){ info_f.at(index).set(1u); }, - {{0u,0u}}, - {{dim_x, dim_y}}, + {}, + info_f.get_dims(), {{1u,1u}} ); @@ -76,12 +76,30 @@ saw::error_or<void> setup_initial_conditions(saw::data<sch::ChunkStruct<T,Desc>> {{dim_x, dim_y}}, {{0u,1u}} ); +/* + // + auto& df_f = fields.template get<"dfs">(); + saw::data<T> rho{1}; + saw::data<sch::FixedArray<T,Desc::D>> vel{}; + auto eq = equilibrium<T,Desc>(rho,vel); + + iterator<Desc::D>::apply( + [&](auto& index){ + auto& df = df_f.at(index); + + df = eq; + }, + {},// 0-index + df_f.get_dims() + ); +*/ return saw::make_void(); } template<typename T, typename Desc> -saw::error_or<void> step(){ +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){ + auto& q = dev.get_handle(); return saw::make_void(); } @@ -111,7 +129,9 @@ saw::error_or<void> lbm_main(int argc, char** argv){ // saw::data<sch::FixedArray<sch::UInt64,Desc::D>> meta{{dim_x,dim_y}}; saw::data<sch::ChunkStruct<T,Desc>> lbm_data{}; - acpp::sycl::queue sycl_q; + device dev; + auto& sycl_q = dev.get_handle(); + sycl_q.wait(); { auto eov = setup_initial_conditions<T,Desc>(lbm_data); @@ -120,6 +140,27 @@ saw::error_or<void> lbm_main(int argc, char** argv){ } } + saw::data<sch::ChunkStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_data; + { + auto eov = dev.allocate_on_device(lbm_sycl_data); + if(eov.is_error()){ + return eov; + } + } + { + auto eov = dev.copy_to_device(lbm_data,lbm_sycl_data); + if(eov.is_error()){ + return eov; + } + } + + for(saw::data<sch::UInt64> i{0u}; i < saw::data<sch::UInt64>{32ul}; ++i){ + auto eov = step<T,Desc>(lbm_sycl_data,i,dev); + if(eov.is_error()){ + return eov; + } + } + /* iterator<Desc::D>::apply( [&](auto& index){ diff --git a/lib/core/c++/chunk.hpp b/lib/core/c++/chunk.hpp index ad7de34..6b028ba 100644 --- a/lib/core/c++/chunk.hpp +++ b/lib/core/c++/chunk.hpp @@ -47,7 +47,7 @@ public: data<ValueSchema, Encode>& ghost_at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index){ return values_.at(index); } - + const data<ValueSchema, Encode>& ghost_at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index) const { return values_.at(index); } @@ -55,7 +55,7 @@ public: static constexpr auto get_ghost_dims() { return data<InnerSchema,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){ @@ -63,7 +63,7 @@ public: } return values_.at(ind); } - + const 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){ @@ -75,6 +75,14 @@ public: static constexpr auto get_dims(){ return data<schema::FixedArray<schema::UInt64, sizeof...(Sides)>,Encode>{{Sides...}}; } + + auto flat_data(){ + return values_.flat_data(); + } + + static constexpr auto flat_size() { + return data<InnerSchema,Encode>::flat_size(); + } }; template<typename Sch, uint64_t Ghost, uint64_t... Sides> diff --git a/lib/core/c++/flatten.hpp b/lib/core/c++/flatten.hpp index 164bb77..1609589 100644 --- a/lib/core/c++/flatten.hpp +++ b/lib/core/c++/flatten.hpp @@ -20,7 +20,6 @@ public: return 1u; } - private: /// 2,3,4 => 2,6,24 /// i + j * 2 + k * 3*2 @@ -35,7 +34,6 @@ private: public: static saw::data<T> apply(const saw::data<sch::FixedArray<T,D>>& index, const saw::data<sch::FixedArray<T,D>>& meta){ saw::data<T> flat_ind{0u}; - apply_i<0u>(flat_ind, index, meta); return flat_ind; } diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp index bb8b4bf..d9976dd 100644 --- a/lib/sycl/c++/data.hpp +++ b/lib/sycl/c++/data.hpp @@ -11,19 +11,115 @@ struct Sycl { }; } +/* namespace impl { template<typename Schema> struct struct_has_only_equal_dimension_array{}; } +*/ } } namespace saw { +template<typename Sch, uint64_t... Dims, typename Encode> +class data<schema::FixedArray<Sch,Dims...>, kel::lbm::encode::Sycl<Encode>> final { +public: + using Schema = schema::FixedArray<Sch,Dims...>; +private: + acpp::sycl::queue* q_; + data<Sch>* values_; + + SAW_FORBID_COPY(data); +public: + data(acpp::sycl::queue& q__): + q_{&q__}, + values_{nullptr} + { + values_ = acpp::sycl::malloc<data<Sch>>(ct_multiply<uint64_t,Dims...>::value,q_); + } + + ~data(){ + if(not values_){ + return; + } + + acpp::sycl::free(values_,q_); + } + + static constexpr data<schema::FixedArray<schema::UInt64, sizeof...(Dims)>> get_dims() { + return {std::array<uint64_t, sizeof...(Dims)>{Dims...}}; + } + + data<Sch>& 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::Native>* 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: + using Schema = kel::lbm::sch::Chunk<Sch,Ghost,Sides...>; +private: + using InnerSchema = typename Schema::InnerSchema; + using ValueSchema = typename InnerSchema::ValueType; + + data<InnerSchema, kel::lbm::encode::Sycl<Encode>> values_; +public: + data(acpp::sycl::queue& q__): + values_{q__} + {} + + data<ValueSchema, kel::lbm::encode::Sycl<Encode>>& ghost_at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index){ + return values_.at(index); + } + + const 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, kel::lbm::encode::Sycl<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); + } + + const data<ValueSchema, kel::lbm::encode::Sycl<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<uint64_t Ghost, uint64_t... Meta, typename... Sch, string_literal... Keys, typename Encode> class data<schema::Struct<schema::Member<kel::lbm::sch::Chunk<Sch,Ghost,Meta...>, Keys>...>, kel::lbm::encode::Sycl<Encode> > final { public: static constexpr data<schema::FixedArray<schema::UInt64,sizeof...(Meta)>> meta = {{Meta...}}; - using StorageT = std::tuple<data<Sch,Encode>*...>; + using StorageT = std::tuple<data<kel::lbm::sch::Chunk<Sch,Ghost,Meta...>,kel::lbm::encode::Sycl<Encode>>...>; + using Schema = schema::Struct<schema::Member<kel::lbm::sch::Chunk<Sch,Ghost,Meta...>, Keys>...>; private: /** @@ -31,43 +127,18 @@ private: * Do it here by specializing. */ StorageT members_; - kel::lbm::sycl::queue* q_; public: - data(): - members_{}, - q_{nullptr} - {} - - data(StorageT members__, kel::lbm::sycl::queue& q__): - members_{members__}, - q_{&q__} + data(acpp::sycl::queue& q__): + members_{{data<kel::lbm::sch::Chunk<Sch,Ghost,Meta...>,kel::lbm::encode::Sycl<Encode>>{q__}...}} {} - ~data(){ - SAW_ASSERT(q_){ - return; - } - std::visit([this](auto arg){ - if(not arg){ - return; - } - acpp::sycl::free(arg,*q_); - arg = nullptr; - },members_); - } - - template<saw::string_literal K> - auto* get_ptr(){ - return std::get<parameter_key_pack_index<K, Keys...>::value>(members_); - } - template<saw::string_literal K> auto& get(){ - auto ptr = get_ptr<K>(); - SAW_ASSERT(ptr); - return *ptr; + return std::get<parameter_key_pack_index<K, Keys...>::value>(members_); } }; + + } namespace kel { @@ -81,7 +152,7 @@ struct sycl_malloc_struct_helper<sch::Struct<Members...>, Encode> final { using Schema = sch::Struct<Members...>; template<uint64_t i> - static saw::error_or<void> allocate_on_device_member(typename saw::data<typename saw::parameter_pack_type<i,Members...>::type::ValueType,encode::Sycl<Encode>>::StorageT& storage, sycl::queue& q){ + static saw::error_or<void> allocate_on_device_member(typename saw::data<Schema,encode::Sycl<Encode>>::StorageT& storage, sycl::queue& q){ if constexpr (i < sizeof...(Members)){ using M = typename saw::parameter_pack_type<i,Members...>::type; auto& ptr = std::get<i>(storage); @@ -103,6 +174,60 @@ struct sycl_malloc_struct_helper<sch::Struct<Members...>, Encode> final { return eov; } }; + +template<typename Sch, typename Encode> +struct sycl_copy_helper; + +template<typename... Members, typename Encode> +struct sycl_copy_helper<sch::Struct<Members...>, Encode> final { + using Schema = sch::Struct<Members...>; + + template<uint64_t i> + static saw::error_or<void> copy_to_device_member(saw::data<Schema,Encode>& host_data, saw::data<Schema,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){ + if constexpr (i < sizeof...(Members)){ + using M = typename saw::parameter_pack_type<i,Members...>::type; + auto& host_member_data = host_data.template get<i>(); + auto& sycl_member_data = sycl_data.template get<i>(); + + auto host_ptr = host_member_data.flat_data(); + auto sycl_ptr = sycl_member_data.flat_data(); + + q.memcpy(host_ptr, sycl_ptr, sizeof(std::decay_t<decltype(host_ptr)>) * host_member_data.flat_size() ); + + return copy_to_device_member<i+1u>(host_data,sycl_data,q); + } + + return saw::make_void(); + } + + static saw::error_or<void> copy_to_device(saw::data<Schema,Encode>& host_data, saw::data<Schema, encode::Sycl<Encode>>& sycl_data, sycl::queue& q){ + + return copy_to_device_member<0u>(host_data, sycl_data, q); + } + + template<uint64_t i> + static saw::error_or<void> copy_to_host_member(saw::data<Schema,encode::Sycl<Encode>>& sycl_data, saw::data<Schema,Encode>& host_data, sycl::queue& q){ + if constexpr (i < sizeof...(Members)){ + using M = typename saw::parameter_pack_type<i,Members...>::type; + auto& host_member_data = host_data.template get<i>(); + auto& sycl_member_data = sycl_data.template get<i>(); + + auto host_ptr = host_member_data.flat_data(); + auto sycl_ptr = sycl_member_data.flat_data(); + + q.memcpy(sycl_ptr, host_ptr, sizeof(std::decay_t<decltype(host_ptr)>) * host_member_data.flat_size() ); + + return copy_to_host_member<i+1u>(sycl_data,host_data,q); + } + + return saw::make_void(); + } + + + static saw::error_or<void> copy_to_host(saw::data<Schema,Encode>& sycl_data, saw::data<Schema,encode::Sycl<Encode>>& host_data, sycl::queue& q){ + return copy_to_host_member<0u>(sycl_data, host_data, q); + } +}; } class device final { private: @@ -120,7 +245,21 @@ public: if(eov.is_error()){ return eov; } - sycl_data.set_queue(q_); + return saw::make_void(); + } + + template<typename Sch, typename Encode> + saw::error_or<void> copy_to_device(saw::data<Sch,Encode>& host_data, saw::data<Sch,encode::Sycl<Encode>>& sycl_data){ + return impl::sycl_copy_helper<Sch,Encode>::copy_to_device(host_data, sycl_data, q_); + } + + template<typename Sch, typename Encode> + saw::error_or<void> copy_to_host(saw::data<Sch,encode::Sycl<Encode>>& sycl_data, saw::data<Sch,Encode>& host_data){ + return impl::sycl_copy_helper<Sch,Encode>::copy_to_host(sycl_data, host_data, q_); + } + + auto& get_handle(){ + return q_; } }; } diff --git a/lib/sycl/tests/data.cpp b/lib/sycl/tests/data.cpp index e57a1e3..3073a22 100644 --- a/lib/sycl/tests/data.cpp +++ b/lib/sycl/tests/data.cpp @@ -1,6 +1,21 @@ #include <forstio/test/suite.hpp> +#include "../c++/lbm.hpp" -SAW_TEST("Sycl Dummy"){ +SAW_TEST("Sycl Data Compilation"){ + acpp::sycl::queue q; + saw::data< + saw::schema::Struct< + saw::schema::Member< + kel::lbm::sch::Chunk<saw::schema::UInt8,1u,1u,1u>, + "test" + > + >, + kel::lbm::encode::Sycl<saw::encode::Native> + > dat{q}; + auto& test_f = dat.template get<"test">(); + + // test_f.at({}).set(1); + // SAW_EXPECT(test_f.at({}).get() == 1, "Value check failed"); } |