diff options
Diffstat (limited to 'lib/sycl/c++/data.hpp')
| -rw-r--r-- | lib/sycl/c++/data.hpp | 596 |
1 files changed, 558 insertions, 38 deletions
diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp index 50adb4f..9f43848 100644 --- a/lib/sycl/c++/data.hpp +++ b/lib/sycl/c++/data.hpp @@ -30,12 +30,23 @@ private: SAW_FORBID_COPY(data); SAW_FORBID_MOVE(data); public: + data(const data<typename meta_schema<Schema>::MetaSchema>& meta__, acpp::sycl::queue& q__): + q_{&q__}, + values_{nullptr} + { + (void) meta__; + SAW_ASSERT(q_); + values_ = acpp::sycl::malloc_device<data<Sch,Encode>>(ct_multiply<uint64_t,Dims...>::value,*q_); + SAW_ASSERT(values_); + } + data(acpp::sycl::queue& q__): q_{&q__}, values_{nullptr} { + SAW_ASSERT(q_); values_ = acpp::sycl::malloc_device<data<Sch,Encode>>(ct_multiply<uint64_t,Dims...>::value,*q_); - SAW_ASSERT(values_ and q_); + SAW_ASSERT(values_); } ~data(){ @@ -60,7 +71,7 @@ public: return values_[kel::lbm::flatten_index<schema::UInt64,sizeof...(Dims)>::apply(index,get_dims()).get()]; } - constexpr data<Sch,Encode>* flat_data() { + constexpr data<Sch,Encode>* flat_data() const { return values_; } }; @@ -81,7 +92,7 @@ public: {} data(data<schema::FixedArray<Sch,Dims...>, kel::lbm::encode::Sycl<Encode>>& values__): - values_{&values__.at({})} + values_{values__.flat_data()} {} data(data<Sch,Encode>* values__): @@ -100,9 +111,141 @@ public: return values_[kel::lbm::flatten_index<schema::UInt64,sizeof...(Dims)>::apply(index,get_dims()).get()]; } - constexpr data<Sch,Encode>* flat_data() { + constexpr data<Sch,Encode>* flat_data() const { + return values_; + } +}; + +template<typename Sch, uint64_t Dims, typename Encode> +class data<schema::Array<Sch,Dims>, kel::lbm::encode::Sycl<Encode>> final { +public: + using Schema = schema::Array<Sch,Dims>; +private: + static_assert(Dims > 0u, "Zero Dim Arrays make no sense here. If you meant to use this for math style approaches then use Tensor instead of Array"); + + data<Sch,Encode>* values_; + data<schema::FixedArray<schema::UInt64,Dims>,Encode> meta_; + acpp::sycl::queue* q_; + + SAW_FORBID_COPY(data); + SAW_FORBID_MOVE(data); +public: + data(const data<typename meta_schema<Schema>::MetaSchema>& meta__, acpp::sycl::queue& q__): + values_{nullptr}, + meta_{meta__}, + q_{&q__} + { + SAW_ASSERT(q_); + values_ = acpp::sycl::malloc_device<data<Sch,Encode>>(flat_size().get(),*q_); + SAW_ASSERT(values_); + } + + data(acpp::sycl::queue& q__): + values_{nullptr}, + meta_{}, + q_{&q__} + { + SAW_ASSERT(q_); + } + + ~data(){ + if(not values_){ + return; + } + SAW_ASSERT(q_); + + acpp::sycl::free(values_,*q_); + values_ = nullptr; + } + + constexpr data<schema::FixedArray<schema::UInt64, Dims>, Encode> meta() const { + return meta_; + } + + constexpr data<Sch,Encode>& at(const data<schema::FixedArray<schema::UInt64,Dims>, Encode>& index){ + return values_[kel::lbm::flatten_index<schema::UInt64,Dims>::apply(index,meta()).get()]; + } + + constexpr data<Sch,Encode>& at(const data<schema::FixedArray<schema::UInt64,Dims>, Encode>& index) const{ + return values_[kel::lbm::flatten_index<schema::UInt64,Dims>::apply(index,meta()).get()]; + } + + constexpr error_or<void> reset_to(const data<typename meta_schema<Schema>::MetaSchema>& meta_arg){ + SAW_ASSERT(q_); + meta_ = meta_arg; + + if(values_){ + acpp::sycl::free(values_,*q_); + } + values_ = acpp::sycl::malloc_device<data<Sch,Encode>>(flat_size().get(),*q_); + SAW_ASSERT(q_); + + return make_void(); + } + + constexpr data<Sch,Encode>* flat_data() const { + return values_; + } + + constexpr data<schema::UInt64,Encode> flat_size() const { + data<schema::UInt64> mult{1u}; + + for(uint64_t i{0u}; i < Dims; ++i){ + mult = mult * meta_.at({i}); + } + + return mult; + } +}; + +template<typename Sch, uint64_t Dims, typename Encode> +class data<schema::Ptr<schema::Array<Sch,Dims>>, kel::lbm::encode::Sycl<Encode>> final { +public: + using Schema = schema::Ptr<schema::Array<Sch,Dims>>; +private: + static_assert(Dims > 0u, "Zero Dim Arrays make no sense here. If you meant to use this for math style approaches then use Tensor instead of Array"); + + data<Sch,Encode>* values_; + data<schema::FixedArray<schema::UInt64,Dims>,Encode> meta_; + +public: + SAW_DEFAULT_COPY(data); + SAW_DEFAULT_MOVE(data); + + data(): + values_{nullptr} + {} + + data(const data<schema::Array<Sch,Dims>, kel::lbm::encode::Sycl<Encode>>& values__): + values_{values__.flat_data()}, + meta_{values__.meta()} + {} + + constexpr data<schema::FixedArray<schema::UInt64, Dims>, Encode> meta() const { + return meta_; + } + + constexpr data<Sch,Encode>& at(const data<schema::FixedArray<schema::UInt64,Dims>, Encode>& index){ + return values_[kel::lbm::flatten_index<schema::UInt64,Dims>::apply(index,meta()).get()]; + } + + constexpr data<Sch,Encode>& at(const data<schema::FixedArray<schema::UInt64,Dims>, Encode>& index) const{ + return values_[kel::lbm::flatten_index<schema::UInt64,Dims>::apply(index,meta()).get()]; + } + + constexpr data<Sch,Encode>* flat_data() const { return values_; } + + constexpr data<schema::UInt64,Encode> flat_size() const { + data<schema::UInt64> mult{1u}; + + for(uint64_t i{0u}; i < Dims; ++i){ + mult = mult * meta_.at({i}); + } + + return mult; + } }; template<typename Sch, uint64_t Ghost, uint64_t... Sides, typename Encode> @@ -115,6 +258,10 @@ private: data<InnerSchema, kel::lbm::encode::Sycl<Encode>> values_; public: + data(const data<typename meta_schema<Schema>::MetaSchema>& meta__, acpp::sycl::queue& q__): + values_{meta__,q__} + {} + data(acpp::sycl::queue& q__): values_{q__} {} @@ -131,6 +278,10 @@ public: return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::get_dims(); } + static constexpr auto ghost_meta() { + return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::meta(); + } + 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){ @@ -151,7 +302,11 @@ public: return data<schema::FixedArray<schema::UInt64, sizeof...(Sides)>,Encode>{{Sides...}}; } - auto flat_data(){ + static constexpr auto meta(){ + return data<schema::FixedArray<schema::UInt64, sizeof...(Sides)>,Encode>{{Sides...}}; + } + + auto flat_data() const { return values_.flat_data(); } @@ -178,6 +333,10 @@ public: values_{nullptr} {} + data(const data<kel::lbm::sch::Chunk<Sch,Ghost,Sides...>, kel::lbm::encode::Sycl<Encode>>& values__): + values_{values__.flat_data()} + {} + data(data<Sch,Encode>* values__): values_{values__} {} @@ -194,6 +353,10 @@ public: return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::get_dims(); } + static constexpr auto ghost_meta() { + return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::meta(); + } + 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){ @@ -210,11 +373,15 @@ public: return values_.at(ind); } + static constexpr auto meta(){ + return data<schema::FixedArray<schema::UInt64, sizeof...(Sides)>,Encode>{{Sides...}}; + } + static constexpr auto get_dims(){ return data<schema::FixedArray<schema::UInt64, sizeof...(Sides)>,Encode>{{Sides...}}; } - auto flat_data(){ + auto flat_data() const { return values_.flat_data(); } @@ -224,6 +391,70 @@ public: }; template<typename... Members, typename Encode> +struct data<schema::Tuple<Members...>, kel::lbm::encode::Sycl<Encode>> final { +public: + using StorageT = std::tuple<data<Members,kel::lbm::encode::Sycl<Encode>>...>; + using Schema = schema::Tuple<Members...>; +private: + StorageT members_; + + /** + * A helper constructor to forward the sycl queue to the inner "default" constructors + */ + template<std::size_t... Is> + constexpr data(acpp::sycl::queue& q, std::index_sequence<Is...>): + members_{(static_cast<void>(Is), q)...} + {} +public: + /* + data(data<typename meta_schema<Schema>::MetaSchema>& meta__, acpp::sycl::queue& q__): + data{q__, std::make_index_sequence<sizeof...(Members)>{}} + {} + */ + + data(acpp::sycl::queue& q__): + data{q__, std::make_index_sequence<sizeof...(Members)>{}} + { + q__.wait(); + } + + template<size_t i> + auto& get(){ + return std::get<i>(members_); + } + + template<size_t i> + auto& get() const { + return std::get<i>(members_); + } +}; + +template<typename... Members, typename Encode> +struct data<schema::Ptr<schema::Tuple<Members...>>, kel::lbm::encode::Sycl<Encode>> final { +public: + using StorageT = std::tuple<data<schema::Ptr<Members>,kel::lbm::encode::Sycl<Encode>>...>; + using Schema = schema::Tuple<Members...>; +private: + StorageT members_; + + /** + * A helper constructor to forward the sycl queue to the inner "default" constructors + */ +public: + data() = default; + + template<size_t i> + auto& get(){ + return std::get<i>(members_); + } + + template<size_t i> + const auto& get() const { + return std::get<i>(members_); + } +}; + +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>>...>; @@ -235,6 +466,9 @@ private: */ StorageT members_; + /** + * A helper constructor to forward the sycl queue to the inner "default" constructors + */ template<std::size_t... Is> constexpr data(acpp::sycl::queue& q, std::index_sequence<Is...>): members_{(static_cast<void>(Is), q)...} @@ -246,13 +480,23 @@ public: q__.wait(); } + template<size_t i> + auto& get(){ + return std::get<i>(members_); + } + + template<size_t i> + auto& get() const { + return std::get<i>(members_); + } + template<saw::string_literal K> auto& get(){ return std::get<parameter_key_pack_index<K, Members::KeyLiteral...>::value>(members_); } template<saw::string_literal K> - const auto& get() const { + auto& get() const { return std::get<parameter_key_pack_index<K, Members::KeyLiteral...>::value>(members_); } }; @@ -268,17 +512,26 @@ private: * Do it here by specializing. */ StorageT members_; - public: data() = default; + template<size_t i> + auto& get(){ + return std::get<i>(members_); + } + + template<size_t i> + auto& get() const { + return std::get<i>(members_); + } + 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 { + auto& get() const { return std::get<parameter_key_pack_index<K, Keys...>::value>(members_); } }; @@ -301,14 +554,10 @@ struct sycl_copy_helper<sch::Struct<Members...>, Encode> final { auto& host_member_data = host_data.template get<M::KeyLiteral>(); auto& sycl_member_data = sycl_data.template get<M::KeyLiteral>(); - auto host_ptr = host_member_data.flat_data(); - auto sycl_ptr = sycl_member_data.flat_data(); - - static_assert(sizeof(std::decay_t<decltype(sycl_ptr)>) == sizeof(std::decay_t<decltype(host_ptr)>), "Unequal size"); - - q.submit([&](acpp::sycl::handler& h){ - h.copy(host_ptr,sycl_ptr, host_member_data.flat_size().get()); - }).wait(); + auto eov = sycl_copy_helper<typename M::ValueType,Encode>::copy_to_device(host_member_data,sycl_member_data,q); + if(eov.is_error()){ + return eov; + } return copy_to_device_member<i+1u>(host_data,sycl_data,q); } @@ -328,12 +577,10 @@ struct sycl_copy_helper<sch::Struct<Members...>, Encode> final { auto& host_member_data = host_data.template get<M::KeyLiteral>(); auto& sycl_member_data = sycl_data.template get<M::KeyLiteral>(); - auto host_ptr = host_member_data.flat_data(); - auto sycl_ptr = sycl_member_data.flat_data(); - - q.submit([&](acpp::sycl::handler& h){ - h.copy(sycl_ptr,host_ptr, host_member_data.flat_size().get()); - }).wait(); + auto eov = sycl_copy_helper<typename M::ValueType,Encode>::copy_to_host(sycl_member_data,host_member_data,q); + if(eov.is_error()){ + return eov; + } return copy_to_host_member<i+1u>(sycl_data,host_data,q); } @@ -345,6 +592,107 @@ struct sycl_copy_helper<sch::Struct<Members...>, Encode> final { static saw::error_or<void> copy_to_host(saw::data<Schema,encode::Sycl<Encode>>& sycl_data, saw::data<Schema,Encode>& host_data, sycl::queue& q){ return copy_to_host_member<0u>(sycl_data, host_data, q); } + + template<uint64_t i> + static saw::error_or<void> malloc_on_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 eov = sycl_copy_helper<typename M::ValueType,Encode>::malloc_on_device(host_member_data,sycl_member_data,q); + if(eov.is_error()){ + return eov; + } + + return malloc_on_device_member<i+1u>(host_data,sycl_data,q); + } + + return saw::make_void(); + } + + static saw::error_or<void> malloc_on_device(saw::data<Schema,Encode>& host_data, saw::data<Schema,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){ + return malloc_on_device_member<0u>(host_data,sycl_data,q); + } +}; + +template<typename... Members, typename Encode> +struct sycl_copy_helper<sch::Tuple<Members...>, Encode> final { + using Schema = sch::Tuple<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 eov = sycl_copy_helper<M,Encode>::copy_to_device(host_member_data,sycl_member_data,q); + if(eov.is_error()){ + return eov; + } + + 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 eov = sycl_copy_helper<M,Encode>::copy_to_host(sycl_member_data,host_member_data,q); + if(eov.is_error()){ + return eov; + } + + 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>& host_data, + saw::data<Schema,encode::Sycl<Encode>>& sycl_data, + sycl::queue& q + ){ + return copy_to_host_member<0u>(sycl_data, host_data, q); + } + + template<uint64_t i> + static saw::error_or<void> malloc_on_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 eov = sycl_copy_helper<M,Encode>::malloc_on_device(host_member_data,sycl_member_data,q); + if(eov.is_error()){ + return eov; + } + + return malloc_on_device_member<i+1u>(host_data,sycl_data,q); + } + + return saw::make_void(); + } + + static saw::error_or<void> malloc_on_device(saw::data<Schema,Encode>& host_data, saw::data<Schema,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){ + return malloc_on_device_member<0u>(host_data,sycl_data,q); + } }; template<typename Sch, uint64_t... Dims, typename Encode> @@ -374,43 +722,205 @@ struct sycl_copy_helper<sch::FixedArray<Sch,Dims...>, Encode> final { }).wait(); return saw::make_void(); } + + static saw::error_or<void> malloc_on_device(saw::data<Schema,Encode>& host_data, saw::data<Schema,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){ + (void) host_data; + (void) sycl_data; + (void) q; + return saw::make_void(); + } +}; + +template<typename Sch, uint64_t Ghost, uint64_t... Dims, typename Encode> +struct sycl_copy_helper<sch::Chunk<Sch,Ghost,Dims...>, Encode> final { + using Schema = sch::Chunk<Sch,Ghost,Dims...>; + + static saw::error_or<void> copy_to_host(saw::data<Schema,encode::Sycl<Encode>>& sycl_data, saw::data<Schema,Encode>& host_data, sycl::queue& q){ + auto host_ptr = host_data.flat_data(); + auto sycl_ptr = sycl_data.flat_data(); + + static_assert(sizeof(std::decay_t<decltype(sycl_ptr)>) == sizeof(std::decay_t<decltype(host_ptr)>), "Unequal size"); + + auto flat_size = host_data.flat_size(); + + q.submit([&](acpp::sycl::handler& h){ + h.copy(sycl_ptr,host_ptr, flat_size.get()); + }).wait(); + 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){ + auto host_ptr = host_data.flat_data(); + auto sycl_ptr = sycl_data.flat_data(); + + static_assert(sizeof(std::decay_t<decltype(sycl_ptr)>) == sizeof(std::decay_t<decltype(host_ptr)>), "Unequal size"); + + auto flat_size = host_data.flat_size(); + + q.submit([&](acpp::sycl::handler& h){ + h.copy(host_ptr,sycl_ptr, flat_size.get()); + }).wait(); + return saw::make_void(); + } +}; + +template<typename Sch, uint64_t Dims, typename Encode> +struct sycl_copy_helper<sch::Array<Sch,Dims>, Encode> final { + using Schema = sch::Array<Sch,Dims>; + + static saw::error_or<void> copy_to_host(saw::data<Schema,encode::Sycl<Encode>>& sycl_data, saw::data<Schema,Encode>& host_data, sycl::queue& q){ + auto host_ptr = host_data.flat_data(); + auto sycl_ptr = sycl_data.flat_data(); + + static_assert(sizeof(std::decay_t<decltype(sycl_ptr)>) == sizeof(std::decay_t<decltype(host_ptr)>), "Unequal size"); + + SAW_ASSERT(host_data.flat_size() == sycl_data.flat_size()); + q.submit([&](acpp::sycl::handler& h){ + h.copy(sycl_ptr,host_ptr, host_data.flat_size().get()); + }).wait(); + 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){ + + { + auto hm = host_data.meta(); + auto sm = sycl_data.meta(); + bool equ{true}; + for(uint64_t i{0u}; i < Dims; ++i){ + equ &= (hm.at({i}).get() == sm.at({i}).get()); + } + if(not equ){ + sycl_data.reset_to(hm); + } + } + + auto host_ptr = host_data.flat_data(); + auto sycl_ptr = sycl_data.flat_data(); + static_assert(sizeof(std::decay_t<decltype(sycl_ptr)>) == sizeof(std::decay_t<decltype(host_ptr)>), "Unequal size"); + + q.submit([&](acpp::sycl::handler& h){ + h.copy(host_ptr,sycl_ptr, host_data.flat_size().get()); + }).wait(); + + return saw::make_void(); + } + + static saw::error_or<void> malloc_on_device(saw::data<Schema,Encode>& host_data, saw::data<Schema,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){ + sycl_data = {host_data.meta(),q}; + return saw::make_void(); + } }; template<typename Schema, typename Encode> -struct make_chunk_struct_view_helper; +struct make_view_helper; + +template<typename Sch, uint64_t Dims, typename Encode> +struct make_view_helper<sch::Array<Sch,Dims>, encode::Sycl<Encode>> { +public: +static saw::error_or<void> apply( + saw::data<sch::Array<Sch,Dims>, encode::Sycl<Encode>>& dat, + saw::data<sch::Ptr<sch::Array<Sch,Dims>>, encode::Sycl<Encode>>& dat_view +){ + dat_view = {dat}; + return saw::make_void(); +} +}; + +template<typename Sch, uint64_t... Dims, typename Encode> +struct make_view_helper<sch::FixedArray<Sch,Dims...>, encode::Sycl<Encode>> { +public: + static saw::error_or<void> apply( + saw::data<sch::FixedArray<Sch,Dims...>, encode::Sycl<Encode>>& dat, + saw::data<sch::Ptr<sch::FixedArray<Sch,Dims...>>, encode::Sycl<Encode>>& dat_view + ){ + dat_view = {dat}; + return saw::make_void(); + } +}; + +template<typename Sch, uint64_t Ghost, uint64_t... Dims, typename Encode> +struct make_view_helper<sch::Chunk<Sch,Ghost,Dims...>, encode::Sycl<Encode>> { +public: + static saw::error_or<void> apply( + saw::data<sch::Chunk<Sch,Ghost,Dims...>, encode::Sycl<Encode>>& dat, + saw::data<sch::Ptr<sch::Chunk<Sch,Ghost,Dims...>>, encode::Sycl<Encode>>& dat_view + ){ + dat_view = {dat}; + return saw::make_void(); + } +}; template<typename... Sch, saw::string_literal... Keys, typename Encode> -struct make_chunk_struct_view_helper<sch::Struct<sch::Member<Sch,Keys>...>, encode::Sycl<Encode>> { +struct make_view_helper<sch::Struct<sch::Member<Sch,Keys>...>, encode::Sycl<Encode>> { private: - template<uint64_t i> - static void apply_i( +template<uint64_t i> +static saw::error_or<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) - { + 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); + auto eov = make_view_helper<typename M::ValueType,encode::Sycl<Encode>>::apply(dat.template get<M::KeyLiteral>(),dat_view.template get<M::KeyLiteral>()); + if(eov.is_error()){ + return eov; + } + + return apply_i<i+1u>(dat,dat_view); } + + return saw::make_void(); } 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; +static saw::error_or<void> 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>>& dat_view +){ + return apply_i<0u>(dat,dat_view); +} +}; + +template<typename... Sch, typename Encode> +struct make_view_helper<sch::Tuple<Sch...>, encode::Sycl<Encode>> { +private: + template<uint64_t i> + static saw::error_or<void> apply_i( + saw::data<sch::Tuple<Sch...>, encode::Sycl<Encode>>& dat, + saw::data<sch::Ptr<sch::Tuple<Sch...>>, encode::Sycl<Encode>>& dat_view + ){ + if constexpr (i < sizeof...(Sch)){ + using M = typename saw::parameter_pack_type<i,Sch...>::type; + + auto eov = make_view_helper<M,encode::Sycl<Encode>>::apply(dat.template get<i>(), dat_view.template get<i>()); + if(eov.is_error()){ + return eov; + } - apply_i<0u>(dat,str_view); + return apply_i<i+1u>(dat,dat_view); + } - return str_view; + return saw::make_void(); + } +public: + static saw::error_or<void> apply( + saw::data<sch::Tuple<Sch...>, encode::Sycl<Encode>>& dat, + saw::data<sch::Ptr<sch::Tuple<Sch...>>, encode::Sycl<Encode>> dat_view + ){ + return apply_i<0u>(dat,dat_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); +auto make_view(saw::data<Schema,Encode>& dat){ + saw::data<sch::Ptr<Schema>,Encode> dat_view; + auto eov = impl::make_view_helper<Schema,Encode>::apply(dat,dat_view); + (void) eov; + + return dat_view; } class device final { @@ -433,6 +943,16 @@ public: return impl::sycl_copy_helper<Sch,Encode>::copy_to_host(sycl_data, host_data, q_); } + template<typename Sch, typename Encode> + saw::error_or<void> malloc_on_device( + saw::data<Sch,Encode>& host_data, + saw::data<Sch,encode::Sycl<Encode>>& sycl_data + ){ + auto eov = impl::sycl_copy_helper<Sch,Encode>::malloc_on_device(host_data, sycl_data, q_); + q_.wait(); + return eov; + } + auto& get_handle(){ return q_; } |