I seem to have no clue what I'm doing

4 files changed, 107 insertions, 68 deletions

diff --git a/default.nix b/default.nix
index 4303f93..d233d90 100644
--- a/default.nix
+++ b/default.nix
@@ -46,7 +46,7 @@ let
 
 		 src = builtins.fetchurl {
 			url = "https://git.keldu.de/forstio-forstio/snapshot/master.tar.gz";
-			sha256 = "sha256:0khnwmrhdric5i701wfbkbdrfzp9vrswg6raan4njmwlx146vbrf";
+			sha256 = "sha256:15iqzmymza47jjx4wpc19mbg3zzwmkabpssf5y968f566n0fnb9a";
 		 };
 
 		 phases = [ "unpackPhase" "installPhase" ];
diff --git a/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp b/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp
index d4bf053..e1bd3ba 100644
--- a/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp
+++ b/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp
@@ -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,61 +335,6 @@ saw::error_or<void> step(
 		// h.depends_on(collision_ev);
 	}).wait();
 	*/
-
-	auto& force_f = macros.template get<"force">();
-
-	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& 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();
-
 	return saw::make_void();
 }
 }
@@ -493,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);
diff --git a/lib/core/c++/collision.hpp b/lib/core/c++/collision.hpp
index 9c76c1a..023f61f 100644
--- a/lib/core/c++/collision.hpp
+++ b/lib/core/c++/collision.hpp
@@ -146,7 +146,8 @@ public:
 
 		saw::data<sch::Scalar<T>> half;
 		half.at({}).set(0.5);
-		saw::data<sch::Vector<T,Descriptor::D>> vel = vel_f.at(index) + total_force * ( half / rho );
+		auto& vel = vel_f.at(index);
+		vel = vel + total_force * ( half / rho );
 
 		compute_rho_u<T,Descriptor>(dfs_old_f.at(index),rho,vel);
 		auto eq = equilibrium<T,Descriptor>(rho,vel);
diff --git a/lib/core/c++/math/n_closest.hpp b/lib/core/c++/math/n_closest.hpp
index 13414e2..ac0fe2f 100644
--- a/lib/core/c++/math/n_closest.hpp
+++ b/lib/core/c++/math/n_closest.hpp
@@ -7,7 +7,7 @@ namespace kel {
 namespace lbm {
 
 template<typename FieldSchema, typename Encode, typename T, uint64_t D>
-saw::data<typename FieldSchema::InnerValueSchema> n_closest_read(const saw::data<sch::Ptr<FieldSchema>,Encode>& f, const saw::data<sch::Vector<T,D>>& frac_ind){
+saw::data<typename FieldSchema::StoredValueSchema> n_closest_read(const saw::data<sch::Ptr<FieldSchema>,Encode>& f, const saw::data<sch::Vector<T,D>>& frac_ind){
 
 	auto shift_frac_ind = frac_ind;
 	for(uint64_t i{0u}; i < D; ++i){
@@ -18,13 +18,16 @@ saw::data<typename FieldSchema::InnerValueSchema> n_closest_read(const saw::data
 		}
 	}
 
-	auto shift_ind = frac_ind.template cast_to<sch::UInt64>();
+	saw::data<sch::FixedArray<sch::UInt64,D>> shift_ind;
+	for(uint64_t i{0u}; i < D; ++i){
+		shift_ind.at({i}) = frac_ind.at({{i}}).template cast_to<sch::UInt64>();
+	}
 
 	return f.at(shift_ind);
 }
 
 template<typename FieldSchema, typename Encode, typename T, uint64_t D>
-void n_closest_add(saw::data<sch::Ptr<FieldSchema>,Encode>& f, const saw::data<sch::Vector<T,D>>& frac_ind, const saw::data<typename FieldSchema::InnerValueSchema>& val){
+void n_closest_add(const saw::data<sch::Ptr<FieldSchema>,Encode>& f, const saw::data<sch::Vector<T,D>>& frac_ind, const saw::data<typename FieldSchema::StoredValueSchema>& val){
 	auto shift_frac_ind = frac_ind;
 	for(uint64_t i{0u}; i < D; ++i){
 		
@@ -34,7 +37,14 @@ void n_closest_add(saw::data<sch::Ptr<FieldSchema>,Encode>& f, const saw::data<s
 		}
 	}
 
-	auto shift_ind = frac_ind.template cast_to<sch::UInt64>();
+	auto f_meta = f.meta();
+	saw::data<sch::FixedArray<sch::UInt64,D>> shift_ind;
+	for(uint64_t i{0u}; i < D; ++i){
+		shift_ind.at({i}) = frac_ind.at({{i}}).template cast_to<sch::UInt64>();
+		if(shift_ind.at({i}) < f_meta.at({i})){
+			shift_ind.at({i}) = f_meta.at({i}) - 1u;
+		}
+	}
 	auto& f_i = f.at(shift_ind);
 
 	f_i = f_i + val;