Merge branch 'dev'

7 files changed, 0 insertions, 845 deletions

diff --git a/examples/one_particle_sedimentation/.nix/derivation.nix b/examples/one_particle_sedimentation/.nix/derivation.nix
deleted file mode 100644
index a062133..0000000
--- a/examples/one_particle_sedimentation/.nix/derivation.nix
+++ /dev/null
@@ -1,43 +0,0 @@
-{ lib
-, stdenv
-, scons
-, clang-tools
-, forstio
-, python3
-, pname
-, version
-, adaptive-cpp
-, kel
-, particle_coupling ? "fplbm"
-}:
-
-let
-in stdenv.mkDerivation {
-	pname = "${pname}-examples-ops_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/one_particle_sedimentation/SConscript b/examples/one_particle_sedimentation/SConscript
deleted file mode 100644
index 40f4052..0000000
--- a/examples/one_particle_sedimentation/SConscript
+++ /dev/null
@@ -1,34 +0,0 @@
-#!/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/ops_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/one_particle_sedimentation/SConstruct b/examples/one_particle_sedimentation/SConstruct
deleted file mode 100644
index 0611b67..0000000
--- a/examples/one_particle_sedimentation/SConstruct
+++ /dev/null
@@ -1,81 +0,0 @@
-#!/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/one_particle_sedimentation/common.hpp b/examples/one_particle_sedimentation/common.hpp
deleted file mode 100644
index be9a538..0000000
--- a/examples/one_particle_sedimentation/common.hpp
+++ /dev/null
@@ -1,71 +0,0 @@
-#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/one_particle_sedimentation/init.hpp b/examples/one_particle_sedimentation/init.hpp
deleted file mode 100644
index af7eb89..0000000
--- a/examples/one_particle_sedimentation/init.hpp
+++ /dev/null
@@ -1,121 +0,0 @@
-#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/one_particle_sedimentation/sim.cpp b/examples/one_particle_sedimentation/sim.cpp
deleted file mode 100644
index b3422a5..0000000
--- a/examples/one_particle_sedimentation/sim.cpp
+++ /dev/null
@@ -1,238 +0,0 @@
-#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/one_particle_sedimentation/step.hpp b/examples/one_particle_sedimentation/step.hpp
deleted file mode 100644
index 07f51d1..0000000
--- a/examples/one_particle_sedimentation/step.hpp
+++ /dev/null
@@ -1,257 +0,0 @@
-#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,{16u});
-			});
-		}).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,{16u});
-			});
-		}).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,{16u});
-			});
-		}).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();
-}
-
-}
-}