5 files changed, 107 insertions, 40 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;
 }
-}