13 files changed, 295 insertions, 45 deletions
diff --git a/examples/cavity_2d_gpu/cavity_2d_gpu.cpp b/examples/cavity_2d_gpu/cavity_2d_gpu.cpp
index 507c2d6..0fc490a 100644
--- a/examples/cavity_2d_gpu/cavity_2d_gpu.cpp
+++ b/examples/cavity_2d_gpu/cavity_2d_gpu.cpp
@@ -319,7 +319,7 @@ int main(){
 	auto& lbm_dir = eo_lbm_dir.get_value();
 	auto out_dir = lbm_dir / "cavity_gpu_2d";
 
-	acpp::sycl::queue sycl_q{acpp::sycl::default_selector_v, acpp::sycl::property::queue::in_order{}};
+	sycl::queue sycl_q{sycl::default_selector_v, sycl::property::queue::in_order{}};
 
 	constexpr size_t num_cells = dim_x * dim_y;
 
diff --git a/examples/heterogeneous_computing/SConscript b/examples/heterogeneous_computing/SConscript
index 1e52d88..226185b 100644
--- a/examples/heterogeneous_computing/SConscript
+++ b/examples/heterogeneous_computing/SConscript
@@ -22,7 +22,7 @@ env.headers += examples_env.headers;
 
 # Cavity2D
 examples_objects = [];
-examples_env.add_source_files(examples_objects, ['heterogeneous_computing.cpp'], shared=False);
+examples_env.add_source_files(examples_objects, ['sim.cpp'], shared=False);
 examples_env.heterogeneous_computing = examples_env.Program('#bin/heterogeneous_computing', [examples_objects]);
 
 # Set Alias
diff --git a/examples/heterogeneous_computing/heterogeneous_computing.cpp b/examples/heterogeneous_computing/heterogeneous_computing.cpp
deleted file mode 100644
index 8a79354..0000000
--- a/examples/heterogeneous_computing/heterogeneous_computing.cpp
+++ /dev/null
@@ -1,36 +0,0 @@
-#include <forstio/error.hpp>
-#include <iostream>
-
-
-namespace kel {
-namespace lbm {
-namespace sch {
-using namespace saw::schema;
-using KelConfig = Struct<
-	Member<String,"resolution">
->;
-}
-}
-
-saw::error_or<void> lbm_main(int argc, char** argv){
-	return saw::make_void();
-}
-}
-
-int main(int argc, char** argv){
-	auto eov = kel::lbm_main(argc, argv);
-	if(eov.is_error()){
-		auto& err = eov.get_error();
-		auto err_msg = err.get_message();
-		std::cerr<<"[Error]: "<<err.get_category();
-
-		if(not err_msg.empty()){
-			std::cerr<<" - "<<err_msg;
-		}
-		std::cerr<<std::endl;
-
-		return err.get_id();
-	}
-
-	return 0;
-}
diff --git a/examples/heterogeneous_computing/sim.cpp b/examples/heterogeneous_computing/sim.cpp
new file mode 100644
index 0000000..d074ad8
--- /dev/null
+++ b/examples/heterogeneous_computing/sim.cpp
@@ -0,0 +1,96 @@
+#include <forstio/error.hpp>
+#include <iostream>
+
+#include <kel/lbm/lbm.hpp>
+
+namespace kel {
+namespace lbm {
+namespace sch {
+using namespace saw::schema;
+
+/**
+ * struct lbm_data {
+ *   std::array<std::array<float,9>, 64u*64u> dfs;
+ *   std::array<uint8_t, 64u*64u> info;
+ * };
+ *
+ * which leads to the form
+ *
+ * template<uint64_t Dim, uint64_t Size>
+ * struct lbm_data {
+ *   std::array<std::array<float,9>, Size*Size> dfs;
+ *   std::array<uint8_t, Size*Size> info;
+ * };
+ *
+ * which transferred into sycl requires us to go to
+ *
+ * template<uint64_t Dim, uint64_t Size>
+ * struct lbm_sycl_data {
+ *   std::array<float,9>* dfs;
+ *   uint8_t* info;
+ * };
+ *
+ * in data form on host
+ *
+ * template<uint64_t Dim, uint64_t Size>
+ * using LbmData = Struct<
+ *  Member<FixedArray<FixedArray<Float32,9u>, Size*Size>, "dfs">,
+ *  Member<FixedArray<UInt8, Size*Size>, "info">
+ * >;
+ *
+ * If we specialize the encode::Sycl<Encoding> data type, then we get
+ * With a helper class we can copy single values back and forth and the whole block is guaranteed
+ * to be allocated. And this one can be dynamic while the host definition might be compile time.
+ *
+ * template<...>
+ * class data<LbmData<...>,encode::Sycl<Encoding>> final {
+ *   saw::data<sch::FixedArray<sch::UInt64,Dim>> meta;
+ *   saw::data<FixedArray<Float32u,9>>* dfs;
+ *   saw::data<UInt8>* info;
+ * };
+ */
+
+template<typename T, typename Desc>
+using CellStruct = Struct<
+	Member<Chunk<FixedArray<T,Desc::Q>,Desc::D, "dfs">,
+	Member<FixedArray<T,Desc::Q>, "dfs_old">,
+	Member<UInt8, "info">
+>;
+
+}
+template<typename T, typename Desc>
+saw::error_or<void> simulate(int argc, char** argv,
+	const saw::data<sch::FixedArray<sch::UInt64,Desc::D>>& meta
+{
+	constexpr auto cell_size = sizeof(saw::data<sch::CellStruct<T,Desc>>);
+
+	auto lbm_data = saw::heap<saw::data<sch::Chunk<CellStruct<>,Desc::D, 64u>>>{meta};
+
+	return saw::make_void();
+}
+
+}
+
+saw::error_or<void> lbm_main(int argc, char** argv){
+	using namespace lbm;
+	return simulate<sch::Float64,sch::D2Q9>(argc, argv);
+}
+}
+
+int main(int argc, char** argv){
+	auto eov = kel::lbm_main(argc, argv);
+	if(eov.is_error()){
+		auto& err = eov.get_error();
+		auto err_msg = err.get_message();
+		std::cerr<<"[Error]: "<<err.get_category();
+
+		if(not err_msg.empty()){
+			std::cerr<<" - "<<err_msg;
+		}
+		std::cerr<<std::endl;
+
+		return err.get_id();
+	}
+
+	return 0;
+}
diff --git a/examples/poiseulle_particles_2d_gpu/sim.cpp b/examples/poiseulle_particles_2d_gpu/sim.cpp
index 21ff253..8bc60c9 100644
--- a/examples/poiseulle_particles_2d_gpu/sim.cpp
+++ b/examples/poiseulle_particles_2d_gpu/sim.cpp
@@ -47,10 +47,18 @@ saw::error_or<void> kel_main(int argc, char** argv){
 	uint64_t x_d = 256u;
 	uint64_t y_d = 64u;
 	saw::data<sch::FixedArray<sch::UInt64,Desc::D>> meta{{x_d,y_d}};
+	saw::data<sch::Array<lbm::sch::CellStruct<T,Desc>,Desc::D>> lbm_data{meta};
 
 	acpp::sycl::queue sycl_q;
-
 	sycl_q.wait();
+	{
+		auto eov = setup_initial_conditions(lbm_data);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+
+	
 
 	return saw::make_void();
 }
@@ -71,4 +79,3 @@ int main(int argc, char** argv){
 	}
 	return 0;
 }
-}
diff --git a/lib/core/c++/chunk.hpp b/lib/core/c++/chunk.hpp
new file mode 100644
index 0000000..bfef358
--- /dev/null
+++ b/lib/core/c++/chunk.hpp
@@ -0,0 +1,38 @@
+#pragma once
+
+#include "common.hpp"
+
+namespace kel {
+namespace lbm {
+namespace sch {
+namespace impl {
+template<typename Sch, uint64_t Dim, uint64_t Side, typename... AddedSides>
+struct chunk_schema_type_helper {
+	using Schema = typename chunk_schema_type_helper<Sch,Dim-1u,Side,AddedSides...,Side>::Schema;
+};
+
+template<typename Sch, uint64_t Side, typename... AddedSides>
+struct chunk_schema_type_helper<Sch, 0u, Side, AddedSides...> {
+	using Schema = FixedArray<Sch,AddedSides...>;
+};
+}
+
+
+template<typename Schema, uint64_t Dim, uint64_t Side, uint64_t Ghost = 0u>
+struct Chunk {
+	using InnerSchema = typename impl::chunk_schema_type_helper<Sch, Dim, Side + (2u*Ghost)>::Schema;
+};
+
+template<typename ChunkSchema, uint64_t Dim>
+using SuperChunk = Array<ChunkSchema,Dim>;
+}
+}
+}
+
+namespace saw {
+template<typename Sch, uint64_t Dim, uint64_t Side, uint64_t Ghost, typename Encode>
+class data<kel::lbm::sch::Chunk<Sch,Dim,Side,Ghost>,Encode> final {
+private:
+public:
+};
+}
diff --git a/lib/core/c++/common.hpp b/lib/core/c++/common.hpp
new file mode 100644
index 0000000..5f7129f
--- /dev/null
+++ b/lib/core/c++/common.hpp
@@ -0,0 +1,11 @@
+#pragma once
+
+#include <forstio/codec/data.hpp>
+
+namespace kel {
+namespace lbm {
+namespace sch {
+using namespace saw::schema;
+}
+}
+}
diff --git a/lib/core/c++/descriptor.hpp b/lib/core/c++/descriptor.hpp
index c6938e3..9cc2591 100644
--- a/lib/core/c++/descriptor.hpp
+++ b/lib/core/c++/descriptor.hpp
@@ -15,6 +15,10 @@ struct Descriptor {
 	static constexpr uint64_t Q = QV;
 };
 
+using D2Q9 = Descriptor<2u,9u>;
+//using D2Q5 = Descriptor<2u,5u>;
+using D3Q27 = Descriptor<3u,27u>;
+
 template<typename Sch, typename Desc, uint64_t SC_V, uint64_t DC_V, uint64_t QC_V>
 struct Cell {
 	using Descriptor = Desc;
diff --git a/lib/core/c++/lbm.hpp b/lib/core/c++/lbm.hpp
index 473ca69..aff38e9 100644
--- a/lib/core/c++/lbm.hpp
+++ b/lib/core/c++/lbm.hpp
@@ -19,9 +19,6 @@
 #include <iostream>
 
 namespace kel {
-namespace sch {
-using namespace saw::schema;
-}
 namespace lbm {
 template<typename T, typename Desc>
 void print_lbm_meta(const converter<T>& conv, const saw::data<sch::SiKinematicViscosity<T>>& kin_vis_si){
diff --git a/lib/sycl/.nix/derivation.nix b/lib/sycl/.nix/derivation.nix
index 68fbab7..02032ba 100644
--- a/lib/sycl/.nix/derivation.nix
+++ b/lib/sycl/.nix/derivation.nix
@@ -5,7 +5,7 @@
 , pname
 , version
 , forstio
-, kel-lbm
+, kel
 , adaptive-cpp
 }:
 
@@ -27,7 +27,7 @@ stdenv.mkDerivation {
 		forstio.codec-json
 		forstio.remote
 		forstio.remote-sycl
-		kel-lbm.core
+		kel.lbm.core
 		adaptive-cpp
 	];
 
diff --git a/lib/sycl/c++/common.hpp b/lib/sycl/c++/common.hpp
new file mode 100644
index 0000000..8ff76fc
--- /dev/null
+++ b/lib/sycl/c++/common.hpp
@@ -0,0 +1,12 @@
+#pragma once
+
+#include <kel/lbm/common.hpp>
+#include <AdaptiveCpp/sycl/sycl.hpp>
+
+namespace kel {
+namespace lbm {
+namespace sycl {
+using namespace acpp::sycl;
+}
+}
+}
diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp
new file mode 100644
index 0000000..67422e2
--- /dev/null
+++ b/lib/sycl/c++/data.hpp
@@ -0,0 +1,117 @@
+#pragma once
+
+#include "common.hpp"
+
+namespace kel {
+namespace lbm {
+namespace encode {
+template<typename Encode>
+struct Sycl {
+};
+}
+
+namespace impl {
+template<typename Schema>
+struct struct_has_only_equal_dimension_array
+}
+}
+}
+
+namespace saw {
+template<uint64_t... Meta, typename... Sch, string_literal... Keys, typename Encode>
+class data<schema::Struct<schema::Member<schema::FixedArray<Sch,Meta...>, Keys>...>, kel::lbm::encode::Sycl<Encode>> final {
+public:
+	static constexpr data<schema::FixedArray meta = {{Meta...}};
+	using StorageT = std::tuple<data<Members::Type::InnerSchema,Encode>*...>;
+private:
+
+	/**
+	 * @todo Check by static assert that the members all have the same dimensions. Alternatively
+	 * Do it here by specializing.
+	 */
+	StorageT members_;
+	kel::lbm::sycl::queue* q_;
+public:
+	data():
+		members_{},
+		q_{nullptr}
+	{}
+
+	~data(){
+		SAW_ASSERT(q_){
+			exit(-1);
+		}
+		std::visit([this](auto arg){
+			if(not arg){
+				return;
+			}
+			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;
+	}
+
+	void set_queue(kel::lbm::sycl::queue& q){
+		q_ = &q;
+	}
+};
+
+}
+
+namespace kel {
+namespace lbm {
+namespace impl {
+template<typename Sch, typename Encode>
+struct sycl_malloc_struct_helper;
+	
+template<typename... Members, typename Encode>
+struct sycl_malloc_struct_helper<sch::Struct<Members...>, Encode> final {
+	template<uint64_t i>
+	static saw::error_or<void> allocate_on_device_member(typename data<Sch,encode::Sycl<Encode>>::StorageT& storage, sycl::queue& q){
+		if constexpr (i < sizeof...(Members)){
+			auto& ptr = std::get<i>(storage);
+
+			return allocate_on_device_member<i+1u>(sycl_data,q);
+		}
+		
+		return saw::make_void();
+	}
+	
+	static saw::error_or<void> allocate_on_device(data<Sch,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){
+		typename data<Sch,encode::Sycl<Encode>>::StorageT storage;
+		return allocate_on_device_member<0u>(storage,q);
+	}
+};
+}
+class device final {
+private:
+	sycl::queue q_;
+
+	SAW_FORBID_COPY(device);
+	SAW_FORBID_MOVE(device);
+public:
+	device() = default;
+	~device() = default;
+
+	template<typename Sch, typename Encode>
+	saw::error_or<void> allocate_on_device(data<Sch,encode::Sycl<Encode>>& sycl_data){
+		auto eov = sycl_malloc_struct_helper<Sch,Encode>::allocate_on_device(sycl_data, q_);
+		if(eov.is_error()){
+			return eov;
+		}
+		sycl_data.set_queue(q_);
+	}
+};
+}
+}
diff --git a/lib/sycl/c++/sycl.hpp b/lib/sycl/c++/sycl.hpp
new file mode 100644
index 0000000..8ddc3cd
--- /dev/null
+++ b/lib/sycl/c++/sycl.hpp
@@ -0,0 +1,4 @@
+#pragma once
+
+#include "common.hpp"
+#include "data.hpp"