From 20e46825069b3974d5cc883163d7e37b4836b2af Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Mon, 13 Apr 2026 19:26:57 +0200 Subject: Progress? --- lib/sycl/c++/data.hpp | 300 +++++++++++++++++++++++++++++++++++++++++++++++++- 1 file changed, 294 insertions(+), 6 deletions(-) (limited to 'lib/sycl/c++') diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp index 50adb4f..661c313 100644 --- a/lib/sycl/c++/data.hpp +++ b/lib/sycl/c++/data.hpp @@ -29,13 +29,15 @@ private: SAW_FORBID_COPY(data); SAW_FORBID_MOVE(data); + public: data(acpp::sycl::queue& q__): q_{&q__}, values_{nullptr} { + SAW_ASSERT(q_); values_ = acpp::sycl::malloc_device>(ct_multiply::value,*q_); - SAW_ASSERT(values_ and q_); + SAW_ASSERT(values_); } ~data(){ @@ -63,6 +65,7 @@ public: constexpr data* flat_data() { return values_; } + }; template @@ -81,7 +84,7 @@ public: {} data(data, kel::lbm::encode::Sycl>& values__): - values_{&values__.at({})} + values_{values__.flat_data()} {} data(data* values__): @@ -105,6 +108,113 @@ public: } }; +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, Encode>& 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(){ + 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 data* flat_data() { + 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_; + + SAW_FORBID_COPY(data); + SAW_FORBID_MOVE(data); +public: + 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() { + 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: @@ -223,6 +333,64 @@ public: } }; +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(acpp::sycl::queue& q__): + data{q__, std::make_index_sequence{}} + { + q__.wait(); + } + + template + auto& get(){ + return std::get(members_); + } + + template + const 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: @@ -235,6 +403,9 @@ 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)...} @@ -246,6 +417,16 @@ public: q__.wait(); } + template + auto& get(){ + return std::get(members_); + } + + template + const auto& get() const { + return std::get(members_); + } + template auto& get(){ return std::get::value>(members_); @@ -347,6 +528,60 @@ struct sycl_copy_helper, Encode> final { } }; +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)){ + auto& host_member_data = host_data.template get(); + auto& sycl_member_data = sycl_data.template get(); + + auto host_ptr = host_member_data.flat_data(); + auto sycl_ptr = sycl_member_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_member_data.flat_size().get()); + }).wait(); + + 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)){ + auto& host_member_data = host_data.template get(); + auto& sycl_member_data = sycl_data.template get(); + + 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(); + + 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 struct sycl_copy_helper, Encode> final { using Schema = sch::FixedArray; @@ -376,12 +611,41 @@ struct sycl_copy_helper, Encode> final { } }; +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"); + + q.submit([&](acpp::sycl::handler& h){ + h.copy(sycl_ptr,host_ptr, host_data.flat_size()); + }).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, host_data.flat_size()); + }).wait(); + return saw::make_void(); + } +}; + template -struct make_chunk_struct_view_helper; +struct make_view_helper; template -struct make_chunk_struct_view_helper...>, encode::Sycl> { +struct make_view_helper...>, encode::Sycl> { private: template static void apply_i( @@ -405,12 +669,36 @@ public: return str_view; } }; + +template +struct make_view_helper, encode::Sycl> { +private: + template + static void apply_i( + saw::data, encode::Sycl>& dat, + saw::data>, encode::Sycl>& dat_view) + { + if constexpr (i < sizeof...(Sch)){ + dat_view.template get() = {dat.template get().flat_data()}; + + apply_i(dat,dat_view); + } + } +public: + static auto apply(saw::data, encode::Sycl>& dat){ + saw::data>, encode::Sycl> str_view; + + apply_i<0u>(dat,str_view); + + return str_view; + } +}; } // Ptr => Ptr>> => Ptr template -auto make_chunk_struct_view(saw::data& dat){ - return impl::make_chunk_struct_view_helper::apply(dat); +auto make_view(saw::data& dat){ + return impl::make_view_helper::apply(dat); } class device final { -- cgit v1.2.3