Merge branch 'dev'

1 files changed, 91 insertions, 46 deletions

diff --git a/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp b/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp
index e68d7da..e1bd3ba 100644
--- a/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp
+++ b/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp
@@ -1,7 +1,7 @@
 #include <kel/lbm/sycl/lbm.hpp>
 #include <kel/lbm/lbm.hpp>
 #include <kel/lbm/particle.hpp>
-#include <kel/lbm/math/n_linear.hpp>
+#include <kel/lbm/math/math.hpp>
 
 #include <forstio/io/io.hpp>
 #include <forstio/remote/filesystem/easy.hpp>
@@ -114,6 +114,7 @@ saw::error_or<void> setup_initial_conditions(
 	auto& df_f = fields.template get<"dfs_old">();
 	auto& rho_f = macros.template get<"density">();
 	auto& vel_f = macros.template get<"velocity">();
+	auto& force_f = macros.template get<"force">();
 	
 	iterator<Desc::D>::apply(
 		[&](auto& index){
@@ -135,6 +136,9 @@ saw::error_or<void> setup_initial_conditions(
 			auto& rho = rho_f.at(index);
 			rho.at({}) = {1};
 			auto& vel = vel_f.at(index);
+			auto& force = force_f.at(index);
+			force = {};
+
 			if(info_f.at(index).get() == 2u){
 				vel.at({{0u}}) = 0.0;
 			}
@@ -195,15 +199,79 @@ saw::error_or<void> step(
 ){
 	auto& q = dev.get_handle();
 	auto& info_f = fields.template get<"info">();
+	auto& force_f = macros.template get<"force">();
+
+	q.submit([&](acpp::sycl::handler& h){
+		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& force = force_f.at(index);
+
+			for(uint64_t i{0u}; i < Desc::D; ++i){
+				force.at({{i}}) = 0.0;
+			}
+		});
+	}).wait();
+			
+	q.submit([&](acpp::sycl::handler& h){
+		h.parallel_for(acpp::sycl::range<1u>{1u}, [=](acpp::sycl::id<1u> idx){
+			auto& vel_f = macros.template get<"velocity">();
+			auto& dense_f = macros.template get<"density">();
+
+			auto& ps = particles;
+			auto& mask = ps.template get<"mask">();
+			auto& mask_step = ps.template get<"mask_step">().at({});
+			auto& p_dense = ps.template get<"density">().at({});
+			auto& com = ps.template get<"center_of_mass">();
+
+			auto& parts = ps.template get<"particles">();
+
+			auto& p_i = parts.at({{idx[0u]}});
+
+			auto& p_i_rb = p_i.template get<"rigid_body">();
+			/// 0. Iterate over mask and calculate position in LBM grid
+			/// In this case it's simple since I'm too lazy to do scaling and rotation
+			/// Technically scale => rotate => translate
+			/// Here it's only translate
+			auto& p_i_rb_pos = p_i_rb.template get<"position">();
+
+			iterator<Desc::D>::apply([&](const auto& index){
+				/// Calculate the shift from the mask
+				saw::data<sch::Vector<T,Desc::D>> index_shift;
+				for(uint64_t i = 0u; i < Desc::D; ++i){
+					index_shift.at({{i}}) = index.at({i}).template cast_to<T>() - com.at({}).at({{i}});
+					// Scale to LBM Grid
+					index_shift.at({{i}}) = index_shift.at({{i}}) * mask_step.at({});
+				}
+
+				// Shift our pos into the index
+				auto p_i_rb_pos_ind = p_i_rb_pos + index_shift;
+							
+				/// Calculate force pickup from neigbouring u_vel cells
+				// auto inter_vel_fluid = n_linear_interpolate(vel_f,p_i_rb_pos_ind);
+				auto inter_vel_fluid = n_closest_read(vel_f,p_i_rb_pos_ind);
+
+				// Technically TODO to use moment
+				auto inter_moment_fluid = inter_vel_fluid;
+
+				// Technically Particles can have more timesteps than the fluid
+				auto force_response = -inter_moment_fluid;
+
+				/// Distribute force to fluid
+				// n_linear_spread(force_f,p_i_rb_pos_ind, force_response);
+				n_closest_add(force_f,p_i_rb_pos_ind, force_response);
+
+			}, {}, mask.meta());
+		});
+	}).wait();
 
 	// auto coll_ev = 
 	q.submit([&](acpp::sycl::handler& h){
 		// Need nicer things to handle the flow. I see improvement here
-		saw::data<sch::Vector<T,Desc::D>> f;
-		f.at({{0u}}) = 0.0;
-		f.at({{1u}}) = -1.0;
-
-		component<T,Desc,cmpt::BGKGuo, encode::Sycl<saw::encode::Native>> collision{0.65,f};
+		component<T,Desc,cmpt::BGKGuo, encode::Sycl<saw::encode::Native>> collision{0.65};
 		component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
 		component<T,Desc,cmpt::AntiBounceBack<0u>,encode::Sycl<saw::encode::Native>> abb;
 
@@ -233,7 +301,7 @@ saw::error_or<void> step(
 			}
 
 			auto info = info_f.at(index);
-			
+						
 			switch(info.get()){
 				case 0u:
 				break;
@@ -256,6 +324,8 @@ saw::error_or<void> step(
 				break;
 			}
 		});
+
+		
 	}).wait();
 
 	
@@ -265,44 +335,6 @@ saw::error_or<void> step(
 		// h.depends_on(collision_ev);
 	}).wait();
 	*/
-	q.submit([&](acpp::sycl::handler& h){
-		h.parallel_for(acpp::sycl::range<1u>{1u}, [=](acpp::sycl::id<1u> idx){
-			auto& vel = macros.template get<"velocity">();
-
-			auto& ps = particles;
-			auto& mask = ps.template get<"mask">();
-			auto& dense = ps.template get<"density">().at({});
-			auto& com = ps.template get<"center_of_mass">();
-
-			auto& parts = ps.template get<"particles">();
-
-			auto& p_i = parts.at({{0u}});
-
-			auto& p_i_rb = p_i.template get<"rigid_body">();
-			/// 0. Iterate over mask and calculate position in LBM grid
-			/// In this case it's simple since I'm too lazy to do scaling and rotation
-			/// Technically scale => rotate => translate
-			/// Here it's only translate
-			auto& p_i_rb_pos = p_i_rb.template get<"position">();
-
-			iterator<Desc::D>::apply([&](const auto& index){
-				/// Calculate the shift from the mask
-				saw::data<sch::Vector<T,Desc::D>> index_shift;
-				for(uint64_t i = 0u; i < Desc::D; ++i){
-					index_shift.at({{i}}) = index.at({i}).template cast_to<T>() - com.at({{i}});
-				}
-
-
-				/// TODO 1. Calculate force pickup from neigbouring u_vel cells
-				auto inter_vel = n_linear_interpolate(vel,index_shift);
-
-				/// TODO 3. Distribute force to fluid
-
-
-			}, {}, mask.meta());
-		});
-	}).wait();
-
 	return saw::make_void();
 }
 }
@@ -476,8 +508,21 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 			});
 		}).wait();
 		wait.poll();
+
+		// PRINT STATUS ON SIGUSR1
 		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;
+			{
+				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,"m",i.get(), *lbm_macro_data_ptr);
+				if(eov.is_error()){
+					return eov;
+				}
+			}
 			print_status = false;
 		}
 		print_progress_bar(i.get(), time_steps.get()-1u);