Dangling

21 files changed, 1637 insertions, 0 deletions

diff --git a/default.nix b/default.nix
index d7aeb9c..21944a8 100644
--- a/default.nix
+++ b/default.nix
@@ -222,6 +222,11 @@ in rec {
 			inherit pname version stdenv forstio adaptive-cpp;
 			inherit kel;
 		};
+		
+		stdkr_2d_bgk = pkgs.callPackage ./examples/stdkr_2d_bgk/.nix/derivation.nix {
+			inherit pname version stdenv forstio adaptive-cpp;
+			inherit kel;
+		};
 
 		heterogeneous_computing = pkgs.callPackage ./examples/heterogeneous_computing/.nix/derivation.nix {
 			inherit pname version stdenv forstio adaptive-cpp;
diff --git a/examples/dkt/.nix/derivation.nix b/examples/dkt/.nix/derivation.nix
new file mode 100644
index 0000000..3a640d8
--- /dev/null
+++ b/examples/dkt/.nix/derivation.nix
@@ -0,0 +1,43 @@
+{ lib
+, stdenv
+, scons
+, clang-tools
+, forstio
+, python3
+, pname
+, version
+, adaptive-cpp
+, kel
+, particle_coupling ? "fplbm"
+}:
+
+let
+in stdenv.mkDerivation {
+	pname = "${pname}-examples-dkt_2d_gpu_${particle_coupling}";
+	inherit version;
+	src = ./..;
+
+	nativeBuildInputs = [
+		scons
+		clang-tools
+		python3
+	];
+
+	buildInputs = [
+		forstio.core
+		forstio.async
+		forstio.codec
+		forstio.codec-unit
+		forstio.io
+    forstio.remote
+    forstio.remote-filesystem
+		forstio.codec-json
+		adaptive-cpp
+		kel.lbm.core
+		kel.lbm.sycl
+	];
+
+	preferLocalBuild = true;
+
+	outputs = [ "out" "dev" ];
+}
diff --git a/examples/dkt/SConscript b/examples/dkt/SConscript
new file mode 100644
index 0000000..b33b7ae
--- /dev/null
+++ b/examples/dkt/SConscript
@@ -0,0 +1,34 @@
+#!/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['CXX'] = 'syclcc-clang';
+examples_env['CXXFLAGS'] += ['-O3'];
+
+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, ['sim.cpp'], shared=False);
+examples_env.poiseulle_2d_gpu = examples_env.Program('#bin/stdkr_2d', [examples_objects]);
+
+# Set Alias
+env.examples = [
+	examples_env.poiseulle_2d_gpu
+];
+env.Alias('examples', env.examples);
+env.targets += ['examples'];
+env.Install('$prefix/bin/', env.examples);
diff --git a/examples/dkt/SConstruct b/examples/dkt/SConstruct
new file mode 100644
index 0000000..0611b67
--- /dev/null
+++ b/examples/dkt/SConstruct
@@ -0,0 +1,81 @@
+#!/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',
+        'forstio-async',
+        'forstio-io'
+    ]
+);
+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/dkt/common.hpp b/examples/dkt/common.hpp
new file mode 100644
index 0000000..be9a538
--- /dev/null
+++ b/examples/dkt/common.hpp
@@ -0,0 +1,71 @@
+#pragma once
+
+#include <kel/lbm/sycl/lbm.hpp>
+#include <kel/lbm/lbm.hpp>
+#include <kel/lbm/particle.hpp>
+#include <kel/lbm/particle/particle.hpp>
+
+#include <forstio/io/io.hpp>
+#include <forstio/remote/filesystem/easy.hpp>
+#include <forstio/codec/json/json.hpp>
+#include <forstio/codec/simple.hpp>
+
+namespace kel {
+namespace lbm {
+
+/**
+ * length 2.2m
+ * height 0.41m
+ * viscosity 10^-3
+ * particle diameter 0.1m
+ * position at (0.16m,0.2m)
+ */
+
+constexpr uint64_t dim_y = 512ul;
+constexpr uint64_t dim_x = static_cast<uint64_t>(dim_y * (220.0 / 41.0));
+
+constexpr uint64_t particle_amount = 1ul;
+
+namespace sch {
+using namespace saw::schema;
+
+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 ScalarChunk = Chunk<Scalar<T>, 0u, dim_x, dim_y>;
+
+template<typename T, typename Desc>
+using VectorChunk = Chunk<Vector<T,Desc::D>, 0u, dim_x, dim_y>;
+
+template<typename T, typename Desc>
+using ChunkStruct = Struct<
+	Member<InfoChunk, "info">,
+	Member<DfChunk<T,Desc>, "dfs">,
+	Member<DfChunk<T,Desc>, "dfs_old">,
+	Member<VectorChunk<T,Desc>, "particle_N">,
+	Member<ScalarChunk<T,Desc>, "particle_D">
+>;
+
+template<typename T, typename Desc>
+using VelChunk = Chunk<Vector<T,Desc::D>, 0u, dim_x, dim_y>;
+
+template<typename T>
+using RhoChunk = Chunk<Scalar<T>, 0u, dim_x, dim_y>;
+
+template<typename T, typename Desc>
+using MacroStruct = Struct<
+	Member<VelChunk<T,Desc>, "momentum">,
+	Member<RhoChunk<T>, "density">,
+	Member<ScalarChunk<T,Desc>, "porosity">,
+	Member<VelChunk<T,Desc>, "force">
+>;
+
+template<typename T, typename Desc>
+using ParticleSpheroidGroup = ParticleGroup<T,Desc::D,1u,coll::Spheroid<T>>;
+}
+
+}
+}
diff --git a/examples/dkt/init.hpp b/examples/dkt/init.hpp
new file mode 100644
index 0000000..af7eb89
--- /dev/null
+++ b/examples/dkt/init.hpp
@@ -0,0 +1,121 @@
+#pragma once
+
+#include "common.hpp"
+
+namespace kel {
+namespace lbm {
+
+template<typename T, typename Desc, typename Coll>
+saw::error_or<void> init(
+		const converter<T>& conv,
+		saw::data<sch::ChunkStruct<T,Desc>>& fields,
+		saw::data<sch::MacroStruct<T,Desc>>& macros,
+		saw::data<sch::ParticleSpheroidGroup<T,Desc>>& particles
+){
+	(void) conv;
+
+	auto& info_f = fields.template get<"info">();
+	auto& porous_f = macros.template get<"porosity">();
+	// Set everything as walls
+	iterator<Desc::D>::apply(
+		[&](auto& index){
+			info_f.at(index).set(1u);
+		},
+		{},
+		info_f.get_dims(),
+		{}
+	);
+	// Fluid
+	iterator<Desc::D>::apply(
+		[&](auto& index){
+			info_f.at(index).set(2u);
+		},
+		{},
+		info_f.get_dims(),
+		{{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}}
+	);
+	//
+	auto& df_f = fields.template get<"dfs_old">();
+	auto& rho_f = macros.template get<"density">();
+	auto& vel_f = macros.template get<"momentum">();
+	auto& por_f = macros.template get<"porosity">();
+	
+	iterator<Desc::D>::apply(
+		[&](auto& index){
+			auto& df = df_f.at(index);
+			auto& rho = rho_f.at(index);
+			por_f.at(index).at({}) = {1};
+			rho.at({}) = {1};
+			auto& vel = vel_f.at(index);
+			auto eq = equilibrium<T,Desc>(rho,vel);
+
+			df = eq;
+		},
+		{},// 0-index
+		df_f.get_dims()
+	);
+
+	iterator<Desc::D>::apply(
+		[&](auto& index){
+			auto& df = df_f.at(index);
+			auto& rho = rho_f.at(index);
+			rho.at({}) = {1};
+			auto& vel = vel_f.at(index);
+			if(info_f.at(index).get() == 2u){
+				vel.at({{0u}}) = 0.0;
+			}
+			auto eq = equilibrium<T,Desc>(rho,vel);
+
+			df = eq;
+		},
+		{},// 0-index
+		df_f.get_dims(),
+		{{1u,1u}}
+	);
+
+	{
+		saw::data<sch::Scalar<T>> radius_p;
+		radius_p.at({}).set(conv.meter_si_to_lbm({{0.05}}).handle().get());
+		std::cout<<"RADIUS: "<<radius_p.at({}).get()<<std::endl;
+		saw::data<sch::Scalar<T>> dense_p;
+		dense_p.at({}).set(1);
+		particles = create_spheroid_particle_group<T,Desc::D,1u>(radius_p, dense_p, {{64u}});
+		auto& parts = particles.template get<"particles">();
+
+		{
+			auto& p = parts.at({0u});
+			auto& prb = p.template get<"rigid_body">();
+			auto& p_pos = prb.template get<"position">();
+			auto& p_posold = prb.template get<"position_old">();
+			p_pos.at({{0u}}) = conv.meter_si_to_lbm({{0.16}}).handle();
+			p_pos.at({{1u}}) = conv.meter_si_to_lbm({{0.2}}).handle();
+
+			p_posold = p_pos;
+		}
+	}
+
+	return saw::make_void();
+}
+
+}
+}
diff --git a/examples/dkt/sim.cpp b/examples/dkt/sim.cpp
new file mode 100644
index 0000000..b3422a5
--- /dev/null
+++ b/examples/dkt/sim.cpp
@@ -0,0 +1,238 @@
+#include "common.hpp"
+#include "init.hpp"
+#include "step.hpp"
+
+#include <forstio/codec/args.hpp>
+
+/**
+ * For deciding what parameters to use maybe?
+ */
+namespace args {
+using namespace saw::schema;
+
+using LbmArgsStruct = Struct<
+	Member<String, "coupling">
+>;
+
+using LbmArgs = Args<
+	LbmArgsStruct,
+	Tuple<>
+>;
+}
+
+template<typename T, typename Desc, typename Coll>
+saw::error_or<void> lbm_main(const saw::data<args::LbmArgs>& args){
+	using namespace kel::lbm;
+
+	auto& an = args.template get<"args">();
+
+	using dfi = df_info<T,Desc>;
+
+	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 / "stdkr_2d_gpu" / an.template get<"coupling">().stl_view();
+
+	{
+		std::error_code ec;
+		std::filesystem::create_directories(out_dir,ec);
+		if(ec != std::errc{}){
+			return saw::make_error<saw::err::critical>("Could not create output directory");
+		}
+	}
+
+	converter<T> conv {
+		// delta_x
+		{{0.41 / dim_y}},
+		// delta_t
+		{{1.067e-4}}
+	};
+	
+	print_lbm_meta<T,Desc>(conv,{1e-3},{1e-4},{0.5 * dim_y});
+
+	// saw::data<sch::FixedArray<sch::UInt64,Desc::D>> meta{{dim_x,dim_y}};
+	auto lbm_data_ptr = saw::heap<saw::data<sch::ChunkStruct<T,Desc>>>();
+	auto lbm_macro_data_ptr = saw::heap<saw::data<sch::MacroStruct<T,Desc>>>();
+	auto lbm_particle_data_ptr = saw::heap<saw::data<sch::ParticleSpheroidGroup<T,Desc>>>();
+
+	std::cout<<"Estimated Bytes: "<<memory_estimate<sch::ChunkStruct<T,Desc>,sch::MacroStruct<T,Desc>>().get()<<std::endl;
+
+	auto eo_aio = saw::setup_async_io();
+	if(eo_aio.is_error()){
+		return std::move(eo_aio.get_error());
+	}
+	auto& aio = eo_aio.get_value();
+	saw::wait_scope wait{aio.event_loop};
+
+	bool krun = true;
+	bool print_status = false;
+	aio.event_port.on_signal(saw::Signal::Terminate).then([&](){
+		krun = false;
+	}).detach();
+	aio.event_port.on_signal(saw::Signal::User1).then([&](){
+		print_status = true;
+	}).detach();
+
+	device dev;
+
+	auto& sycl_q = dev.get_handle();
+
+	sycl_q.wait();
+	{
+		auto eov = init<T,Desc,Coll>(conv,*lbm_data_ptr,*lbm_macro_data_ptr,*lbm_particle_data_ptr);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+	{
+		auto eov = write_vtk_file(out_dir,"initial_state",0u,*lbm_data_ptr);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+
+	saw::data<sch::ChunkStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_data{sycl_q};
+	saw::data<sch::MacroStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_macro_data{sycl_q};
+	saw::data<sch::ParticleSpheroidGroup<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_particle_data{sycl_q};
+	sycl_q.wait();
+
+	{
+		auto eov = dev.copy_to_device(*lbm_data_ptr,lbm_sycl_data);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+	{
+		auto eov = dev.copy_to_device(*lbm_macro_data_ptr,lbm_sycl_macro_data);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+	{
+		auto eov = dev.copy_to_device(*lbm_particle_data_ptr,lbm_sycl_particle_data);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+	sycl_q.wait();
+	auto lsd_view = make_view(lbm_sycl_data);
+	auto lsdm_view = make_view(lbm_sycl_macro_data);
+	auto lsdp_view = make_view(lbm_sycl_particle_data);
+	
+	saw::data<sch::UInt64> time_steps{64u*1024ul};
+
+	auto& info_f = lsd_view.template get<"info">();
+
+	for(saw::data<sch::UInt64> i{0u}; i < time_steps and krun; ++i){
+		// BC + Collision
+		{
+			auto eov = step<T,Desc,Coll>(conv,lsd_view,lsdm_view,lsdp_view,i,dev);
+			if(eov.is_error()){
+				return eov;
+			}
+		}
+
+		sycl_q.wait();
+		if(i.get() % 64u == 0u){
+			{
+				auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr);
+				if(eov.is_error()){
+					return eov;
+				}
+			}
+			{
+				auto eov = write_csv_file(out_dir,"stdkr_2d",i.get(), *lbm_macro_data_ptr);
+				if(eov.is_error()){
+					return eov;
+				}
+			}
+		}
+		// Stream
+		sycl_q.submit([&](acpp::sycl::handler& h){
+			component<T,Desc,cmpt::Stream,encode::Sycl<saw::encode::Native>> stream;
+
+			h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
+				saw::data<sch::FixedArray<sch::UInt64,Desc::D>> index;
+				for(uint64_t i = 0u; i < Desc::D; ++i){
+					index.at({{i}}).set(idx[i]);
+				}
+
+				auto info = info_f.at(index).get();
+
+				if(info > 0u){
+					stream.apply(lsd_view,index,i);
+				}
+			});
+		}).wait();
+		wait.poll();
+
+		if(print_status){
+			std::cout<<"Status: "<<i.get()<<" of "<<time_steps.get()<<" - "<<(i.template cast_to<sch::Float64>().get() * 100 / time_steps.get())<<"%"<<std::endl;
+			print_status = false;
+		}
+		print_progress_bar(i.get(), time_steps.get()-1u);
+	}
+	// After Loop
+	sycl_q.wait();
+	{
+		auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+	{
+		auto eov = write_vtk_file(out_dir,std::string{"stdkr_2d_"}+an.template get<"coupling">().stl_string(),time_steps.get(), *lbm_macro_data_ptr);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+
+	sycl_q.wait();
+	return saw::make_void();
+}
+
+using FloatT = kel::lbm::sch::Float32;
+using DescT = kel::lbm::sch::D2Q9;
+
+saw::error_or<void> k_main(int argc, char** argv){
+	using namespace kel::lbm;
+	
+	saw::data<args::LbmArgs> args;
+	{
+		auto eov = saw::parse_args(args,argc,argv);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+	auto& an = args.template get<"args">();
+	auto& coupling = an.template get<"coupling">();
+
+	if(coupling.stl_view() == "hlbm"){
+		return lbm_main<FloatT,DescT,method::Hlbm>(args);
+	} else if(coupling.stl_view() == "fplbm"){
+		return lbm_main<FloatT,DescT,method::FpLbm>(args);
+	} else if(coupling.stl_view() == "psm"){
+		return lbm_main<FloatT,DescT,method::Psm>(args);
+	}
+
+	return saw::make_error<saw::err::critical>("Invalid coupling");
+}
+
+
+int main(int argc, char** argv){
+	auto eov = k_main(argc, argv);
+	if(eov.is_error()){
+		auto& err = eov.get_error();
+		std::cerr<<"[Error] "<<err.get_category();
+		auto err_msg = err.get_message();
+		if(err_msg.size() > 0u){
+			std::cerr<<" - "<<err_msg;
+		}
+		std::cerr<<std::endl;
+		return err.get_id();
+	}
+	return 0;
+}
diff --git a/examples/dkt/step.hpp b/examples/dkt/step.hpp
new file mode 100644
index 0000000..cbdfd10
--- /dev/null
+++ b/examples/dkt/step.hpp
@@ -0,0 +1,257 @@
+#pragma once
+
+#include "common.hpp"
+
+namespace kel {
+namespace lbm {
+
+template<typename T, typename Desc, typename Coll>
+saw::error_or<void> step(
+		const converter<T>& conv,
+		saw::data<sch::Ptr<sch::ChunkStruct<T,Desc>>,encode::Sycl<saw::encode::Native>>& fields,
+		saw::data<sch::Ptr<sch::MacroStruct<T,Desc>>,encode::Sycl<saw::encode::Native>>& macros,
+		saw::data<sch::Ptr<sch::ParticleSpheroidGroup<T,Desc>>,encode::Sycl<saw::encode::Native>>& particles,
+		saw::data<sch::UInt64> t_i,
+		device& dev
+){
+	static constexpr saw::data<sch::UInt64> sub_steps{0u};
+
+	auto& q = dev.get_handle();
+	auto& info_f = fields.template get<"info">();
+	auto& porous_f = macros.template get<"porosity">();
+
+
+	
+	if constexpr ( std::is_same_v<Coll, method::Hlbm> ){
+		q.submit([&](acpp::sycl::handler& h){
+			component<T,Desc,cmpt::HlbmReset,encode::Sycl<saw::encode::Native>> hlbm_reset;
+			h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
+				saw::data<sch::FixedArray<sch::UInt64,Desc::D>> index;
+				for(uint64_t i = 0u; i < Desc::D; ++i){
+					index.at({{i}}).set(idx[i]);
+				}
+
+				hlbm_reset.apply(fields,macros,index,t_i);
+			});
+		}).wait();
+		q.submit([&](acpp::sycl::handler& h){
+			component<T,Desc,cmpt::HlbmOneParticle,encode::Sycl<saw::encode::Native>> hlbm_one_part;
+
+			h.parallel_for(acpp::sycl::range<1u>{particle_amount}, [=](acpp::sycl::id<1u> idx){
+				saw::data<sch::FixedArray<sch::UInt64,1u>> index;
+				for(uint64_t i = 0u; i < 1u; ++i){
+					index.at({{i}}).set(idx[i]);
+				}
+
+				hlbm_one_part.apply(fields,macros,particles,index,t_i,{0u});
+			});
+		}).wait();
+
+		q.submit([&](acpp::sycl::handler& h){
+			component<T,Desc,cmpt::Hlbm,encode::Sycl<saw::encode::Native>> collision{1.0};
+			component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
+			component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0};
+
+			h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
+				saw::data<sch::FixedArray<sch::UInt64,Desc::D>> index;
+				for(uint64_t i = 0u; i < Desc::D; ++i){
+					index.at({{i}}).set(idx[i]);
+				}
+
+				auto info = info_f.at(index);
+				
+				switch(info.get()){
+					case 0u:
+					break;
+					case 1u:
+						bb.apply(fields,index,t_i);
+					break;
+					case 2u:
+						collision.apply(fields,macros,index,t_i);
+						break;
+					case 3u:
+						{
+							component<T,Desc,cmpt::ZouHeVelocityX<true>,encode::Sycl<saw::encode::Native>> flow_in{
+								[&]() -> saw::data<sch::Vector<T,Desc::D>> {
+									saw::data<sch::Vector<T,Desc::D>> vel;
+									{
+										auto y_si = conv.meter_lbm_to_si(index.at({{1u}}).template cast_to<T>()).handle();
+										vel.at({{0u}}) = conv.velocity_si_to_lbm({{static_cast<typename saw::native_data_type<T>::type>((1.2 / 0.41) * y_si.get() - (1.2 / (0.41*0.41)) * y_si.get() * y_si.get())}}).handle();
+										vel.at({{1u}}) = 0.0f;
+									}
+									return vel;
+								}()
+							};
+							flow_in.apply(fields,index,t_i);
+						}
+						collision.apply(fields,macros,index,t_i);
+						break;
+					case 4u:
+						flow_out.apply(fields,index,t_i);
+						collision.apply(fields,macros,index,t_i);
+					break;
+					default:
+					break;
+				}
+			});
+		}).wait();
+	}else if constexpr ( std::is_same_v<Coll, method::FpLbm> ){
+		q.submit([&](acpp::sycl::handler& h){
+			component<T,Desc,cmpt::FpLbmReset,encode::Sycl<saw::encode::Native>> fplbm_reset;
+			h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
+				saw::data<sch::FixedArray<sch::UInt64,Desc::D>> index;
+				for(uint64_t i = 0u; i < Desc::D; ++i){
+					index.at({{i}}).set(idx[i]);
+				}
+				
+				fplbm_reset.apply(fields,macros,index,t_i);
+			});
+		}).wait();
+
+		q.submit([&](acpp::sycl::handler& h){
+			component<T,Desc,cmpt::FpLbmOneParticle,encode::Sycl<saw::encode::Native>> fplbm_one_part;
+
+			h.parallel_for(acpp::sycl::range<1u>{particle_amount}, [=](acpp::sycl::id<1u> idx){
+				saw::data<sch::FixedArray<sch::UInt64,1u>> index;
+				for(uint64_t i = 0u; i < 1u; ++i){
+					index.at({{i}}).set(idx[i]);
+				}
+
+				fplbm_one_part.apply(fields,macros,particles,index,t_i,{0u});
+			});
+		}).wait();
+
+		// auto coll_ev = 
+		q.submit([&](acpp::sycl::handler& h){
+			component<T,Desc,cmpt::BGK,encode::Sycl<saw::encode::Native>> bgk{1.0};
+			component<T,Desc,cmpt::FpLbm,encode::Sycl<saw::encode::Native>> collision{1.0};
+			component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
+			component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0};
+
+			h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
+				saw::data<sch::FixedArray<sch::UInt64,Desc::D>> index;
+				for(uint64_t i = 0u; i < Desc::D; ++i){
+					index.at({{i}}).set(idx[i]);
+				}
+
+				auto info = info_f.at(index);
+				
+				switch(info.get()){
+					case 0u:
+					break;
+					case 1u:
+						bb.apply(fields,index,t_i);
+					break;
+					case 2u:
+						collision.apply(fields,macros,index,t_i);
+						break;
+					case 3u:
+						{
+							component<T,Desc,cmpt::ZouHeVelocityX<true>,encode::Sycl<saw::encode::Native>> flow_in{
+								[&]() -> saw::data<sch::Vector<T,Desc::D>> {
+									saw::data<sch::Vector<T,Desc::D>> vel;
+									{
+										auto y_si = conv.meter_lbm_to_si(index.at({{1u}}).template cast_to<T>()).handle();
+										vel.at({{0u}}) = conv.velocity_si_to_lbm({{static_cast<typename saw::native_data_type<T>::type>((1.2 / 0.41) * y_si.get() - (1.2 / (0.41*0.41)) * y_si.get() * y_si.get())}}).handle();
+										vel.at({{1u}}) = 0.0f;
+									}
+									return vel;
+								}()
+							};
+							flow_in.apply(fields,index,t_i);
+						}
+						bgk.apply(fields,macros,index,t_i);
+						break;
+					case 4u:
+						flow_out.apply(fields,index,t_i);
+						bgk.apply(fields,macros,index,t_i);
+					break;
+					default:
+					break;
+				}
+			});
+		}).wait();
+	}else if constexpr ( std::is_same_v<Coll, method::Psm> ){
+		q.submit([&](acpp::sycl::handler& h){
+			component<T,Desc,cmpt::PsmReset,encode::Sycl<saw::encode::Native>> fplbm_reset;
+			h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
+				saw::data<sch::FixedArray<sch::UInt64,Desc::D>> index;
+				for(uint64_t i = 0u; i < Desc::D; ++i){
+					index.at({{i}}).set(idx[i]);
+				}
+				
+				fplbm_reset.apply(fields,macros,index,t_i);
+			});
+		}).wait();
+
+		q.submit([&](acpp::sycl::handler& h){
+			component<T,Desc,cmpt::PsmOneParticle,encode::Sycl<saw::encode::Native>> fplbm_one_part;
+
+			h.parallel_for(acpp::sycl::range<1u>{particle_amount}, [=](acpp::sycl::id<1u> idx){
+				saw::data<sch::FixedArray<sch::UInt64,1u>> index;
+				for(uint64_t i = 0u; i < 1u; ++i){
+					index.at({{i}}).set(idx[i]);
+				}
+
+				fplbm_one_part.apply(fields,macros,particles,index,t_i,{0u});
+			});
+		}).wait();
+
+		// auto coll_ev = 
+		q.submit([&](acpp::sycl::handler& h){
+			component<T,Desc,cmpt::BGK,encode::Sycl<saw::encode::Native>> bgk{1.0};
+			component<T,Desc,cmpt::Psm,encode::Sycl<saw::encode::Native>> collision{1.0};
+			component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
+			component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0};
+
+			h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
+				saw::data<sch::FixedArray<sch::UInt64,Desc::D>> index;
+				for(uint64_t i = 0u; i < Desc::D; ++i){
+					index.at({{i}}).set(idx[i]);
+				}
+
+				auto info = info_f.at(index);
+				
+				switch(info.get()){
+					case 0u:
+					break;
+					case 1u:
+						bb.apply(fields,index,t_i);
+					break;
+					case 2u:
+						collision.apply(fields,macros,index,t_i);
+						break;
+					case 3u:
+						{
+							component<T,Desc,cmpt::ZouHeVelocityX<true>,encode::Sycl<saw::encode::Native>> flow_in{
+								[&]() -> saw::data<sch::Vector<T,Desc::D>> {
+									saw::data<sch::Vector<T,Desc::D>> vel;
+									{
+										auto y_si = conv.meter_lbm_to_si(index.at({{1u}}).template cast_to<T>()).handle();
+										vel.at({{0u}}) = conv.velocity_si_to_lbm({{static_cast<typename saw::native_data_type<T>::type>((1.2 / 0.41) * y_si.get() - (1.2 / (0.41*0.41)) * y_si.get() * y_si.get())}}).handle();
+										vel.at({{1u}}) = 0.0f;
+									}
+									return vel;
+								}()
+							};
+							flow_in.apply(fields,index,t_i);
+						}
+						bgk.apply(fields,macros,index,t_i);
+						break;
+					case 4u:
+						flow_out.apply(fields,index,t_i);
+						bgk.apply(fields,macros,index,t_i);
+					break;
+					default:
+					break;
+				}
+			});
+		}).wait();
+
+	}
+
+	return saw::make_void();
+}
+
+}
+}
diff --git a/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp b/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp
index b41f850..b3422a5 100644
--- a/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp
+++ b/examples/schaefer_turek_durst_krause_rannbacher/sim.cpp
@@ -214,6 +214,8 @@ saw::error_or<void> k_main(int argc, char** argv){
 		return lbm_main<FloatT,DescT,method::Hlbm>(args);
 	} else if(coupling.stl_view() == "fplbm"){
 		return lbm_main<FloatT,DescT,method::FpLbm>(args);
+	} else if(coupling.stl_view() == "psm"){
+		return lbm_main<FloatT,DescT,method::Psm>(args);
 	}
 
 	return saw::make_error<saw::err::critical>("Invalid coupling");
