From c0549d71b2109f10c1238db8b22362e7826ba61b Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Sun, 5 Jul 2026 15:59:23 +0200 Subject: Just rename from lib to modules --- lib/sycl/c++/data.hpp | 961 -------------------------------------------------- 1 file changed, 961 deletions(-) delete mode 100644 lib/sycl/c++/data.hpp (limited to 'lib/sycl/c++/data.hpp') 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 - -namespace kel { -namespace lbm { -namespace encode { -template -struct Sycl { -}; -} - -/* -namespace impl { -template -struct struct_has_only_equal_dimension_array{}; -} -*/ -} -} - -namespace saw { -template -class data, kel::lbm::encode::Sycl> final { -public: - using Schema = schema::FixedArray; -private: - acpp::sycl::queue* q_; - data* values_; - - SAW_FORBID_COPY(data); - SAW_FORBID_MOVE(data); -public: - data(const data::MetaSchema>& meta__, acpp::sycl::queue& q__): - q_{&q__}, - values_{nullptr} - { - (void) meta__; - SAW_ASSERT(q_); - values_ = acpp::sycl::malloc_device>(ct_multiply::value,*q_); - SAW_ASSERT(values_); - } - - data(acpp::sycl::queue& q__): - q_{&q__}, - values_{nullptr} - { - SAW_ASSERT(q_); - values_ = acpp::sycl::malloc_device>(ct_multiply::value,*q_); - SAW_ASSERT(values_); - } - - ~data(){ - if(not values_){ - return; - } - SAW_ASSERT(q_); - - acpp::sycl::free(values_,*q_); - values_ = nullptr; - } - - static constexpr data> get_dims() { - return saw::data>{{Dims...}}; - } - - constexpr data& at(const data>& index){ - return values_[kel::lbm::flatten_index::apply(index,get_dims()).get()]; - } - - constexpr data& at(const data>& index) const{ - return values_[kel::lbm::flatten_index::apply(index,get_dims()).get()]; - } - - constexpr data* flat_data() const { - return values_; - } -}; - -template -class data>, kel::lbm::encode::Sycl> final { -public: - using Schema = schema::Ptr>; -private: - data* values_; - -public: - SAW_DEFAULT_COPY(data); - SAW_DEFAULT_MOVE(data); - - data(): - values_{nullptr} - {} - - data(data, kel::lbm::encode::Sycl>& values__): - values_{values__.flat_data()} - {} - - data(data* values__): - values_{values__} - {} - - static constexpr data> get_dims() { - return saw::data>{{Dims...}}; - } - - constexpr data& at(const data>& index){ - return values_[kel::lbm::flatten_index::apply(index,get_dims()).get()]; - } - - constexpr data& at(const data>& index) const{ - return values_[kel::lbm::flatten_index::apply(index,get_dims()).get()]; - } - - constexpr data* flat_data() const { - return values_; - } -}; - -template -class data, kel::lbm::encode::Sycl> final { -public: - using Schema = schema::Array; -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* values_; - data,Encode> meta_; - acpp::sycl::queue* q_; - - SAW_FORBID_COPY(data); - SAW_FORBID_MOVE(data); -public: - data(const data::MetaSchema>& meta__, acpp::sycl::queue& q__): - values_{nullptr}, - meta_{meta__}, - q_{&q__} - { - SAW_ASSERT(q_); - values_ = acpp::sycl::malloc_device>(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, Encode> meta() const { - return meta_; - } - - constexpr data& at(const data, Encode>& index){ - return values_[kel::lbm::flatten_index::apply(index,meta()).get()]; - } - - constexpr data& at(const data, Encode>& index) const{ - return values_[kel::lbm::flatten_index::apply(index,meta()).get()]; - } - - constexpr error_or reset_to(const data::MetaSchema>& meta_arg){ - SAW_ASSERT(q_); - meta_ = meta_arg; - - if(values_){ - acpp::sycl::free(values_,*q_); - } - values_ = acpp::sycl::malloc_device>(flat_size().get(),*q_); - SAW_ASSERT(q_); - - return make_void(); - } - - constexpr data* flat_data() const { - return values_; - } - - constexpr data flat_size() const { - data mult{1u}; - - for(uint64_t i{0u}; i < Dims; ++i){ - mult = mult * meta_.at({i}); - } - - return mult; - } -}; - -template -class data>, kel::lbm::encode::Sycl> final { -public: - using Schema = schema::Ptr>; -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* values_; - data,Encode> meta_; - -public: - SAW_DEFAULT_COPY(data); - SAW_DEFAULT_MOVE(data); - - data(): - values_{nullptr} - {} - - data(const data, kel::lbm::encode::Sycl>& values__): - values_{values__.flat_data()}, - meta_{values__.meta()} - {} - - constexpr data, Encode> meta() const { - return meta_; - } - - constexpr data& at(const data, Encode>& index){ - return values_[kel::lbm::flatten_index::apply(index,meta()).get()]; - } - - constexpr data& at(const data, Encode>& index) const{ - return values_[kel::lbm::flatten_index::apply(index,meta()).get()]; - } - - constexpr data* flat_data() const { - return values_; - } - - constexpr data flat_size() const { - data mult{1u}; - - for(uint64_t i{0u}; i < Dims; ++i){ - mult = mult * meta_.at({i}); - } - - return mult; - } -}; - -template -class data,kel::lbm::encode::Sycl> final { -public: - using Schema = kel::lbm::sch::Chunk; -private: - using InnerSchema = typename Schema::InnerSchema; - using ValueSchema = typename InnerSchema::ValueType; - - data> values_; -public: - data(const data::MetaSchema>& meta__, acpp::sycl::queue& q__): - values_{meta__,q__} - {} - - data(acpp::sycl::queue& q__): - values_{q__} - {} - - constexpr data& ghost_at(const data>& index){ - return values_.at(index); - } - - constexpr data& ghost_at(const data>& index) const { - return values_.at(index); - } - - static constexpr auto get_ghost_dims() { - return data>::get_dims(); - } - - static constexpr auto ghost_meta() { - return data>::meta(); - } - - 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,Encode>{{Sides...}}; - } - - static constexpr auto meta(){ - return data,Encode>{{Sides...}}; - } - - auto flat_data() const { - return values_.flat_data(); - } - - static constexpr auto flat_size() { - return data>::flat_size(); - } -}; - -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_; - -public: - SAW_DEFAULT_MOVE(data); - SAW_DEFAULT_COPY(data); - - data(): - values_{nullptr} - {} - - data(const data, kel::lbm::encode::Sycl>& values__): - values_{values__.flat_data()} - {} - - data(data* values__): - values_{values__} - {} - - constexpr data& ghost_at(const data>& index){ - return values_.at(index); - } - - constexpr data& ghost_at(const data>& index) const { - return values_.at(index); - } - - static constexpr auto get_ghost_dims() { - return data>::get_dims(); - } - - static constexpr auto ghost_meta() { - return data>::meta(); - } - - 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 meta(){ - return data,Encode>{{Sides...}}; - } - - static constexpr auto get_dims(){ - return data,Encode>{{Sides...}}; - } - - auto flat_data() const { - return values_.flat_data(); - } - - static constexpr auto flat_size() { - return data>::flat_size(); - } -}; - -template -struct data, kel::lbm::encode::Sycl> final { -public: - using StorageT = std::tuple>...>; - using Schema = schema::Tuple; -private: - StorageT members_; - - /** - * A helper constructor to forward the sycl queue to the inner "default" constructors - */ - template - constexpr data(acpp::sycl::queue& q, std::index_sequence): - members_{(static_cast(Is), q)...} - {} -public: - /* - data(data::MetaSchema>& meta__, acpp::sycl::queue& q__): - data{q__, std::make_index_sequence{}} - {} - */ - - data(acpp::sycl::queue& q__): - data{q__, std::make_index_sequence{}} - { - q__.wait(); - } - - template - auto& get(){ - return std::get(members_); - } - - template - auto& get() const { - return std::get(members_); - } -}; - -template -struct data>, kel::lbm::encode::Sycl> final { -public: - using StorageT = std::tuple,kel::lbm::encode::Sycl>...>; - using Schema = schema::Tuple; -private: - StorageT members_; - - /** - * A helper constructor to forward the sycl queue to the inner "default" constructors - */ -public: - data() = default; - - template - auto& get(){ - return std::get(members_); - } - - template - const auto& get() const { - return std::get(members_); - } -}; - -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. - */ - StorageT members_; - - /** - * A helper constructor to forward the sycl queue to the inner "default" constructors - */ - template - constexpr data(acpp::sycl::queue& q, std::index_sequence): - members_{(static_cast(Is), q)...} - {} -public: - data(acpp::sycl::queue& q__): - data{q__, std::make_index_sequence{}} - { - q__.wait(); - } - - template - auto& get(){ - return std::get(members_); - } - - template - auto& get() const { - return std::get(members_); - } - - template - auto& get(){ - return std::get::value>(members_); - } - - template - auto& get() const { - 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(members_); - } - - template - auto& get() const { - return std::get(members_); - } - - template - auto& get(){ - return std::get::value>(members_); - } - - template - auto& get() const { - return std::get::value>(members_); - } -}; -} - -namespace kel { -namespace lbm { -namespace impl { -template -struct sycl_copy_helper; - -template -struct sycl_copy_helper, Encode> final { - using Schema = sch::Struct; - - template - static saw::error_or copy_to_device_member(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ - if constexpr (i < sizeof...(Members)){ - using M = typename saw::parameter_pack_type::type; - auto& host_member_data = host_data.template get(); - auto& sycl_member_data = sycl_data.template get(); - - auto eov = sycl_copy_helper::copy_to_device(host_member_data,sycl_member_data,q); - if(eov.is_error()){ - return eov; - } - - return copy_to_device_member(host_data,sycl_data,q); - } - - return saw::make_void(); - } - - static saw::error_or copy_to_device(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ - - return copy_to_device_member<0u>(host_data, sycl_data, q); - } - - template - static saw::error_or copy_to_host_member(saw::data>& sycl_data, saw::data& host_data, sycl::queue& q){ - if constexpr (i < sizeof...(Members)){ - using M = typename saw::parameter_pack_type::type; - auto& host_member_data = host_data.template get(); - auto& sycl_member_data = sycl_data.template get(); - - auto eov = sycl_copy_helper::copy_to_host(sycl_member_data,host_member_data,q); - if(eov.is_error()){ - return eov; - } - - return copy_to_host_member(sycl_data,host_data,q); - } - - return saw::make_void(); - } - - - static saw::error_or copy_to_host(saw::data>& sycl_data, saw::data& host_data, sycl::queue& q){ - return copy_to_host_member<0u>(sycl_data, host_data, q); - } - - template - static saw::error_or malloc_on_device_member(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ - - if constexpr (i < sizeof...(Members)){ - using M = typename saw::parameter_pack_type::type; - auto& host_member_data = host_data.template get(); - auto& sycl_member_data = sycl_data.template get(); - - auto eov = sycl_copy_helper::malloc_on_device(host_member_data,sycl_member_data,q); - if(eov.is_error()){ - return eov; - } - - return malloc_on_device_member(host_data,sycl_data,q); - } - - return saw::make_void(); - } - - static saw::error_or malloc_on_device(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ - return malloc_on_device_member<0u>(host_data,sycl_data,q); - } -}; - -template -struct sycl_copy_helper, Encode> final { - using Schema = sch::Tuple; - - template - static saw::error_or copy_to_device_member(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ - if constexpr (i < sizeof...(Members)){ - using M = typename saw::parameter_pack_type::type; - auto& host_member_data = host_data.template get(); - auto& sycl_member_data = sycl_data.template get(); - - auto eov = sycl_copy_helper::copy_to_device(host_member_data,sycl_member_data,q); - if(eov.is_error()){ - return eov; - } - - return copy_to_device_member(host_data,sycl_data,q); - } - - return saw::make_void(); - } - - static saw::error_or copy_to_device(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ - return copy_to_device_member<0u>(host_data, sycl_data, q); - } - - template - static saw::error_or copy_to_host_member(saw::data>& sycl_data, saw::data& host_data, sycl::queue& q){ - if constexpr (i < sizeof...(Members)){ - using M = typename saw::parameter_pack_type::type; - - auto& host_member_data = host_data.template get(); - auto& sycl_member_data = sycl_data.template get(); - - auto eov = sycl_copy_helper::copy_to_host(sycl_member_data,host_member_data,q); - if(eov.is_error()){ - return eov; - } - - return copy_to_host_member(sycl_data,host_data,q); - } - - return saw::make_void(); - } - - - static saw::error_or copy_to_host( - saw::data& host_data, - saw::data>& sycl_data, - sycl::queue& q - ){ - return copy_to_host_member<0u>(sycl_data, host_data, q); - } - - template - static saw::error_or malloc_on_device_member(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ - - if constexpr (i < sizeof...(Members)){ - using M = typename saw::parameter_pack_type::type; - auto& host_member_data = host_data.template get(); - auto& sycl_member_data = sycl_data.template get(); - - auto eov = sycl_copy_helper::malloc_on_device(host_member_data,sycl_member_data,q); - if(eov.is_error()){ - return eov; - } - - return malloc_on_device_member(host_data,sycl_data,q); - } - - return saw::make_void(); - } - - static saw::error_or malloc_on_device(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ - return malloc_on_device_member<0u>(host_data,sycl_data,q); - } -}; - -template -struct sycl_copy_helper, Encode> final { - using Schema = sch::FixedArray; - - static saw::error_or copy_to_host(saw::data>& sycl_data, saw::data& 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) == sizeof(std::decay_t), "Unequal size"); - - q.submit([&](acpp::sycl::handler& h){ - h.copy(sycl_ptr,host_ptr, saw::ct_multiply::value); - }).wait(); - return saw::make_void(); - } - - static saw::error_or copy_to_device(saw::data& host_data, saw::data>& 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) == sizeof(std::decay_t), "Unequal size"); - - q.submit([&](acpp::sycl::handler& h){ - h.copy(host_ptr,sycl_ptr, saw::ct_multiply::value); - }).wait(); - return saw::make_void(); - } - - static saw::error_or malloc_on_device(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ - (void) host_data; - (void) sycl_data; - (void) q; - return saw::make_void(); - } -}; - -template -struct sycl_copy_helper, Encode> final { - using Schema = sch::Chunk; - - static saw::error_or copy_to_host(saw::data>& sycl_data, saw::data& 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) == sizeof(std::decay_t), "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 copy_to_device(saw::data& host_data, saw::data>& 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) == sizeof(std::decay_t), "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 -struct sycl_copy_helper, Encode> final { - using Schema = sch::Array; - - static saw::error_or copy_to_host(saw::data>& sycl_data, saw::data& 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) == sizeof(std::decay_t), "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 copy_to_device(saw::data& host_data, saw::data>& 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) == sizeof(std::decay_t), "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 malloc_on_device(saw::data& host_data, saw::data>& sycl_data, sycl::queue& q){ - sycl_data = {host_data.meta(),q}; - return saw::make_void(); - } -}; - - -template -struct make_view_helper; - -template -struct make_view_helper, encode::Sycl> { -public: -static saw::error_or apply( - saw::data, encode::Sycl>& dat, - saw::data>, encode::Sycl>& dat_view -){ - dat_view = {dat}; - return saw::make_void(); -} -}; - -template -struct make_view_helper, encode::Sycl> { -public: - static saw::error_or apply( - saw::data, encode::Sycl>& dat, - saw::data>, encode::Sycl>& dat_view - ){ - dat_view = {dat}; - return saw::make_void(); - } -}; - -template -struct make_view_helper, encode::Sycl> { -public: - static saw::error_or apply( - saw::data, encode::Sycl>& dat, - saw::data>, encode::Sycl>& dat_view - ){ - dat_view = {dat}; - return saw::make_void(); - } -}; - -template -struct make_view_helper...>, encode::Sycl> { -private: -template -static saw::error_or 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; - - auto eov = make_view_helper>::apply(dat.template get(),dat_view.template get()); - if(eov.is_error()){ - return eov; - } - - return apply_i(dat,dat_view); - } - - return saw::make_void(); - } -public: -static saw::error_or apply( - saw::data...>, encode::Sycl>& dat, - saw::data...>>, encode::Sycl>& dat_view -){ - return apply_i<0u>(dat,dat_view); -} -}; - -template -struct make_view_helper, encode::Sycl> { -private: - template - static saw::error_or 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; - - auto eov = make_view_helper>::apply(dat.template get(), dat_view.template get()); - if(eov.is_error()){ - return eov; - } - - return apply_i(dat,dat_view); - } - - return saw::make_void(); - } -public: - static saw::error_or apply( - saw::data, encode::Sycl>& dat, - saw::data>, encode::Sycl> dat_view - ){ - return apply_i<0u>(dat,dat_view); - } -}; -} - -// Ptr => Ptr>> => Ptr -template -auto make_view(saw::data& dat){ - saw::data,Encode> dat_view; - auto eov = impl::make_view_helper::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 - saw::error_or copy_to_device(saw::data& host_data, saw::data>& sycl_data){ - return impl::sycl_copy_helper::copy_to_device(host_data, sycl_data, q_); - } - - template - saw::error_or copy_to_host(saw::data>& sycl_data, saw::data& host_data){ - return impl::sycl_copy_helper::copy_to_host(sycl_data, host_data, q_); - } - - template - saw::error_or malloc_on_device( - saw::data& host_data, - saw::data>& sycl_data - ){ - auto eov = impl::sycl_copy_helper::malloc_on_device(host_data, sycl_data, q_); - q_.wait(); - return eov; - } - - auto& get_handle(){ - return q_; - } -}; -} -} -- cgit v1.2.3