summaryrefslogtreecommitdiff
path: root/lib
diff options
context:
space:
mode:
Diffstat (limited to 'lib')
-rw-r--r--lib/core/c++/abstract/data.hpp26
-rw-r--r--lib/core/c++/particle.hpp5
-rw-r--r--lib/core/c++/particle/particle.hpp35
-rw-r--r--lib/core/c++/write_vtk.hpp4
-rw-r--r--lib/sycl/c++/data.hpp596
-rw-r--r--lib/sycl/tests/data.cpp35
6 files changed, 632 insertions, 69 deletions
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<typename StorageT, typename InterfaceT>
struct MixedPrecision {
+ using Meta = Void;
using StorageType = StorageT;
using InterfaceType = InterfaceT;
};
template<typename PrimType, uint64_t N>
struct Primitive {
- using PrimitiveType = PrimType;
+ using Meta = Void;
+ using Type = PrimType;
static constexpr uint64_t Bytes = N;
};
template<typename T, uint64_t... Dims>
-struct Array {
- using InnerType = T;
+struct FixedArray {
+ using Meta = Void;
+ using Inner = T;
static constexpr std::array<uint64_t,sizeof...(Dims)> Dimensions{Dims...};
};
+template<typename T, uint64_t Dims>
+struct Array {
+ using Meta = FixedArray<UInt64,Dims>;
+ using Inner = T;
+ static constexpr std::array<uint64_t,sizeof...(Dims)> Dimensions{Dims};
+};
+
+template<typename... T>
+struct Tuple {
+};
}
+
+template<typename Sch>
+struct schema {
+ using Type = Sch;
+};
}
diff --git a/lib/core/c++/particle.hpp b/lib/core/c++/particle.hpp
index b098ecc..691a74b 100644
--- a/lib/core/c++/particle.hpp
+++ b/lib/core/c++/particle.hpp
@@ -11,7 +11,12 @@ struct Particle {};
template<typename T, typename Descriptor, typename Encode>
class component<T, Descriptor, cmpt::Particle, Encode> {
+private:
+ saw::data<sch::Scalar<T>> dt_;
public:
+ component(saw::data<sch::Scalar<T>> dt__):
+ dt_{dt__}
+ {}
template<typename ParticleSchema, typename MacroFieldSchema>
void apply(const saw::data<ParticleSchema, Encode>& particles, const saw::data<MacroFieldSchema,Encode>& macros, saw::data<sch::UInt64> index, saw::data<sch::UInt64> time_step) const {
diff --git a/lib/core/c++/particle/particle.hpp b/lib/core/c++/particle/particle.hpp
index c0d115f..7af43dc 100644
--- a/lib/core/c++/particle/particle.hpp
+++ b/lib/core/c++/particle/particle.hpp
@@ -40,9 +40,8 @@ using ParticleRigidBody = Struct<
Member<typename impl::rotation_type_helper<T,D>::Schema, "angular_acceleration">
>;
-template<typename T>
+template<typename T, typename saw::native_data_type<T>::type radius = 1.0f>
using ParticleCollisionSpheroid = Struct<
- Member<Scalar<T>, "radius">
>;
template<typename T, uint64_t D>
@@ -52,42 +51,26 @@ using Particle = Struct<
// Member<Array<Float64,D>, "mask">,
>;
-template<typename T, uint64_t D, typename CollisionType = ParticleCollisionSpheroid<T>>
+template<typename T, uint64_t D, typename CollisionType = ParticleCollisionSpheroid<T,2.0f>>
using ParticleGroup = Struct<
Member<Array<T,D>, "mask">,
- Member<CollisionType, "collision">,
- Member<Scalar<T>, "mass">,
+ Member<FixedArray<Scalar<T>,1u>, "density">,
Member<Array<Particle<T,D>>, "particles">
>;
}
-template<typename T, uint64_t D>
-saw::data<sch::ParticleGroup<T,D, sch::ParticleCollisionSpheroid<T>>> create_spheroid_particle_group(
- saw::data<sch::Scalar<T>> rad_p,
+template<typename T, uint64_t D, typename saw::native_data_type<T>::type radius>
+saw::data<sch::ParticleGroup<T,D, sch::ParticleCollisionSpheroid<T,radius>>> create_spheroid_particle_group(
saw::data<sch::Scalar<T>> density_p,
const saw::data<sch::UInt64>& mask_resolution
){
- saw::data<sch::ParticleGroup<T,D,sch::ParticleCollisionSpheroid<T>>> part;
+ saw::data<sch::ParticleGroup<T,D,sch::ParticleCollisionSpheroid<T,2.0f>>> part;
auto& mask = part.template get<"mask">();
- auto& collision = part.template get<"collision">();
- auto& mass = part.template get<"mass">();
+ auto& density = part.template get<"density">().at({{0u}});
- if constexpr ( D == 1u ){
- saw::data<sch::Scalar<T>> c;
- c.at({}).set(2.0);
- mass = rad_p * c * density_p;
- } else if constexpr ( D == 2u){
- saw::data<sch::Scalar<T>> pi;
- pi.at({}).set(3.14159);
- mass = rad_p * rad_p * pi * density_p;
- } else if constexpr ( D == 3u ){
- saw::data<sch::Scalar<T>> c;
- c.at({}).set(3.14159 * 4.0 / 3.0);
- mass = rad_p * rad_p * rad_p * c * density_p;
- } else {
- static_assert(D == 0u or D > 3u, "Dimensions only supported for Dim 1,2 & 3.");
- }
+ static_assert(D >= 1u and D <= 3u, "Dimensions only supported for Dim 1,2 & 3.");
+ density = density_p;
saw::data<sch::FixedArray<sch::UInt64,D>> mask_dims;
for(uint64_t i = 0u; i < D; ++i){
diff --git a/lib/core/c++/write_vtk.hpp b/lib/core/c++/write_vtk.hpp
index f925361..e852172 100644
--- a/lib/core/c++/write_vtk.hpp
+++ b/lib/core/c++/write_vtk.hpp
@@ -136,7 +136,7 @@ struct lbm_vtk_writer<sch::Vector<T,D>> {
template<typename T>
struct lbm_vtk_writer<sch::Scalar<T>> {
static saw::error_or<void> apply_header(std::ostream& vtk_file, std::string_view name){
- vtk_file<<"SCALARS "<<name<<" float\n";
+ vtk_file<<"SCALARS "<<name<<" float 1\n";
vtk_file<<"LOOKUP_TABLE default\n";
return saw::make_void();
}
@@ -206,7 +206,7 @@ struct lbm_vtk_writer<sch::Struct<sch::Member<sch::Chunk<MemberT,Ghost,Dims...>,
if constexpr (sizeof...(MemberT) > 0u){
// POINT DATA
{
- vtk_file << "POINT_DATA " << pd_size.get() <<"\n";
+ vtk_file << "\nPOINT_DATA " << pd_size.get() <<"\n";
}
// HEADER TO BODY
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_;
}
diff --git a/lib/sycl/tests/data.cpp b/lib/sycl/tests/data.cpp
index 3073a22..4321a0d 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<FixedArray<UInt64,2u,2u>, "foo">,
+ Member<Array<Float32>, "bar">,
+ Member<
+ Array<
+ Struct<
+ Member<FixedArray<Float32,2u>,"pos">
+ >
+ >,
+ "baz"
+ >
+>;
+}
+
SAW_TEST("Sycl Data Compilation"){
acpp::sycl::queue q;
saw::data<
@@ -19,3 +37,20 @@ SAW_TEST("Sycl Data Compilation"){
// test_f.at({}).set(1);
// SAW_EXPECT(test_f.at({}).get() == 1, "Value check failed");
}
+
+SAW_TEST("Sycl Data Array of Struct"){
+ acpp::sycl::queue q;
+
+ saw::data<sch::Array<sch::Float64>, kel::lbm::encode::Sycl<saw::encode::Native>> a{{{2u}},q};
+}
+
+/*
+SAW_TEST("Sycl Data Compilation for Particle Similacrum"){
+ acpp::sycl::queue q;
+
+ saw::data<
+ sch::TestObjSchema
+ > a;
+}
+*/
+}
generated by cgit v1.2.3 (git 2.25.1) at 2026-05-11 13:26:00 +0000