From eba2f61ddfaa888f6208aff5caed0c7b112c95b1 Mon Sep 17 00:00:00 2001 From: Omar Sobh Date: Wed, 16 Sep 2026 07:34:28 -0500 Subject: [PATCH] =?UTF-8?q?rtx-cfd=20CUDA=20Poisson=20kernels:=20solve=20?= =?UTF-8?q?=E2=88=87=C2=B2=CF=86=20=3D=20f=20as=20documented=20(the=20upda?= =?UTF-8?q?tes=20subtracted=20the=20source=20with=20the=20wrong=20sign);?= =?UTF-8?q?=20the=20Jacobi=20test's=20stop=20is=20absolute=20on=20a=20sour?= =?UTF-8?q?ce=20of=20size=202=CF=80=C2=B2?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- crates/specialized/rtx-cfd/src/kernels/cuda/poisson.cu | 8 ++++---- crates/specialized/rtx-cfd/tests/kernels_tests.rs | 2 +- 2 files changed, 5 insertions(+), 5 deletions(-) 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,