189 Commits

Include renames

Author	SHA1	Message	Date
Omar SobhandClaude Fable 5.1	d337afa8f9	rtx-cfd embedded3 item 9: wall.rs (binary ghost mask: three face families, trilinear stencil with periodic-z wrap, z-weighted least-squares ghost fit, flux compatibility correction; slip walls allowed as touched sides) + loads.rs (surface-stress route with probes, control-volume route with full-span z faces skipped); the step wired (body/mask, predicates, anchor, ghost re-imposition). Gates HELD: CFD1 ny 41 nz 1 CV 15.6156 / surface 15.7126 both to 1e-6 of the 2D record; nz 4 periodic CV 1.1e-6 / surface 4.2e-4; sphere MMS order 0.89, div 1e-8, ghost correction 1.0e-4 → 2.0e-5, both load routes' errors falling (0.153 → 0.115 surface, 0.214 → 0.139 CV) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 15:14:40 -05:00
Omar SobhandClaude Fable 5.1	d4ffac9ac7	rtx-cfd embedded3 item 8: body.rs (SDF, surface velocity, normal, samplers: sphere Fibonacci lattice, z-cylinder, extruded 2D body) + cut.rs (Kuhn six-tet apertures/volumes, closure, wall distances); gate 8 HELD: volume/area orders 1.94–2.01, closure ≤ 3e-17, Lipschitz ratio 0.100, Σ ΔV per step 6.2e-4 of the swept volume ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 14:58:06 -05:00
Omar SobhandClaude Fable 5.1	54911b4db3	rtx-cfd embedded3 item 7: e3_step.cu + step::device::{DeviceStep, StepTimers} on the shared runtime (FMA off); gate 7 HELD: device = host ≤ 7e-12 under tight tolerances on MMS/Beltrami/Poiseuille, equal CG counts on every step, periodic planes within 2e-16; the default-tolerance differences are the projection's inner stop ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 14:56:46 -05:00
Omar SobhandClaude Fable 5.1	8821e18520	rtx-cfd embedded3 items 4–6: field.rs + step/{mod, predictor, projection} (364/700/419 lines) — the host PISO step re-laid from the verified three_d code; gates HELD: MMS + Poiseuille marches value-identical to the 2D embedded solver at nz=1 over 200 steps; 3D MMS orders 0.88 upwind / 1.61 TVD, div ≤ 5e-9; Beltrami 1.08 / 1.25 with face-averaged data, div−mean ≤ 7e-8; Poiseuille |u−û| ≤ 8e-10 at nz 1 and periodic nz 4 ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 14:55:18 -05:00
Omar SobhandClaude Fable 5.1	765ba6d3f6	rtx-cfd embedded3 item 3: e3_cg.cu (FMA off) + poisson::device_cg::{DeviceCg, DevicePcgCache, solve_pcg_device_cached}; gate 3 HELD: device CG = host PCG (9 / 9 iterations, |Δp| ≤ 1.5e-14 relative, run-to-run bit-identical, anchor exact, mean 9e-16) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 14:53:20 -05:00
Omar SobhandClaude Fable 5.1	5e2b565971	rtx-cfd embedded3 item 2: e3_mg.cu + poisson::{export, device} (one shared CUDA runtime and module loader for embedded3); gate 2 HELD: device V-cycle = host f32 to 4e-7 relative on 96×40×{1,8} and 378×62×62, 4.97 ms per V-cycle at 1.45 M cells incl. transfers ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 14:52:18 -05:00
Omar SobhandClaude Fable 5.1	6526a3bd38	rtx-cfd embedded3 item 1: the Poisson core as poisson/{problem, hierarchy, pcg} (≤ 385 lines each; the periodic wrap built once into neighbour arrays); gate 1 HELD: nz=1 bit-identical to the 2D solver (lex + red-black, cached/uncached, iterations) and every extrusion case bit-identical to the three_d oracle, planes within 1.4e-11 ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 14:50:18 -05:00
Omar SobhandClaude Fable 5.1	13d30c3ce6	rtx-cfd 3D Stage 1 item 6: three_d::geometry::{Body3, CutGeometry3} — nodal signed distance, face apertures and cell volumes exact for the linear interpolant on a fixed Kuhn six-tet split (consistent across faces), the wall polygon by closure, face-centre wall distances. Gate 7 HELD: sphere/z-cylinder volume and area orders 1.94–2.01, closure 1e-17; continuity: the largest neighbour change falls 10× for a 10× finer sweep (ratio 0.100); translating-sphere Σ ΔV per step MEASURED at 6e-4 of the swept volume (the 'to rounding' clause was wrong and is replaced by the number) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 11:49:04 -05:00
Omar SobhandClaude Fable 5.1	4747d9e6b1	rtx-cfd three_d piso_device: fix the by-reference launch args (the previous commit did not build with the cuda feature) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 11:44:49 -05:00
Omar SobhandClaude Fable 5.1	bd43b34850	rtx-cfd three_d piso_device: pass the pointer table by reference (clippy) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 11:44:20 -05:00
Omar SobhandClaude Fable 5.1	d9fd849239	rtx-cfd 3D Stage 1 item 5: piso3.cu + three_d::piso_device::Piso3Device — the PISO step device-resident (predictors compiled without FMA contraction = the host's arithmetic, per-side boundary tables, divergence/correct/add_p kernels with fixed-order partials, the device CG); poisson_operator/boundary_tables/source_tables/anchor_cell/inner_stop split out of the host solver; StepTimers3 (RTX_PROFILE). Gates 4–6 on the device HELD: MMS 9e-16 (upwind) / 7e-16 (TVD), Beltrami 3.3e-12, Poiseuille 2e-15 with periodic planes equal to 2e-16 (tight tolerances; default-tolerance differences = the projection's inner stop); 55 ms per step at 378×62×62 (predictor 22, poisson 32, apply 2) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 11:44:04 -05:00
Omar SobhandClaude Fable 5.1	e0cbb99343	rtx-cfd three_d: the predictors split into piso_predictor.rs (file-size rule), the w predictor's index helpers cleaned; tests unchanged and green ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 11:33:13 -05:00
Omar SobhandClaude Fable 5.1	616d2a3394	rtx-cfd 3D Stage 1 item 4: three_d::{FlowField3D, piso_host::Piso3Solver} — the 2D embedded predictor/projection transcribed with the z terms appended (u/v/w predictors, six sides incl. periodic z, TVD, the apertured-ready projection on PoissonProblem3D); gates 4/5/6 HELD: MMS + Poiseuille marches value-identical to the 2D embedded solver (multigrid) over 200 steps at nz=1; 3D MMS orders 0.88 upwind / 1.61 TVD, div 1e-9; Beltrami orders 1.08/1.25 with face-averaged data (box compatible to 1e-12); Poiseuille |u−û| ≤ 2e-10, |v|,|w| ≤ 4e-10, p spread ≤ 5e-8 at nz 1 and periodic nz 4 ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 11:31:30 -05:00
Omar SobhandClaude Fable 5.1	2f476a38d5	rtx-cfd 3D Stage 1 item 3: cg3.cu + DeviceCg3 — the f64 CG entirely on the device (run_pcg's control flow: true-residual stop and resync, breakdown guard, mean projection, anchor shift; fixed-order block reductions, no atomics; the device V-cycle as preconditioner; DevicePcgCache3); gate 3 HELD: iterations = host, |Δp| ≤ 2e-14 relative, run-to-run bit-identical, singular box anchor exact / mean 2e-15 ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 11:18:52 -05:00
Omar SobhandClaude Fable 5.1	05ed96f383	rtx-cfd 3D Stage 1 item 2: mg3_vcycle.cu (seven-point red-black half-sweep, residual, coarsest; top/bot neighbour arrays carry the periodic wrap), LevelExport3/export_hierarchy3/vcycle_f32_reference3, DeviceVcycle3 (the K=1 sequence); gate 2 HELD: device = host f32 V-cycle to 4e-7 relative on 96×40×{1,8} (periodic and walls) and 378×62×62; 4.88 ms per V-cycle at 1.45 M cells incl. transfers ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 11:15:03 -05:00
Omar SobhandClaude Fable 5.1	c13b07b68d	rtx-cfd three_d: ignore reasons on the probes (clippy) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 11:12:23 -05:00
Omar SobhandClaude Fable 5.1	0794e1b6db	rtx-cfd 3D Stage 1 item 1: three_d::{Grid3, poisson} — the 2D Poisson stack transcribed to a seven-point operator (sanitised coefficients, 2×2×2 Galerkin aggregation, (i+j+k)%2 colouring, periodic z, run_pcg line for line, PcgCache3); gate 1 HELD: nz=1 bit-identical to the 2D solver (solution + iterations, lex + red-black, cached/uncached); extrusion z-invariant to the solve's accuracy (bit-identical planes for lexicographic decoupled); probes for the aggregation/colouring interaction ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 11:11:52 -05:00
Omar SobhandClaude Fable 5.1	7e135893a2	rtx-cfd/rtx-fsi PERF-2 P3-iii K-stream mode: each lane thread's device operator lives on its own CUDA stream (no rendezvous; set_plane_streams, RTX_FSI2O_PLANE_MODE=streams default | batch); stream-mode unit test digit-identical to the solo apply ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 06:40:43 -05:00
Omar SobhandClaude Fable 5.1	95ffde9591	rtx-cfd/rtx-fsi PERF-2 P3-iii: the MarchPlane — lane-table V-cycle kernels (ml_*), the batched apply over K lanes' own operators, the plane rendezvous inside the CG's preconditioner (lane thread-local, CG guard), set_plane_lane/plane_counters; the fsi2_overset_plane driver (K lane threads, one process); G-P0 unit tests: K distinct operators batched = solo apply bit for bit, uneven lanes served exactly ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-16 17:03:51 -05:00
Omar Sobh	49ccebc863	rtx-cfd legacy GPU tests: the two point-source Poisson tests expect the sign of ∇²φ = f	2026-09-16 07:36:07 -05:00
Omar Sobh	eba2f61ddf	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π²	2026-09-16 07:34:28 -05:00
Omar Sobh	52fd40f8f2	rtx-cfd legacy GPU tests: the Jacobi Poisson test stops at 1e-4 (the accuracy check is the pin); rustfmt	2026-09-16 07:32:30 -05:00
Omar Sobh	1f77ec637a	rtx-cfd legacy GPU: per-block buffer mutability; the D3Q19 getters extract before borrowing	2026-09-16 07:30:44 -05:00
Omar Sobh	8999548aca	rtx-cfd legacy GPU: buffers re-borrowed after nested device calls; writing and delegating methods take &mut self	2026-09-16 07:30:02 -05:00
Omar Sobh	fd5972c949	rtx-cfd legacy GPU: the kernel-writing methods take &mut self; read-only ones keep a shared borrow	2026-09-16 07:29:05 -05:00
Omar Sobh	0de750bad1	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	2026-09-16 07:28:42 -05:00
Omar Sobh	f3f6ac21fb	rtx-cfd CUDA kernels: no host includes at all in the NVRTC sources; the float limits defined inline	2026-09-16 07:24:54 -05:00
Omar Sobh	64faa095a2	rtx-cfd CUDA kernels: no host headers in the NVRTC sources (cuda_runtime.h was never found: every legacy GPU test failed at manager creation)	2026-09-16 07:23:02 -05:00
Omar Sobh	f429bdd415	rtx-cfd legacy GPU modules: the launches that already sat in unsafe blocks are not wrapped twice	2026-09-16 07:20:51 -05:00
Omar Sobh	7f4d437199	rtx-cfd: rustfmt over the repaired GPU modules and tests	2026-09-16 07:19:50 -05:00
Omar Sobh	3dadf1ab5c	rtx-cfd legacy GPU tests: one base struct per literal	2026-09-16 07:19:27 -05:00
Omar Sobh	147e0c0422	rtx-cfd legacy GPU tests: struct-update commas and the Poisson in-out buffer	2026-09-16 07:19:15 -05:00
Omar Sobh	534801e16a	rtx-cfd legacy GPU tests: the kernels' output buffers are &mut	2026-09-16 07:18:58 -05:00
Omar Sobh	1381b5b633	rtx-cfd legacy GPU tests: the current CudaKernelManager API (Arc-held manager, allocate-on-copy, copy_from_device returning the host vector), MatrixOpsKernel::tridiagonal_matvec added over the existing kernel, PoissonKernel::solve_jacobi_2d kept as the older name; CfdConfig literals take ..Default	2026-09-16 07:18:34 -05:00
Omar Sobh	af19860c04	rtx-cfd k-ε GPU: memcpy_dtoh argument order	2026-09-16 07:16:52 -05:00
Omar Sobh	5e402918bb	rtx-cfd legacy GPU modules: cudarc 0.18 API — unsafe launches, the stream accessor, memcpy_htod argument order; the renamed k-ε constants and ::standard(); get_module clones; CfdConfig test literals take ..Default	2026-09-16 07:16:39 -05:00
Omar Sobh	de435eb91c	rtx-cfd: the legacy GPU modules back in the cuda build (repair in progress)	2026-09-16 07:15:24 -05:00
Omar Sobh	3fd4fb7b8c	rtx-fsi: the interface noise probe reads RTX_FSI2_MG_DEVICE (the device V-cycle's regime gate)	2026-09-16 01:24:14 -05:00
Omar SobhandClaude Fable 5.1	fc556f8a88	PERF-2 P3-ii: the device V-cycle as the CG's preconditioner — poisson/device.rs (one CUDA runtime per process, persistent per-operator buffers, the mg_vcycle.cu kernels at K = 1; upload r, run the V-cycle, download z; the f64 CG unchanged), MultigridParameters::device, Prepared holds the device hierarchy, the CG driver destructures the prepared operator instead of cloning it; EmbeddedPisoSolver::set_poisson_device, overset pass-through, harness knob RTX_FSI2O_MG_DEVICE=1; export_levels factored out; the quarantine's dangling cfg attribute fixed ... Co-Authored-By: Claude Fable 5.1 <[email protected]> Claude-Session: https://claude.ai/code/session_01YJPeT6WA2e7YvAnS875AHL	2026-09-16 01:16:55 -05:00
Omar Sobh	3b7f8fb362	PERF-2 P3 bench: the coarsest level uses the CPU's red-black symmetric ordering (a nearly singular coarsest system settles on a smoother-dependent constant, which the prolongation doubled up every level: the 19 % disagreement)	2026-09-16 01:10:12 -05:00
Omar Sobh	06a60b5fed	PERF-2 P3 bench: per-level comparison of the device V-cycle's work vectors against the CPU's (vcycle_f32_work) to locate the disagreement	2026-09-16 01:08:31 -05:00
Omar Sobh	cb1647c0ef	rtx-cfd: the legacy GPU solver modules (piso_gpu, simple_gpu, d2q9_gpu, d3q19_gpu, smagorinsky_gpu) kept out of the cuda build — broken since the initial commit; the PERF-2 benchmark uses cudarc directly	2026-09-16 01:05:39 -05:00
Omar Sobh	956e9d38e2	rtx-cfd: the cuda feature builds again — the duplicate CudaKernelManager::stream removed, the legacy k-epsilon GPU model (broken since the initial commit) kept out of the build	2026-09-16 01:03:33 -05:00
Omar SobhandClaude Fable 5.1	0dc95a8de5	PERF-2 P3: the batched device V-cycle microbenchmark — mg_vcycle.cu (masked variable-coefficient five-point red-black half-sweeps, residual, CSR restriction in a fixed order, prolongation, coarsest sweeps; [K][n] layout, blockIdx.y = march), LevelExport/export_hierarchy and vcycle_f32_reference on the CPU side, and the ignored cuda-feature test that checks the K = 1 device V-cycle against the CPU f32 reference and times K = 1 / 8 / 16 per march ... Co-Authored-By: Claude Fable 5.1 <[email protected]> Claude-Session: https://claude.ai/code/session_01YJPeT6WA2e7YvAnS875AHL	2026-09-16 01:03:02 -05:00
Omar Sobh	53078c2aa7	rtx-cfd: get_features compiles under --features cuda (the vector must be mutable)	2026-09-16 01:00:34 -05:00
Omar SobhandClaude Fable 5.1	8c130c55c8	PERF-2 P2-c: RTX_FSI2O_INNER_STOP knob — the background Poisson's inner stop as a multiple of the Schwarz tolerance (default 0.1; the band-gated ladder on the CG iteration count) ... Co-Authored-By: Claude Fable 5.1 <[email protected]> Claude-Session: https://claude.ai/code/session_01YJPeT6WA2e7YvAnS875AHL	2026-09-16 00:50:52 -05:00
Omar Sobh	822041fb3f	PERF-2 P2: the parallel maps run only on levels with ≥ 8,192 active cells and in tasks of ≥ 2,048 cells (a fork-join over a small level costs more than the map); values unchanged	2026-09-16 00:10:42 -05:00
Omar SobhandClaude Fable 5.1	7f144e0005	PERF-2 P2: the red-black colour maps, the residual and the V-cycle's per-level maps on rayon threads (MultigridParameters::threads; RTX_THREADS in the harness; rtx_cfd::configure_threads) — each colour's values computed from the unchanged other colour into a scratch and written back, the L1 sum in the serial order: bit-identical to the serial red-black (pin: 4 right-hand sides at 4 threads); the coarsest level stays serial; no effect on the lexicographic regime ... Co-Authored-By: Claude Fable 5.1 <[email protected]> Claude-Session: https://claude.ai/code/session_01YJPeT6WA2e7YvAnS875AHL	2026-09-15 23:59:29 -05:00
Omar Sobh	364766a54a	rtx-fsi: the interface noise probe reads RTX_FSI2_MG_RB itself (its solver is built outside the march's knob site; the first red-black probe run was vacuous)	2026-09-15 23:53:57 -05:00
Omar Sobh	f9d38bee55	rtx-cfd/rtx-fsi: every EmbeddedParameters literal carries the smoother field (all test targets of both crates compile)	2026-09-15 23:50:19 -05:00