Merge branch 'dev'

6 files changed, 212 insertions, 36 deletions

diff --git a/examples/poiseulle_particles_2d_gpu/sim.cpp b/examples/poiseulle_particles_2d_gpu/sim.cpp
index 0f7c97b..4bc9dfb 100644
--- a/examples/poiseulle_particles_2d_gpu/sim.cpp
+++ b/examples/poiseulle_particles_2d_gpu/sim.cpp
@@ -220,9 +220,9 @@ saw::error_or<void> step(
 			[&](){
 				uint64_t target_t_i = 64u;
 				if(t_i.get() < target_t_i){
-					return 1.0 + (0.015 / target_t_i) * t_i.get();
+					return 1.0 + (0.0015 / target_t_i) * t_i.get();
 				}
-				return 1.015;
+				return 1.0015;
 			}()
 		};
 		component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0};
@@ -240,10 +240,16 @@ saw::error_or<void> step(
 				case 0u:
 				break;
 				case 1u:
-					// bb.apply(fields,index,t_i);
+					bb.apply(fields,index,t_i);
 				break;
 				case 2u:
+					flow_in.apply(fields,index,t_i);
+					collision.apply(fields,macros,index,t_i);
+					break;
 				case 3u:
+					flow_out.apply(fields,index,t_i);
+					collision.apply(fields,macros,index,t_i);
+					break;
 				case 4u:
 					collision.apply(fields,macros,index,t_i);
 				break;
@@ -270,24 +276,6 @@ saw::error_or<void> step(
 		});
 	}).wait();
 	
