459 Commits

This Branch

This Branch

All Branches

Author	SHA1	Message	Date
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
Omar Sobh	43d5abde8d	rtx-fsi harness: the embedded harness's EmbeddedParameters carries the smoother field (the previous commit did not compile the rtx-fsi tests)	2026-09-15 23:49:31 -05:00
Omar SobhandClaude Fable 5.1	cb5771fa71	PERF-2: red-black symmetric Gauss–Seidel smoother behind a knob (MgSmoother::RedBlack; default lexicographic = the recorded regime, bit-identical) — red/black cell lists per level, the symmetric pair red,black,black,red, the operator cache keyed on the smoother; EmbeddedParameters::poisson_smoother, harness knobs RTX_FSI2O_MG_RB (overset) and RTX_FSI2_MG_RB (embedded, the noise probe); pin: solves the masked problem to the same stop, agrees with lexicographic to 2e-12, cached = uncached bit for bit ... Co-Authored-By: Claude Fable 5.1 <[email protected]> Claude-Session: https://claude.ai/code/session_01YJPeT6WA2e7YvAnS875AHL	2026-09-15 23:48:55 -05:00
Omar Sobh	26904e37d8	PERF-2 P1.b: the polyline index prunes from the start with an upper bound from each chunk's first vertex (never the answer; the strict-minimum scan rule is unchanged), so a query near the polygon's far end no longer scans every chunk; pin unchanged, 0 mismatches	2026-09-15 23:35:40 -05:00
Omar SobhandClaude Fable 5.1	e0b5983436	PERF-2 P1.2: the overlap map's cell classification skips the point location for background cells outside the patch's node bounding box (no patch cell can contain them and the body lies inside the patch: Active, the class the search returns); PatchMesh::bounding_box ... Co-Authored-By: Claude Fable 5.1 <[email protected]> Claude-Session: https://claude.ai/code/session_01YJPeT6WA2e7YvAnS875AHL	2026-09-15 23:34:40 -05:00
Omar SobhandClaude Fable 5.1	79c18def32	PERF-2 P1.1: the multigrid-PCG's prepared operator (hierarchy, f64 fine level, active cells, components) cached on the embedded solver and reused while the operator is bit-identical (an exact key over the coefficient bit patterns, the mask and the hierarchy parameters); solve_multigrid_pcg_cached gives the uncached solve's answer bit for bit (pin poisson_cache_exact.rs: five right-hand sides on one operator, a one-coefficient miss, both precisions); the CG driver split into Prepared::build + run_pcg ... Co-Authored-By: Claude Fable 5.1 <[email protected]> Claude-Session: https://claude.ai/code/session_01YJPeT6WA2e7YvAnS875AHL	2026-09-15 23:30:53 -05:00
Omar SobhandClaude Fable 5.1	01830e5c8d	PERF-2 P1: exact chunked polyline index for the O-grid's sliding outer ring — PolylineIndex::nearest returns nearest_on_polyline's point bit for bit (same segments, same order, same strict minimum; chunks skipped only when their box is farther than the best by more than 1e-12 relative), used by the Winslow sweep and the two ring projections; pin: 0 mismatches over 39,864 queries on a 3,166-point rounded outline including vertices, mid-points and ties ... Co-Authored-By: Claude Fable 5.1 <[email protected]> Claude-Session: https://claude.ai/code/session_01YJPeT6WA2e7YvAnS875AHL	2026-09-15 23:08:09 -05:00
Omar Sobh	11f71b832c	rtx-cfd: the regeneration profile block sits above the too_many_arguments attribute, not between it and its function	2026-09-15 22:53:40 -05:00
Omar SobhandClaude Fable 5.1	1547d14ff1	PERF-2 P0-b: CG iterations per Poisson solve in the profile (ms per V-cycle) and sub-timers inside the patch regeneration (outline, hull + offset, ring projection, Winslow sweeps, respace, warm bookkeeping, mesh finalisation) ... Co-Authored-By: Claude Fable 5.1 <[email protected]> Claude-Session: https://claude.ai/code/session_01YJPeT6WA2e7YvAnS875AHL	2026-09-15 22:53:26 -05:00
Omar SobhandClaude Fable 5.1	172a26dee4	PERF-2 P0: intra-step profiler — StepTimers on the overset solver (RTX_PROFILE; overlap build / predictors / patch BiCGSTAB / background Poisson with its setup-vs-iterate split / round exchange / apply / end exchange), PoissonSolution carries setup_ns and iterate_ns, the harness times advance / restore / snapshot / load sampling / force / FEA predictor and prints the wall split of the coupled phase; no clock is read when profiling is off; the reclassified/step progress denominator fixed (was ×4) ... Co-Authored-By: Claude Fable 5.1 <[email protected]> Claude-Session: https://claude.ai/code/session_01YJPeT6WA2e7YvAnS875AHL	2026-09-15 22:25:52 -05:00
Omar SobhandClaude Fable 5.1	2d2fed7181	rtx-fsi harness: 'fixed' coupler (constant relaxation OMEGA0) for the P6-c Robin–Neumann test — at ω = 1 the candidate is the FEA's own output, so the Robin wall's previous-pass datum is the structure's traction at the candidate; design 2's FSI3 divergence is the datum's inconsistency under IQN's ω0 = 0.05 ... Co-Authored-By: Claude Fable 5.1 <[email protected]> Claude-Session: https://claude.ai/code/session_01YJPeT6WA2e7YvAnS875AHL	2026-09-15 19:22:36 -05:00
Omar SobhandClaude Fable 5.1	96a7f1c700	rtx-cfd: Robin wall sign pin — a still closed annulus under a uniform pressure above a zero datum: the explicit offset recedes into the body at exactly −p₀/(α(1+g)) (3e-16), the implicit compliant term answers with a pressure DROP (mean 0.17 of p₀, max 0.36) that absorbs the recession (net wall flux 6e-19 of 0.51); the steady MMS pin is blind to the sign of both parts ... Co-Authored-By: Claude Fable 5.1 <[email protected]> Claude-Session: https://claude.ai/code/session_01YJPeT6WA2e7YvAnS875AHL	2026-09-15 11:17:38 -05:00
Omar SobhandClaude Fable 5.1	1c98156653	rtx-cfd: the deformed-tip outline chooses the flat / semicircle branch against the UNDEFORMED half-thickness — at the semicircle setting the deformed tip's radius breathes by nanometres and flipped the topology (143 → 144 points) within the first steps of the P6-b ny 41 cap-12 release at every α; pin: the point count is fixed under ± 1e-6 thickness breathing at h 10 / 6.6 mm and ny 41, semicircle and flat ... Co-Authored-By: Claude Fable 5.1 <[email protected]> Claude-Session: https://claude.ai/code/session_01YJPeT6WA2e7YvAnS875AHL	2026-09-15 11:13:55 -05:00
Omar SobhandClaude Fable 5.1	c3dbd5040a	P6-b design 2: the Robin wall on the patch's Inner side — RobinWall { alpha, datum } with the wall velocity u_s + (t_f − datum)/α (explicit, damped by the wall's own viscous gain μ/(α d)) and the compliant-wall pressure term |S| p'/α implicit in the projection (the added-mass operator in the fluid's own response); Dirichlet bit for bit when off; MMS pin on the skewed annulus (orders 2.46/2.13 at α = 10 and 100 μ/h, the Dirichlet values; a 1e12 wall reproduces Dirichlet to 2e-6); OversetPisoSolver::patch_mut; harness knob RTX_FSI2O_ROBIN_ALPHA with the previous pass's tractions as the datum, printed in the header ... Co-Authored-By: Claude Fable 5.1 <[email protected]> Claude-Session: https://claude.ai/code/session_01YJPeT6WA2e7YvAnS875AHL	2026-09-15 10:46:35 -05:00
Omar SobhandClaude Fable 5.1	c144a5f733	P6-b: the compensating load uses the candidate's own Newmark acceleration a(d) = (d − u_pred)/(β Δt²) (consistent with the relaxed candidate the fluid saw; cancels exactly at the fixed point) — the lagged structure-output acceleration of the first try made the loop slower (1.8 → 5.2 subit/step at α 0 → 2) ... Co-Authored-By: Claude Fable 5.1 <[email protected]> Claude-Session: https://claude.ai/code/session_01YJPeT6WA2e7YvAnS875AHL	2026-09-15 08:37:35 -05:00
Omar SobhandClaude Fable 5.1	a2086a59de	P6-b: the fictitious added mass for the partitioned loop — rtx-fea NonlinearDynamicStepper::set_added_lumped_mass (a lumped per-DOF mass in the Newmark inertial residual and effective tangent, never the consistent mass or the rest state; zero = the plain stepper bit for bit) with its pin (compensated step reproduces the plain step to 2.5e-9, uncompensated moves it 11 %); the FSI2 overset harness carries RTX_FSI2O_FICT_MASS=α (α × ρ_f π (c/2)² spread over the wetted nodes) and adds the compensating load M_f ü_k of the previous subiterate to every structure solve (predictor and passes), printed in the header ... Co-Authored-By: Claude Fable 5.1 <[email protected]> Claude-Session: https://claude.ai/code/session_01YJPeT6WA2e7YvAnS875AHL	2026-09-15 08:30:55 -05:00