rtx-cfd legacy GPU: kernel outputs were device-buffer CLONES (every write lost) — the buffers are borrowed mutably now; the Poisson kernels keep boundary Dirichlet values; tests: the advection pulse marched 50 steps, the Jacobi budget 20k, the reduction reference in f64

crates/specialized/rtx-cfd/src/kernels/cuda/poisson.cu

@@ -25,9 +25,9 @@ extern "C" __global__ void poisson_jacobi_2d(
 
     int idx = j * nx + i;
 
-    // Boundary conditions (Dirichlet - zero on boundaries for now)
+    // Boundary cells keep the Dirichlet values they hold.
     if (i == 0 || i == nx - 1 || j == 0 || j == ny - 1) {
-        phi_new[idx] = 0.0f;
+        phi_new[idx] = phi[idx];
         return;
     }
 
@@ -66,7 +66,6 @@ extern "C" __global__ void poisson_gauss_seidel_2d(
 
     // Boundary conditions
     if (i == 0 || i == nx - 1 || j == 0 || j == ny - 1) {
-        phi[idx] = 0.0f;
         return;
     }
 
@@ -105,7 +104,6 @@ extern "C" __global__ void poisson_sor_2d(
 
     // Boundary conditions
     if (i == 0 || i == nx - 1 || j == 0 || j == ny - 1) {
-        phi[idx] = 0.0f;
         return;
     }