Final run

6 files changed, 92 insertions, 122 deletions

diff --git a/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp b/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp
index 73964b7..ef4639e 100644
--- a/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp
+++ b/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp
@@ -171,7 +171,7 @@ saw::error_or<void> step(
 	// auto coll_ev = 
 	q.submit([&](acpp::sycl::handler& h){
 		// Need nicer things to handle the flow. I see improvement here
-		component<T,Desc,cmpt::BGK, encode::Sycl<saw::encode::Native>> collision{0.65};
+		component<T,Desc,cmpt::BGK, encode::Sycl<saw::encode::Native>> collision{0.8};
 		component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
 		component<T,Desc,cmpt::AntiBounceBack<0u>,encode::Sycl<saw::encode::Native>> abb;
 
diff --git a/examples/poiseulle_particles_2d_fplbm_gpu/sim.cpp b/examples/poiseulle_particles_2d_fplbm_gpu/sim.cpp
index 93d4b46..3d16884 100644
--- a/examples/poiseulle_particles_2d_fplbm_gpu/sim.cpp
+++ b/examples/poiseulle_particles_2d_fplbm_gpu/sim.cpp
@@ -51,6 +51,7 @@ template<typename T, typename Desc>
 using MacroStruct = Struct<
 	Member<VelChunk<T,Desc>, "velocity">,
 	Member<RhoChunk<T>, "density">,
+	Member<VelChunk<T,Desc>, "force">,
 	Member<ScalarChunk<T,Desc>, "porosity">
 >;
 
@@ -144,6 +145,7 @@ saw::error_or<void> setup_initial_conditions(
 		{{1u,1u}}
 	);
 
+	/*
 	iterator<Desc::D>::apply(
 		[&](auto& index){
 			saw::data<sch::Vector<T,Desc::D>> middle, ind_vec;
@@ -162,6 +164,7 @@ saw::error_or<void> setup_initial_conditions(
 		{},// 0-index
 		df_f.get_dims()
 	);
+	*/
 
 	return saw::make_void();
 }
@@ -179,16 +182,20 @@ saw::error_or<void> step(
 
 	// auto coll_ev =
 	q.submit([&](acpp::sycl::handler& h){
-		component<T,Desc,cmpt::FpLbmOneParticleNoVelocity,encode::Sycl<saw::encode::Native>> collision{0.65};
+		component<T,Desc,cmpt::BGK, encode::Sycl<saw::encode::Native>> bgk{0.8};
+		component<T,Desc,cmpt::FpLbmOneParticleNoVelocity,encode::Sycl<saw::encode::Native>> collision{0.8};
 		component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
-		component<T,Desc,cmpt::AntiBounceBack<0u>,encode::Sycl<saw::encode::Native>> abb;
-
-		saw::data<sch::Scalar<T>> rho_b;
-		rho_b.at({}) = 1.0;
-		saw::data<sch::Vector<T,Desc::D>> vel_b;
-		vel_b.at({{0u}}) = 0.015;
 
-		component<T,Desc,cmpt::Equilibrium,encode::Sycl<saw::encode::Native>> equi{rho_b,vel_b};
+		saw::data<sch::Scalar<T>> eps;
+		eps.at({}).set(1.5f);
+		saw::data<sch::Scalar<T>> rad;
+		rad.at({}).set(dim_y*0.1);
+		saw::data<sch::Vector<T,Desc::D>> p_pos;
+		{
+			p_pos.at({{0u}}) = dim_x * 0.25;
+			p_pos.at({{1u}}) = dim_y * 0.5;
+		}
+		component<T,Desc,cmpt::OneParticleAt,encode::Sycl<saw::encode::Native>> opa{p_pos,rad,eps};
 
 		component<T,Desc,cmpt::ZouHeHorizontal<true>,encode::Sycl<saw::encode::Native>> flow_in{
 			[&](){
@@ -217,17 +224,18 @@ saw::error_or<void> step(
 					bb.apply(fields,index,t_i);
 				break;
 				case 2u:
+					opa.apply(macros,index,t_i);
 					collision.apply(fields,macros,index,t_i);
 					break;
 				case 3u:
 					flow_in.apply(fields,index,t_i);
 					// equi.apply(fields,index,t_i);
-					collision.apply(fields,macros,index,t_i);
+					bgk.apply(fields,macros,index,t_i);
 					break;
 				case 4u:
 					flow_out.apply(fields,index,t_i);
 					// equi.apply(fields,index,t_i);
-					collision.apply(fields,macros,index,t_i);
+					bgk.apply(fields,macros,index,t_i);
 				break;
 				default:
 				break;
@@ -349,20 +357,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 			}
 		}
 		sycl_q.wait();
-		if(i.get() % 32u == 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,"m",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;
@@ -402,6 +397,12 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 			return eov;
 		}
 	}
+	{
+		auto eov = write_csv_file(out_dir,"m",time_steps.get(), *lbm_macro_data_ptr);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
 
 	sycl_q.wait();
 	return saw::make_void();
diff --git a/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp b/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp
index b09aa64..b1bd37d 100644
--- a/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp
+++ b/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp
@@ -48,6 +48,7 @@ template<typename T, typename Desc>
 using MacroStruct = Struct<
 	Member<VelChunk<T,Desc>, "velocity">,
 	Member<RhoChunk<T>, "density">,
+	Member<VelChunk<T,Desc>, "force">,
 	Member<ScalarChunk<T,Desc>, "porosity">
 >;
 
@@ -140,7 +141,8 @@ saw::error_or<void> setup_initial_conditions(
 		df_f.get_dims(),
 		{{1u,1u}}
 	);
-	
+
+	/*
 	iterator<Desc::D>::apply(
 		[&](auto& index){
 			saw::data<sch::Vector<T,Desc::D>> middle, ind_vec;
@@ -159,6 +161,7 @@ saw::error_or<void> setup_initial_conditions(
 		{},// 0-index
 		df_f.get_dims()
 	);
+	*/
 
 	return saw::make_void();
 }
@@ -176,16 +179,23 @@ saw::error_or<void> step(
 
 	// auto coll_ev = 
 	q.submit([&](acpp::sycl::handler& h){
-		component<T,Desc,cmpt::Hlbm,encode::Sycl<saw::encode::Native>> collision{0.65};
+		component<T,Desc,cmpt::BGK, encode::Sycl<saw::encode::Native>> bgk{0.8};
+		component<T,Desc,cmpt::Hlbm,encode::Sycl<saw::encode::Native>> collision{0.8};
 		component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
 		component<T,Desc,cmpt::AntiBounceBack<0u>,encode::Sycl<saw::encode::Native>> abb;
 
-		saw::data<sch::Scalar<T>> rho_b;
-		rho_b.at({}) = 1.0;
-		saw::data<sch::Vector<T,Desc::D>> vel_b;
-		vel_b.at({{0u}}) = 0.015;
-
-		component<T,Desc,cmpt::Equilibrium,encode::Sycl<saw::encode::Native>> equi{rho_b,vel_b};
+		saw::data<sch::Scalar<T>> eps;
+		eps.at({}).set(1.5f);
+		saw::data<sch::Scalar<T>> rad;
+		rad.at({}).set(dim_y*0.1);
+		saw::data<sch::Vector<T,Desc::D>> p_pos;
+		{
+			p_pos.at({{0u}}) = dim_x * 0.25;
+			p_pos.at({{1u}}) = dim_y * 0.5;
+		}
+		component<T,Desc,cmpt::OneParticleAt,encode::Sycl<saw::encode::Native>> opa{p_pos,rad,eps};
+		
+		component<T,Desc,cmpt::ForceGather, encode::Sycl<saw::encode::Native>> fg;
 
 		component<T,Desc,cmpt::ZouHeHorizontal<true>,encode::Sycl<saw::encode::Native>> flow_in{
 			[&](){
@@ -214,32 +224,19 @@ saw::error_or<void> step(
 					bb.apply(fields,index,t_i);
 				break;
 				case 2u:
+					opa.apply(macros,index,t_i);
 					collision.apply(fields,macros,index,t_i);
-					{
-						saw::data<sch::Vector<T,Desc::D>> middle;
-						middle.at({{0u}}) = dim_x * 0.25;
-						middle.at({{1u}}) = dim_y * 0.5;
-
-						saw::data<sch::Vector<T,Desc::D>> ind_vec;
-						ind_vec.at({{0u}}) = index.at({{0u}}).template cast_to<T>();
-						ind_vec.at({{1u}}) = index.at({{1u}}).template cast_to<T>();
-
-						auto dist = middle - ind_vec;
-						auto dist_2 = saw::math::dot(dist,dist);
-						if(dist_2.at({}).get() < dim_y*dim_y*0.01){
-							porous_f.at(index).at({}) = 0.0;
-						}
-					}
+					fg.apply(fields,macros,index,t_i);
 					break;
 				case 3u:
 					flow_in.apply(fields,index,t_i);
 					// equi.apply(fields,index,t_i);
-					collision.apply(fields,macros,index,t_i);
+					bgk.apply(fields,macros,index,t_i);
 					break;
 				case 4u:
 					flow_out.apply(fields,index,t_i);
 					// equi.apply(fields,index,t_i);
-					collision.apply(fields,macros,index,t_i);
+					bgk.apply(fields,macros,index,t_i);
 				break;
 				default:
 				break;
diff --git a/examples/poiseulle_particles_2d_psm_gpu/sim.cpp b/examples/poiseulle_particles_2d_psm_gpu/sim.cpp
index b3ea7ae..7d3c6f5 100644
--- a/examples/poiseulle_particles_2d_psm_gpu/sim.cpp
+++ b/examples/poiseulle_particles_2d_psm_gpu/sim.cpp
@@ -49,6 +49,7 @@ template<typename T, typename Desc>
 using MacroStruct = Struct<
 	Member<VelChunk<T,Desc>, "velocity">,
 	Member<RhoChunk<T>, "density">,
+	Member<VelChunk<T,Desc>, "force">,
 	Member<ScalarChunk<T,Desc>, "porosity">
 >;
 
@@ -216,17 +217,9 @@ saw::error_or<void> step(
 	auto& porous_f = macros.template get<"porosity">();
 
 	q.submit([&](acpp::sycl::handler& h){
+		component<T,Desc,cmpt::BGK, encode::Sycl<saw::encode::Native>> bgk{0.8};
 		component<T,Desc,cmpt::PSM,encode::Sycl<saw::encode::Native>> collision{0.8};
 		component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
-		component<T,Desc,cmpt::AntiBounceBack<0u>,encode::Sycl<saw::encode::Native>> abb;
-
-		saw::data<sch::Scalar<T>> rho_b;
-		rho_b.at({}) = 1.0;
-		saw::data<sch::Vector<T,Desc::D>> vel_b;
-		vel_b.at({{0u}}) = 0.015;
-
-		component<T,Desc,cmpt::Equilibrium,encode::Sycl<saw::encode::Native>> equi{rho_b,vel_b};
-
 
 		saw::data<sch::Scalar<T>> eps;
 		eps.at({}).set(1.5f);
@@ -250,6 +243,7 @@ saw::error_or<void> step(
 		};
 		component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0};
 
+		component<T,Desc,cmpt::ForceGather, encode::Sycl<saw::encode::Native>> fg;
 
 		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;
@@ -268,16 +262,17 @@ saw::error_or<void> step(
 				case 2u:
 					opa.apply(macros,index,t_i);
 					collision.apply(fields,macros,index,t_i);
+					fg.apply(fields,macros,index,t_i);
 					break;
 				case 3u:
 					flow_in.apply(fields,index,t_i);
 					//equi.apply(fields,index,t_i);
-					collision.apply(fields,macros,index,t_i);
+					bgk.apply(fields,macros,index,t_i);
 					break;
 				case 4u:
 					flow_out.apply(fields,index,t_i);
 					// equi.apply(fields,index,t_i);
-					collision.apply(fields,macros,index,t_i);
+					bgk.apply(fields,macros,index,t_i);
 					break;
 				case 5u:
 					// Corners
diff --git a/examples/stokes_drag_particle_2d_hlbm_gpu/sim.cpp b/examples/stokes_drag_particle_2d_hlbm_gpu/sim.cpp
index a8ba21f..81498e2 100644
--- a/examples/stokes_drag_particle_2d_hlbm_gpu/sim.cpp
+++ b/examples/stokes_drag_particle_2d_hlbm_gpu/sim.cpp
@@ -13,7 +13,7 @@ namespace lbm {
 constexpr uint64_t dim_y = 1024ul;
 constexpr uint64_t dim_x = dim_y * 2ul;
 
-constexpr uint64_t particle_amount = 1ul;
+// constexpr uint64_t particle_amount = 1ul;
 
 namespace sch {
 using namespace saw::schema;
@@ -47,6 +47,7 @@ using RhoChunk = Chunk<Scalar<T>, 0u, dim_x, dim_y>;
 template<typename T, typename Desc>
 using MacroStruct = Struct<
 	Member<VelChunk<T,Desc>, "velocity">,
+	Member<VelChunk<T,Desc>, "force">,
 	Member<RhoChunk<T>, "density">,
 	Member<ScalarChunk<T,Desc>, "porosity">
 >;
@@ -158,25 +159,6 @@ saw::error_or<void> setup_initial_conditions(
 		{{1u,1u}}
 	);
 	
-	iterator<Desc::D>::apply(
-		[&](auto& index){
-			saw::data<sch::Vector<T,Desc::D>> middle, ind_vec;
-			middle.at({{0u}}) = dim_x * 0.5;
-			middle.at({{1u}}) = dim_y * 0.5;
-
-			ind_vec.at({{0u}}) = index.at({{0u}}).template cast_to<T>();
-			ind_vec.at({{1u}}) = index.at({{1u}}).template cast_to<T>();
-
-			auto dist = middle - ind_vec;
-			auto dist_2 = saw::math::dot(dist,dist);
-			if(dist_2.at({}).get() < dim_y*dim_y*0.01){
-				porous_f.at(index).at({}) = 0.0;
-			}
-		},
-		{},// 0-index
-		df_f.get_dims()
-	);
-
 	return saw::make_void();
 }
 
@@ -196,12 +178,16 @@ saw::error_or<void> step(
 		component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
 		component<T,Desc,cmpt::AntiBounceBack<0u>,encode::Sycl<saw::encode::Native>> abb;
 
-		saw::data<sch::Scalar<T>> rho_b;
-		rho_b.at({}) = 1.0;
-		saw::data<sch::Vector<T,Desc::D>> vel_b;
-		vel_b.at({{0u}}) = 0.015;
-
-		component<T,Desc,cmpt::Equilibrium,encode::Sycl<saw::encode::Native>> equi{rho_b,vel_b};
+		saw::data<sch::Scalar<T>> eps;
+		eps.at({}).set(1.5f);
+		saw::data<sch::Scalar<T>> rad;
+		rad.at({}).set(dim_y*0.1*0.25);
+		saw::data<sch::Vector<T,Desc::D>> p_pos;
+		{
+			p_pos.at({{0u}}) = dim_x * 0.5;
+			p_pos.at({{1u}}) = dim_y * 0.5;
+		}
+		component<T,Desc,cmpt::OneParticleAt,encode::Sycl<saw::encode::Native>> opa{p_pos,rad,eps};
 
 		component<T,Desc,cmpt::ZouHeHorizontal<true>,encode::Sycl<saw::encode::Native>> flow_in{
 			[&](){
@@ -230,6 +216,7 @@ saw::error_or<void> step(
 					abb.apply(fields,index,t_i);
 				break;
 				case 2u:
+					opa.apply(macros,index,t_i);
 					collision.apply(fields,macros,index,t_i);
 					break;
 				case 3u:
@@ -385,20 +372,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 			}
 		}
 		*/
-		if(i.get() % 32u == 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,"m",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;
@@ -438,6 +412,12 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 			return eov;
 		}
 	}
+	{
+		auto eov = write_csv_file(out_dir,"m",time_steps.get(), *lbm_macro_data_ptr);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
 
 	sycl_q.wait();
 	return saw::make_void();
diff --git a/examples/stokes_drag_particle_2d_psm_gpu/sim.cpp b/examples/stokes_drag_particle_2d_psm_gpu/sim.cpp
index 0ff6e6b..6942cc7 100644
--- a/examples/stokes_drag_particle_2d_psm_gpu/sim.cpp
+++ b/examples/stokes_drag_particle_2d_psm_gpu/sim.cpp
@@ -48,6 +48,7 @@ template<typename T, typename Desc>
 using MacroStruct = Struct<
 	Member<VelChunk<T,Desc>, "velocity">,
 	Member<RhoChunk<T>, "density">,
+	Member<VelChunk<T,Desc>, "force">,
 	Member<ScalarChunk<T,Desc>, "porosity">
 >;
 
@@ -196,12 +197,16 @@ saw::error_or<void> step(
 		component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
 		component<T,Desc,cmpt::AntiBounceBack<0u>,encode::Sycl<saw::encode::Native>> abb;
 
-		saw::data<sch::Scalar<T>> rho_b;
-		rho_b.at({}) = 1.0;
-		saw::data<sch::Vector<T,Desc::D>> vel_b;
-		vel_b.at({{0u}}) = 0.015;
-
-		component<T,Desc,cmpt::Equilibrium,encode::Sycl<saw::encode::Native>> equi{rho_b,vel_b};
+		saw::data<sch::Scalar<T>> eps;
+		eps.at({}).set(1.5f);
+		saw::data<sch::Scalar<T>> rad;
+		rad.at({}).set(dim_y*0.1*0.25);
+		saw::data<sch::Vector<T,Desc::D>> p_pos;
+		{
+			p_pos.at({{0u}}) = dim_x * 0.5;
+			p_pos.at({{1u}}) = dim_y * 0.5;
+		}
+		component<T,Desc,cmpt::OneParticleAt,encode::Sycl<saw::encode::Native>> opa{p_pos,rad,eps};
 
 		component<T,Desc,cmpt::ZouHeHorizontal<true>,encode::Sycl<saw::encode::Native>> flow_in{
 			[&](){
@@ -230,6 +235,7 @@ saw::error_or<void> step(
 					abb.apply(fields,index,t_i);
 				break;
 				case 2u:
+					opa.apply(macros,index,t_i);
 					collision.apply(fields,macros,index,t_i);
 					break;
 				case 3u:
@@ -385,20 +391,6 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 			}
 		}
 		*/
-		if(i.get() % 32u == 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,"m",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;
@@ -438,7 +430,12 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 			return eov;
 		}
 	}
-
+	{
+		auto eov = write_csv_file(out_dir,"m",time_steps.get(), *lbm_macro_data_ptr);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
 	sycl_q.wait();
 	return saw::make_void();
 }