Merge branch 'dev'

4 files changed, 55 insertions, 48 deletions

diff --git a/examples/poiseulle_particles_2d_gpu/common.hpp b/examples/poiseulle_particles_2d_gpu/common.hpp
index 6c05b64..90b7687 100644
--- a/examples/poiseulle_particles_2d_gpu/common.hpp
+++ b/examples/poiseulle_particles_2d_gpu/common.hpp
@@ -55,7 +55,7 @@ using MacroStruct = Struct<
 >;
 
 template<typename T, typename Desc>
-using ParticleSpheroidGroup = ParticleGroup<T,Desc::D,coll::Spheroid<T>>;
+using ParticleSpheroidGroup = ParticleGroup<T,Desc::D,1u,coll::Spheroid<T>>;
 }
 
 }
diff --git a/examples/poiseulle_particles_2d_gpu/init.hpp b/examples/poiseulle_particles_2d_gpu/init.hpp
index 617b296..cdce25e 100644
--- a/examples/poiseulle_particles_2d_gpu/init.hpp
+++ b/examples/poiseulle_particles_2d_gpu/init.hpp
@@ -5,8 +5,8 @@
 namespace kel {
 namespace lbm {
 
-template<typename T, typename Desc>
-saw::error_or<void> setup_initial_conditions(
+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,
@@ -117,7 +117,7 @@ saw::error_or<void> setup_initial_conditions(
 		radius_p.at({}).set(2);
 		saw::data<sch::Scalar<T>> dense_p;
 		dense_p.at({}).set(1);
-		particles = create_spheroid_particle_group<T,Desc::D>(radius_p, dense_p, {{16u}});
+		particles = create_spheroid_particle_group<T,Desc::D,1u>(radius_p, dense_p, {{16u}});
 	}
 
 	return saw::make_void();
diff --git a/examples/poiseulle_particles_2d_gpu/sim.cpp b/examples/poiseulle_particles_2d_gpu/sim.cpp
index 865babf..fd6cdca 100644
--- a/examples/poiseulle_particles_2d_gpu/sim.cpp
+++ b/examples/poiseulle_particles_2d_gpu/sim.cpp
@@ -22,18 +22,11 @@ using LbmArgs = Args<
 }
 
 template<typename T, typename Desc, typename Coll>
-saw::error_or<void> lbm_main(int argc, char** argv){
+saw::error_or<void> lbm_main(const saw::data<args::LbmArgs>& args){
 	using namespace kel::lbm;
 
 	using dfi = df_info<T,Desc>;
 
-	auto eo_args = saw::parse_args<args::LbmArgs>(argc,argv);
-	if(eo_args.is_error()){
-		return std::move(eo_args.get_error());
-	}
-	auto& args = eo_args.get_value();
-	(void)args;
-
 	auto eo_lbm_dir = output_directory();
 	if(eo_lbm_dir.is_error()){
 		return std::move(eo_lbm_dir.get_error());
@@ -88,7 +81,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 
 	sycl_q.wait();
 	{
-		auto eov = setup_initial_conditions<T,Desc>(conv,*lbm_data_ptr,*lbm_macro_data_ptr,*lbm_particle_data_ptr);
+		auto eov = init<T,Desc,Coll>(conv,*lbm_data_ptr,*lbm_macro_data_ptr,*lbm_particle_data_ptr);
 		if(eov.is_error()){
 			return eov;
 		}
@@ -118,26 +111,25 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 		}
 	}
 	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{16u*4096ul};
+	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>(conv,lsd_view,lsdm_view,lsdp_view,i,dev);
+			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() % 32u == 0u){
+		if(i.get() % 64u == 0u){
 			{
 				auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr);
 				if(eov.is_error()){
@@ -195,17 +187,32 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 	return saw::make_void();
 }
 
-
 using FloatT = kel::lbm::sch::Float32;
+using DescT = kel::lbm::sch::D2Q9;
 
-saw::error_or<void> kmain(int argc, char** argv){
+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);
+		return eov;
+	}
+	auto& an = args.template get<"args">();
+	auto& coupling = an.template get<"coupling">();
 
-	// auto eo_args = saw::parse_args<
+	if(coupling.stl_view() == "hlbm"){
+		return lbm_main<FloatT,DescT,method::Hlbm>(args);
+	//}else if(coupling.view() == "psm"){
+	//	return kmain<FloatT,DescT,coupling::Psm>(args);
+	}
+
+	return saw::make_error<saw::err::critical>("Invalid coupling");
 }
 
+
 int main(int argc, char** argv){
-	auto eov = lbm_main<FloatT,kel::lbm::sch::D2Q9, >(argc, argv);
+	auto eov = k_main(argc, argv);
 	if(eov.is_error()){
 		auto& err = eov.get_error();
 		std::cerr<<"[Error] "<<err.get_category();
diff --git a/examples/poiseulle_particles_2d_gpu/step.hpp b/examples/poiseulle_particles_2d_gpu/step.hpp
index aa0e382..37eb4e7 100644
--- a/examples/poiseulle_particles_2d_gpu/step.hpp
+++ b/examples/poiseulle_particles_2d_gpu/step.hpp
@@ -5,7 +5,7 @@
 namespace kel {
 namespace lbm {
 
-template<typename T, typename Desc>
+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,
@@ -17,29 +17,44 @@ saw::error_or<void> step(
 	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::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]);
+			}
 
-	component<T,Desc,cmpt::HlbmParticle,encode::Sycl<saw::encode::Native>> particle;
+			hlbm_one_part.apply(fields,macros,particles,index,t_i,{16u});
+		});
+	}).wait();
 
 	// auto coll_ev = 
 	q.submit([&](acpp::sycl::handler& h){
 		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};
 
 		component<T,Desc,cmpt::ZouHeHorizontal<true>,encode::Sycl<saw::encode::Native>> flow_in{
 			[&](){
 				uint64_t target_t_i = 16u;
 				if(t_i.get() < target_t_i){
-					return 1.0 + (0.0002 / target_t_i) * t_i.get();
+					return 1.0 + (0.01 / target_t_i) * t_i.get();
 				}
-				return 1.0002;
+				return 1.01;
 			}()
 		};
 		component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0};
@@ -63,12 +78,10 @@ saw::error_or<void> step(
 					break;
 				case 3u:
 					flow_in.apply(fields,index,t_i);
-					// equi.apply(fields,index,t_i);
 					collision.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);
 				break;
 				default:
@@ -77,19 +90,6 @@ saw::error_or<void> step(
 		});
 	}).wait();
 
-	q.submit([&](acpp::sycl::handler& h){
-		h.parallel_for(acpp::sycl::range<1u>{1u}, [=](acpp::sycl::id<1u> idx){
-			particle.apply(fields,macros,particles,{{0u}},t_i);
-		});
-	}).wait();
-	
-	// Step
-	/*
-	q.submit([&](acpp::sycl::handler& h){
-		// h.depends_on(collision_ev);
-	}).wait();
-	*/
-
 	return saw::make_void();
 }