Fixing compilation errors

10 files changed, 117 insertions, 38 deletions

diff --git a/default.nix b/default.nix
index 461335c..a8ec430 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:0ckfcygk4f9x51y45hs5jpi95fjdfqfhl6bx6y374f7x98ng9xx3";
+		sha256 = "sha256:09xxs9dlnd52k1082lb5p034ik4jfg9zd5zc1x5da2clqhdb6s4m";
 	}) + "/default.nix"){
 		inherit stdenv;
 		inherit clang-tools;
diff --git a/examples/poiseulle_2d_gpu/poiseulle_2d_gpu.cpp b/examples/poiseulle_2d_gpu/poiseulle_2d_gpu.cpp
index 0f0d219..1aaa5c0 100644
--- a/examples/poiseulle_2d_gpu/poiseulle_2d_gpu.cpp
+++ b/examples/poiseulle_2d_gpu/poiseulle_2d_gpu.cpp
@@ -338,9 +338,7 @@ saw::error_or<void> kel_main(int argc, char** argv){
 	}
 	auto& lbm_dir = eo_lbm_dir.get_value();
 	auto out_dir = lbm_dir / "poiseulle_channel_2d_gpu";
-
 	std::filesystem::create_directories(out_dir);
-
 	// 
 	lbm::converter<lbm::sch::Float64> conv {
 		// delta_x
diff --git a/examples/poiseulle_particles_2d_gpu/.nix/derivation.nix b/examples/poiseulle_particles_2d_gpu/.nix/derivation.nix
index 127243d..e08f42a 100644
--- a/examples/poiseulle_particles_2d_gpu/.nix/derivation.nix
+++ b/examples/poiseulle_particles_2d_gpu/.nix/derivation.nix
@@ -26,11 +26,12 @@ stdenv.mkDerivation {
 		forstio.async
 		forstio.codec
 		forstio.codec-unit
-		forstio.remote
+      forstio.remote
+      forstio.remote-filesystem
 		forstio.codec-json
 		adaptive-cpp
 		kel.lbm.core
-#		kel-lbm.sycl
+		kel.lbm.sycl
 	];
 
 	preferLocalBuild = true;
diff --git a/examples/poiseulle_particles_2d_gpu/sim.cpp b/examples/poiseulle_particles_2d_gpu/sim.cpp
index 8bc60c9..be6b4f0 100644
--- a/examples/poiseulle_particles_2d_gpu/sim.cpp
+++ b/examples/poiseulle_particles_2d_gpu/sim.cpp
@@ -1,8 +1,15 @@
 #include <kel/lbm/lbm.hpp>
-#include <AdaptiveCpp/sycl/sycl.hpp>
+#include <kel/lbm/sycl/lbm.hpp>
+
+#include <forstio/remote/filesystem/easy.hpp>
+#include <forstio/codec/json/json.hpp>
 
 namespace kel {
 namespace lbm {
+
+constexpr uint64_t dim_x = 16u;
+constexpr uint64_t dim_y = 4u;
+
 namespace sch {
 using namespace saw::schema;
 
@@ -14,6 +21,63 @@ using CellStruct = Struct<
 	Member<Vector<T,Desc::D>, "velocity">,
 	Member<Vector<T,Desc::D>, "force">
 >;
+
+
+using InfoChunk = Chunk<UInt8, 0u, dim_x, dim_y>;
+
+template<typename T, typename Desc>
+using DfChunk = Chunk<FixedArray<T,Desc::Q>, 1u, dim_x, dim_y>;
+
+template<typename T, typename Desc>
+using ChunkStruct = Struct<
+	Member<InfoChunk, "info">
+>;
+
+}
+
+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">();
+	// Set everything as walls
+	iterator<Desc::D>::apply(
+		[&](auto& index){
+			info_f.at(index).set(2u);
+		},
+		{{0u,0u}},
+		{{dim_x, dim_y}},
+		{{0u,0u}}
+	);
+	// Fluid
+	iterator<Desc::D>::apply(
+		[&](auto& index){
+			info_f.at(index).set(1u);
+		},
+		{{0u,0u}},
+		{{dim_x, dim_y}},
+		{{1u,1u}}
+	);
+	
+	// Inflow
+	iterator<Desc::D>::apply(
+		[&](auto& index){
+			info_f.at(index).set(3u);
+		},
+		{{0u,0u}},
+		{{1u, dim_y}},
+		{{0u,1u}}
+	);
+	
+	// Outflow
+	iterator<Desc::D>::apply(
+		[&](auto& index){
+			info_f.at(index).set(4u);
+		},
+		{{dim_x-1u,0u}},
+		{{dim_x, dim_y}},
+		{{0u,1u}}
+	);
+
+	return saw::make_void();
 }
 
 template<typename T, typename Desc>
@@ -25,48 +89,54 @@ saw::error_or<void> step(){
 }
 
 template<typename T, typename Desc>
-saw::error_or<void> kel_main(int argc, char** argv){
-	using namespace kel;
+saw::error_or<void> lbm_main(int argc, char** argv){
+	using namespace kel::lbm;
 
-	using dfi = lbm::df_info<T,Desc>;
+	using dfi = df_info<T,Desc>;
 
-	auto eo_lbm_dir = lbm::output_directory();
+	auto eo_lbm_dir = output_directory();
 	if(eo_lbm_dir.is_error()){
 		return std::move(eo_lbm_dir.get_error());
 	}
 	auto& lbm_dir = eo_lbm_dir.get_value();
 	auto out_dir = lbm_dir / "poiseulle_particles_2d_gpu";
 
-	lbm::converter<lbm::sch::Float64> conv {
+	converter<sch::Float64> conv {
 		// delta_x
 		{{1.0}},
 		// delta_t
 		{{1.0}}
 	};
 
-	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};
+	// 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;
 	sycl_q.wait();
 	{
-		auto eov = setup_initial_conditions(lbm_data);
+		auto eov = setup_initial_conditions<T,Desc>(lbm_data);
 		if(eov.is_error()){
 			return eov;
 		}
 	}
 
-	
+	/*
+	iterator<Desc::D>::apply(
+		[&](auto& index){
+			std::cout<<index.at({0u}).get()<<" "<<index.at({1u}).get()<<" "<<lbm_data.get<"info">().at(index).template cast_to<sch::UInt16>().get()<<"\n";
+		},
+		{{0u,0u}},
+		{{dim_x, dim_y}}
+	);
+	*/
 
 	return saw::make_void();
 }
 
-using FloatT = kel::sch::Float32;
+using FloatT = kel::lbm::sch::Float32;
 
 int main(int argc, char** argv){
-	auto eov = kel_main<FloatT,kel::lbm::sch::D2Q9>(argc, argv);
+	auto eov = lbm_main<FloatT,kel::lbm::sch::D2Q9>(argc, argv);
 	if(eov.is_error()){
 		auto& err = eov.get_error();
 		std::cerr<<"[Error] "<<err.get_category();
diff --git a/lib/core/c++/chunk.hpp b/lib/core/c++/chunk.hpp
index 607bec0..ad7de34 100644
--- a/lib/core/c++/chunk.hpp
+++ b/lib/core/c++/chunk.hpp
@@ -76,4 +76,10 @@ public:
 		return data<schema::FixedArray<schema::UInt64, sizeof...(Sides)>,Encode>{{Sides...}};
 	}
 };
+
+template<typename Sch, uint64_t Ghost, uint64_t... Sides>
+struct meta_schema<kel::lbm::sch::Chunk<Sch,Ghost,Sides...>> {
+	using MetaSchema = typename meta_schema<Sch>::MetaSchema;
+};
+
 }
diff --git a/lib/core/c++/geometry.hpp b/lib/core/c++/geometry.hpp
index 6875e90..c8a48a6 100644
--- a/lib/core/c++/geometry.hpp
+++ b/lib/core/c++/geometry.hpp
@@ -21,7 +21,9 @@ class component<Schema, Desc, cmpt::PoiseulleChannel, Encode> final {
 private:
 public:
 	template<typename CellFieldSchema>
-	void apply(saw::data<CellFieldSchema,Encode>& field, const saw::data<sch::FixedArraysch::UInt64,Desc::D>)
+	void apply(saw::data<CellFieldSchema,Encode>& field, const saw::data<sch::FixedArraysch::UInt64,Desc::D>){
+		auto& info_f = field.template get<"info">();
+	}
 };
 
 // Ghost - 0
diff --git a/lib/core/c++/lbm.hpp b/lib/core/c++/lbm.hpp
index aff38e9..00f153a 100644
--- a/lib/core/c++/lbm.hpp
+++ b/lib/core/c++/lbm.hpp
@@ -2,6 +2,7 @@
 
 #include "schema.hpp"
 #include "flatten.hpp"
+#include "chunk.hpp"
 #include "descriptor.hpp"
 #include "boundary.hpp"
 #include "converter.hpp"
diff --git a/lib/sycl/c++/SConscript b/lib/sycl/c++/SConscript
index 85a078f..2ed63ba 100644
--- a/lib/sycl/c++/SConscript
+++ b/lib/sycl/c++/SConscript
@@ -15,18 +15,13 @@ core_env = env.Clone();
 core_env.sources = sorted(glob.glob(dir_path + "/*.cpp"));
 core_env.headers = sorted(glob.glob(dir_path + "/*.hpp"));
 
-core_env.particle_headers = sorted(glob.glob(dir_path + "/particle/*.hpp"));
-core_env.particle_geometry_headers = sorted(glob.glob(dir_path + "/particle/geometry/*.hpp"));
-
 env.sources += core_env.sources;
 env.headers += core_env.headers;
 
 ## Static lib
 objects = []
 core_env.add_source_files(objects, core_env.sources, shared=False);
-env.library_static = core_env.StaticLibrary('#build/kel-lbm', [objects]);
+env.library_static = core_env.StaticLibrary('#build/kel-lbm-sycl', [objects]);
 
 env.Install('$prefix/lib/', env.library_static);
-env.Install('$prefix/include/kel/lbm/', core_env.headers);
-env.Install('$prefix/include/kel/lbm/particle/', core_env.particle_headers);
-env.Install('$prefix/include/kel/lbm/particle/geometry/', core_env.particle_geometry_headers);
+env.Install('$prefix/include/kel/lbm/sycl/', core_env.headers);
diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp
index 44bc5dc..bb8b4bf 100644
--- a/lib/sycl/c++/data.hpp
+++ b/lib/sycl/c++/data.hpp
@@ -1,6 +1,7 @@
 #pragma once
 
 #include "common.hpp"
+#include <kel/lbm/lbm.hpp>
 
 namespace kel {
 namespace lbm {
@@ -19,9 +20,9 @@ struct struct_has_only_equal_dimension_array{};
 
 namespace saw {
 template<uint64_t Ghost, uint64_t... Meta, typename... Sch, string_literal... Keys, typename Encode>
-class data<schema::Struct<schema::Member<schema::Chunk<Sch,Ghost,Meta...>, Keys>...>, kel::lbm::encode::Sycl<Encode>> final {
+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 meta = {{Meta...}};
+	static constexpr data<schema::FixedArray<schema::UInt64,sizeof...(Meta)>> meta = {{Meta...}};
 	using StorageT = std::tuple<data<Sch,Encode>*...>;
 private:
 
@@ -50,14 +51,14 @@ public:
 			if(not arg){
 				return;
 			}
-			sycl::free(arg,*q_);
+			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_);
+		return std::get<parameter_key_pack_index<K, Keys...>::value>(members_);
 	}
 
 	template<saw::string_literal K>
@@ -67,6 +68,7 @@ public:
 		return *ptr;
 	}
 };
+}
 
 namespace kel {
 namespace lbm {
@@ -76,14 +78,15 @@ struct sycl_malloc_struct_helper;
 	
 template<typename... Members, typename Encode>
 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 data<Sch,encode::Sycl<Encode>>::StorageT& storage, sycl::queue& q){
+	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){
 		if constexpr (i < sizeof...(Members)){
 			using M = typename saw::parameter_pack_type<i,Members...>::type;
 			auto& ptr = std::get<i>(storage);
 
-			
-			ptr = sycl::malloc_device<M::ValueType::InnerSchema>(,q);
+			ptr = sycl::malloc_device<M::ValueType::InnerSchema>(1u,q);
 
 			return allocate_on_device_member<i+1u>(storage,q);
 		}
@@ -91,8 +94,8 @@ struct sycl_malloc_struct_helper<sch::Struct<Members...>, Encode> final {
 		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;
+	static saw::error_or<void> allocate_on_device(saw::data<Schema,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){
+		typename saw::data<Schema,encode::Sycl<Encode>>::StorageT storage;
  
 		auto eov = allocate_on_device_member<0u>(storage,q);
 		sycl_data = {storage, q};
@@ -112,8 +115,8 @@ public:
 	~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_);
+	saw::error_or<void> allocate_on_device(saw::data<Sch,encode::Sycl<Encode>>& sycl_data){
+		auto eov = impl::sycl_malloc_struct_helper<Sch,Encode>::allocate_on_device(sycl_data, q_);
 		if(eov.is_error()){
 			return eov;
 		}
diff --git a/lib/sycl/c++/lbm.hpp b/lib/sycl/c++/lbm.hpp
new file mode 100644
index 0000000..3e67ae6
--- /dev/null
+++ b/lib/sycl/c++/lbm.hpp
@@ -0,0 +1,3 @@
+#pragma once
+
+#include "data.hpp"