rtx-cfd CUDA Poisson kernels: solve ∇²φ = f as documented (the updates subtracted the source with the wrong sign); the Jacobi test's stop is absolute on a source of size 2π²

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

@@ -38,7 +38,7 @@ extern "C" __global__ void poisson_jacobi_2d(
 
     // Jacobi iteration: φ^{k+1}_ij = factor * (source_ij + (1/dx²)(φ_{i+1,j} + φ_{i-1,j}) + (1/dy²)(φ_{i,j+1} + φ_{i,j-1}))
     phi_new[idx] = factor * (
-        source[idx] +
+        -source[idx] +
         dx2_inv * (phi[idx_ip1] + phi[idx_im1]) +
         dy2_inv * (phi[idx_jp1] + phi[idx_jm1])
     );
@@ -76,7 +76,7 @@ extern "C" __global__ void poisson_gauss_seidel_2d(
 
     // Gauss-Seidel update (in-place)
     phi[idx] = factor * (
-        source[idx] +
+        -source[idx] +
         dx2_inv * (phi[idx_ip1] + phi[idx_im1]) +
         dy2_inv * (phi[idx_jp1] + phi[idx_jm1])
     );
@@ -116,7 +116,7 @@ extern "C" __global__ void poisson_sor_2d(
 
     // Compute Gauss-Seidel update
     float phi_gs = factor * (
-        source[idx] +
+        -source[idx] +
         dx2_inv * (phi[idx_ip1] + phi[idx_im1]) +
         dy2_inv * (phi[idx_jp1] + phi[idx_jm1])
     );
@@ -195,7 +195,7 @@ extern "C" __global__ void poisson_jacobi_3d(
 
     // 3D Jacobi iteration
     phi_new[idx] = factor * (
-        source[idx] +
+        -source[idx] +
         dx2_inv * (phi[idx_ip1] + phi[idx_im1]) +
         dy2_inv * (phi[idx_jp1] + phi[idx_jm1]) +
         dz2_inv * (phi[idx_kp1] + phi[idx_km1])

crates/specialized/rtx-cfd/tests/kernels_tests.rs

@@ -213,7 +213,7 @@ mod cuda_tests {
         // Jacobi needs O(n²) sweeps on this grid and its residual floor in f32 sits
         // above 1e-6; the accuracy check below is the pin, the stop is loose.
         let max_iterations = 20_000;
-        let tolerance = 1e-4;
+        let tolerance = 1e-2; // absolute, on a source of size 2π²
         let iterations = poisson_kernel.solve_2d(
             &mut d_phi,
             &d_source,