diff --git a/examples/schaefer_turek_durst_krause_rannbacher/step.hpp b/examples/schaefer_turek_durst_krause_rannbacher/step.hpp
index 554a6ea..cbdfd10 100644
--- a/examples/schaefer_turek_durst_krause_rannbacher/step.hpp
+++ b/examples/schaefer_turek_durst_krause_rannbacher/step.hpp
@@ -172,6 +172,82 @@ saw::error_or<void> step(
 			});
 		}).wait();
 	}else if constexpr ( std::is_same_v<Coll, method::Psm> ){
+		q.submit([&](acpp::sycl::handler& h){
+			component<T,Desc,cmpt::PsmReset,encode::Sycl<saw::encode::Native>> fplbm_reset;
+			h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
+				saw::data<sch::FixedArray<sch::UInt64,Desc::D>> index;
+				for(uint64_t i = 0u; i < Desc::D; ++i){
+					index.at({{i}}).set(idx[i]);
+				}
+				
+				fplbm_reset.apply(fields,macros,index,t_i);
+			});
+		}).wait();
+
+		q.submit([&](acpp::sycl::handler& h){
+			component<T,Desc,cmpt::PsmOneParticle,encode::Sycl<saw::encode::Native>> fplbm_one_part;
+
+			h.parallel_for(acpp::sycl::range<1u>{particle_amount}, [=](acpp::sycl::id<1u> idx){
+				saw::data<sch::FixedArray<sch::UInt64,1u>> index;
+				for(uint64_t i = 0u; i < 1u; ++i){
+					index.at({{i}}).set(idx[i]);
+				}
+
+				fplbm_one_part.apply(fields,macros,particles,index,t_i,{0u});
+			});
+		}).wait();
+
+		// auto coll_ev = 
+		q.submit([&](acpp::sycl::handler& h){
+			component<T,Desc,cmpt::BGK,encode::Sycl<saw::encode::Native>> bgk{1.0};
+			component<T,Desc,cmpt::Psm,encode::Sycl<saw::encode::Native>> collision{1.0};
+			component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
+			component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0};
+
+			h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
+				saw::data<sch::FixedArray<sch::UInt64,Desc::D>> index;
+				for(uint64_t i = 0u; i < Desc::D; ++i){
+					index.at({{i}}).set(idx[i]);
+				}
+
+				auto info = info_f.at(index);
+				
+				switch(info.get()){
+					case 0u:
+					break;
+					case 1u:
+						bb.apply(fields,index,t_i);
+					break;
+					case 2u:
+						collision.apply(fields,macros,index,t_i);
+						break;
+					case 3u:
+						{
+							component<T,Desc,cmpt::ZouHeVelocityX<true>,encode::Sycl<saw::encode::Native>> flow_in{
+								[&]() -> saw::data<sch::Vector<T,Desc::D>> {
+									saw::data<sch::Vector<T,Desc::D>> vel;
+									{
+										auto y_si = conv.meter_lbm_to_si(index.at({{1u}}).template cast_to<T>()).handle();
+										vel.at({{0u}}) = conv.velocity_si_to_lbm({{static_cast<typename saw::native_data_type<T>::type>((1.2 / 0.41) * y_si.get() - (1.2 / (0.41*0.41)) * y_si.get() * y_si.get())}}).handle();
+										vel.at({{1u}}) = 0.0f;
+									}
+									return vel;
+								}()
+							};
+							flow_in.apply(fields,index,t_i);
+						}
+						bgk.apply(fields,macros,index,t_i);
+						break;
+					case 4u:
+						flow_out.apply(fields,index,t_i);
+						bgk.apply(fields,macros,index,t_i);
+					break;
+					default:
+					break;
+				}
+			});
+		}).wait();
+
 	}
 
 	return saw::make_void();
