diff --git a/crates/specialized/rtx-cfd/src/kernels/cuda/poisson.cu b/crates/specialized/rtx-cfd/src/kernels/cuda/poisson.cu
index 2cc6937..2eaada6 100644
--- a/crates/specialized/rtx-cfd/src/kernels/cuda/poisson.cu
+++ b/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])
diff --git a/crates/specialized/rtx-cfd/tests/kernels_tests.rs b/crates/specialized/rtx-cfd/tests/kernels_tests.rs
index f4b9f7f..9e2a57c 100644
--- a/crates/specialized/rtx-cfd/tests/kernels_tests.rs
+++ b/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,