From 0a8dd2541e20f59812db21e8bad069b50cf8ebaf Mon Sep 17 00:00:00 2001
From: "Claudius \"keldu\" Holeksa" <mail@keldu.de>
Date: Fri, 23 Jan 2026 12:04:35 +0100
Subject: Preparing for SYCL accesors

---
 default.nix                                 |   2 +-
 examples/poiseulle_particles_2d_gpu/sim.cpp |  63 +++++++--
 lib/core/c++/chunk.hpp                      |  14 +-
 lib/core/c++/flatten.hpp                    |   2 -
 lib/sycl/c++/data.hpp                       | 205 +++++++++++++++++++++++-----
 lib/sycl/tests/data.cpp                     |  17 ++-
 6 files changed, 252 insertions(+), 51 deletions(-)

diff --git a/default.nix b/default.nix
index a8ec430..fd36323 100644
--- a/default.nix
+++ b/default.nix
@@ -39,7 +39,7 @@ let
 
 	forstio = (import ((builtins.fetchTarball {
 		url = "https://git.keldu.de/forstio-forstio/snapshot/master.tar.gz";
-		sha256 = "sha256:09xxs9dlnd52k1082lb5p034ik4jfg9zd5zc1x5da2clqhdb6s4m";
+		sha256 = "sha256:0078544k33h6ihhzz1150xs9zcywdmb6y8p16038v488kblvh3nc";
 	}) + "/default.nix"){
 		inherit stdenv;
 		inherit clang-tools;
diff --git a/examples/poiseulle_particles_2d_gpu/sim.cpp b/examples/poiseulle_particles_2d_gpu/sim.cpp
index be6b4f0..9646b0a 100644
--- a/examples/poiseulle_particles_2d_gpu/sim.cpp
+++ b/examples/poiseulle_particles_2d_gpu/sim.cpp
@@ -7,8 +7,8 @@
 namespace kel {
 namespace lbm {
 
-constexpr uint64_t dim_x = 16u;
-constexpr uint64_t dim_y = 4u;
+constexpr uint64_t dim_x = 256u;
+constexpr uint64_t dim_y = 16u;
 
 namespace sch {
 using namespace saw::schema;
@@ -31,29 +31,29 @@ using DfChunk = Chunk<FixedArray<T,Desc::Q>, 1u, dim_x, dim_y>;
 template<typename T, typename Desc>
 using ChunkStruct = Struct<
 	Member<InfoChunk, "info">
+	//,Member<DfChunk<T,Desc>, "dfs">
 >;
-
 }
 
 template<typename T, typename Desc>
 saw::error_or<void> setup_initial_conditions(saw::data<sch::ChunkStruct<T,Desc>>& fields){
-	auto& info_f = fields.get<"info">();
+	auto& info_f = fields.template get<"info">();
 	// Set everything as walls
 	iterator<Desc::D>::apply(
 		[&](auto& index){
 			info_f.at(index).set(2u);
 		},
-		{{0u,0u}},
-		{{dim_x, dim_y}},
-		{{0u,0u}}
+		{},
+		info_f.get_dims(),
+		{}
 	);
 	// Fluid
 	iterator<Desc::D>::apply(
 		[&](auto& index){
 			info_f.at(index).set(1u);
 		},
-		{{0u,0u}},
-		{{dim_x, dim_y}},
+		{},
+		info_f.get_dims(),
 		{{1u,1u}}
 	);
 	
@@ -76,12 +76,30 @@ saw::error_or<void> setup_initial_conditions(saw::data<sch::ChunkStruct<T,Desc>>
 		{{dim_x, dim_y}},
 		{{0u,1u}}
 	);
+/*
+	//
+	auto& df_f = fields.template get<"dfs">();
+	saw::data<T> rho{1};
+	saw::data<sch::FixedArray<T,Desc::D>> vel{};
+	auto eq = equilibrium<T,Desc>(rho,vel);
+
+	iterator<Desc::D>::apply(
+		[&](auto& index){
+			auto& df = df_f.at(index);
+
+			df = eq;
+		},
+		{},// 0-index
+		df_f.get_dims()
+	);
+*/
 
 	return saw::make_void();
 }
 
 template<typename T, typename Desc>
-saw::error_or<void> step(){
+saw::error_or<void> step(saw::data<sch::ChunkStruct<T,Desc>,encode::Sycl<saw::encode::Native>>& fields, saw::data<sch::UInt64> t_i, device& dev){
+	auto& q = dev.get_handle();
 
 	return saw::make_void();
 }
@@ -111,7 +129,9 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 	// saw::data<sch::FixedArray<sch::UInt64,Desc::D>> meta{{dim_x,dim_y}};
 	saw::data<sch::ChunkStruct<T,Desc>> lbm_data{};
 
-	acpp::sycl::queue sycl_q;
+	device dev;
+	auto& sycl_q = dev.get_handle();
+
 	sycl_q.wait();
 	{
 		auto eov = setup_initial_conditions<T,Desc>(lbm_data);
@@ -120,6 +140,27 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 		}
 	}
 
+	saw::data<sch::ChunkStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_data;
+	{
+		auto eov = dev.allocate_on_device(lbm_sycl_data);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+	{
+		auto eov = dev.copy_to_device(lbm_data,lbm_sycl_data);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+
+	for(saw::data<sch::UInt64> i{0u}; i < saw::data<sch::UInt64>{32ul}; ++i){
+		auto eov = step<T,Desc>(lbm_sycl_data,i,dev);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+
 	/*
 	iterator<Desc::D>::apply(
 		[&](auto& index){
diff --git a/lib/core/c++/chunk.hpp b/lib/core/c++/chunk.hpp
index ad7de34..6b028ba 100644
--- a/lib/core/c++/chunk.hpp
+++ b/lib/core/c++/chunk.hpp
@@ -47,7 +47,7 @@ public:
 	data<ValueSchema, Encode>& ghost_at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index){
 		return values_.at(index);
 	}
-	
+
 	const data<ValueSchema, Encode>& ghost_at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index) const {
 		return values_.at(index);
 	}
@@ -55,7 +55,7 @@ public:
 	static constexpr auto get_ghost_dims() {
 		return data<InnerSchema,Encode>::get_dims();
 	}
-	
+
 	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){
@@ -63,7 +63,7 @@ public:
 		}
 		return values_.at(ind);
 	}
-	
+
 	const 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){
@@ -75,6 +75,14 @@ public:
 	static constexpr auto get_dims(){
 		return data<schema::FixedArray<schema::UInt64, sizeof...(Sides)>,Encode>{{Sides...}};
 	}
+
+	auto flat_data(){
+		return values_.flat_data();
+	}
+
+	static constexpr auto flat_size() {
+		return data<InnerSchema,Encode>::flat_size();
+	}
 };
 
 template<typename Sch, uint64_t Ghost, uint64_t... Sides>
diff --git a/lib/core/c++/flatten.hpp b/lib/core/c++/flatten.hpp
index 164bb77..1609589 100644
--- a/lib/core/c++/flatten.hpp
+++ b/lib/core/c++/flatten.hpp
@@ -20,7 +20,6 @@ public:
 
 		return 1u;
 	}
-
 private:
 	/// 2,3,4 => 2,6,24
 	/// i + j * 2 + k * 3*2
@@ -35,7 +34,6 @@ private:
 public:
 	static saw::data<T> apply(const saw::data<sch::FixedArray<T,D>>& index, const saw::data<sch::FixedArray<T,D>>& meta){
 		saw::data<T> flat_ind{0u};
-
 		apply_i<0u>(flat_ind, index, meta);
 		return flat_ind;
 	}
diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp
index bb8b4bf..d9976dd 100644
--- a/lib/sycl/c++/data.hpp
+++ b/lib/sycl/c++/data.hpp
@@ -11,19 +11,115 @@ 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>* values_;
+
+	SAW_FORBID_COPY(data);
+public:
+	data(acpp::sycl::queue& q__):
+		q_{&q__},
+		values_{nullptr}
+	{
+		values_ = acpp::sycl::malloc<data<Sch>>(ct_multiply<uint64_t,Dims...>::value,q_);
+	}
+
+	~data(){
+		if(not values_){
+			return;
+		}
+
+		acpp::sycl::free(values_,q_);
+	}
+
+	static constexpr data<schema::FixedArray<schema::UInt64, sizeof...(Dims)>> get_dims() {
+		return {std::array<uint64_t, sizeof...(Dims)>{Dims...}};
+	}
+
+	data<Sch>& at(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::Native>* flat_data() {
+		return values_;
+	}
+};
+
+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(acpp::sycl::queue& q__):
+		values_{q__}
+	{}
+
+	data<ValueSchema, kel::lbm::encode::Sycl<Encode>>& ghost_at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index){
+		return values_.at(index);
+	}
+
+	const data<ValueSchema, kel::lbm::encode::Sycl<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();
+	}
+
+	data<ValueSchema, kel::lbm::encode::Sycl<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);
+	}
+
+	const data<ValueSchema, kel::lbm::encode::Sycl<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)>,kel::lbm::encode::Sycl<Encode>>{{Sides...}};
+	}
+
+	auto flat_data(){
+		return values_.flat_data();
+	}
+
+	static constexpr auto flat_size() {
+		return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::flat_size();
+	}
+};
+
 template<uint64_t Ghost, uint64_t... Meta, typename... Sch, string_literal... Keys, typename Encode>
 class data<schema::Struct<schema::Member<kel::lbm::sch::Chunk<Sch,Ghost,Meta...>, Keys>...>, kel::lbm::encode::Sycl<Encode> > final {
 public:
 	static constexpr data<schema::FixedArray<schema::UInt64,sizeof...(Meta)>> meta = {{Meta...}};
-	using StorageT = std::tuple<data<Sch,Encode>*...>;
+	using StorageT = std::tuple<data<kel::lbm::sch::Chunk<Sch,Ghost,Meta...>,kel::lbm::encode::Sycl<Encode>>...>;
+	using Schema = schema::Struct<schema::Member<kel::lbm::sch::Chunk<Sch,Ghost,Meta...>, Keys>...>;
 private:
 
 	/**
@@ -31,43 +127,18 @@ private:
 	 * Do it here by specializing.
 	 */
 	StorageT members_;
-	kel::lbm::sycl::queue* q_;
 public:
-	data():
-		members_{},
-		q_{nullptr}
-	{}
-
-	data(StorageT members__, kel::lbm::sycl::queue& q__):
-		members_{members__},
-		q_{&q__}
+	data(acpp::sycl::queue& q__):
+		members_{{data<kel::lbm::sch::Chunk<Sch,Ghost,Meta...>,kel::lbm::encode::Sycl<Encode>>{q__}...}}
 	{}
 
-	~data(){
-		SAW_ASSERT(q_){
-			return;
-		}
-		std::visit([this](auto arg){
-			if(not arg){
-				return;
-			}
-			acpp::sycl::free(arg,*q_);
-			arg = nullptr;
-		},members_);
-	}
-
-	template<saw::string_literal K>
-	auto* get_ptr(){
-		return std::get<parameter_key_pack_index<K, Keys...>::value>(members_);
-	}
-
 	template<saw::string_literal K>
 	auto& get(){
-		auto ptr = get_ptr<K>();
-		SAW_ASSERT(ptr);
-		return *ptr;
+		return std::get<parameter_key_pack_index<K, Keys...>::value>(members_);
 	}
 };
+
+
 }
 
 namespace kel {
@@ -81,7 +152,7 @@ struct sycl_malloc_struct_helper<sch::Struct<Members...>, Encode> final {
 	using Schema = sch::Struct<Members...>;
 
 	template<uint64_t i>
-	static saw::error_or<void> allocate_on_device_member(typename saw::data<typename saw::parameter_pack_type<i,Members...>::type::ValueType,encode::Sycl<Encode>>::StorageT& storage, sycl::queue& q){
+	static saw::error_or<void> allocate_on_device_member(typename saw::data<Schema,encode::Sycl<Encode>>::StorageT& storage, sycl::queue& q){
 		if constexpr (i < sizeof...(Members)){
 			using M = typename saw::parameter_pack_type<i,Members...>::type;
 			auto& ptr = std::get<i>(storage);
@@ -103,6 +174,60 @@ struct sycl_malloc_struct_helper<sch::Struct<Members...>, Encode> final {
 		return eov;
 	}
 };
+
+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<i>();
+			auto& sycl_member_data = sycl_data.template get<i>();
+
+			auto host_ptr = host_member_data.flat_data();
+			auto sycl_ptr = sycl_member_data.flat_data();
+
+			q.memcpy(host_ptr, sycl_ptr, sizeof(std::decay_t<decltype(host_ptr)>) * host_member_data.flat_size() );
+
+			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 host_ptr = host_member_data.flat_data();
+			auto sycl_ptr = sycl_member_data.flat_data();
+
+			q.memcpy(sycl_ptr, host_ptr, sizeof(std::decay_t<decltype(host_ptr)>) * host_member_data.flat_size() );
+
+			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_data, saw::data<Schema,encode::Sycl<Encode>>& host_data, sycl::queue& q){
+		return copy_to_host_member<0u>(sycl_data, host_data, q);
+	}
+};
 }
 class device final {
 private:
@@ -120,7 +245,21 @@ public:
 		if(eov.is_error()){
 			return eov;
 		}
-		sycl_data.set_queue(q_);
+		return saw::make_void();
+	}
+
+	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_);
+	}
+
+	auto& get_handle(){
+		return q_;
 	}
 };
 }
diff --git a/lib/sycl/tests/data.cpp b/lib/sycl/tests/data.cpp
index e57a1e3..3073a22 100644
--- a/lib/sycl/tests/data.cpp
+++ b/lib/sycl/tests/data.cpp
@@ -1,6 +1,21 @@
 #include <forstio/test/suite.hpp>
 
+#include "../c++/lbm.hpp"
 
-SAW_TEST("Sycl Dummy"){
+SAW_TEST("Sycl Data Compilation"){
+	acpp::sycl::queue q;
+	saw::data<
+		saw::schema::Struct<
+			saw::schema::Member<
+				kel::lbm::sch::Chunk<saw::schema::UInt8,1u,1u,1u>,
+				"test"
+			>
+		>,
+		kel::lbm::encode::Sycl<saw::encode::Native>
+	> dat{q};
 
+	auto& test_f = dat.template get<"test">();
+
+	// test_f.at({}).set(1);
+	// SAW_EXPECT(test_f.at({}).get() == 1, "Value check failed");
 }
-- 
cgit v1.2.3