From d01f8faadcec62593e0af5304bcc2db3d22ed0c5 Mon Sep 17 00:00:00 2001
From: "Claudius \"keldu\" Holeksa" <mail@keldu.de>
Date: Thu, 16 Apr 2026 16:20:23 +0200
Subject: Brain fried

---
 examples/settling_cubes_2d_ibm_gpu/sim.cpp | 21 ++++++++++++++-------
 lib/core/c++/particle.hpp                  |  5 +++++
 lib/sycl/c++/data.hpp                      | 16 ++++++++++++++++
 3 files changed, 35 insertions(+), 7 deletions(-)

diff --git a/examples/settling_cubes_2d_ibm_gpu/sim.cpp b/examples/settling_cubes_2d_ibm_gpu/sim.cpp
index 80ba90c..dceb156 100644
--- a/examples/settling_cubes_2d_ibm_gpu/sim.cpp
+++ b/examples/settling_cubes_2d_ibm_gpu/sim.cpp
@@ -201,7 +201,7 @@ saw::error_or<void> step(
 				[&](auto& m_ind){
 					saw::data<sch::Vector<T,Desc::D>> index_shift;
 					for(uint64_t i{0u}; i < Desc::D; ++i){
-						index_shift.at({{i}}) = m_ind.at({i}).template cast_to<T>() - (p_mask.meta().at({i})+1u).template cast_to<T>() * 0.5;
+						index_shift.at({{i}}) = m_ind.at({i}).template cast_to<T>() - (vels.meta().at({i})+1u).template cast_to<T>() * 0.5;
 					}
 
 					saw::data<sch::Vector<T,Desc::D>> transformed_pos;
@@ -213,12 +213,12 @@ saw::error_or<void> step(
 					// Lagrange indicator position
 					auto p_pos_lag = p_pos + transformed_pos;
 
-					// Pick the closest velocity
+					// Pick the closest velocity and clamp it
 					saw::data<sch::FixedArray<sch::UInt64,Desc::D>> p_cell_pos;
 					saw::data<sch::Vector<sch::UInt64,Desc::D>> p_cell_pos_vec;
 					for(uint64_t i{0u}; i < Desc::D; ++i){
 						p_cell_pos.at({{i}}) = (p_pos_lag.at({{i}}) + 0.5).template cast_to<sch::UInt64>();
-						p_cell_pos.at({{i}}).set(std::max(1ul,std::min(p_cell_pos.at({{i}}).get(), p_meta.at({{i}}).get() - 2ul)));
+						p_cell_pos.at({{i}}).set(std::max(0ul,std::min(p_cell_pos.at({{i}}).get(), vels.meta().at({{i}}).get() - 2ul)));
 						p_cell_pos_vec.at({{i}}) = p_cell_pos.at({{i}});
 					}
 
@@ -228,14 +228,16 @@ saw::error_or<void> step(
 					auto rel_cell_to_part_pos = p_cell_pos_vec.template cast_to<T>() - p_pos;
 
 					auto p_vel = (p_pos - p_rb.template get<"position_old">()) * delta_t;
-					auto u_solid = p_vel + saw::math::cross(p_rot,rel_cell_to_part_pos);
 
+					auto u_solid = p_vel + saw::math::cross(p_rot,rel_cell_to_part_pos);
 
 					// Force
 					auto force = (u_solid - u_fluid) / delta_t;
 
 					// TODO HERE ATOMIC! !!!!
 					forces.at(p_cell_pos) = forces.at(p_cell_pos) + force;
+
+					// TODO APPLY FORCE TO PARTICLE
 				},
 				{},
 				p_mask.meta()
@@ -248,7 +250,7 @@ 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
-		component<T,Desc,cmpt::BGKGuo, encode::Sycl<saw::encode::Native>> collision{0.8};
+		component<T,Desc,cmpt::BGK, encode::Sycl<saw::encode::Native>> collision{0.8};
 		component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
 
 		h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
@@ -271,7 +273,6 @@ saw::error_or<void> step(
 				default:
 				break;
 			}
-
 		});
 	}).wait();
 
@@ -389,6 +390,12 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 	auto lsdm_view = make_view(lbm_sycl_macro_data);
 	auto lsdp_view = make_view(lbm_sycl_particle_data);
 
+	{
+		auto eov = write_vtk_file(out_dir,"ms",0u, *lbm_macro_data_ptr);
+		if(eov.is_error()){
+			return eov;
+		}
+	}
 	saw::data<sch::UInt64> time_steps{16u*4096ul};
 
 	auto& info_f = lsd_view.template get<"info">();
@@ -402,7 +409,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 			}
 		}
 		sycl_q.wait();
-		if(i.get()%32u ==0u){
+		if(i.get()%1u ==0u){
 			{
 				auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr);
 				if(eov.is_error()){
diff --git a/lib/core/c++/particle.hpp b/lib/core/c++/particle.hpp
index b098ecc..691a74b 100644
--- a/lib/core/c++/particle.hpp
+++ b/lib/core/c++/particle.hpp
@@ -11,7 +11,12 @@ struct Particle {};
 
 template<typename T, typename Descriptor, typename Encode>
 class component<T, Descriptor, cmpt::Particle, Encode> {
+private:
+	saw::data<sch::Scalar<T>> dt_;
 public:
+	component(saw::data<sch::Scalar<T>> dt__):
+		dt_{dt__}
+	{}
 
 	template<typename ParticleSchema, typename MacroFieldSchema>
 	void apply(const saw::data<ParticleSchema, Encode>& particles, const saw::data<MacroFieldSchema,Encode>& macros, saw::data<sch::UInt64> index, saw::data<sch::UInt64> time_step) const {
diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp
index 0206833..71627de 100644
--- a/lib/sycl/c++/data.hpp
+++ b/lib/sycl/c++/data.hpp
@@ -278,6 +278,10 @@ public:
 		return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::get_dims();
 	}
 
+	static constexpr auto ghost_meta() {
+		return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::meta();
+	}
+
 	data<ValueSchema, Encode>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index){
 		std::decay_t<decltype(index)> ind;
 		for(uint64_t i = 0u; i < sizeof...(Sides); ++i){
@@ -298,6 +302,10 @@ public:
 		return data<schema::FixedArray<schema::UInt64, sizeof...(Sides)>,Encode>{{Sides...}};
 	}
 
+	static constexpr auto meta(){
+		return data<schema::FixedArray<schema::UInt64, sizeof...(Sides)>,Encode>{{Sides...}};
+	}
+
 	auto flat_data() const {
 		return values_.flat_data();
 	}
@@ -345,6 +353,10 @@ public:
 		return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::get_dims();
 	}
 
+	static constexpr auto ghost_meta() {
+		return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::meta();
+	}
+
 	data<ValueSchema, Encode>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index){
 		std::decay_t<decltype(index)> ind;
 		for(uint64_t i = 0u; i < sizeof...(Sides); ++i){
@@ -361,6 +373,10 @@ public:
 		return values_.at(ind);
 	}
 
+	static constexpr auto meta(){
+		return data<schema::FixedArray<schema::UInt64, sizeof...(Sides)>,Encode>{{Sides...}};
+	}
+
 	static constexpr auto get_dims(){
 		return data<schema::FixedArray<schema::UInt64, sizeof...(Sides)>,Encode>{{Sides...}};
 	}
-- 
cgit v1.2.3