IBM reimplement

4 files changed, 72 insertions, 32 deletions

diff --git a/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp b/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp
index 9905906..ea49d51 100644
--- a/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp
+++ b/examples/poiseulle_particles_2d_ibm_gpu/sim.cpp
@@ -54,12 +54,21 @@ using MacroStruct = Struct<
 //using ParticleArray = Array<
 //	Particle<T,Desc::D>
 //>;
+
+template<typename T, typename Desc>
+using ParticleSpheroidGroup = ParticleGroup<
+	T,
+	Desc::D,
+	sch::ParticleCollisionSpheroid<T,2.0f>
+>;
+
 }
 
 template<typename T, typename Desc>
 saw::error_or<void> setup_initial_conditions(
 		saw::data<sch::ChunkStruct<T,Desc>>& fields,
-		saw::data<sch::MacroStruct<T,Desc>>& macros
+		saw::data<sch::MacroStruct<T,Desc>>& macros,
+		saw::data<sch::ParticleSpheroidGroup<T,Desc>>& particles
 ){
 	auto& info_f = fields.template get<"info">();
 	// Set everything as walls
@@ -136,25 +145,41 @@ 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;
-			middle.at({{0u}}) = dim_x * 0.25;
-			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>();
+	// Particles
+	{
+		saw::data<sch::Scalar<T>> dense_p;
+		dense_p.at({}).set(1);
 
-			auto dist = middle - ind_vec;
-			auto dist_2 = saw::math::dot(dist,dist);
-			if(dist_2.at({}).get() < dim_y*dim_y*0.01){
-				info_f.at(index) = 1u;
-			}
-		},
-		{},// 0-index
-		df_f.get_dims()
-	);
+		auto& spheroid_grp = particles;
+
+		spheroid_grp = create_spheroid_particle_group<T,Desc::D,2.0f>(
+			dense_p,
+			{32u}
+		);
+
+		{
+			auto& p = spheroid_grp.template get<"particles">();
+			p = {{{1u}}};
+
+			iterator<1u>::apply(
+				[&](auto& index){
+					auto& p_ind = p.at(index);
+					auto& p_rb = p_ind.template get<"rigid_body">();
+					auto& p_pos = p_rb.template get<"position">();
+					auto& p_pos_old = p_rb.template get<"position_old">();
+
+					// Set position
+					p_pos.at({{0u}}) = dim_x * 0.25;
+					p_pos.at({{1u}}) = dim_y * 0.5;
+
+					p_pos_old = p_pos;
+				},
+				{},
+				p.meta()
+			);
+		}
+	}
 
 	return saw::make_void();
 }
@@ -163,6 +188,7 @@ template<typename T, typename Desc>
 saw::error_or<void> step(
 		saw::data<sch::Ptr<sch::ChunkStruct<T,Desc>>,encode::Sycl<saw::encode::Native>>& fields,
 		saw::data<sch::Ptr<sch::MacroStruct<T,Desc>>,encode::Sycl<saw::encode::Native>>& macros,
+		saw::data<sch::Ptr<sch::ParticleSpheroidGroup<T,Desc>>,encode::Sycl<saw::encode::Native>>& particles,
 		saw::data<sch::UInt64> t_i,
 		device& dev
 ){
@@ -172,9 +198,10 @@ 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
-			saw::data<sch::Vector<T,Desc::D>> f;
-			f.at({{0u}}) = 0.0;
-			f.at({{1u}}) = -1.0;
+		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::BounceBack,encode::Sycl<saw::encode::Native>> bb;
 		component<T,Desc,cmpt::AntiBounceBack<0u>,encode::Sycl<saw::encode::Native>> abb;
@@ -237,6 +264,11 @@ 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){
+				
+		});
+	}).wait();
 
 	return saw::make_void();
 }
@@ -285,6 +317,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 	// saw::data<sch::FixedArray<sch::UInt64,Desc::D>> meta{{dim_x,dim_y}};
 	auto lbm_data_ptr = saw::heap<saw::data<sch::ChunkStruct<T,Desc>>>();
 	auto lbm_macro_data_ptr = saw::heap<saw::data<sch::MacroStruct<T,Desc>>>();
