24 Commits

Include renames

Author	SHA1	Message	Date
Omar SobhandClaude Fable 5.1	9c3fac0755	embedded3 S2-6: the transverse gradient from the face at least 0.2 h off the wall (full neighbours over their own distance along the cut partner's normal) — the first form blew up at DFG ny 61 with the fine floor (coefficient ~ 1/d_f); host = device 7.7e-14 with the closures on; Couette linear exactness kept (<= 0.01 h); RTX_E3_OBLIQUE_N rung knob ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-18 16:54:52 -05:00
Omar SobhandClaude Fable 5.1	e9b2e7887b	embedded3 S2-6: the oblique-wall instrument (tests/embedded3_wall_position_oblique.rs: z-flow / in-plane Poiseuille + Couette linear exactness; gate oblique_wall_position_is_second_order) and the closures it found, host + device, default OFF — the transverse centroid correction (RTX_E3_DIFFUSION_TRANSVERSE=1, wall_order bit 7), the fine distance floor (RTX_E3_DISTANCE_FLOOR=fine, bit 8); kernel geometry factored into cut_cv; DFG 2D-1 x-shift knob (RTX_E3_DFG_SHIFT_X); FLAG_X0 0.6 -> 0.25 in the three flag tests (the flag of every record was detached) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-18 16:22:38 -05:00
Omar SobhandClaude Fable 5.1	fdfb6da769	embedded3 S2-5: the cut wall sat ½(1−α)h inside the body — cross diffusion over the open-part centroid spacing (RTX_E3_DIFFUSION_CENTROID; host + e3_cut.cu, shift tables, point-implicit excess); flat-wall effective-position instrument; DFG 2D-1 ladder tests (device + host); knobs tried and refuted along the way (oblique distance, axis exchange, centroid pressure gradient) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-18 10:54:52 -05:00
Omar SobhandClaude Fable 5.1	cda973df26	embedded3 S2-4: quadratic wall gradient behind Parameters::wall_order (host predictor, operator route, e3_cut.cu; RTX_E3_WALL_ORDER); density pin holds at order 2 ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-18 07:21:57 -05:00
Omar SobhandClaude Fable 5.1	f6add276c0	embedded3: cut predictor convection carries ρ (host + e3_cut.cu; density-scaling pin); operator load route includes the wall exchange (exchange.rs); reconstructed_parts, probe aperture floor knob; dfg_split diagnostic test ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-18 03:59:57 -05:00
Omar SobhandClaude Fable 5.1	aca77fe1ac	embedded3: the cut mask's imposition and the device surface-velocity tables restricted to a 4-cell band (deeper prescribed faces are never read; the surface velocity was the moving step's cost); the operator rebuild counted in the rebuild share ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 19:33:33 -05:00
Omar SobhandClaude Fable 5.1	aa096465a6	embedded3 S2-2a/S2-3: moving bodies on the device (phased tables, host rebuild per step, impose kernel; gate test), the per-span cut wall route, the flag-wake driver and its geometry pre-flight ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 18:53:00 -05:00
Omar SobhandClaude Fable 5.1	3b3d6c84c0	embedded3 item 9b: the device step carries a static cut-cell mask (e3_cut.cu: cut predictor, apertured merged continuity with the fold, owner-read corrections; device CG off-stencil links) — host = device to 2e-10 (CFD1 cylinder nz 4) and 4e-14 (sphere) under tight tolerances ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 17:22:52 -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	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	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	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	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	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	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	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	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	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
redclawsystems	4d88dc0584	Initial commit	2026-03-04 00:08:42 +00:00