Cleaning up examples

16 files changed, 552 insertions, 274 deletions

diff --git a/default.nix b/default.nix
index 6b79f2f..c408acd 100644
--- a/default.nix
+++ b/default.nix
@@ -32,6 +32,21 @@ in rec {
 			inherit pname version stdenv forstio adaptive-cpp-dev;
 			kel-lbm = lbm;
 		};
+		
+		cavity_2d = pkgs.callPackage ./examples/cavity_2d/.nix/derivation.nix {
+			inherit pname version stdenv forstio;
+			kel-lbm = lbm;
+		};
+		
+		poiseulle_2d = pkgs.callPackage ./examples/poiseulle_2d/.nix/derivation.nix {
+			inherit pname version stdenv forstio;
+			kel-lbm = lbm;
+		};
+		
+		poiseulle_3d = pkgs.callPackage ./examples/poiseulle_3d/.nix/derivation.nix {
+			inherit pname version stdenv forstio;
+			kel-lbm = lbm;
+		};
 
 		poiseulle_particles_channel_2d = pkgs.callPackage ./examples/poiseulle_particles_channel_2d/.nix/derivation.nix {
 			inherit pname version stdenv forstio;
diff --git a/examples/cavity_2d/.nix/derivation.nix b/examples/cavity_2d/.nix/derivation.nix
new file mode 100644
index 0000000..cab8b8d
--- /dev/null
+++ b/examples/cavity_2d/.nix/derivation.nix
@@ -0,0 +1,33 @@
+{ lib
+, stdenv
+, scons
+, clang-tools
+, forstio
+, pname
+, version
+, kel-lbm
+}:
+
+stdenv.mkDerivation {
+	pname = pname + "-examples-" + "cavity_2d";
+	inherit version;
+	src = ./..;
+
+	nativeBuildInputs = [
+		scons
+		clang-tools
+	];
+
+	buildInputs = [
+		forstio.core
+		forstio.async
+		forstio.codec
+		forstio.codec-unit
+		forstio.codec-json
+		kel-lbm
+	];
+
+	preferLocalBuild = true;
+
+	outputs = [ "out" "dev" ];
+}
diff --git a/examples/cavity_2d/SConscript b/examples/cavity_2d/SConscript
new file mode 100644
index 0000000..3fabe4b
--- /dev/null
+++ b/examples/cavity_2d/SConscript
@@ -0,0 +1,32 @@
+#!/bin/false
+
+import os
+import os.path
+import glob
+
+
+Import('env')
+
+dir_path = Dir('.').abspath
+
+# Environment for base library
+examples_env = env.Clone();
+
+examples_env.sources = sorted(glob.glob(dir_path + "/*.cpp"))
+examples_env.headers = sorted(glob.glob(dir_path + "/*.hpp"))
+
+env.sources += examples_env.sources;
+env.headers += examples_env.headers;
+
+# Cavity2D
+examples_objects = [];
+examples_env.add_source_files(examples_objects, ['cavity_2d.cpp'], shared=False);
+examples_env.cavity_2d = examples_env.Program('#bin/cavity_2d', [examples_objects]);
+
+# Set Alias
+env.examples = [
+	examples_env.cavity_2d
+];
+env.Alias('examples', env.examples);
+env.targets += ['examples'];
+env.Install('$prefix/bin/', env.examples);
diff --git a/examples/cavity_2d/SConstruct b/examples/cavity_2d/SConstruct
new file mode 100644
index 0000000..a7201cb
--- /dev/null
+++ b/examples/cavity_2d/SConstruct
@@ -0,0 +1,79 @@
+#!/usr/bin/env python3
+
+import sys
+import os
+import os.path
+import glob
+import re
+
+
+if sys.version_info < (3,):
+    def isbasestring(s):
+        return isinstance(s,basestring)
+else:
+    def isbasestring(s):
+        return isinstance(s, (str,bytes))
+
+def add_kel_source_files(self, sources, filetype, lib_env=None, shared=False, target_post=""):
+
+    if isbasestring(filetype):
+        dir_path = self.Dir('.').abspath
+        filetype = sorted(glob.glob(dir_path+"/"+filetype))
+
+    for path in filetype:
+        target_name = re.sub( r'(.*?)(\.cpp|\.c\+\+)', r'\1' + target_post, path )
+        if shared:
+            target_name+='.os'
+            sources.append( self.SharedObject( target=target_name, source=path ) )
+        else:
+            target_name+='.o'
+            sources.append( self.StaticObject( target=target_name, source=path ) )
+    pass
+
+def isAbsolutePath(key, dirname, env):
+	assert os.path.isabs(dirname), "%r must have absolute path syntax" % (key,)
+
+env_vars = Variables(
+	args=ARGUMENTS
+)
+
+env_vars.Add('prefix',
+	help='Installation target location of build results and headers',
+	default='/usr/local/',
+	validator=isAbsolutePath
+)
+
+env_vars.Add('build_examples',
+	help='If examples should be built',
+	default="true"
+)
+
+env=Environment(ENV=os.environ, variables=env_vars, CPPPATH=[],
+    CPPDEFINES=['SAW_UNIX'],
+    CXXFLAGS=[
+        '-std=c++20',
+        '-g',
+        '-Wall',
+        '-Wextra'
+    ],
+    LIBS=[
+        'forstio-core'
+    ]
+);
+env.__class__.add_source_files = add_kel_source_files
+env.Tool('compilation_db');
+env.cdb = env.CompilationDatabase('compile_commands.json');
+
+env.objects = [];
+env.sources = [];
+env.headers = [];
+env.targets = [];
+
+Export('env')
+SConscript('SConscript')
+
+env.Alias('cdb', env.cdb);
+env.Alias('all', [env.targets]);
+env.Default('all');
+
+env.Alias('install', '$prefix')
diff --git a/examples/cavity_2d.cpp b/examples/cavity_2d/cavity_2d.cpp
index e228a06..702661a 100644
--- a/examples/cavity_2d.cpp
+++ b/examples/cavity_2d/cavity_2d.cpp
@@ -1,9 +1,4 @@
-#include "../c++/descriptor.hpp"
-#include "../c++/macroscopic.hpp"
-#include "../c++/lbm.hpp"
-#include "../c++/component.hpp"
-#include "../c++/collision.hpp"
-#include "../c++/boundary.hpp"
+#include <kel/lbm/lbm.hpp>
 
 #include <forstio/codec/data.hpp>
 
diff --git a/examples/cavity_2d_gpu/cavity_2d_gpu.cpp b/examples/cavity_2d_gpu/cavity_2d_gpu.cpp
index 5efe7a3..964bfde 100644
--- a/examples/cavity_2d_gpu/cavity_2d_gpu.cpp
+++ b/examples/cavity_2d_gpu/cavity_2d_gpu.cpp
@@ -6,6 +6,7 @@
 #include <iostream>
 #include <fstream>
 #include <vector>
+#include <chrono>
 #include <cmath>
 
 namespace saw {
@@ -153,18 +154,19 @@ public:
 constexpr size_t dim_x = 128;
 constexpr size_t dim_y = 128;
 
+template<typename Desc>
 void set_geometry(
-	saw::data<sch::CellInfo<Desc>>& info_field
+	saw::data<sch::CellInfo<Desc>>* info_field
 ){
 	using namespace kel::lbm;
 	using namespace acpp;
-
-	auto meta = lattice.meta();
+	saw::data<sch::FixedArray<sch::UInt64,Desc::D>> meta{{dim_x,dim_y}};
 	/**
 	 * Set ghost
 	 */
 	iterate_over([&](const saw::data<sch::FixedArray<sch::UInt64,2u>>& index){
-		auto& info = info_field.at(index);
+		size_t i = index.at({0u}).get() * dim_x + index.at({1u}).get();
+		auto& info = info_field[i];
 
 		info({0u}).set(0u);
 
@@ -174,9 +176,10 @@ void set_geometry(
 	 * Set wall
 	 */
 	iterate_over([&](const saw::data<sch::FixedArray<sch::UInt64,2u>>& index){
-		auto& info = info_field.at(index);
+		size_t i = index.at({0u}).get() * dim_x + index.at({1u}).get();
+		auto& info = info_field[i];
 
-		info({0u}).set(2u);
+		info({0u}).set(1u);
 
 	}, {{0u,0u}}, meta, {{1u,1u}});
 
@@ -184,9 +187,10 @@ void set_geometry(
 	 * Set fluid
 	 */
 	iterate_over([&](const saw::data<sch::FixedArray<sch::UInt64,2u>>& index){
-		auto& info = info_field.at(index);
+		size_t i = index.at({0u}).get() * dim_x + index.at({1u}).get();
+		auto& info = info_field[i];
 
-		info({0u}).set(1u);
+		info({0u}).set(2u);
 
 	}, {{0u,0u}}, meta, {{2u,2u}});
 
@@ -194,12 +198,14 @@ void set_geometry(
 	 * Set top lid
 	 */
 	iterate_over([&](const saw::data<sch::FixedArray<sch::UInt64,2u>>& index){
-		auto& info = info_field.at(index);
+		size_t i = index.at({0u}).get() * dim_x + index.at({1u}).get();
+		auto& info = info_field[i];
 		info({0u}).set(3u);
 
 	}, {{0u,1u}}, {{meta.at({0u}), 2u}}, {{2u,0u}});
 }
 
+template<typename Desc>
 void set_initial_conditions(
 	saw::data<sch::CellInfo<Desc>>* info_field,
 	saw::data<sch::DfCell<Desc>>* dfs_field,
@@ -209,19 +215,19 @@ void set_initial_conditions(
 	using namespace acpp;
 
 	saw::data<sch::T> rho{1.0};
-	saw::data<sch::FixedArray<sch::T,sch::D2Q9::D>> vel{{0.0,0.0}};
-	auto eq = equilibrium<sch::T,sch::D2Q9>(rho, vel);
-	saw::data<sch::FixedArray<sch::UInt64,sch::D2Q9::D>> meta{{dim_x,dim_y}};
+	saw::data<sch::FixedArray<sch::T,Desc::D>> vel{{0.0,0.0}};
+	auto eq = equilibrium<sch::T,Desc>(rho, vel);
+	saw::data<sch::FixedArray<sch::UInt64,Desc::D>> meta{{dim_x,dim_y}};
 
 	/**
 	 * Set distribution
 	 */
-	iterate_over([&](const saw::data<sch::FixedArray<sch::UInt64,sch::D2Q9::D>>& index){
+	iterate_over([&](const saw::data<sch::FixedArray<sch::UInt64,Desc::D>>& index){
 		size_t i = index.at({0u}).get() * dim_x + index.at({1u}).get();
 		auto& dfs = dfs_field[i];
 		auto& dfs_old = dfs_old_field[i];
 
-		for(saw::data<sch::UInt64> k = 0; k < saw::data<sch::UInt64>{sch::D2Q9::Q}; ++k){
+		for(saw::data<sch::UInt64> k = 0; k < saw::data<sch::UInt64>{Desc::Q}; ++k){
 			dfs(k) = eq.at(k);
 			dfs_old(k) = eq.at(k);
 		}
@@ -249,11 +255,6 @@ void lbm_step(
 	component<T, Desc, cmpt::MovingWall> bb_lid;
 	bb_lid.lid_vel = {0.1, 0.0};
 
-	/*
-	const size_t Nx = latt.template get_dim_size<0>().get();
-	const size_t Ny = latt.template get_dim_size<1>().get();
-	*/
-
 	// Submit collision kernel
 	sycl_q.submit([&](sycl::handler& cgh) {
 	 // Accessor for latt with read/write
@@ -272,20 +273,18 @@ void lbm_step(
 			 case 1u: {
 				// bb.apply(latt_acc, {i, j}, time_step);
 
-						auto& dfs_old = (is_even) ? dfs_r_old[acc_id] : dfs_r[acc_id];
-						// auto& dfs = (not is_even) ? cell.template get<"dfs_old">() : cell.template get<"dfs">();
-						auto df_cpy = dfs_old.copy();
+				auto& dfs_old = is_even ? dfs_r_old[acc_id] : dfs_r[acc_id];
+				auto df_cpy = dfs_old.copy();
 
-						for(uint64_t i = 1u; i < Desc::Q; ++i){
-							dfs_old({i}) = df_cpy({dfi::opposite_index.at(i)});
-						}
+				for(uint64_t i = 1u; i < Desc::Q; ++i){
+					dfs_old({i}) = df_cpy({dfi::opposite_index.at(i)});
+				}
 
 				break;
 			 }
 			 case 2u: {
 				// coll.apply(latt_acc, {i, j}, time_step);
-
-				auto& dfs_old = (is_even) ? dfs_r_old[acc_id] : dfs_r[acc_id];
+				auto& dfs_old = is_even ? dfs_r_old[acc_id] : dfs_r[acc_id];
 
 				saw::data<T> rho;
 				saw::data<sch::FixedArray<T,Desc::D>> vel;
@@ -294,7 +293,6 @@ void lbm_step(
 
 				for(uint64_t i = 0u; i < Desc::Q; ++i){
 					dfs_old({i}) = dfs_old({i}) + frequency * (eq.at(i) - dfs_old({i}));
-					// dfs_old({i}).set(dfs_old({i}).get() + (1.0 / relaxation_) * (eq.at(i).get() - dfs_old({i}).get()));
 				}
 				break;
 			}
@@ -315,26 +313,27 @@ void lbm_step(
 
 			size_t acc_id = i * dim_x + j;
 
-			auto dfs_new = dfs_r[acc_id];
+			auto& dfs_new = is_even ? dfs_r[acc_id] : dfs_r_old[acc_id];
 			auto& info = info_r[acc_id];
 
+
 			if (info({0u}).get() > 0u && info({0u}).get() != 3u) {
-			 for (uint64_t k = 0u; k < Desc::Q; ++k) {
-				auto dir = dfi::directions[dfi::opposite_index[k]];
-				// auto& cell_dir_old = latt_acc({i + dir[0], j + dir[1]});
-				//
-				size_t acc_old_id = (i+dir[0]) * dim_x + (j+dir[1]);
-
-				auto& dfs_old = dfs_r_old[acc_old_id];
-				auto& info_old = info_r[acc_old_id];
-
-				if (info_old({0}).get() == 3u) {
-					auto& dfs_old_loc = dfs_r_old[acc_old_id];
-					dfs_new({k}) = dfs_old_loc({dfi::opposite_index.at(k)}) - 2.0 * dfi::inv_cs2 * dfi::weights.at(k) * 1.0 * (bb_lid.lid_vel[0] * dir[0] + bb_lid.lid_vel[1] * dir[1]);
-				} else {
-					dfs_new({k}) = dfs_old({k});
+				for (uint64_t k = 0u; k < Desc::Q; ++k) {
+					auto dir = dfi::directions[dfi::opposite_index[k]];
+					// auto& cell_dir_old = latt_acc({i + dir[0], j + dir[1]});
+					//
+					size_t acc_old_id = (i+dir[0]) * dim_x + (j+dir[1]);
+
+					auto& dfs_old = is_even ? dfs_r_old[acc_old_id] : dfs_r[acc_old_id];
+					auto& info_old = info_r[acc_old_id];
+
+					if (info_old({0}).get() == 3u) {
+						auto& dfs_old_loc = is_even ? dfs_r_old[acc_id] : dfs_r[acc_id];
+						dfs_new({k}) = dfs_old_loc({dfi::opposite_index.at(k)}) - 2.0 * dfi::inv_cs2 * dfi::weights.at(k) * 1.0 * (bb_lid.lid_vel[0] * dir[0] + bb_lid.lid_vel[1] * dir[1]);
+					} else {
+						dfs_new({k}) = dfs_old({k});
+					}
 				}
-			 }
 			}
 		});
 	});
@@ -346,7 +345,7 @@ int main(){
 	using namespace kel::lbm;
 	using namespace acpp;
 
-	saw::data<sch::FixedArray<sch::UInt64,sch::D2Q9::D>> dim{{dim_x, dim_y}};
+	saw::data<sch::FixedArray<sch::UInt64,sch::D2Q9::D>> meta{{dim_x, dim_y}};
 	constexpr size_t dim_size = dim_x * dim_y;
 
 	converter<sch::T> conv{
@@ -391,30 +390,35 @@ int main(){
 	uint64_t print_every = 256u;
 	uint64_t file_no = 0u;
 
-	saw::data<sch::Array<sch::MacroStruct<sch::T,sch::D2Q9::D>,sch::D2Q9::D>> macros{dim};
+	saw::data<sch::Array<sch::MacroStruct<sch::T,sch::D2Q9::D>,sch::D2Q9::D>> macros{meta};
 
+	std::cout<<"Start"<<std::endl;
 
-	for(uint64_t i = 0u; i < 256u; ++i){
-		{
+	auto start = std::chrono::steady_clock::now();
+	sycl_q.wait();
+	for(uint64_t i = 0u; i < lattice_steps; ++i){
+		lbm_step<sch::T,sch::D2Q9>(info, dfs, dfs_old, even_step, i, sycl_q);
+		if(i % 1024u == 0u){
 			sycl_q.wait();
-			{
-				//auto acc = dfs_sycl.template get_access<acpp::sycl::access::mode::read>(dfs_field.internal_data());
-
-			}
-			apply_for_cells([&](auto& cell, std::size_t i, std::size_t j){
-				auto dfs = even_step ? cell.template get<"dfs_old">() : cell.template get<"dfs">();
-
-				auto& rho = macros.at({{i,j}}).template get<"pressure">();
-				auto& vel = macros.at({{i,j}}).template get<"velocity">();
-				compute_rho_u
-			dfs_sycl.synchronize();<sch::T,sch::D2Q9>(dfs,rho,vel);
-			}, lattice);
-			std::string vtk_f_name{"tmp/poiseulle_2d_"};
+			iterate_over([&](const saw::data<sch::FixedArray<sch::UInt64,sch::D2Q9::D>>& index){
+				size_t j = index.at({0u}).get() * dim_x + index.at({1u}).get();
+				auto dfs_field = even_step ? dfs_old : dfs;
+
+				auto& rho = macros.at(index).template get<"pressure">();
+				auto& vel = macros.at(index).template get<"velocity">();
+				compute_rho_u<sch::T,sch::D2Q9>(dfs_field[j],rho,vel);
+			}, {{0u,0u}}, meta);
+			std::string vtk_f_name{"tmp/cavity_2d_gpu_"};
 			vtk_f_name += std::to_string(i) + ".vtk";
 			write_vtk_file(vtk_f_name, macros);
 		}
 
-		lbm_step<sch::T,sch::D2Q9>(info_sycl, dfs_sycl, dfs_old_sycl, (i%2u == 0u), i, sycl_q);
+		even_step = not even_step;
 	}
+	auto stop = std::chrono::steady_clock::now();
+	std::cout<<std::format("{:%H:%M:%S}",(stop-start))<<std::endl;
+	sycl::free(info,sycl_q);
+	sycl::free(dfs,sycl_q);
+	sycl::free(dfs_old,sycl_q);
 	return 0;
 }
diff --git a/examples/particle_ibm.cpp b/examples/particle_ibm.cpp
deleted file mode 100644
index d54dbdc..0000000
--- a/examples/particle_ibm.cpp
+++ /dev/null
@@ -1,201 +0,0 @@
-#include "../c++/descriptor.hpp"
-#include "../c++/equilibrium.hpp"
-
-#include <forstio/codec/data.hpp>
-
-#include <iostream>
-
-namespace kel {
-namespace lbm {
-namespace sch {
-using namespace saw::schema;
-
-/**
- * Basic distribution function
- * Base type
- * D
- * Q
- * Scalar factor
- * D factor
- * Q factor
- */
-using T = Float32;
-using D2Q5 = Descriptor<2,5>;
-using D2Q9 = Descriptor<2,9>;
-
-template<typename Desc>
-using DfCell = Cell<T, Desc, 0, 0, 1>;
-
-template<typename Desc>
-using CellInfo = Cell<UInt8, D2Q9, 1, 0, 0>;
-
-/**
- * Basic type for simulation
- */
-template<typename Desc>
-using CellStruct = Struct<
-	Member<DfCell<Desc>, "dfs">,
-	Member<DfCell<Desc>, "dfs_old">,
-	Member<CellInfo<Desc>, "info">,
-	Member<UInt8, "geometry">,
-	Member<UInt32, "particle">
->;
-
-template<typename T, uint64_t D>
-using MacroStruct = Struct<
-	Member<FixedArray<T,D>, "velocity">,
-	Member<T, "pressure">
->;
-
-using CavityFieldD2Q9 = CellField<D2Q9, CellStruct<D2Q9>>;
-}
-}
-}
-
-void set_geometry(saw::data<kel::lbm::sch::CavityFieldD2Q9>& latt){
-	using namespace kel::lbm;
-	/*
-	apply_for_cells([](auto& cell, std::size_t i, std::size_t j){
-		uint8_t val = 0;
-		if(j == 1){
-			val = 2u;
-		}
-		if(i == 1 || (i+2) == dim_x || (j+2) == dim_y){
-			val = 3u;
-		}
-		if(i > 1 && (i+2) < dim_x && j > 1 && (j+2) < dim_y){
-			val = 1u;
-		}
-		if(i == 0 || j == 0 || (i+1) == dim_x || (j+1) == dim_y){
-			val = 0u;
-		}
-		cell.template get<"info">()(0u).set(val);
-	}, latt);
-	*/
-}
-
-void lbm_step(
-	saw::data<kel::lbm::sch::CavityFieldD2Q9>& latt,
-	bool even_step
-){
-	using namespace kel::lbm;
-  using dfi = df_info<sch::T,sch::D2Q9>;
-
-	component<cmpt::BGK,sch::D2Q9> coll;
-	coll.relaxation_ = 0.5384;
-	component<cmpt::ConstRhoBGK,sch::D2Q9> rho_coll;
-	rho_coll.relaxation_ = 0.5384;
-	rho_coll.rho_ = 1.0;
-
-	component<cmpt::BounceBack,sch::D2Q9> bb;
-	component<cmpt::MovingWall,sch::D2Q9> bb_lid;
-	bb_lid.lid_vel = {0.1,0.0};
-
-	// Collide
-	apply_for_cells([&](auto& cell, std::size_t i, std::size_t j){
-		auto& df = even_step ? cell.template get<"dfs_old">() : cell.template get<"dfs">();
-		auto& info = cell.template get<"info">();
-
-		auto info_val = info({0u}).get();
-		switch(info_val){
-		case 1u:
-			coll.apply(df);
-			break;
-		case 2u:
-			// bb.apply(df);
-			bb_lid.apply(df);
-			break;
-		case 3u:
-			bb.apply(df);
-			break;
-		}
-	}, latt);
-
-	// Stream
-	for(uint64_t i = 1; (i+1) < latt.template get_dim_size<0>().get(); ++i){
-		for(uint64_t j = 1; (j+1) < latt.template get_dim_size<1>().get(); ++j){
-			auto& cell = latt({{i,j}});
-			auto& df_new = even_step ? cell.template get<"dfs">() : cell.template get<"dfs_old">();
-			auto& info_new = cell.template get<"info">();
-
-			if(info_new({0u}).get() > 0u && info_new({0u}).get() != 2u){
-				for(uint64_t k = 0u; k < sch::D2Q9::Q; ++k){
-					auto dir = dfi::directions[dfi::opposite_index[k]];
-					auto& cell_dir_old = latt({{i+dir[0],j+dir[1]}});
-
-					auto& df_old = even_step ? cell_dir_old.template get<"dfs_old">() : cell_dir_old.template get<"dfs">();
-					auto& info_old = cell_dir_old.template get<"info">();
-
-					if( info_old({0}).get() == 2u ){
-						auto& df_old_loc = even_step ? latt({{i,j}}).template get<"dfs_old">() : latt({{i,j}}).template get<"dfs">();
-						df_new({k}) = df_old_loc({dfi::opposite_index.at(k)}) - 2.0 * dfi::inv_cs2 * dfi::weights.at(k) * 1.0 * ( bb_lid.lid_vel[0] * dir[0] + bb_lid.lid_vel[1] * dir[1]);
-						// dfs({dfi::opposite_index.at(i)}) = dfs_cpy({i}) - 2.0 * dfi::weights[i] * 1.0 * ( lid_vel[0] * dfi::directions[i][0] + lid_vel[1] * dfi::directions[i][1]) * dfi::inv_cs2;
-					} else {
-						df_new({k}) = df_old({k});
-					}
-				}
-			}
-		}
-	}
-}
-
-int main(int argc, char** argv){
-	using namespace kel::lbm;
-
-	std::string_view cfg_file_name = "config.json";
-	if(argc > 1){
-		cfg_file_name = argv[1];
-	}
-
-	auto eo_conf = load_lbm_config<sch::Float64,sch::Descriptor<2,9>>(cfg_file_name);
-	if(eo_conf.is_error()){
-		auto& err = eo_conf.get_error();
-		std::cerr<<"[Error]: "<<err.get_category();
-		auto err_msg = err.get_message();
-		if(!err_msg.empty()){
-			std::cerr<<" - "<<err_msg;
-		}
-		std::cerr<<std::endl;
-
-		return err.get_id();
-	}
-
-	auto& conf = eo_conf.get_value();
-
-	converter<sch::Float64> conv{
-		{conf.template get<"delta_x">()},
-		{conf.template get<"delta_t">()}
-	};
-
-	print_lbm_meta<sch::Float64,sch::Descriptor<2,9>>(conv, {conf.template get<"kinematic_viscosity">()});
-
-
-	saw::data<sch::FixedArray<sch::UInt64,sch::D2Q9::D>> dim{{128, 128}};
-
-	saw::data<sch::CavityFieldD2Q9, saw::encode::Native> lattice{dim};
-	lbm::particle_system<sch::T,2,sch::Particle<sch::T,2>> system;
-
-	/**
-	 * Set meta information describing what this cell is
-	 */
-	set_geometry(lattice);
-
-
-	uint64_t lattice_steps{128u};
-	uint64_t print_every{64u};
-	bool even_step{true};
-
-	uint64_t file_num{0u};
-
-	saw::data<sch::Array<sch::MacroStruct<sch::T,sch::D2Q9::D>,sch::D2Q9::D>> macros{dim};
-
-	for(uint64_t step{0u}; step < lattice_steps; ++step){
-
-		lbm_step(lattice, even_step);
-
-		even_step = not even_step;
-
-	}
-
-	return 0;
-}
diff --git a/examples/planetary_3d/planetary_3d.cpp b/examples/planetary_3d/planetary_3d.cpp
index e100f48..990652a 100644
--- a/examples/planetary_3d/planetary_3d.cpp
+++ b/examples/planetary_3d/planetary_3d.cpp
@@ -2,6 +2,12 @@
 #include <iostream>
 
 namespace kel {
+namespace sch {
+using namespace saw::schema;
+using KelConfig = Struct<
+	Member<String,"resolution">
+>;
+}
 
 saw::error_or<void> real_main(int argc, char** argv){
 	return saw::make_void();
diff --git a/examples/poiseulle_2d/SConscript b/examples/poiseulle_2d/SConscript
new file mode 100644
index 0000000..041f14c
--- /dev/null
+++ b/examples/poiseulle_2d/SConscript
@@ -0,0 +1,32 @@
+#!/bin/false
+
+import os
+import os.path
+import glob
+
+
+Import('env')
+
+dir_path = Dir('.').abspath
+
+# Environment for base library
+examples_env = env.Clone();
+
+examples_env.sources = sorted(glob.glob(dir_path + "/*.cpp"))
+examples_env.headers = sorted(glob.glob(dir_path + "/*.hpp"))
+
+env.sources += examples_env.sources;
+env.headers += examples_env.headers;
+
+# Cavity2D
+examples_objects = [];
+examples_env.add_source_files(examples_objects, ['poiseulle_2d.cpp'], shared=False);
+examples_env.poiseulle_2d = examples_env.Program('#bin/poiseulle_2d', [examples_objects]);
+
+# Set Alias
+env.examples = [
+	examples_env.poiseulle_2d
+];
+env.Alias('examples', env.examples);
+env.targets += ['examples'];
+env.Install('$prefix/bin/', env.examples);
diff --git a/examples/poiseulle_2d/SConstruct b/examples/poiseulle_2d/SConstruct
new file mode 100644
index 0000000..a7201cb
--- /dev/null
+++ b/examples/poiseulle_2d/SConstruct
@@ -0,0 +1,79 @@
+#!/usr/bin/env python3
+
+import sys
+import os
+import os.path
+import glob
+import re
+
+
+if sys.version_info < (3,):
+    def isbasestring(s):
+        return isinstance(s,basestring)
+else:
+    def isbasestring(s):
+        return isinstance(s, (str,bytes))
+
+def add_kel_source_files(self, sources, filetype, lib_env=None, shared=False, target_post=""):
+
+    if isbasestring(filetype):
+        dir_path = self.Dir('.').abspath
+        filetype = sorted(glob.glob(dir_path+"/"+filetype))
+
+    for path in filetype:
+        target_name = re.sub( r'(.*?)(\.cpp|\.c\+\+)', r'\1' + target_post, path )
+        if shared:
+            target_name+='.os'
+            sources.append( self.SharedObject( target=target_name, source=path ) )
+        else:
+            target_name+='.o'
+            sources.append( self.StaticObject( target=target_name, source=path ) )
+    pass
+
+def isAbsolutePath(key, dirname, env):
+	assert os.path.isabs(dirname), "%r must have absolute path syntax" % (key,)
+
+env_vars = Variables(
+	args=ARGUMENTS
+)
+
+env_vars.Add('prefix',
+	help='Installation target location of build results and headers',
+	default='/usr/local/',
+	validator=isAbsolutePath
+)
+
+env_vars.Add('build_examples',
+	help='If examples should be built',
+	default="true"
+)
+
+env=Environment(ENV=os.environ, variables=env_vars, CPPPATH=[],
+    CPPDEFINES=['SAW_UNIX'],
+    CXXFLAGS=[
+        '-std=c++20',
+        '-g',
+        '-Wall',
+        '-Wextra'
+    ],
+    LIBS=[
+        'forstio-core'
+    ]
+);
+env.__class__.add_source_files = add_kel_source_files
+env.Tool('compilation_db');
+env.cdb = env.CompilationDatabase('compile_commands.json');
+
+env.objects = [];
+env.sources = [];
+env.headers = [];
+env.targets = [];
+
+Export('env')
+SConscript('SConscript')
+
+env.Alias('cdb', env.cdb);
+env.Alias('all', [env.targets]);
+env.Default('all');
+
+env.Alias('install', '$prefix')
diff --git a/examples/poiseulle_2d.cpp b/examples/poiseulle_2d/poiseulle_2d.cpp
index 8c31cb8..8c31cb8 100644
--- a/examples/poiseulle_2d.cpp
+++ b/examples/poiseulle_2d/poiseulle_2d.cpp
diff --git a/examples/poiseulle_3d/.nix/derivation.nix b/examples/poiseulle_3d/.nix/derivation.nix
new file mode 100644
index 0000000..a98e736
--- /dev/null
+++ b/examples/poiseulle_3d/.nix/derivation.nix
@@ -0,0 +1,33 @@
+{ lib
+, stdenv
+, scons
+, clang-tools
+, forstio
+, pname
+, version
+, kel-lbm
+}:
+
+stdenv.mkDerivation {
+	pname = pname + "-examples-" + "poiseulle_3d";
+	inherit version;
+	src = ./..;
+
+	nativeBuildInputs = [
+		scons
+		clang-tools
+	];
+
+	buildInputs = [
+		forstio.core
+		forstio.async
+		forstio.codec
+		forstio.codec-unit
+		forstio.codec-json
+		kel-lbm
+	];
+
+	preferLocalBuild = true;
+
+	outputs = [ "out" "dev" ];
+}
diff --git a/examples/poiseulle_3d/SConscript b/examples/poiseulle_3d/SConscript
new file mode 100644
index 0000000..02338cd
--- /dev/null
+++ b/examples/poiseulle_3d/SConscript
@@ -0,0 +1,32 @@
+#!/bin/false
+
+import os
+import os.path
+import glob
+
+
+Import('env')
+
+dir_path = Dir('.').abspath
+
+# Environment for base library
+examples_env = env.Clone();
+
+examples_env.sources = sorted(glob.glob(dir_path + "/*.cpp"))
+examples_env.headers = sorted(glob.glob(dir_path + "/*.hpp"))
+
+env.sources += examples_env.sources;
+env.headers += examples_env.headers;
+
+# Cavity2D
+examples_objects = [];
+examples_env.add_source_files(examples_objects, ['poiseulle_3d.cpp'], shared=False);
+examples_env.poiseulle_3d = examples_env.Program('#bin/poiseulle_3d', [examples_objects]);
+
+# Set Alias
+env.examples = [
+	examples_env.poiseulle_3d
+];
+env.Alias('examples', env.examples);
+env.targets += ['examples'];
+env.Install('$prefix/bin/', env.examples);
diff --git a/examples/poiseulle_3d/SConstruct b/examples/poiseulle_3d/SConstruct
new file mode 100644
index 0000000..a7201cb
--- /dev/null
+++ b/examples/poiseulle_3d/SConstruct
@@ -0,0 +1,79 @@
+#!/usr/bin/env python3
+
+import sys
+import os
+import os.path
+import glob
+import re
+
+
+if sys.version_info < (3,):
+    def isbasestring(s):
+        return isinstance(s,basestring)
+else:
+    def isbasestring(s):
+        return isinstance(s, (str,bytes))
+
+def add_kel_source_files(self, sources, filetype, lib_env=None, shared=False, target_post=""):
+
+    if isbasestring(filetype):
+        dir_path = self.Dir('.').abspath
+        filetype = sorted(glob.glob(dir_path+"/"+filetype))
+
+    for path in filetype:
+        target_name = re.sub( r'(.*?)(\.cpp|\.c\+\+)', r'\1' + target_post, path )
+        if shared:
+            target_name+='.os'
+            sources.append( self.SharedObject( target=target_name, source=path ) )
+        else:
+            target_name+='.o'
+            sources.append( self.StaticObject( target=target_name, source=path ) )
+    pass
+
+def isAbsolutePath(key, dirname, env):
+	assert os.path.isabs(dirname), "%r must have absolute path syntax" % (key,)
+
+env_vars = Variables(
+	args=ARGUMENTS
+)
+
+env_vars.Add('prefix',
+	help='Installation target location of build results and headers',
+	default='/usr/local/',
+	validator=isAbsolutePath
+)
+
+env_vars.Add('build_examples',
+	help='If examples should be built',
+	default="true"
+)
+
+env=Environment(ENV=os.environ, variables=env_vars, CPPPATH=[],
+    CPPDEFINES=['SAW_UNIX'],
+    CXXFLAGS=[
+        '-std=c++20',
+        '-g',
+        '-Wall',
+        '-Wextra'
+    ],
+    LIBS=[
+        'forstio-core'
+    ]
+);
+env.__class__.add_source_files = add_kel_source_files
+env.Tool('compilation_db');
+env.cdb = env.CompilationDatabase('compile_commands.json');
+
+env.objects = [];
+env.sources = [];
+env.headers = [];
+env.targets = [];
+
+Export('env')
+SConscript('SConscript')
+
+env.Alias('cdb', env.cdb);
+env.Alias('all', [env.targets]);
+env.Default('all');
+
+env.Alias('install', '$prefix')
diff --git a/examples/poiseulle_3d/poiseulle_3d.cpp b/examples/poiseulle_3d/poiseulle_3d.cpp
new file mode 100644
index 0000000..9e7bba7
--- /dev/null
+++ b/examples/poiseulle_3d/poiseulle_3d.cpp
@@ -0,0 +1,39 @@
+#include <kel/lbm/lbm.hpp>
+
+namespace kel {
+namespace lbm {
+namespace sch {
+using T = Float32;
+using D3Q27 = Descriptor<3u,27u>;
+
+
+}
+saw::error_or<void> real_main(int argc, char** argv){
+
+
+	return saw::make_void();
+}
+
+} // lbm
+} // kel
+
+/**
+ * main, but I don't like the error handling
+ */
+int main(int argc, char** argv){
+	auto eov = kel::lbm::real_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/lib/c++/descriptor.hpp b/lib/c++/descriptor.hpp
index e81738e..18ea440 100644
--- a/lib/c++/descriptor.hpp
+++ b/lib/c++/descriptor.hpp
@@ -164,6 +164,7 @@ public:
 	static constexpr typename saw::native_data_type<T>::type inv_cs2 = 3.0;
 	static constexpr typename saw::native_data_type<T>::type cs2 = 1./3.;
 };
+
 /*
 template<typename T>
 class df_info<T,sch::Descriptor<3, 27>> {
@@ -177,16 +178,36 @@ public:
 		{ 0, 0, 0},
 		{-1, 0, 0},
 		{ 1, 0, 0},
-		{ 0,-1, 0},
+		// Expand into 2D
+		{ 0, -1, 0},
+		{-1, -1, 0},
+		{ 1, -1, 0},
 		{ 0, 1, 0},
-		{-1,-1, 0},
-		{ 1, 1, 0},
 		{-1, 1, 0},
-		{ 1,-1, 0}
+		{ 1, 1, 0},
+		// Expand into 3D
+		{ 0, 0, -1},
+		{-1, 0, -1},
+		{ 1, 0, -1},
+		{ 0, -1, -1},
+		{-1, -1, -1},
+		{ 1, -1, -1},
+		{ 0, 1, -1},
+		{-1, 1, -1},
+		{ 1, 1, -1},
+		{ 0, 0, 1},
+		{-1, 0, 1},
+		{ 1, 0, 1},
+		{ 0, -1, 1},
+		{-1, -1, 1},
+		{ 1, -1, 1},
+		{ 0, 1, 1},
+		{-1, 1, 1},
+		{ 1, 1, 1}
 	}};
 
 	static constexpr std::array<typename saw::native_data_type<T>::type,Q> weights = {
-		4./9.,
+		8./27.,
 		1./9.,
 		1./9.,
 		1./9.,