100 Commits

This Branch

This Branch

All Branches

Author	SHA1	Message	Date
Omar SobhandClaude Fable 5.1	44caeb8110	embedded3 PERF-3 P1-3 (f): the mask build's whole-grid scans as parallel per-entry maps (cell classification with the anchor as the smallest fluid index, the three face-kind tables, the merging scan's small flags, the centroid-shift table per face) — slab flag ny 62 CSV byte-identical, device moving/cg green; ny 124 rebuild block 2,680 → 2,569 ms per step (build_mask 409 → 309) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-20 12:31:00 -05:00
Omar SobhandClaude Fable 5.1	5a376630a7	embedded3 PERF-3 P1-3 (c)(d)(e): the moving path's host mirror persists between steps and only the arrays the rebuild reads come down (u, v, w, p, the previous velocities) and only those it writes go up (u, v, w, p — the correction stays on the device; the first form re-uploaded a stale one and the byte gate caught it); the operator assembly in parallel by plane with the predicates lifted off the solver; the fine export's casts as parallel maps — slab flag ny 62 CSV byte-identical, device moving/cg green; ny 124 rebuild block 3,386 → 2,680 ms per step ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-20 12:23:01 -05:00
Omar SobhandClaude Fable 5.1	651f75d96a	embedded3 PERF-3 P1-3 (a): the operator refresh moves the operator into the level (no eight-array clone) and takes its components/links from the level's problem; Components::find_parallel (hook-and-shortcut union-find, same numbering) on the refresh and the hierarchy rebuild; the S2-7b device carry's shift-table clone (830 MB per build at ny 124) replaced by a borrow — slab flag ny 62 CSV byte-identical, device moving/cg tests green; ny 124 rebuild block 4,570 → 3,353 ms per step ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-20 11:57:13 -05:00
Omar SobhandClaude Fable 5.1	f6b34f924d	embedded3 S2-7b: the moving-wall forms on the device — the ub table at the open part's centroid foot (RTX_E3_WALL_FOOT=centroid, table only), the solid exchange's wall velocity at the axis feet (RTX_E3_WALL_EXCHANGE=axisfoot: per-face six-entry foot table, bit 12 in e3_cut.cu), the true-normal wall flux through the shared host table (RTX_E3_WALL_FLUX=true); refusals lifted (conv sides stay host-only); host = device green on the sphere (varying surface velocity), the cylinder and the moving circle with each knob and all three ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-20 08:59:30 -05:00
Omar SobhandClaude Fable 5.1	7fffbb1301	embedded3 S2-7b: the curved instrument's per-wall operator probe (curved_operator_probe) and RTX_E3_CURVED_SCHEME; host prototype RTX_E3_WALL_FOOT=centroid (the shear closure's wall velocity at the foot from the open part's CENTROID, not the face centre — on a rotating wall the two feet are ½h(1−α) apart: rigid-mode probe residual 4.5× smaller, the rotating wall's offset −0.022 → ±0.002 h, its cut pressure zeroth → first order); default off, device refuses it ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-20 08:48:41 -05:00
Omar SobhandClaude Fable 5.1	599023d786	embedded3 S2-7b A1: the moving wall's mass flux from the body's true normal at the wall foot (RTX_E3_WALL_FLUX=true, host prototype, device refuses it) — the wall flux of a rotating cylinder 5e-3 → 1e-9 of ΩR h², the rotating wall's cut-cell pressure error 0.33 → 0.11 (n 16) / 0.30 → 0.09 (n 32) of ρ(ΩR)², full cells second order; the rotating wall's offset unchanged (−0.02 h) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-20 07:58:10 -05:00
Omar SobhandClaude Fable 5.1	d25a36d908	embedded3 S2-7b: the curved-wall instrument (tests/embedded3_wall_position_curved.rs: Taylor–Couette between embedded cylinders, rigid-rotation linear mode, outer-driven mode; walls' effective radii, pressure error by cell class against the exact p(r), wall-flux and ghost diagnostics); the sphere operator probe's sub-bands and force units; two host prototypes, default off, device refuses them: RTX_E3_WALL_EXCHANGE=axisfoot (the axis exchange takes the wall velocity at the axis foot — the S2-7 default read the solid face's own value, exact only for a uniform wall velocity: rigid rotation A 1.012 → 1.003) and RTX_E3_CONV_SIDES=exact (convective mass fluxes from the sides' own apertures; moves nothing) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-20 07:54:01 -05:00
Omar SobhandClaude Fable 5.1	c430510802	embedded3 S2-7b instruments: the quadratic / full-cells-only pressure probe (Mask::pressure_at_quadratic[_from]) and the DFG 2D-1 test's RTX_E3_DFG_DP_PROBE line; the manufactured sphere's pressure-error read by cell class + six signed wall-point reads, RTX_E3_MMS_SCHEME=tvd, the cut_cell_pressure_ladder and sphere_operator_probe tests (defaults untouched; MMS gates green) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-20 07:34:40 -05:00
Omar SobhandClaude Fable 5.1	44ecc164f1	embedded3 S2-7 VERDICT: the diffusive exchange with SOLID neighbour faces over the axis distance (wall_exchange_axis, S2-5's knob) is the DEFAULT — the oblique wall's h-independent offset closes (four-registration means ≤ 0.005 h at n 32, p_cut halving per rung), DFG 2D generic ladder monotone +0.68/+0.10/+0.014 %, DFG 3D ny 62 / 124 inside the bar (6.1877 / 0.00907 / 0.1691 at ny 124); RTX_E3_WALL_EXCHANGE=h restores every earlier record digit for digit; host suite 21/21, device cut + moving identities green on the new default ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-19 23:25:50 -05:00
Omar SobhandClaude Fable 5.1	b80af59ca4	embedded3 S2-7: oblique instrument gains registration sweep / slope subset / z-flow skip / merged-cell count (RTX_E3_OBLIQUE_{C0_SHIFTS,SLOPES,ZFLOW,NS}) and the operator probe (oblique_operator_probe: predictor acceleration on the exact centroid-valued field by aperture band, cut-cell divergence under four valuations, spurious pressure; z-flow control); two host prototypes, default off, device refuses them: RTX_E3_CV_SIDES=exact (the momentum CV's side apertures from the interpolant on the half faces / cell-centre planes) and RTX_E3_WALL_ORDER2=centroid (the quadratic wall gradient's second point at the neighbour's centroid distance); quad_fraction / tri_area_fraction lifted to module fns (default path digit-identical: oblique record, host suite 21/21, device cut tests) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-19 22:26:37 -05:00
Omar SobhandClaude Fable 5.1	0cf9d20afb	embedded3 A1-b: the tiled momentum volume as a host prototype (RTX_E3_MOMENTUM_VOLUME=tiled; pressure force, source and inertia on ½(V−+V+)hA; default off; the device refuses it) — REFUTED on the oblique instrument (level moves, the h-independent remainder and the near-wall pressure residual do not) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-19 20:38:29 -05:00
Omar SobhandClaude Fable 5.1	b126915759	embedded3 PERF-3 P3-1: the device predictor's interior fast path (bit 11, default ON; RTX_E3_PREDICT_FAST=0 = uniform path), the face update templated on the component + unrolled (392-byte local frame -> 64), exact-zero divisions/sqrt skipped on the fast path; all byte-identical (DFG 2D ny 122, slab flag ny 62, DFG 3D ny 62; device tests); predictor 22.4 -> 12.5 ms/step at DFG 3D ny 62 ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-19 20:29:57 -05:00
Omar SobhandClaude Fable 5.1	77ecc54b24	embedded3: GuessBasis re-export behind the cuda feature (the host-only build broke); A2 build: no convective exchange with prescribed faces (RTX_E3_EXCHANGE_CONVECTION=off, host + device bit 10, load route consistent) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-19 19:22:08 -05:00
Omar SobhandClaude Fable 5.1	d7250bc4cf	embedded3 A3-i: the advancing-wall friction closure (RTX_E3_WALL_ADVANCING=1, host + device, per-face normal-velocity table; f(0) = 1 exactly) — REFUTED as built on the slab ladder (flag shear 34.2 / 33.6 / 31.1 vs 28.9 / 32.3 / 30.1: spread 3.4 -> 3.1, now from above; worst mass residual 1e-8 -> 1e-4); knob stays off, default digit-identical ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-19 19:19:47 -05:00
Omar SobhandClaude Fable 5.1	d4cd7d9545	embedded3 PERF-3 P2: the Poisson initial guess by projection onto the last K corrections of the same corrector index (GuessBasis, DeviceCg::project_guess: Gram matrix rebuilt per solve, Cholesky with pivot dropping; e3_cg_axpy) behind RTX_E3_POISSON_PROJECT=K (default 0, digit-identical); static DFG 2D ny 122: CG 9.5 -> 2.5 per step, loads identical to 4 digits; the moving flag: -4 % only (consecutive corrections nearly uncorrelated); RTX_E3_POISSON_PROJECT_DEBUG prints the residual before / after ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-19 18:57:50 -05:00
Omar SobhandClaude Fable 5.1	845e0ae01a	embedded3 PERF-3 P1-2: the flag test's centreline polyline cached per thread and time (the 41-point rebuild with four transcendental evaluations each ran per surface-velocity call, ~10^6 calls per step): table faces 1,620 -> 54 ms, impose 622 -> 12 ms, build_mask 898 -> 404 ms, rebuild block 6.9 -> 4.3 s per step at 11.6 M cells; slab CSV byte-identical; table sub-laps ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-19 18:42:55 -05:00
Omar SobhandClaude Fable 5.1	d9a2752cb9	embedded3 PERF-3 P1-2: the predictor's device tables derived from the projection's (shared buffers move over; 282 -> 48 ms), the operator key lazy after a refresh (240 -> 0 ms), parallel active / owner maps — slab flag CSV byte-identical; RTX_E3_TABLES_REBUILD=1 = the full-rebuild reference ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-19 18:35:04 -05:00
Omar SobhandClaude Fable 5.1	f6f1fe048b	embedded3: export_fine_from is crate-private (it takes a crate-private level) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-19 16:42:13 -05:00
Omar SobhandClaude Fable 5.1	73e43d597d	embedded3 PERF-3 P1-1: the level constructor's masking pass, cell lists and the diagonal sum as parallel per-entry maps (level 0.90 -> 0.60 s per step at 11.6 M cells), digit-identical (slab flag CSV byte-identical; host Poisson / MMS / identity and device cg / moving / vcycle tests green) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-19 16:41:28 -05:00
Omar SobhandClaude Fable 5.1	40268ad19b	embedded3 PERF-3 P1-1: the moving path's host work, digit-identical — impose_from's solid faces by z plane in parallel (1.95 -> 0.62 s per step at 11.6 M cells), the f32 fine export derived from the f64 level the refresh already builds (export_fine 1.04 -> 0.34 s); sub-laps in the mask rebuild; rebuild block 9.4 -> 7.4 s ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-19 15:59:19 -05:00
Omar SobhandClaude Fable 5.1	103e2576ae	embedded3 S2-8: RTX_E3_PROBE_D (diagnostic: the reconstructed route's probes at D h and 2 D h) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-19 10:08:57 -05:00
Omar SobhandClaude Fable 5.1	5e1b3e0731	embedded3 S2-8 instrument: the load routes in PARTS per body — exchange split into diffusive / convective (cut_wall_exchange_parts), a diagnostic x window on every load route (exchange::set_load_window), the flag test as a 2D periodic slab (RTX_E3_FLAG_NZ) with an amplitude knob (RTX_E3_FLAG_AMP, 0 = frozen) and a PARTS summary; cut_wall_force / _per_span moved to exchange.rs (cutwall.rs line cap) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-19 08:53:40 -05:00
Omar SobhandClaude Fable 5.1	3f0bd0650a	embedded3 oblique instrument: RTX_E3_OBLIQUE_DRIVE=pressure (no body force; the sides pin the flow rate) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-18 19:12:29 -05:00
Omar SobhandClaude Fable 5.1	dbc5dbaefb	embedded3 S2-6: the three closures default ON (oblique distance, transverse centroid correction, fine floor) — RTX_E3_WALL_DISTANCE=face RTX_E3_DIFFUSION_TRANSVERSE=0 RTX_E3_DISTANCE_FLOOR=coarse reproduce every earlier record digit for digit (checked: DFG 2D ny 62 host); DFG 2D D/h 30 c_D -0.19 % -> -0.01 %, 3D ny 124 6.1750 / 0.00844 / 0.1692 inside the bar on both routes ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-18 19:05:51 -05:00
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	9ae0e42dcd	embedded3 S2-3d: face-momentum unsteady term for the box (Mask::box_momentum_rates, Solver::previous_apertures); falsifier prints the residual against it ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-18 12:58:30 -05:00
Omar SobhandClaude Fable 5.1	cdd5ad0b66	embedded3 S2-2b: the moving path's host work 5.0 s → 0.6 s per step at 1.45 M cells, bit-identical — Problem::diagonals (the per-cell link diagonal was cells × links: 1.8 s twice per step), counting CSR + rayon in the device tables, rayon in the cut geometry builder; RTX_E3_MOVING_PROFILE laps ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-18 12:48:45 -05:00
Omar SobhandClaude Fable 5.1	99e4a7214b	embedded3 S2-5: centroid diffusion default ON (RTX_E3_DIFFUSION_CENTROID=0 restores the records); flat_wall_position_is_second_order gate; closure.rs split (cutwall.rs under 700) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-18 12:37:56 -05:00
Omar SobhandClaude Fable 5.1	37086300ba	embedded3 S2-5: the exchange route mirrors the predictor's centroid spacing toward solid neighbours (operator = box back to the box's own error) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-18 11:01:27 -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	4c3e58fa27	embedded3: DFG 2D-1 ladder test (periodic-z slab on the device cut wall) — S2-5's premise check ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-18 09:41:43 -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	53b1babb91	embedded3: moving-body box route (Δ(Vu) unsteady term, Solver::previous_volumes), floor-source / closure-lag / applied-force instruments, moving.rs split; falsifier prints the route residual under each ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-18 06:27:04 -05:00
Omar SobhandClaude Fable 5.1	62b46194dd	embedded3: reconstructed route moved to reconstruct.rs (cutwall.rs back under 700 lines) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-18 04:34:54 -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	680041d63d	embedded3: the flag driver's dt-scale knob (the loads' dt ladder) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-18 03:06:18 -05:00
Omar SobhandClaude Fable 5.1	9c58fdf602	embedded3 S2-2b-ii: the per-step refresh exports the fine level alone (export_fine) instead of the whole hierarchy ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-18 01:12:33 -05:00
Omar SobhandClaude Fable 5.1	3eb40f98ce	embedded3: the flag driver logs the reconstructed wall route per sample and in the final read ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 23:19:07 -05:00
Omar SobhandClaude Fable 5.1	b34f6ad966	embedded3 S2-1 remedy: the reconstructed wall route on the cut geometry's polygons (two probes along the interpolant normal) — sphere MMS 11.3/8.3 % (the best route); wired into the DFG and flag drivers ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 23:18:21 -05:00
Omar SobhandClaude Fable 5.1	f2cdda4691	embedded3 S2-3c: the 2D solver on the flag wake's kinematics (the reference for the 3D full-span run) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 23:14:42 -05:00
Omar SobhandClaude Fable 5.1	6d2872e7d0	embedded3 S2-2b-ii HELD: the device V-cycle's fine level refreshed with the operator (coarser levels kept K steps), the fine rhs zeroed before the gather — moving circle identity 1.7e-14 at K = 10, CG +9 % ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 23:10:15 -05:00
Omar SobhandClaude Fable 5.1	c713bbbbb2	embedded3: study knobs — aperture_substeps (space-time sub-sampling), momentum_volume_cell_mean, RTX_E3_MERGE_FRACTION; all default off (refuted as the fresh-row residual's source) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 21:38:46 -05:00
Omar SobhandClaude Fable 5.1	33de5d0080	embedded3 S2-2b-ii: DeviceCg::refresh (fine operator per step, hierarchy every K steps, z zeroed before the scatter) + gate test; the flag driver's correctors 3 / inner 1e-3 (residual remedy measured 7.6e-9 on the moving circle); residual study test ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 21:15:02 -05:00
Omar SobhandClaude Fable 5.1	fc8d69af29	embedded3 S2-3b: the flag driver's span knob (full width = the 2D geometry extruded) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 21:11:16 -05:00
Omar SobhandClaude Fable 5.1	58a5501cef	embedded3: the cut mask's surface velocity uses the interpolant's distance and normal (one body call per face instead of eight); the device shares the surface-velocity tables between a step's phases; sphere gate re-read 0.965/0.961 ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 19:48:15 -05:00
Omar SobhandClaude Fable 5.1	11d856b3d4	embedded3: operator export profile at ny 62 (poisson_operator 0.03 s, export_hierarchy 1.44 s); export and poisson_operator public for profiling ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 19:44:17 -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	c30e3dba34	embedded3 S2-2b (host lever): narrow-band φ re-evaluation for moving bodies (max_surface_speed) — bit-identical, the flag body's ny 62 rebuild 3.54 → 0.30 s; the flag driver uses it ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 19:30:54 -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	2c94ae0bb1	embedded3 S2-1: the traction sampler on the cut field (pressure fit excludes merged cells; sampler route in the sphere MMS and DFG drivers) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 18:47:46 -05:00
Omar SobhandClaude Fable 5.1	32546b1222	embedded3 loads: the control-volume route between no-slip z walls carries the walls' shear (control_volume_force_with_walls); the DFG driver uses it ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 17:27:19 -05:00
Omar SobhandClaude Fable 5.1	4fecb3098c	embedded3: the device download carries the previous step's field (the control-volume route's unsteady term read the initial profile) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 17:25:32 -05:00
Omar SobhandClaude Fable 5.1	338d66f85e	embedded3 item 12: the DFG 3D-2Z device driver (loads by both routes, Δp, settle criterion, CSV) and the VTK instant export ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 17:24:38 -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	0fa05f2056	embedded3 item 10b: virtual merging of small cells in the projection (fraction < 0.1 → master = largest active face neighbour; off-stencil links on the fine Poisson level; merged rhs, anchor, mass residual and source scale); static sphere rows within 0.02 %, loads 9.4/10.4 %; stadium falsifier spikes 39–54× below the binary wall (circle 6–8×), energy per event 6–9× lower ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 16:41:29 -05:00
Omar SobhandClaude Fable 5.1	5b1621e6ad	embedded3 item 11: moving bodies (end-of-step mask, fresh-cell refill, space-time cut cell: step-averaged apertures, GCL wall flux, Reynolds-transport momentum), the 3D fresh-cell falsifier (plate / circle / stadium, wall + control-volume routes) and the Lipschitz sweep; ghost wall reproduces the 2D falsifier to the digit; cut wall 5–14× smoother on the circle, gates not met (fresh cell's first step); wall.rs split (impose.rs) ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 16:20:35 -05:00
Omar SobhandClaude Fable 5.1	0e4c97ed24	embedded3 item 10: the apertured cut-cell wall (AM-wall) — cutwall.rs classification, apertured projection with the compatible wall flux, cut predictor (V_u = αhA, averaged mass fluxes, implicit wall shear, inertia floor), cut load route; sphere MMS CutCell ≤ GhostBinary at n 12/24 (ratio 0.96), loads 9.3/10.7 % at n 24 ... Co-Authored-By: Claude Fable 5.1 <[email protected]>	2026-09-17 15:34:07 -05:00
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
quantumandClaude Fable 5	c36cf2f8a7	rtx-backend-metal: fix swap_dims returning a corrupt strided view ... swap_dims copied the buffer but wrote elements back at their ORIGINAL positions (new_idx was computed with the swapped strides), returning a stride-swapped non-contiguous tensor. Every other op in this backend — elementwise kernels, MPS matmul, to_vec — reads raw buffers and ignores strides, so any transpose consumer (notably the autograd matmul backward, grad_a = grad_c @ b^T) silently computed on untransposed data. Found via CPU-vs-Metal gradient parity on the DigiGraph HetGAT: forward matched, gradients were ~2x off. swap_dims now physically permutes into a contiguous result (reading through the input's strides + offset), and reshape asserts contiguity instead of silently reinterpreting a non-contiguous buffer. 3 new parity tests incl. matmul-after-transpose (25 total pass on-device). Co-Authored-By: Claude Fable 5 <[email protected]>	2026-08-21 06:11:38 -07:00
quantumandClaude Fable 5	9297976929	rtx-backend-metal: GPU index_select / index_add via one-hot CSR SpMM ... Override the Backend trait's host-round-trip defaults: gather is S @ X with S the [E x N] one-hot selection CSR; scatter-add is the adjoint S^T @ X, whose CSR is built directly by counting sort so duplicate indices land in one row and the spmm kernel (one thread per output element) accumulates them without atomics. CSR matrices are cached per thread keyed by the exact index list + dims, so a static graph topology (GNN message passing) builds each matrix once. Host fallback on degenerate shapes or any sparse-pipeline failure. 13 new parity tests vs CPU reference: duplicates, unreferenced rows, D=1/2/3, 15k x 5k x 64 gather/scatter, cache reuse, adjoint roundtrip. Verified on-device that the SpMM path (not the fallback) serves all 13. Co-Authored-By: Claude Fable 5 <[email protected]>	2026-08-21 05:56:23 -07:00
quantumandClaude Fable 5	2e23d0f4c6	rtx-backend: gate CUDA parity tests to x86_64 Linux ... The rtx-backend-cuda dev-dependency is already gated to cfg(all(linux, x86_64)) in Cargo.toml, but the test file imported it unconditionally, so cargo test -p rtx-backend failed to compile on macOS. Co-Authored-By: Claude Fable 5 <[email protected]>	2026-08-21 05:52:03 -07:00
quantumandClaude Fable 5	72b41e3167	rtx-backend-metal: align to current rtx-metal API ... rtx-metal dropped tensor_ops::{sin,cos,pow,clamp,gt_scalar,var} and never had nn::{max_pool2d,avg_pool2d}; nn::conv2d grew a scalar 14-arg signature. Give the missing ops correct host fallbacks (the sum_dim pattern), and dispatch conv2d to the Metal kernel when its restricted signature applies (symmetric stride/padding, dilation 1, groups 1), host fallback otherwise. cargo test -p rtx-backend-metal: 22/22 parity tests pass on-device (M5). Co-Authored-By: Claude Fable 5 <[email protected]>	2026-08-21 05:52:03 -07:00
osobhandClaude Sonnet 5	5155c081ca	feat(mamba): GPU-accelerated backward pass (backward_cuda) ... MambaBlock::forward already had a working, tested CUDA dispatch (forward_cuda: cuBLAS matmuls for projections, CPU for the scan). backward() had none — it silently ran entirely CPU-serial on GPU tensors via to_vec()/from_vec() D2H/H2D round-trips. This adds the missing acceleration, mirroring forward_cuda's hybrid split: the four large projection-parameter gradients (in_proj, x_proj, dt_proj, out_proj) now go through batched GPU matmuls; the inherently sequential scan reverse-pass and small per-channel grads stay CPU. Extracted CpuWeights::pull and recompute_forward_cpu out of the old inline per-batch forward-recompute block inside backward() (pure refactor, gradient-checked unchanged by real_selective_scan.rs's existing 6 tests) so CPU backward and the new CUDA backward share identical forward math and can never numerically diverge on it. New CUDA-vs-CPU gradient-check test (mamba_cuda_backward_matches_cpu, #[ignore]-gated, GPU-only) caught a real bug during development: Tensor::contiguous() is a no-op stub in this rtx-tensor version, and cuda_matmul reads raw GPU storage by shape.dims() ignoring strides/offset, so .transpose(..).matmul(..) on a GPU tensor silently computed garbage (80-200x relative error on 3 of 4 accelerated gradients). Fixed by building already-transposed [dim, b*l] buffers on CPU before upload instead of transposing GPU-side. All 9 gradients now match CPU backward within ~2.2e-5 max relative error (tolerance 1e-4). Co-Authored-By: Claude Sonnet 5 <[email protected]>	2026-08-10 07:11:20 -07:00
osobhandClaude Sonnet 5	4aaa36a57a	style: cargo fmt --workspace (whitespace/wrapping only, no semantic change) ... Whole-workspace rustfmt pass picked up while iterating on Mamba GPU backward work. Verified formatting-only via diff sampling; no logic changed. Co-Authored-By: Claude Sonnet 5 <[email protected]>	2026-08-10 07:09:36 -07:00
osobhandClaude Fable 5	ad6405663f	fix(streaming): sane DynamicBatchingConfig default; worker lifecycle regression tests ... The new lifecycle tests caught that the batch-optimizer worker crashed at spawn: DynamicBatchingConfig derived Default (all zeros), and tokio's interval() panics on a zero period. Default is now a usable config (batch 32 in [1,128], step 4, 1ms latency target, 100-sample window, 1s optimization interval, AIMD adaptation). New regression tests in AdaptiveProcessor, EdgeComputingManager, and MonitoringSystem assert that all workers are still alive shortly after start() (catches workers dying at startup) and that stop() completes via the graceful control-channel path, not the 5s abort backstop (catches shutdown hangs). cargo test -p rtx-streaming: 58 lib + 8 integration + 6 aux, all green. Co-Authored-By: Claude Fable 5 <[email protected]>	2026-07-10 17:16:04 -07:00
osobhandClaude Fable 5	c83e0fb22d	fix(streaming): wire worker control planes for real graceful shutdown ... Follow-up to a0bf294, which tolerated dead control channels; this makes them functional: - AdaptiveProcessor: mpsc control channel (single consumer behind a mutex, broke on ANY message including Start) replaced with broadcast; all three workers (resource monitor, batch optimizer, pressure monitor) subscribe and exit only on ControlCommand::Stop - EdgeComputingManager / MonitoringSystem: their 7 interval-loop workers now subscribe to the existing broadcast control channels and exit on Stop instead of looping forever - stop() in all three: graceful join with 5s timeout, abort only as a backstop (previously unconditional abort mid-tick) - benches: criterion needs async_tokio for Bencher::to_async — bench target now compiles (clippy --all-targets clean) cargo test -p rtx-streaming: 55 lib + 8 integration + 6 aux, all green. Co-Authored-By: Claude Fable 5 <[email protected]>	2026-07-10 17:04:35 -07:00
osobhandClaude Fable 5	a0bf29461b	fix(streaming): real inference backend wiring and lifecycle fixes; full suite green ... - token_generator: backend is now an optional real rtx-inference engine (RwLock<Option<Arc<InferenceEngine>>>) with ServingTokenizer support; set_backend/set_tokenizer plumbing through StreamingServer - connection_manager: ConnectionPool::acquire no longer errors when the idle cache is full — creates fresh connections up to max_connections - streaming_server: ServerState::Running on construction; stream_inference generates one token per step (chunk_size semantics) - lifecycle bugs surfaced by the newly-compiling integration tests: * start(): broadcast control-channel send with zero subscribers was treated as fatal ("channel closed") in RealtimePipeline, EdgeComputingManager, MonitoringSystem — now tolerated * stop(): AdaptiveProcessor/EdgeComputingManager/MonitoringSystem awaited worker interval loops that never exit (test hung 5h) — workers are now aborted with cancellation-aware join - integration_tests: removed stale .await on now-synchronous methods cargo test -p rtx-streaming: 55 lib + 8 integration + 6 aux, all passing. Co-Authored-By: Claude Fable 5 <[email protected]>	2026-07-10 15:48:33 -07:00
osobhandClaude Fable 5	73102b71cf	feat(transformers): real compute in orchestrator modalities and attention planner ... - orchestrator_core: execute_classical is genuine seeded QKV self-attention; execute_flash_attention delegates to rtx_flash_attention::flash_attention_forward (8-head reshape, clear Err on indivisible hidden); execute_hybrid composes the two; execute_edge/execute_distributed return explicit "modality not implemented" errors instead of fake success — which makes the previously-dead fallback_modalities retry loop real. - tensor_core_kernels: AttentionComputationOptimizer::execute dispatches Standard/Online to scaled_dot_product_attention, Flash to rtx-flash-attention, Approximated to an explicit Err (no approximation kernel exists; refuses to compute exact attention under an approximated label). - 10 new always-on tests incl. flash-vs-standard 1e-3 agreement and an Edge->Distributed->Classical fallback end-to-end. - docs/consolidation.md updated: both entries moved from scaffolding/no-op-stub status to real dispatch descriptions. 982 rtx-transformers lib tests pass. Co-Authored-By: Claude Fable 5 <[email protected]>	2026-07-10 07:06:50 -07:00
osobhandClaude Fable 5	fce6cef262	docs: JEPA roadmap — GPU resume re-upload done; remaining items need multi-GPU ... Co-Authored-By: Claude Fable 5 <[email protected]>	2026-07-10 03:59:48 -07:00
osobhandClaude Fable 5	a755627269	feat(jepa): GPU weight re-upload on checkpoint resume ... GpuViTEncoder gains upload_weights (extracted from construction), cpu_weights_mut, and reupload_weights; JepaTrainerV2 exposes context_encoder_as_any_mut for backend-specific downcasts. The runner resume path now restores checkpoint fields into the GPU encoder's host copy and pushes them back to the device buffers — previously GPU resume restored only the step counter with a warning. If re-upload fails after host restore, the run aborts rather than training on stale device weights. Verified live on the RTX 5060 Ti: train 20 steps -> resume from the .jepa binary with total_steps=30 -> "Resumed from step 20", exactly 10 further steps, eval runs, no warnings. New tests: GPU output changes after host mutation + re-upload; CPU-target re-upload is a no-op Ok. 972 CPU tests / 37 GPU jepa_gpu tests pass. Co-Authored-By: Claude Fable 5 <[email protected]>	2026-07-10 03:59:36 -07:00
osobhandClaude Fable 5	74d3db7ee7	docs: refresh honesty notes for rewired demos, consolidation audit, ServingTokenizer ... Co-Authored-By: Claude Fable 5 <[email protected]>	2026-07-10 03:52:13 -07:00
osobhandClaude Fable 5	68481ef314	docs: JEPA roadmap round-2 done items; queue GPU-resume re-upload and TP/PP ... Co-Authored-By: Claude Fable 5 <[email protected]>	2026-07-10 03:51:50 -07:00
osobhandClaude Fable 5	ac3f2af06b	feat(jepa): eval in the training loop, NCCL GPU AllReduce, GPU checkpointing ... - Eval: run_jepa_training now runs k-NN (k=5) + linear-probe evaluation every eval_every steps and at the end (deduped when aligned); JepaEvalResult recorded in JepaTrainingSummary (final_knn_acc / final_probe_acc), printed by the CLI, appended as an # eval section to the metrics CSV. Probe set is deterministic LCG synthetic (offset seed, never aliases training batches) or real shard labels when loaded. - NCCL: real GPU-direct AllReduce backend behind the new `nccl` feature (cudarc/nccl, dlopen-based so builds don't need libnccl). NCCL unique id is bootstrapped over the existing TCP rendezvous (master_port+137); data path is htod -> ncclAllReduce(Sum) -> dtoh -> mean. catch_unwind guards cudarc's panic-on-missing-lib so training falls back instead of aborting. Verified for real on the RTX 5060 Ti: single-rank GPU all_reduce identity test passes (26/26 with --features nccl). - GPU checkpointing/eval: JepaTrainerV2::context_encoder_cpu_weights() exposes host-side weights for both CPU and GPU encoders (GpuViTEncoder::cpu_weights); checkpoint save and eval now work for GPU training runs (verified: .jepa binaries written and 2 eval passes during a live GPU CLI run). Resume with a GPU encoder restores the step counter and warns that weight re-upload is not yet implemented rather than silently training on stale weights. 125 runner/distributed/vit tests pass; CLI 8/8; cuda check clean. Co-Authored-By: Claude Fable 5 <[email protected]>	2026-07-10 03:51:36 -07:00
osobhandClaude Fable 5	f6b2308381	docs: JEPA roadmap — 2026-07-10 items done, next tier queued ... Co-Authored-By: Claude Fable 5 <[email protected]>	2026-07-10 00:04:52 -07:00
osobhandClaude Fable 5	19b6581f9c	feat(jepa): extended GPU training, data pipeline, integration, and cargo config ... - jepa_gpu: remove the bring-up 2-block cap; full-depth GPU-resident ViT verified against a full-depth CPU reference on the RTX 5060 Ti (depth-12 ViT-Tiny max_rel_err <= 6.6e-5). CpuViTEncoder's own hidden min(depth,2) cap removed too — CPU-path callers now get the model they configured. - jepa_distributed: real TCP parameter-server AllReduce backend (rendezvous handshake with world-size/rank validation, length- prefixed f32 payloads, connect/read/accept timeouts, connect retry until deadline so early peers survive rank 0 still computing); jepa_runner wires it for world_size > 1 and fails hard on collective errors. Two-rank loopback training run covered by test. - rtx-jepa-cli (new crate): rtx-jepa binary with train/bench/plan/ validate subcommands driving JepaRunConfig, run_jepa_training, run_jepa_benchmark, and ClusterTrainingPlan (plan --emit-config round-trips through a config serializer). GPU bench on this node: 86k patches/sec vs 1.3k CPU (~64x). - ViTSizeStr::Micro (d=32, depth=2) added as an explicit test/smoke size now that no hidden caps keep full-size configs cheap; heavy tests moved onto it (rtx-transformers suite: 367s -> 5s, and the runner subset had ballooned to 35min at full depth before this). 966 lib tests pass; 35/35 jepa_gpu with cuda; 8/8 CLI tests. Co-Authored-By: Claude Fable 5 <[email protected]>	2026-07-10 00:04:21 -07:00
osobhandClaude Fable 5	b6440905e7	feat(jepa): gzip-compressed WebDataset shard support ... read_webdataset_shard detects the gzip magic bytes (1F 8B, not extension) and decompresses via flate2 before tar parsing; WebDatasetShard::load no longer rejects .tar.gz/.tgz. Round-trip test writes a real gzipped tar and loads it back. Co-Authored-By: Claude Fable 5 <[email protected]>	2026-07-09 22:37:05 -07:00
osobhandClaude Fable 5	4ba1b78215	docs(consolidation): audit outcome — flagged MoE/flash-attn duplicates are not duplicates ... Call-site rewrite pass audited every flagged site; none needed consolidation: metal_moe is an API-consistent backend specialization, modular/router.rs is module-level (not expert-token) routing, glam.rs is disabled dead code with a pre-existing bug (noted for whoever re-enables it), and the three flash-attention "reimplementations" turn out to be planner scaffolding, no-op stubs, and a doc comment — no attention math exists to delegate. jepa_gpu's attention is documented as part of the fused GPU ViT block by design. Verified no regressions: rtx-transformers 961 lib tests pass, jepa_gpu 34/34 with cuda, rtx-training cuda check clean. Co-Authored-By: Claude Fable 5 <[email protected]>	2026-07-09 22:30:35 -07:00
osobhandClaude Fable 5	fdc1432072	feat(jepa): full GPU-resident ViT block on CUDA ... jepa_gpu.rs rewrite: all GEMMs (QKV, attention scores via a single cuBLAS call with folded 1/sqrt(dk) scale, weighted sum, projections, FFN) plus layernorm/GELU/bias/residual kernels now operate on device-resident CudaSlice buffers; new nvrtc kernels for numerically stable row softmax and head extract/scatter replace the CPU reorder loops. Two host transfers remain per encode(): patch-token upload and final output download. Fallback is now genuine-unavailability only (no device / feature off / kernel compile failure); per-op errors in live GPU mode are hard errors instead of silent per-op CPU downgrades. Parity vs CpuViTEncoder verified on RTX 5060 Ti / CUDA 13.1: max_rel_err <= 3.1e-5 across tiny/Tiny-192 configs (tolerance 1e-3). 34/34 jepa_gpu tests pass with --features cuda; 33/33 CPU-only. CLAUDE.md JEPA "Next" list updated to reflect completed GPU wiring, WebDataset reading, and cluster-plan consumption. Co-Authored-By: Claude Fable 5 <[email protected]>	2026-07-09 22:10:27 -07:00
osobhandClaude Fable 5	e080748d88	feat(demos,inference): wire simulation demos to real compute; fix embedding lookup and weight-name aliases ... Demos: - rtx-distllm-demo: real rtx-tensor weights per shard, real scaled-dot-product attention forward, metrics measured (Instant) instead of hardcoded constants; network topology remains a documented simulation fed by real tensor byte sizes. - rtx-model-zoo: MockInferenceEngine deleted; RealInferenceEngine loads a tiny real transformer into rtx_inference::InferenceEngine and runs genuine engine.infer per request; domain outputs are explicitly- labeled toy proxies derived from real output tokens. - rtx-inference-profiler: mock models deleted; profiles real matmul/softmax pipelines on rtx-tensor with measured latency/memory. Inference-path bugs the demos surfaced (fixed here): - ForwardPass::apply_embedding misused Tensor::gather for the embedding lookup — gather returns the indices' shape, silently dropping the hidden dim and breaking every downstream broadcast. Now uses the existing Tensor::embedding_lookup ([vocab,hidden] x [batch,seq] -> [batch,seq,hidden]). - Attention weight lookup accepts both self_attn. (HF-LLaMA) and attention. prefixes; final layer norm accepts norm.weight / model.norm.weight / ln_f.weight aliases. - Integration fixture gains the final norm weight; the previously always-failing engine tests now pass (8/8 model_loading_test). End-to-end inference through the real engine now works for the first time — verified via model_zoo_demo producing real forward-pass outputs across all categories. Co-Authored-By: Claude Fable 5 <[email protected]>	2026-07-09 22:06:29 -07:00
osobhandClaude Fable 5	733b02cd8b	feat(inference): concrete EAGLE draft model + real tokenizer at the serving boundary ... EAGLE (rtx-inference/src/eagle.rs, ~610 lines, mirrors medusa.rs conventions): EagleDraftHead autoregressive FFN with Concat/Add/ Attention feature fusion, EagleHeads draft model with draft/ draft_steps (per-step top-k for candidate trees) and teacher-forced training_loss; implements the speculative::EagleDraftModel trait so it plugs into the orchestration layer. 38 unit tests. Tokenizer (rtx-inference/src/tokenizer.rs): ServingTokenizer enum — Vocab (HuggingFace tokenizers, loadable from tokenizer.json) or ByteLevel fallback preserving previous behavior. rtx-serving-api's AppState and rtx-streaming's token generator now encode/decode through it (with_engine_and_tokenizer / set_tokenizer added; existing signatures unchanged). Also fixes two pre-existing compile errors in rtx-streaming (missing import, stray .await) that blocked its lib tests entirely. Tests: rtx-inference 328 pass, rtx-serving-api 193 pass, rtx-streaming 53 pass (2 pre-existing mock-server connection failures unrelated to these changes). Co-Authored-By: Claude Fable 5 <[email protected]>	2026-07-09 21:54:05 -07:00
osobhandClaude Fable 5	0cbfc1a739	fix(tests): repair rtx-onnx-codegen build and all pre-existing test failures in rtx-serving-api and rtx-runtime ... - rtx-onnx-codegen: re-export AttributeValue from ir (private-module import broke the whole crate; remaining errors were knock-ons). - rtx-runtime: gate test_kernel_launch/test_kernel_statistics behind the cuda feature (they need a real CUDA stream; verified passing with --features cuda on the RTX 5060 Ti); non-cuda stream_to_cuda_handle error message now says "not supported" so error-propagation tests are valid in both build modes. - rtx-serving-api (31 failures → 0, 192 pass): per-instance Prometheus registries (macros were silently registering into the global one), kv-cache eviction scoring at microsecond precision + memory_bytes actually reported, #[serde(default)] on cache config for partial TOML, radix-tree capacity/cleanup/prefix-length fixes, sliding-window context-carry fixes, speculative beam-search early-stop fix, CacheValue::is_expired off-by-one, n-gram double-append fix, grammar validation fix, deterministic health status, streaming no-subscriber send no longer treated as an error, websocket messages switched to adjacently-tagged serde (internally-tagged could not serialize the newtype variants at all — the old wire format errored at runtime for those messages; no external consumers existed since the serving layer was mock until this sweep), plus a handful of test-side numerical/formula corrections. Co-Authored-By: Claude Fable 5 <[email protected]>	2026-07-09 19:49:01 -07:00