From c61ba8f8eb86f66915a54551fcc39dfbeab1fad9 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Tue, 14 Apr 2026 21:09:42 +0200 Subject: Working on meta schema instantiation, because I'm stupid --- lib/core/c++/abstract/data.hpp | 26 ++++++++-- lib/sycl/c++/data.hpp | 112 +++++++++++++++++++++++++++++++++++++++-- lib/sycl/tests/data.cpp | 27 ++++++++++ 3 files changed, 157 insertions(+), 8 deletions(-) (limited to 'lib') diff --git a/lib/core/c++/abstract/data.hpp b/lib/core/c++/abstract/data.hpp index e8f1757..0075718 100644 --- a/lib/core/c++/abstract/data.hpp +++ b/lib/core/c++/abstract/data.hpp @@ -4,28 +4,48 @@ namespace kel { namespace sch { +struct Void {}; + struct UnsignedInteger {}; struct SignedInteger {}; struct FloatingPoint {}; template struct MixedPrecision { + using Meta = Void; using StorageType = StorageT; using InterfaceType = InterfaceT; }; template struct Primitive { - using PrimitiveType = PrimType; + using Meta = Void; + using Type = PrimType; static constexpr uint64_t Bytes = N; }; template -struct Array { - using InnerType = T; +struct FixedArray { + using Meta = Void; + using Inner = T; static constexpr std::array Dimensions{Dims...}; }; +template +struct Array { + using Meta = FixedArray; + using Inner = T; + static constexpr std::array Dimensions{Dims}; +}; + +template +struct Tuple { +}; } + +template +struct schema { + using Type = Sch; +}; } diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp index 4e5129b..3ac51e0 100644 --- a/lib/sycl/c++/data.hpp +++ b/lib/sycl/c++/data.hpp @@ -29,8 +29,17 @@ private: 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} @@ -65,7 +74,6 @@ public: constexpr data* flat_data() const { return values_; } - }; template @@ -122,6 +130,20 @@ private: SAW_FORBID_COPY(data); SAW_FORBID_MOVE(data); public: + data(const data::MetaSchema>& meta__, acpp::sycl::queue& q__): + q_{&q__}, + values_{nullptr} + { + SAW_ASSERT(q_); + /// TODO use meta + data m{1u}; + for(uint64_t i = 0u; i < Dims; ++i){ + m = m * meta__.at({i}); + } + values_ = acpp::sycl::malloc_device>(m.get(),*q_); + SAW_ASSERT(values_); + } + data(acpp::sycl::queue& q__): values_{nullptr}, meta_{}, @@ -198,8 +220,7 @@ public: data(const data, kel::lbm::encode::Sycl>& values__): values_{values__.flat_data()}, meta_{values__.meta()} - { - } + {} constexpr data, Encode> meta() const { return meta_; @@ -238,6 +259,10 @@ private: data> values_; public: + data(const data::MetaSchema>& meta__, acpp::sycl::queue& q__): + values_{meta__,q__} + {} + data(acpp::sycl::queue& q__): values_{q__} {} @@ -366,6 +391,11 @@ private: 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{}} { @@ -546,6 +576,29 @@ struct sycl_copy_helper, Encode> final { 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 @@ -594,9 +647,36 @@ struct sycl_copy_helper, Encode> final { } - static saw::error_or copy_to_host(saw::data>& sycl_data, saw::data& host_data, sycl::queue& q){ + 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 @@ -626,6 +706,13 @@ struct sycl_copy_helper, Encode> final { }).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 @@ -684,6 +771,11 @@ struct sycl_copy_helper, Encode> final { }).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(); + } }; @@ -817,6 +909,16 @@ public: 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_; } diff --git a/lib/sycl/tests/data.cpp b/lib/sycl/tests/data.cpp index 3073a22..6b17622 100644 --- a/lib/sycl/tests/data.cpp +++ b/lib/sycl/tests/data.cpp @@ -2,6 +2,24 @@ #include "../c++/lbm.hpp" +namespace { + +namespace sch { +using namespace kel::lbm::sch; +using TestObjSchema = Tuple< + Member, "foo">, + Member, "bar">, + Member< + Array< + Struct< + Member,"pos"> + > + >, + "baz" + > +>; +} + SAW_TEST("Sycl Data Compilation"){ acpp::sycl::queue q; saw::data< @@ -19,3 +37,12 @@ SAW_TEST("Sycl Data Compilation"){ // test_f.at({}).set(1); // SAW_EXPECT(test_f.at({}).get() == 1, "Value check failed"); } + +SAW_TEST("Sycl Data Compilation for Particle Similacrum"){ + acpp::sycl::queue q; + + saw::data< + sch::TestObjSchema + > a; +} +} -- cgit v1.2.3