+	auto lbm_particle_ptr = saw::heap<saw::data<sch::ParticleSpheroidGroup<T,Desc>>>();
 	
 	std::cout<<"Estimated Bytes: "<<memory_estimate<sch::ChunkStruct<T,Desc>,sch::MacroStruct<T,Desc>>().get()<<std::endl;
 
@@ -310,7 +343,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 
 	sycl_q.wait();
 	{
-		auto eov = setup_initial_conditions<T,Desc>(*lbm_data_ptr,*lbm_macro_data_ptr);
+		auto eov = setup_initial_conditions<T,Desc>(*lbm_data_ptr,*lbm_macro_data_ptr,*lbm_particle_ptr);
 		if(eov.is_error()){
 			return eov;
 		}
@@ -324,6 +357,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 
 	saw::data<sch::ChunkStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_data{sycl_q};
 	saw::data<sch::MacroStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_macro_data{sycl_q};
+	saw::data<sch::ParticleSpheroidGroup<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_particle{sycl_q};
 	sycl_q.wait();
 
 	{
@@ -338,9 +372,16 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 			return eov;
 		}
 	}
+	{
+		auto eov = dev.copy_to_device(*lbm_particle_ptr,lbm_sycl_particle);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
 	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);
 	saw::data<sch::UInt64> time_steps{16u*4096ul};
 
 	auto& info_f = lsd_view.template get<"info">();
@@ -348,7 +389,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 	for(saw::data<sch::UInt64> i{0u}; i < time_steps and krun; ++i){
 		// BC + Collision
 		{
-			auto eov = step<T,Desc>(lsd_view,lsdm_view,i,dev);
+			auto eov = step<T,Desc>(lsd_view,lsdm_view,lsdp_view,i,dev);
 			if(eov.is_error()){
 				return eov;
 			}
diff --git a/examples/settling_cubes_2d_ibm_gpu/sim.cpp b/examples/settling_cubes_2d_ibm_gpu/sim.cpp
index d7b402a..9fdea8c 100644
--- a/examples/settling_cubes_2d_ibm_gpu/sim.cpp
+++ b/examples/settling_cubes_2d_ibm_gpu/sim.cpp
@@ -145,8 +145,7 @@ saw::error_or<void> setup_initial_conditions(
 		}
 	}
 	// Particle in hacky flavour
-	{
-	}
+	{}
 
 	return saw::make_void();
 }
@@ -197,6 +196,8 @@ saw::error_or<void> step(
 			auto& p_pos = p_rb.template get<"position">();
 			auto& p_rot = p_rb.template get<"rotation">();
 
+			auto& p_acc = p_rb.template get<"acceleration">();
+
 			iterator<Desc::D>::apply(
 				[&](auto& m_ind){
 					saw::data<sch::Vector<T,Desc::D>> index_shift;
@@ -238,6 +239,7 @@ saw::error_or<void> step(
 					forces.at(p_cell_pos) = forces.at(p_cell_pos) + force;
 
 					// TODO APPLY FORCE TO PARTICLE
+					p_acc = p_acc + force; // TODO divide by mass
 				},
 				{},
 				p_mask.meta()
diff --git a/lib/core/c++/iterator.hpp b/lib/core/c++/iterator.hpp
index 696c521..7fd6f58 100644
--- a/lib/core/c++/iterator.hpp
+++ b/lib/core/c++/iterator.hpp
@@ -50,8 +50,11 @@ public:
 	}
 };
 
+template<typename Tupl>
+struct ct_tuple_iterator;
+
 template<typename... Tup>
-struct ct_tuple_iterator final {
+struct ct_tuple_iterator<sch::Tuple<Tup...>> final {
 private:
 	template<typename Func, uint64_t i>
 	static constexpr saw::error_or<void> apply_i(
diff --git a/lib/core/c++/term_renderer.hpp b/lib/core/c++/term_renderer.hpp
deleted file mode 100644
index 5cbb551..0000000
--- a/lib/core/c++/term_renderer.hpp
+++ /dev/null
@@ -1,6 +0,0 @@
-#pragma once
-
-namespace kel {
-namespace lbm {
-}
-}