-	q.submit([&](acpp::sycl::handler& h){
-		h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
-			switch(info.get()){
-				case 1u:
-					bb.apply(fields,index,t_i);
-				break;
-				case 3u:
-					// equi.apply(fields,index,t_i);
-					flow_in.apply(fields,index,t_i);
-				break;
-				case 4u:
-					// equi.apply(fields,index,t_i);
-					flow_out.apply(fields,index,t_i);
-				break;
-			}
-		});
-	}).wait();
-
 	// Step
 	/*
 	q.submit([&](acpp::sycl::handler& h){
@@ -396,6 +384,12 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 
 	for(saw::data<sch::UInt64> i{0u}; i < time_steps and krun; ++i){
 		{
+			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,"t",i.get(), *lbm_macro_data_ptr);
 				if(eov.is_error()){
@@ -409,13 +403,6 @@ saw::error_or<void> lbm_main(int argc, char** argv){
 				return eov;
 			}
 		}
-		{
-			auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr);
-			if(eov.is_error()){
-				return eov;
-			}
-		}
-
 		wait.poll();
 		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;
diff --git a/examples/poiseulle_particles_channel_2d/poiseulle_particles_channel_2d.cpp b/examples/poiseulle_particles_channel_2d/poiseulle_particles_channel_2d.cpp
index e4c6de6..7ac663f 100644
--- a/examples/poiseulle_particles_channel_2d/poiseulle_particles_channel_2d.cpp
+++ b/examples/poiseulle_particles_channel_2d/poiseulle_particles_channel_2d.cpp
@@ -603,7 +603,7 @@ void lbm_step(
 }
 
 /**
- * 2-Way coupling through trinlinearity
+ * 2-Way coupling through bilinearity (n-linearity)
  * o -- o
  * |    |
  * |    |
diff --git a/lib/core/c++/math/n_linear.hpp b/lib/core/c++/math/n_linear.hpp
new file mode 100644
index 0000000..62906cd
--- /dev/null
+++ b/lib/core/c++/math/n_linear.hpp
@@ -0,0 +1,29 @@
+#pragma once
+
+#include "../common.hpp"
+
+namespace kel {
+namespace lbm {
+
+template<typename FieldSchema, typename T, uint64_t D>
+void n_linear_interpolate(const saw::data<FieldSchema>& field, const saw::data<sch::FixedArray<T,D>>& pos){
+
+	auto pos_bound = pos;
+	auto meta = field.dims();
+	saw::data<sch::UInt64> ind;
+	for(saw::data<sch::UInt64> i{0u}; i < saw::data<sch::UInt64>{D}; ++i){
+		pos_bound.at(i).set(
+			std::max(
+				std::min(
+					meta.at(i).template cast_to<T>().get(),
+					pos.at(i).get() + 1.5
+				),
+				1,
+			)
+			-1
+		);
+		ind.at(i) = pos_bound.at(i).template cast_to<sch::UInt64>();
+	}
+}
+}
+}
diff --git a/lib/core/c++/particle/blur.hpp b/lib/core/c++/particle/blur.hpp
new file mode 100644
index 0000000..a304e8d
--- /dev/null
+++ b/lib/core/c++/particle/blur.hpp
@@ -0,0 +1,37 @@
+#pragma once
+
+namespace kel {
+namespace lbm {
+template<typename T, uint64_t D>
+void blur_mask(saw::data<sch::Array<T,D>>& p_mask){
+
+	saw::data<T> mid{2.0/3.0};
+	saw::data<T> edge{1.0/6.0};
+
+	// saw::data<sch::FixedArray<sch::UInt64>,2u> blur_weights{{2.0/3.0},{1.0/6.0}};
+
+	auto meta = p_mask.dims();
+	saw::data<sch::Array<T,D>> blurred_mask{meta};
+
+	for(saw::data<sch::UInt64> i{0u}; i < saw::data<sch::UInt64>{D}; ++i){
+		iterator<D>::apply([&](const auto& index){
+			blurred_mask.at(index) = p_mask.at(index) * mid;
+
+			if(index.at({i}).get() > 0u){
+				auto l_ind = index;
+				l_ind.at({i}) = ind.at({i}) - 1u;
+				blurred_mask.at(index) = blurred_mask.at(index) + p_mask.at(l_ind) * edge;
+			}
+			if((index.at({i}).get() + 1u) < meta.at({i})){
+				auto r_ind = index;
+				r_ind.at({i}) = ind.at({i}) + 1u;
+				blurred_mask.at(index) = blurred_mask.at(index) + p_mask.at(r_ind) * edge;
+			}
+		},{},meta);
+
+		p_mask = blurred_mask;
+	}
+
+}
+}
+}
diff --git a/lib/core/c++/particle/geometry/circle.hpp b/lib/core/c++/particle/geometry/circle.hpp
index 331f06d..cf9b87e 100644
--- a/lib/core/c++/particle/geometry/circle.hpp
+++ b/lib/core/c++/particle/geometry/circle.hpp
@@ -5,7 +5,7 @@
 namespace kel {
 namespace lbm {
 
-template<typename T>
+template<typename T, uint64_t D>
 class particle_circle_geometry {
 private:
 public:
@@ -13,8 +13,8 @@ public:
 	{}
 
 	template<typename MT = T>
-	saw::data<sch::ParticleMask<MT,2>> generate_mask(uint64_t resolution, uint64_t boundary_nodes = 0) const {
-		saw::data<sch::ParticleMask<MT,2>> mask;
+	saw::data<sch::ParticleMask<MT,D>> generate_mask(uint64_t resolution, uint64_t boundary_nodes = 0) const {
+		saw::data<sch::ParticleMask<MT,D>> mask;
 
 		auto& grid = mask.template get<"grid">();
 		auto& com = mask.template get<"center_of_mass">();
@@ -28,6 +28,24 @@ public:
 		saw::data<T> delta_x{static_cast<typename saw::native_data_type<T>::type>(2.0  / static_cast<double>(diameter_i))};
 		saw::data<T> center = (saw::data<T>{1.0} + saw::data<T>{2.0} * saw::data<T>{static_cast<saw::native_data_type<T>::type>(boundary_nodes)/diameter_i});
 
+		saw::data<sch::FixedArray<sch::UInt64,D>> boundary_arr;
+		for(uint64_t i = 0u; i < D; ++i){
+			boundary_arr.set(boundary_nodes);
+		}
+
+		iterator<D>::apply([&](const auto& index){
+			grid.at(index).set(0);
+		}, {}, grid.dims());
+		
+		iterator<D>::apply([&](const auto& index){
+			saw::data<sch::Vector<T,D>> f_ind;
+			for(uint64_t i = 0u; i < D; ++i){
+				f_ind.at({{i}}) = (index.at({{i}}) + 0.5) * delta_x - center;
+			}
+
+			auto norm_f_ind = saw::math::norm(f_ind);
+		}, {}, grid.dims(),boundary_arr);
+
 		for(uint64_t i = 0; i < size; ++i){
 			for(uint64_t j = 0; j < size; ++j){
 				grid.at({{i,j}}).set(0);
@@ -43,18 +61,19 @@ public:
 			}
 		}
 
+		// I don't fully remember what this did
 		saw::data<T> total_mass{};
-		iterate_over([&](const saw::data<sch::FixedArray<sch::UInt64,2u>>& index){
+		iterator::apply<D>::apply([&](const saw::data<sch::FixedArray<sch::UInt64,D>>& index){
 			auto ind_vec = saw::math::vectorize_data(index).template cast_to<T>();
-			for(uint64_t i{0u}; i < 2u; ++i){
+			for(uint64_t i{0u}; i < D; ++i){
 				ind_vec.at({{i}}) = ind_vec.at({{i}}) * grid.at(index);
 			}
 			com = com + ind_vec;
 
 			total_mass = total_mass + grid.at(index);
-		},{{0u,0u}}, grid.dims());
+		},{}, grid.dims());
 
-		for(uint64_t i{0u}; i < 2u; ++i){
+		for(uint64_t i{0u}; i < D; ++i){
 			com.at({{i}}) = com.at({{i}}) / total_mass;
 		}
 
diff --git a/util/podman/norce_prefetch_build_and_run.sh b/util/podman/norce_prefetch_build_and_run.sh
new file mode 100644
index 0000000..318ec4a
--- /dev/null
+++ b/util/podman/norce_prefetch_build_and_run.sh
@@ -0,0 +1,104 @@
+#!/usr/bin/env bash
+set -euo pipefail
+
+# --- Usage ---
+# ./build.sh <url> <sha256>
+if [[ $# -ne 2 ]]; then
+    echo "Usage: $0 <url> <sha256>"
+    exit 1
+fi
+
+SRC_URL="$1"
+SRC_HASH="$2"
+
+NIX_STORE_VOL="nix-store"
+NIX_STATE_VOL="nix-state"
+
+# --- Ensure Podman volumes exist ---
+ensure_volume() {
+    local vol="$1"
+    if ! podman volume inspect "$vol" >/dev/null 2>&1; then
+        echo "📦 Creating Podman volume: $vol"
+        podman volume create "$vol" >/dev/null
+    fi
+}
+
+ensure_volume "$NIX_STORE_VOL"
+ensure_volume "$NIX_STATE_VOL"
+
+# --- Detect host CA bundle ---
+HOST_CA_BUNDLE=""
+if [[ -f /etc/ssl/certs/ca-certificates.crt ]]; then
+    HOST_CA_BUNDLE="/etc/ssl/certs/ca-certificates.crt"
+elif [[ -f /etc/pki/tls/certs/ca-bundle.crt ]]; then
+    HOST_CA_BUNDLE="/etc/pki/tls/certs/ca-bundle.crt"
+fi
+
+# --- Run inside Nix container ---
+podman run --rm -it \
+    -v "$NIX_STORE_VOL":/nix/store \
+    -v "$NIX_STATE_VOL":/nix/var \
+    -v /etc/ssl/certs:/etc/ssl/certs:ro \
+    -v /etc/pki:/etc/pki:ro \
+    ${HOST_CA_BUNDLE:+-v "$HOST_CA_BUNDLE:$HOST_CA_BUNDLE:ro"} \
+    -e SRC_URL="$SRC_URL" \
+    -e SRC_HASH="$SRC_HASH" \
+    -e SSL_CERT_FILE="$HOST_CA_BUNDLE" \
+    -e NIX_SSL_CERT_FILE="$HOST_CA_BUNDLE" \
+    docker.io/nixos/nix:latest \
+    nix-shell -p bash nix-prefetch-url --run "
+        set -euo pipefail
+
+        echo \"⬇ Fetching into nix store...\"
+
+        STORE_PATH=\$(nix-prefetch-url --unpack --type sha256 \"\$SRC_URL\" \"\$SRC_HASH\")
+
+        echo \"📦 Source stored at: \$STORE_PATH\"
+
+        if [[ ! -d \"\$STORE_PATH\" ]]; then
+            echo \"❌ Expected unpacked directory in nix store\"
+            exit 1
+        fi
+
+        cd \"\$STORE_PATH\"
+
+        echo \"📂 Entered: \$(pwd)\"
+
+        if [[ ! -f default.nix ]]; then
+            echo \"❌ No default.nix found in source\"
+            exit 1
+        fi
+
+        echo \"🔨 Running nix-build...\"
+        nix-build default.nix --out-link result
+
+        BIN_DIR=./result/bin
+        if [[ ! -d \"\$BIN_DIR\" ]]; then
+            echo \"ℹ No binaries produced.\"
+            exit 0
+        fi
+
+        mapfile -t BINARIES < <(ls -1 \"\$BIN_DIR\")
+        if (( \${#BINARIES[@]} == 0 )); then
+            echo \"ℹ No binaries found in result/bin\"
+            exit 0
+        fi
+
+        echo \"Available binaries:\"
+        select CHOSEN_BIN in \"\${BINARIES[@]}\" \"Quit\"; do
+            if [[ \"\$CHOSEN_BIN\" == \"Quit\" ]]; then
+                echo \"Exiting.\"
+                break
+            elif [[ -n \"\$CHOSEN_BIN\" ]]; then
+                echo \"▶ Running \$CHOSEN_BIN...\"
+                \"\$BIN_DIR/\$CHOSEN_BIN\"
+                break
+            else
+                echo \"Invalid selection, try again.\"
+            fi
+        done
+    "
+
+echo "✅ Done!"
+echo "   • Persistent Nix store: $NIX_STORE_VOL"
+echo "   • Persistent Nix state: $NIX_STATE_VOL"