Rewriting parts to handle different ghost layers

8 files changed, 188 insertions, 5 deletions

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"