diff options
Diffstat (limited to 'lib/sycl/c++/data.hpp')
| -rw-r--r-- | lib/sycl/c++/data.hpp | 961 |
1 files changed, 0 insertions, 961 deletions
diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp deleted file mode 100644 index 9f43848..0000000 --- a/lib/sycl/c++/data.hpp +++ /dev/null @@ -1,961 +0,0 @@ -#include "common.hpp" -#include <kel/lbm/lbm.hpp> - -namespace kel { -namespace lbm { -namespace encode { -template<typename Encode> -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,Encode>* values_; - - 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_); - } - - ~data(){ - if(not values_){ - return; - } - SAW_ASSERT(q_); - - acpp::sycl::free(values_,*q_); - values_ = nullptr; - } - - 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(const 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(const 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() const { - return values_; - } -}; - -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_; - -public: - SAW_DEFAULT_COPY(data); - SAW_DEFAULT_MOVE(data); - - data(): - values_{nullptr} - {} - - data(data<schema::FixedArray<Sch,Dims...>, kel::lbm::encode::Sycl<Encode>>& values__): - values_{values__.flat_data()} - {} - - 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(const 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(const 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() 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> -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(const data<typename meta_schema<Schema>::MetaSchema>& meta__, acpp::sycl::queue& q__): - values_{meta__,q__} - {} - - data(acpp::sycl::queue& q__): - values_{q__} - {} - - constexpr data<ValueSchema, Encode>& ghost_at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index){ - return values_.at(index); - } - - constexpr data<ValueSchema, 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(); - } - - 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){ - 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)>,Encode>{{Sides...}}; - } - - static constexpr auto meta(){ - return data<schema::FixedArray<schema::UInt64, sizeof...(Sides)>,Encode>{{Sides...}}; - } - - auto flat_data() const { - return values_.flat_data(); - } - - static constexpr auto flat_size() { - return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::flat_size(); - } -}; - -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_; - -public: - SAW_DEFAULT_MOVE(data); - SAW_DEFAULT_COPY(data); - - data(): - 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__} - {} - - constexpr data<ValueSchema, Encode>& ghost_at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index){ - return values_.at(index); - } - - constexpr data<ValueSchema, 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(); - } - - 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){ - 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 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() const { - return values_.flat_data(); - } - - static constexpr auto flat_size() { - return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::flat_size(); - } -}; - -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>>...>; - 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. - */ - 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(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<saw::string_literal K> - auto& get(){ - return std::get<parameter_key_pack_index<K, Members::KeyLiteral...>::value>(members_); - } - - template<saw::string_literal K> - auto& get() const { - 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<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> - auto& get() const { - return std::get<parameter_key_pack_index<K, Keys...>::value>(members_); - } -}; -} - -namespace kel { -namespace lbm { -namespace impl { -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<M::KeyLiteral>(); - auto& sycl_member_data = sycl_data.template get<M::KeyLiteral>(); - - 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); - } - - 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<M::KeyLiteral>(); - auto& sycl_member_data = sycl_data.template get<M::KeyLiteral>(); - - 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); - } - - return saw::make_void(); - } - - - 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> -struct sycl_copy_helper<sch::FixedArray<Sch,Dims...>, Encode> final { - using Schema = sch::FixedArray<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"); - - q.submit([&](acpp::sycl::handler& h){ - h.copy(sycl_ptr,host_ptr, saw::ct_multiply<uint64_t,Dims...>::value); - }).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"); - - q.submit([&](acpp::sycl::handler& h){ - h.copy(host_ptr,sycl_ptr, saw::ct_multiply<uint64_t,Dims...>::value); - }).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_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_view_helper<sch::Struct<sch::Member<Sch,Keys>...>, encode::Sycl<Encode>> { -private: -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 - ){ - if constexpr (i < sizeof...(Sch)){ - using M = typename saw::parameter_pack_type<i,sch::Member<Sch,Keys>...>::type; - - 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 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; - } - - return apply_i<i+1u>(dat,dat_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_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 { -private: - sycl::queue q_; - - SAW_FORBID_COPY(device); - SAW_FORBID_MOVE(device); -public: - device() = default; - ~device() = default; - - 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_); - } - - 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_; - } -}; -} -} |