diff --git a/crates/specialized/rtx-cfd/src/solvers/lbm/d3q19_gpu.rs b/crates/specialized/rtx-cfd/src/solvers/lbm/d3q19_gpu.rs index 447d675..31d9ae7 100644 --- a/crates/specialized/rtx-cfd/src/solvers/lbm/d3q19_gpu.rs +++ b/crates/specialized/rtx-cfd/src/solvers/lbm/d3q19_gpu.rs @@ -671,7 +671,8 @@ __global__ void d3q19_bounce_back_boundaries( } /// Get macroscopic variables at a specific point (copy from GPU) - pub fn gpu_macroscopic_variables_at(&mut self, + pub fn gpu_macroscopic_variables_at( + &mut self, x: usize, y: usize, z: usize, diff --git a/crates/specialized/rtx-cfd/tests/kernels_tests.rs b/crates/specialized/rtx-cfd/tests/kernels_tests.rs index a84975c..f4b9f7f 100644 --- a/crates/specialized/rtx-cfd/tests/kernels_tests.rs +++ b/crates/specialized/rtx-cfd/tests/kernels_tests.rs @@ -48,7 +48,13 @@ mod cuda_tests { // Run advection kernel // CFL = v dt / dx ≈ 0.1: march 50 steps so the pulse moves several cells. for _ in 0..50 { - advection_kernel.apply(&d_phi, &mut d_phi_new, velocity as f32, dt as f32, dx as f32)?; + advection_kernel.apply( + &d_phi, + &mut d_phi_new, + velocity as f32, + dt as f32, + dx as f32, + )?; std::mem::swap(&mut d_phi, &mut d_phi_new); } @@ -204,8 +210,10 @@ mod cuda_tests { d_source = kernel_manager.copy_to_device(&source)?; // Solve Poisson equation - let max_iterations = 20_000; // Jacobi needs O(n²) sweeps on this grid - let tolerance = 1e-6; + // 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 iterations = poisson_kernel.solve_2d( &mut d_phi, &d_source,