Fixed Stokes example (yikes)

4 files changed, 103 insertions, 33 deletions

diff --git a/examples/stokes_drag_particle_2d_fplbm_gpu/sim.cpp b/examples/stokes_drag_particle_2d_fplbm_gpu/sim.cpp
index fb31f47..5de963d 100644
--- a/examples/stokes_drag_particle_2d_fplbm_gpu/sim.cpp
+++ b/examples/stokes_drag_particle_2d_fplbm_gpu/sim.cpp
@@ -150,7 +150,7 @@ saw::error_or<void> setup_initial_conditions(
 			rho.at({}) = {1};
 			auto& vel = vel_f.at(index);
 			if(info_f.at(index).get() == 2u){
-				vel.at({{0u}}) = 0.0;
+				vel.at({{0u}}) = 0.0002f;
 			}
 			auto eq = equilibrium<T,Desc>(rho,vel);
 
@@ -214,17 +214,30 @@ saw::error_or<void> step(
 		}
 		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{
-			[&](){
-				uint64_t target_t_i = 16u;
+		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;
+				uint64_t target_t_i = 8u;
 				if(t_i.get() < target_t_i){
-					return 1.0 + (0.00002 / target_t_i) * t_i.get();
+					vel.at({{0u}}) = (0.0002 / target_t_i) * t_i.get();
+				}else{
+					vel.at({{0u}}) = 0.0002;
 				}
-				return 1.00002;
+				return vel;
+			}()
+		};
+		component<T,Desc,cmpt::ZouHeVelocityX<false>,encode::Sycl<saw::encode::Native>> flow_out{
+			[&]() -> saw::data<sch::Vector<T,Desc::D>>{
+				saw::data<sch::Vector<T,Desc::D>> vel;
+				uint64_t target_t_i = 8u;
+				if(t_i.get() < target_t_i){
+					vel.at({{0u}}) = (0.0002 / target_t_i) * t_i.get();
+				}else{
+					vel.at({{0u}}) = 0.0002;
+				}
+				return vel;
 			}()
 		};
-		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;
@@ -365,7 +378,9 @@ 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);
-	saw::data<sch::UInt64> time_steps{16u*4096ul};
+	
+	saw::data<sch::UInt64> time_steps{32u*4096ul};
+	
 	auto& info_f = lsd_view.template get<"info">();
 
 	for(saw::data<sch::UInt64> i{0u}; i < time_steps and krun; ++i){
diff --git a/examples/stokes_drag_particle_2d_hlbm_gpu/sim.cpp b/examples/stokes_drag_particle_2d_hlbm_gpu/sim.cpp
index 9a04211..68c0b35 100644
--- a/examples/stokes_drag_particle_2d_hlbm_gpu/sim.cpp
+++ b/examples/stokes_drag_particle_2d_hlbm_gpu/sim.cpp
@@ -147,7 +147,7 @@ saw::error_or<void> setup_initial_conditions(
 			rho.at({}) = {1};
 			auto& vel = vel_f.at(index);
 			if(info_f.at(index).get() == 2u){
-				vel.at({{0u}}) = 0.0;
+				vel.at({{0u}}) = 0.0002f;
 			}
 			auto eq = equilibrium<T,Desc>(rho,vel);
 
@@ -214,17 +214,30 @@ saw::error_or<void> step(
 		
 		component<T,Desc,cmpt::ForceGather, encode::Sycl<saw::encode::Native>> fg;
 
-		component<T,Desc,cmpt::ZouHeHorizontal<true>,encode::Sycl<saw::encode::Native>> flow_in{
-			[&](){
-				uint64_t target_t_i = 16u;
+		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;
+				uint64_t target_t_i = 8u;
 				if(t_i.get() < target_t_i){
-					return 1.0 + (0.00002 / target_t_i) * t_i.get();
+					vel.at({{0u}}) = (0.0002 / target_t_i) * t_i.get();
+				}else{
+					vel.at({{0u}}) = 0.0002;
 				}
-				return 1.00002;
+				return vel;
+			}()
+		};
+		component<T,Desc,cmpt::ZouHeVelocityX<false>,encode::Sycl<saw::encode::Native>> flow_out{
+			[&]() -> saw::data<sch::Vector<T,Desc::D>>{
+				saw::data<sch::Vector<T,Desc::D>> vel;
+				uint64_t target_t_i = 8u;
+				if(t_i.get() < target_t_i){
+					vel.at({{0u}}) = (0.0002 / target_t_i) * t_i.get();
+				}else{
+					vel.at({{0u}}) = 0.0002;
+				}
+				return vel;
 			}()
 		};
-		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;
@@ -367,7 +380,9 @@ 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);
-	saw::data<sch::UInt64> time_steps{16u*4096ul};
+	
+	saw::data<sch::UInt64> time_steps{32u*4096ul};
+	
 	auto& info_f = lsd_view.template get<"info">();
 
 	for(saw::data<sch::UInt64> i{0u}; i < time_steps and krun; ++i){
diff --git a/examples/stokes_drag_particle_2d_psm_gpu/sim.cpp b/examples/stokes_drag_particle_2d_psm_gpu/sim.cpp
index 9781691..3ab3dda 100644
--- a/examples/stokes_drag_particle_2d_psm_gpu/sim.cpp
+++ b/examples/stokes_drag_particle_2d_psm_gpu/sim.cpp
@@ -134,6 +134,7 @@ saw::error_or<void> setup_initial_conditions(
 			por_f.at(index).at({}) = {1};
 			rho.at({}) = {1};
 			auto& vel = vel_f.at(index);
+
 			auto eq = equilibrium<T,Desc>(rho,vel);
 
 			df = eq;
@@ -149,7 +150,7 @@ saw::error_or<void> setup_initial_conditions(
 			rho.at({}) = {1};
 			auto& vel = vel_f.at(index);
 			if(info_f.at(index).get() == 2u){
-				vel.at({{0u}}) = 0.0;
+				vel.at({{0u}}) = 0.0002f;
 			}
 			auto eq = equilibrium<T,Desc>(rho,vel);
 
@@ -233,16 +234,30 @@ saw::error_or<void> step(
 		}
 		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{
-			[&](){
-				uint64_t target_t_i = 16u;
+		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;
+				uint64_t target_t_i = 8u;
+				if(t_i.get() < target_t_i){
+					vel.at({{0u}}) = (0.0002 / target_t_i) * t_i.get();
+				}else{
+					vel.at({{0u}}) = 0.0002;
+				}
+				return vel;
+			}()
+		};
+		component<T,Desc,cmpt::ZouHeVelocityX<false>,encode::Sycl<saw::encode::Native>> flow_out{
+			[&]() -> saw::data<sch::Vector<T,Desc::D>>{
+				saw::data<sch::Vector<T,Desc::D>> vel;
+				uint64_t target_t_i = 8u;
 				if(t_i.get() < target_t_i){
-					return 1.0 + (0.00002 / target_t_i) * t_i.get();
+					vel.at({{0u}}) = (0.0002 / target_t_i) * t_i.get();
+				}else{
+					vel.at({{0u}}) = 0.0002;
 				}
-				return 1.00002;
+				return vel;
 			}()
 		};
-		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;
 
@@ -389,7 +404,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 	auto lsd_view = make_view(lbm_sycl_data);
 	auto lsdm_view = make_view(lbm_sycl_macro_data);
 
-	saw::data<sch::UInt64> time_steps{16u*4096ul};
+	saw::data<sch::UInt64> time_steps{32u*4096ul};
 
 	auto& info_f = lsd_view.template get<"info">();
 
diff --git a/modules/core/c++/boundary.hpp b/modules/core/c++/boundary.hpp
index 01ae7b5..22a6f34 100644
--- a/modules/core/c++/boundary.hpp
+++ b/modules/core/c++/boundary.hpp
@@ -221,14 +221,14 @@ public:
 	}
 };
 
-/*
-template<typename FP, typename Descriptor, bool East, typename Encode>
-class component<FP, Descriptor, cmpt::ZouHeVelocityX<East>, Encode> final {
+
+template<typename T, typename Descriptor, bool East, typename Encode>
+class component<T, Descriptor, cmpt::ZouHeVelocityX<East>, Encode> final {
 private:
-	saw::data<sch::Vector<FP,Descriptor::D>> velocity_;
+	saw::data<sch::Vector<T,Descriptor::D>> velocity_;
 public:
 	component(
-		saw::data<sch::Vector<FP,Descriptor::D>> velocity__
+		saw::data<sch::Vector<T,Descriptor::D>> velocity__
 	):
 		velocity_{velocity__}
 	{}
@@ -237,13 +237,38 @@ public:
 	void apply(const saw::data<CellFieldSchema, Encode>& field, const saw::data<sch::FixedArray<sch::UInt64,Descriptor::D>>& index, saw::data<sch::UInt64> time_step) const {
 		
 		bool is_even = ((time_step.get() % 2u) == 0u);
-		using dfi = df_info<FP,Descriptor>;
+		using dfi = df_info<T,Descriptor>;
 		
 		auto& dfs_old_f = (is_even) ? field.template get<"dfs_old">() : field.template get<"dfs">();
 
+
+		auto& dfs = dfs_old_f.at(index);
+
+		saw::data<T> dir_sum;
+		if constexpr (East) {
+			dir_sum = dfs.at({1u}) + dfs.at({5u}) + dfs.at({7u});
+		} else if constexpr(not East){
+			dir_sum = dfs.at({2u}) + dfs.at({6u}) + dfs.at({8u});
+		}
+
+		saw::data<sch::Scalar<T>> rho;
+		rho.at({}) = (dfs.at({0u}) + dfs.at({4u}) + dfs.at({3u}) + dir_sum * 2) / (velocity_.at({{0u}}) + saw::data<T>{static_cast<saw::native_data_type<T>::type>(1.0)});
+
+		if constexpr (East) {
+			dfs.at({2u}) = dfs.at({1u}) + saw::data<T>{static_cast<saw::native_data_type<T>::type>(2.0 / 3.0)} * rho.at({}) * velocity_.at({{0u}});
+			dfs.at({6u}) = dfs.at({5u}) + saw::data<T>{static_cast<saw::native_data_type<T>::type>(1.0 / 6.0)} * rho.at({}) * velocity_.at({{0u}})
+				+ (dfs.at({3u}) - dfs.at({4u}) + rho.at({}) * velocity_.at({{1u}})) * 0.5f;
+			dfs.at({8u}) = dfs.at({7u}) + saw::data<T>{static_cast<saw::native_data_type<T>::type>(1.0 / 6.0)} * rho.at({}) * velocity_.at({{0u}})
+				+ (dfs.at({4u}) - dfs.at({3u}) + rho.at({}) * velocity_.at({{1u}})) * 0.5f;
+		}else if constexpr ( not East ){
+			dfs.at({1u}) = dfs.at({2u}) + saw::data<T>{static_cast<saw::native_data_type<T>::type>(2.0 / 3.0)} * rho.at({}) * velocity_.at({{0u}});
+			dfs.at({5u}) = dfs.at({6u}) + saw::data<T>{static_cast<saw::native_data_type<T>::type>(1.0 / 6.0)} * rho.at({}) * velocity_.at({{0u}})
+				+ (dfs.at({4u}) - dfs.at({3u}) + rho.at({}) * velocity_.at({{1u}}))*0.5f;
+			dfs.at({7u}) = dfs.at({8u}) + saw::data<T>{static_cast<saw::native_data_type<T>::type>(1.0 / 6.0)} * rho.at({}) * velocity_.at({{0u}})
+				+ (dfs.at({3u}) - dfs.at({4u}) + rho.at({}) * velocity_.at({{1u}}))*0.5f;
+		}
 	}
 };
-*/
 
 }
 }