diff --git a/examples/stdkr_2d_bgk/.nix/derivation.nix b/examples/stdkr_2d_bgk/.nix/derivation.nix
new file mode 100644
index 0000000..766e849
--- /dev/null
+++ b/examples/stdkr_2d_bgk/.nix/derivation.nix
@@ -0,0 +1,42 @@
+{ lib
+, stdenv
+, scons
+, clang-tools
+, forstio
+, python3
+, pname
+, version
+, adaptive-cpp
+, kel
+}:
+
+let
+in stdenv.mkDerivation {
+	pname = "${pname}-examples-stdkr_2d_gpu_bgk";
+	inherit version;
+	src = ./..;
+
+	nativeBuildInputs = [
+		scons
+		clang-tools
+		python3
+	];
+
+	buildInputs = [
+		forstio.core
+		forstio.async
+		forstio.codec
+		forstio.codec-unit
+		forstio.io
+    forstio.remote
+    forstio.remote-filesystem
+		forstio.codec-json
+		adaptive-cpp
+		kel.lbm.core
+		kel.lbm.sycl
+	];
+
+	preferLocalBuild = true;
+
+	outputs = [ "out" "dev" ];
+}
diff --git a/examples/stdkr_2d_bgk/SConscript b/examples/stdkr_2d_bgk/SConscript
new file mode 100644
index 0000000..b33b7ae
--- /dev/null
+++ b/examples/stdkr_2d_bgk/SConscript
@@ -0,0 +1,34 @@
+#!/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['CXX'] = 'syclcc-clang';
+examples_env['CXXFLAGS'] += ['-O3'];
+
+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, ['sim.cpp'], shared=False);
+examples_env.poiseulle_2d_gpu = examples_env.Program('#bin/stdkr_2d', [examples_objects]);
+
+# Set Alias
+env.examples = [
+	examples_env.poiseulle_2d_gpu
+];
+env.Alias('examples', env.examples);
+env.targets += ['examples'];
+env.Install('$prefix/bin/', env.examples);
diff --git a/examples/stdkr_2d_bgk/SConstruct b/examples/stdkr_2d_bgk/SConstruct
new file mode 100644
index 0000000..0611b67
--- /dev/null
+++ b/examples/stdkr_2d_bgk/SConstruct
@@ -0,0 +1,81 @@
+#!/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',
+        'forstio-async',
+        'forstio-io'
+    ]
+);
+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/stdkr_2d_bgk/common.hpp b/examples/stdkr_2d_bgk/common.hpp
new file mode 100644
index 0000000..eff9d84
--- /dev/null
+++ b/examples/stdkr_2d_bgk/common.hpp
@@ -0,0 +1,70 @@
+#pragma once
+
+#include <kel/lbm/sycl/lbm.hpp>
+#include <kel/lbm/lbm.hpp>
+#include <kel/lbm/particle.hpp>
+#include <kel/lbm/particle/particle.hpp>
+
+#include <forstio/io/io.hpp>
+#include <forstio/remote/filesystem/easy.hpp>
+#include <forstio/codec/json/json.hpp>
+#include <forstio/codec/simple.hpp>
+
+namespace kel {
+namespace lbm {
+namespace method {
+struct Bgk {};
+}
+
+/**
+ * length 2.2m
+ * height 0.41m
+ * viscosity 10^-3
+ * particle diameter 0.1m
+ * position at (0.16m,0.2m)
+ */
+
+constexpr uint64_t dim_y = 512ul;
+constexpr uint64_t dim_x = static_cast<uint64_t>(dim_y * (220.0 / 41.0));
+
+constexpr uint64_t particle_amount = 1ul;
+
+namespace sch {
+using namespace saw::schema;
+
+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 ScalarChunk = Chunk<Scalar<T>, 0u, dim_x, dim_y>;
+
+template<typename T, typename Desc>
+using VectorChunk = Chunk<Vector<T,Desc::D>, 0u, dim_x, dim_y>;
+
+template<typename T, typename Desc>
+using ChunkStruct = Struct<
+	Member<InfoChunk, "info">,
+	Member<DfChunk<T,Desc>, "dfs">,
+	Member<DfChunk<T,Desc>, "dfs_old">
+>;
+
+template<typename T, typename Desc>
+using VelChunk = Chunk<Vector<T,Desc::D>, 0u, dim_x, dim_y>;
+
+template<typename T>
+using RhoChunk = Chunk<Scalar<T>, 0u, dim_x, dim_y>;
+
+template<typename T, typename Desc>
+using MacroStruct = Struct<
+	Member<VelChunk<T,Desc>, "momentum">,
+	Member<RhoChunk<T>, "density">
+>;
+
+template<typename T, typename Desc>
+using ParticleSpheroidGroup = ParticleGroup<T,Desc::D,1u,coll::Spheroid<T>>;
+}
+
+}
+}
diff --git a/examples/stdkr_2d_bgk/init.hpp b/examples/stdkr_2d_bgk/init.hpp
new file mode 100644
index 0000000..abc2a26
--- /dev/null
+++ b/examples/stdkr_2d_bgk/init.hpp
@@ -0,0 +1,124 @@
+#pragma once
+
+#include "common.hpp"
+
+namespace kel {
+namespace lbm {
+
+template<typename T, typename Desc, typename Coll>
+saw::error_or<void> init(
+		const converter<T>& conv,
+		saw::data<sch::ChunkStruct<T,Desc>>& fields,
+		saw::data<sch::MacroStruct<T,Desc>>& macros
+){
+	(void) conv;
+
+	auto& info_f = fields.template get<"info">();
+	auto& porous_f = macros.template get<"porosity">();
+	// Set everything as walls
+	iterator<Desc::D>::apply(
+		[&](auto& index){
+			info_f.at(index).set(1u);
+		},
+		{},
+		info_f.get_dims(),
+		{}
+	);
+	// Fluid
+	iterator<Desc::D>::apply(
+		[&](auto& index){
+			info_f.at(index).set(2u);
+		},
+		{},
+		info_f.get_dims(),
+		{{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}}
+	);
+	//
+	auto& df_f = fields.template get<"dfs_old">();
+	auto& rho_f = macros.template get<"density">();
+	auto& vel_f = macros.template get<"momentum">();
+	auto& por_f = macros.template get<"porosity">();
+	
+	iterator<Desc::D>::apply(
+		[&](auto& index){
+			auto& df = df_f.at(index);
+			auto& rho = rho_f.at(index);
+			por_f.at(index).at({}) = {1};
+			rho.at({}) = {1};
+			auto& vel = vel_f.at(index);
+			auto eq = equilibrium<T,Desc>(rho,vel);
+
+			df = eq;
+		},
+		{},// 0-index
+		df_f.get_dims()
+	);
+
+	iterator<Desc::D>::apply(
+		[&](auto& index){
+			auto& df = df_f.at(index);
+			auto& rho = rho_f.at(index);
+			rho.at({}) = {1};
+			auto& vel = vel_f.at(index);
+			if(info_f.at(index).get() == 2u){
+				vel.at({{0u}}) = 0.0;
+			}
+			auto eq = equilibrium<T,Desc>(rho,vel);
+
+			df = eq;
+		},
+		{},// 0-index
+		df_f.get_dims(),
+		{{1u,1u}}
+	);
+
+	{
+
+	}
+	{
+		saw::data<sch::Scalar<T>> radius_p;
+		radius_p.at({}).set(conv.meter_si_to_lbm({{0.05}}).handle().get());
+		std::cout<<"RADIUS: "<<radius_p.at({}).get()<<std::endl;
+		saw::data<sch::Scalar<T>> dense_p;
+		dense_p.at({}).set(1);
+		/*
+		auto& parts = particles.template get<"particles">();
+
+		{
+			auto& p = parts.at({0u});
+			auto& prb = p.template get<"rigid_body">();
+			auto& p_pos = prb.template get<"position">();
+			auto& p_posold = prb.template get<"position_old">();
+			p_pos.at({{0u}}) = conv.meter_si_to_lbm({{0.16}}).handle();
+			p_pos.at({{1u}}) = conv.meter_si_to_lbm({{0.2}}).handle();
+
+			p_posold = p_pos;
+		}
+		*/
+	}
+
+	return saw::make_void();
+}
+
+}
+}
diff --git a/examples/stdkr_2d_bgk/sim.cpp b/examples/stdkr_2d_bgk/sim.cpp
new file mode 100644
index 0000000..558a219
--- /dev/null
+++ b/examples/stdkr_2d_bgk/sim.cpp
@@ -0,0 +1,230 @@
+#include "common.hpp"
+#include "init.hpp"
+#include "step.hpp"
+
+#include <forstio/codec/args.hpp>
+
+
+/**
+ * For deciding what parameters to use maybe?
+ */
+namespace args {
+using namespace saw::schema;
+
+using LbmArgsStruct = Struct<
+>;
+
+using LbmArgs = Args<
+	LbmArgsStruct,
+	Tuple<>
+>;
+}
+
+template<typename T, typename Desc, typename Coll>
+saw::error_or<void> lbm_main(const saw::data<args::LbmArgs>& args){
+	using namespace kel::lbm;
+
+	auto& an = args.template get<"args">();
+
+	using dfi = df_info<T,Desc>;
+
+	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 / "stdkr_2d_gpu" / "bgk";
+
+	{
+		std::error_code ec;
+		std::filesystem::create_directories(out_dir,ec);
+		if(ec != std::errc{}){
+			return saw::make_error<saw::err::critical>("Could not create output directory");
+		}
+	}
+
+	converter<T> conv {
+		// delta_x
+		{{0.41 / dim_y}},
+		// delta_t
+		{{1.067e-4}}
+	};
+	
+	print_lbm_meta<T,Desc>(conv,{1e-3},{1e-4},{0.5 * dim_y});
+
+	// saw::data<sch::FixedArray<sch::UInt64,Desc::D>> meta{{dim_x,dim_y}};
+	auto lbm_data_ptr = saw::heap<saw::data<sch::ChunkStruct<T,Desc>>>();
+	auto lbm_macro_data_ptr = saw::heap<saw::data<sch::MacroStruct<T,Desc>>>();
+	auto lbm_particle_data_ptr = saw::heap<saw::data<sch::ParticleSpheroidGroup<T,Desc>>>();
+
+	std::cout<<"Estimated Bytes: "<<memory_estimate<sch::ChunkStruct<T,Desc>,sch::MacroStruct<T,Desc>>().get()<<std::endl;
+
+	auto eo_aio = saw::setup_async_io();
+	if(eo_aio.is_error()){
+		return std::move(eo_aio.get_error());
+	}
+	auto& aio = eo_aio.get_value();
+	saw::wait_scope wait{aio.event_loop};
+
+	bool krun = true;
+	bool print_status = false;
+	aio.event_port.on_signal(saw::Signal::Terminate).then([&](){
+		krun = false;
+	}).detach();
+	aio.event_port.on_signal(saw::Signal::User1).then([&](){
+		print_status = true;
+	}).detach();
+
+	device dev;
+
+	auto& sycl_q = dev.get_handle();
+
+	sycl_q.wait();
+	{
+		auto eov = init<T,Desc,Coll>(conv,*lbm_data_ptr,*lbm_macro_data_ptr,*lbm_particle_data_ptr);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+	{
+		auto eov = write_vtk_file(out_dir,"initial_state",0u,*lbm_data_ptr);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+
+	saw::data<sch::ChunkStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_data{sycl_q};
+	saw::data<sch::MacroStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_macro_data{sycl_q};
+	saw::data<sch::ParticleSpheroidGroup<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_particle_data{sycl_q};
+	sycl_q.wait();
+
+	{
+		auto eov = dev.copy_to_device(*lbm_data_ptr,lbm_sycl_data);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+	{
+		auto eov = dev.copy_to_device(*lbm_macro_data_ptr,lbm_sycl_macro_data);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+	{
+		auto eov = dev.copy_to_device(*lbm_particle_data_ptr,lbm_sycl_particle_data);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+	sycl_q.wait();
+	auto lsd_view = make_view(lbm_sycl_data);
+	auto lsdm_view = make_view(lbm_sycl_macro_data);
+	auto lsdp_view = make_view(lbm_sycl_particle_data);
+	
+	saw::data<sch::UInt64> time_steps{64u*1024ul};
+
+	auto& info_f = lsd_view.template get<"info">();
+
+	for(saw::data<sch::UInt64> i{0u}; i < time_steps and krun; ++i){
+		// BC + Collision
+		{
+			auto eov = step<T,Desc,Coll>(conv,lsd_view,lsdm_view,lsdp_view,i,dev);
+			if(eov.is_error()){
+				return eov;
+			}
+		}
+
+		sycl_q.wait();
+		if(i.get() % 64u == 0u){
+			{
+				auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr);
+				if(eov.is_error()){
+					return eov;
+				}
+			}
+			{
+				auto eov = write_csv_file(out_dir,"stdkr_2d",i.get(), *lbm_macro_data_ptr);
+				if(eov.is_error()){
+					return eov;
+				}
+			}
+		}
+		// Stream
+		sycl_q.submit([&](acpp::sycl::handler& h){
+			component<T,Desc,cmpt::Stream,encode::Sycl<saw::encode::Native>> stream;
+
+			h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
+				saw::data<sch::FixedArray<sch::UInt64,Desc::D>> index;
+				for(uint64_t i = 0u; i < Desc::D; ++i){
+					index.at({{i}}).set(idx[i]);
+				}
+
+				auto info = info_f.at(index).get();
+
+				if(info > 0u){
+					stream.apply(lsd_view,index,i);
+				}
+			});
+		}).wait();
+		wait.poll();
+
+		if(print_status){
+			std::cout<<"Status: "<<i.get()<<" of "<<time_steps.get()<<" - "<<(i.template cast_to<sch::Float64>().get() * 100 / time_steps.get())<<"%"<<std::endl;
+			print_status = false;
+		}
+		print_progress_bar(i.get(), time_steps.get()-1u);
+	}
+	// After Loop
+	sycl_q.wait();
+	{
+		auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+	{
+		auto eov = write_vtk_file(out_dir,"stdkr_2d_bgk",time_steps.get(), *lbm_macro_data_ptr);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+
+	sycl_q.wait();
+	return saw::make_void();
+}
+
+using FloatT = kel::lbm::sch::Float32;
+using DescT = kel::lbm::sch::D2Q9;
+
+saw::error_or<void> k_main(int argc, char** argv){
+	using namespace kel::lbm;
+	
+	saw::data<args::LbmArgs> args;
+	{
+		auto eov = saw::parse_args(args,argc,argv);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
+
+	return lbm_main<FloatT,DescT,method::Bgk>(args);
+
+	return saw::make_error<saw::err::critical>("Invalid coupling");
+}
+
+
+int main(int argc, char** argv){
+	auto eov = k_main(argc, argv);
+	if(eov.is_error()){
+		auto& err = eov.get_error();
+		std::cerr<<"[Error] "<<err.get_category();
+		auto err_msg = err.get_message();
+		if(err_msg.size() > 0u){
+			std::cerr<<" - "<<err_msg;
+		}
+		std::cerr<<std::endl;
+		return err.get_id();
+	}
+	return 0;
+}
diff --git a/examples/stdkr_2d_bgk/step.hpp b/examples/stdkr_2d_bgk/step.hpp
new file mode 100644
index 0000000..667f5a3
--- /dev/null
+++ b/examples/stdkr_2d_bgk/step.hpp
@@ -0,0 +1,77 @@
+#pragma once
+
+#include "common.hpp"
+
+namespace kel {
+namespace lbm {
+
+template<typename T, typename Desc, typename Coll>
+saw::error_or<void> step(
+		const converter<T>& conv,
+		saw::data<sch::Ptr<sch::ChunkStruct<T,Desc>>,encode::Sycl<saw::encode::Native>>& fields,
+		saw::data<sch::Ptr<sch::MacroStruct<T,Desc>>,encode::Sycl<saw::encode::Native>>& macros,
+		saw::data<sch::UInt64> t_i,
+		device& dev
+){
+	static constexpr saw::data<sch::UInt64> sub_steps{0u};
+
+	auto& q = dev.get_handle();
+	auto& info_f = fields.template get<"info">();
+	auto& porous_f = macros.template get<"porosity">();
+
+	{
+		q.submit([&](acpp::sycl::handler& h){
+			component<T,Desc,cmpt::BGK,encode::Sycl<saw::encode::Native>> collision{1.0};
+			component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
+			component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0};
+
+			h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
+				saw::data<sch::FixedArray<sch::UInt64,Desc::D>> index;
+				for(uint64_t i = 0u; i < Desc::D; ++i){
+					index.at({{i}}).set(idx[i]);
+				}
+
+				auto info = info_f.at(index);
+				
+				switch(info.get()){
+					case 0u:
+					break;
+					case 1u:
+						bb.apply(fields,index,t_i);
+					break;
+					case 2u:
+						collision.apply(fields,macros,index,t_i);
+						break;
+					case 3u:
+						{
+							component<T,Desc,cmpt::ZouHeVelocityX<true>,encode::Sycl<saw::encode::Native>> flow_in{
+								[&]() -> saw::data<sch::Vector<T,Desc::D>> {
+									saw::data<sch::Vector<T,Desc::D>> vel;
+									{
+										auto y_si = conv.meter_lbm_to_si(index.at({{1u}}).template cast_to<T>()).handle();
+										vel.at({{0u}}) = conv.velocity_si_to_lbm({{static_cast<typename saw::native_data_type<T>::type>((1.2 / 0.41) * y_si.get() - (1.2 / (0.41*0.41)) * y_si.get() * y_si.get())}}).handle();
+										vel.at({{1u}}) = 0.0f;
+									}
+									return vel;
+								}()
+							};
+							flow_in.apply(fields,index,t_i);
+						}
+						collision.apply(fields,macros,index,t_i);
+						break;
+					case 4u:
+						flow_out.apply(fields,index,t_i);
+						collision.apply(fields,macros,index,t_i);
+					break;
+					default:
+					break;
+				}
+			});
+		}).wait();
+	}
+
+	return saw::make_void();
+}
+
+}
+}
diff --git a/poc/component_generator/ops.hpp b/poc/component_generator/ops.hpp
new file mode 100644
index 0000000..f4449d1
--- /dev/null
+++ b/poc/component_generator/ops.hpp
@@ -0,0 +1,13 @@
+#pragma once
+
+namespace kel {
+/**
+ * With this I technically could do code reduction in a cell step
+ */
+
+template<typename A, typename B>
+struct Add {};
+
+template<typename L, typename R>
+struct Multiply {};
+}
diff --git a/scripts/python/csv_magnitude_of_sum.py b/scripts/python/csv_magnitude_of_sum.py
new file mode 100755
index 0000000..7ec25a1
--- /dev/null
+++ b/scripts/python/csv_magnitude_of_sum.py
@@ -0,0 +1,19 @@
+#!/usr/bin/env python3
+
+import argparse
+import numpy as np
+
+
+def main():
+    parser = argparse.ArgumentParser(description="Compute the sum of magnitude");
+    parser.add_argument("one");
+    args = parser.parse_args();
+
+    a = np.loadtxt(args.one, delimiter=",");
+
+    print(np.linalg.norm(np.sum(a,axis=0)));
+#enddef
+
+if __name__ == "__main__":
+    main();
+#enddef
diff --git a/scripts/python/csv_sum.py b/scripts/python/csv_sum.py
new file mode 100755
index 0000000..bc16006
--- /dev/null
+++ b/scripts/python/csv_sum.py
@@ -0,0 +1,19 @@
+#!/usr/bin/env python3
+
+import argparse
+import numpy as np
+
+
+def main():
+    parser = argparse.ArgumentParser(description="Compute the sum of magnitude");
+    parser.add_argument("one");
+    args = parser.parse_args();
+
+    a = np.loadtxt(args.one, delimiter=",");
+
+    print(np.sum(a,axis=0));
+#enddef
+
+if __name__ == "__main__":
+    main();
+#enddef
diff --git a/scripts/python/csv_magnitude_sum.py b/scripts/python/csv_sum_of_magnitude.py
index 773ce52..773ce52 100755
--- a/scripts/python/csv_magnitude_sum.py
+++ b/scripts/python/csv_sum_of_magnitude.py