fp-cuda: device-resident F₂ row reduction (RREF) on Hopper by JoeyBF · Pull Request #274 · SpectralSequences/sseq

JoeyBF · 2026-07-17T21:06:22Z

Stacked on #273 (the fp-cuda Hopper GEMM backend) — review/merge #273 first;
until it lands this PR's diff includes #273's commits. Rebases to blas3-only once
#273 merges.

Adds a fully device-resident F₂ (GF(2)) row reduction (reduced row echelon form)
on Hopper, dispatched from fp::Matrix::row_reduce for large p = 2 matrices,
and bit-exact vs the CPU M4RI oracle.

What's here

CPU foundation: a blocked GEMM-based row-reduction design + prototype and
benchmarks establishing the algorithm and the bit-exact oracle.
Device port (one commit per phase): DeviceMatrix + on-device GEMM packing
(Phase 3), panel_factor base kernel (Phase 4), device-resident forward
row-echelon pass (Phase 5a), back-substitution to canonical RREF (Phase 5b),
and dispatch of fp::Matrix::row_reduce to the device above a size threshold
(Phase 7). The whole reduction stays on the GPU — uploaded once, downloaded
once.
2^n scaling benchmarks + an effective binary-op throughput metric (Tbop/s).
Optimization campaign (each pass its own commit): cooperative multi-CTA
panel_factor / promote / block_reduce, wide adaptive forward panels,
active-row compaction, blocked-TRSM back-substitution (now default) with a
small-grid reduce, grid-cap tuning, and a low-barrier grid_sync. Flag-gated
net-loss/wash experiments (deep recursion, blocked-TRSM promote, recursive
panel) are kept as distinct commits documenting the exploration.
Cleanup: retires the experiments + FP_CUDA_PROF instrumentation, drops
the A/B kill-switches (commits to the proven defaults), keeps the tuning knobs,
and re-validates the dispatch threshold at 8192.

Testing (H200)

cargo test -p fp --features gpu --test cuda_dispatch — gpu_row_reduce_matches_cpu
and gpu_dispatch_matches_cpu pass (bit-exact vs CPU).
All fp-cuda demos (row_reduce_demo, forward_reduce_demo, panel_factor_demo,
gemm_xor_into_demo, matmul_b1_dev_demo) print all-OK.
Throughput: ~234 Tbop/s at 2^16, ~488 at 2^17 (half-rank square).

🤖 Generated with Claude Code

Prototype CUDA kernel and the Phase 2 wgmma.b1 core: both operands pre-arranged on the host as K-major tiles, one m64n128k256 binary MMA. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

…zation Steps (a)-(d): hoist the A TMA load out of the column-group loop, double-buffer A in the K pipeline, load B via TMA, and split into producer/consumer warpgroups. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

Both b1 operands are K-major; move them to CU_TENSOR_MAP_SWIZZLE_128B so wgmma operand reads avoid bank conflicts. The TMA applies the swizzle on load, so the host now emits plain row-major K-major tiles (the hand-rolled cm() interleave is gone) and the wgmma matrix descriptors carry the matching layout bits (layout=1, LBO=16B, SBO=1024B), derived from CUTLASS make_gmma_desc<Major::K> / LayoutType::B128. The SMEM K-tile grows to 1024 bits (one full 128B K-major swizzle atom = 4 k256 sub-chunks) and moves to dynamic shared memory (~82KB, opt-in via CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES). Per-stage wgmmas now run behind a single commit_group/wait_group and accumulate popcounts in-hardware (scale-D=1) into one resident accumulator per column group, replacing the previous one-wgmma-per-commit/wait serialization. Compile-verified only (nvcc + rustc); not yet validated on H100. Validate with a 64x256x64 identity product first, then matmul_b1_demo, then the bench. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

Replace the four m64n64k256 wgmmas per k-step (one per output limb) with a single m64n256k256 covering all NG=4 limbs at once. Binary wgmma is k256-only, so N is the throughput lever; n256 is the max. Same accumulator register count (128 s32/thread) and same SMEM as the 4x n64 version, but 1/4 the wgmma instructions, B descriptors, and fence/commit churn. B is now arranged as one contiguous 256-column tile per CTA (host transpose_b packs 4 limbs side by side; the B TMA box is 256 rows tall) instead of four separate 64-column tiles, and is zero-padded to whole 256-column groups. A and the K tiling are unchanged. The output bit-pack splits the single acc[128] into the 4 output limbs: the m64n256 fragment is the m64n64 layout tiled along N, so register group gi in 0..32 maps to columns [gi*8, gi*8+8). Compile-verified (PTX emits m64n256k256); fragment->limb bit-pack still needs the 64x256x64 identity check on H100. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

The producer warpgroup needs few registers; the consumer holds the 128-reg m64n256 accumulator. Issue setmaxnreg.dec(40) in the producer and setmaxnreg.inc(216) in the consumer so the consumer claims the producer's surplus. 128*(40+216)=32768 regs/CTA, leaving room for 2 CTAs/SM. Both are warpgroup-aligned and executed by all 128 threads of their warpgroup (the wg branch splits at warpgroup granularity). Compile-verified; PTX emits the dec/inc. Also fixes a stale m64n64k256 mention in the lib doc comment. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

Bump the producer-consumer pipeline depth from 2 to 3 so the producer can run two K-chunks ahead, hiding more TMA latency. SMEM grows to ~122 KB/CTA (one CTA/SM vs two at STAGES=2); the host smem_bytes/cuFuncSetAttribute track it automatically. STAGES is the latency-vs-occupancy knob to sweep on hardware. Phase-array inits made depth-agnostic ({0}). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

Replace the per-thread global stores with a single cp.async.bulk.tensor.2d.global.shared::cta per CTA. sC is now packed row-major ([row][limb]) so the 64×NG tile is contiguous for the store; thread 0 issues it after a __syncthreads + fence.proxy.async (making the atomicXor writes visible to the async proxy), then cp.async.bulk.commit_group/wait_group, then a final __syncthreads keeps sC alive until the store lands. The kernel now takes a C tensor map instead of a raw pointer (nlim/C* params dropped). C is padded on the host to whole NG-limb column groups (n_padded_lim = n_groups*NG) so every stored tile is complete; padded columns carry zeros from the zero-padded B and are trimmed on readback. README roadmap updated to reflect Phases 4-7. Compile-verified (PTX emits the bulk store + fence.proxy.async); store path and n256 bit-pack still need the H100 identity check. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

Replace the cuda-oxide `cuda-core` dependency with `cudarc`. cuda-core pulled in nightly via an incidental `#![feature(f16)]` (a DeviceCopy impl for the unstable f16 primitive we never use); cudarc is stable Rust, mainstream/maintained, and dynamically loads the driver at runtime — so the crate builds with no CUDA present and the whole crate (lib + examples + benches) now compiles AND links locally without libcuda (previously examples failed at link with -lcuda). Host glue rewritten against cudarc: CudaContext/load_module(Ptx)/load_function, clone_htod/alloc_zeros/clone_dtoh, device_ptr for the TMA descriptor addresses, CudaFunction::set_attribute for the >48 KB dynamic-SMEM opt-in, and the typed launch builder. The three CUtensorMaps are passed by value as grid-constant args via a #[repr(transparent)] TmaArg: DeviceRepr wrapper. cuTensorMapEncodeTiled is called through cudarc::driver::sys. Kernel (.cu/PTX) and build.rs are unchanged. cudarc features: driver + nvrtc (for Ptx) + cuda-12080 + dynamic-loading; the version only selects pre-generated bindings, the real driver (12.8+/13.x) is resolved at runtime. Also delete examples/bench_breakdown.rs: it was stale (called the pre-TMA 7-pointer kernel signature with the old cm() transpose) and the only remaining cuda-core user; bench_kernel_only supersedes it. Compile + link verified locally (lib, examples, benches). Stable-toolchain build to be confirmed on the server (dev box only has Nix nightly). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

The >16384 throughput cliff is L2-residency bound on B: with the old (grid_x=n_groups, grid_y=m_tiles) launch (N fastest), every B column-panel is re-touched only once per full sweep of B, so its reuse distance is the whole matrix and it spills from L2 once K*N/8 > ~50 MB, getting re-streamed from HBM per M-tile. Convert to a persistent 1-D grid of ~SM-count CTAs that sweep all output tiles in a grouped-along-M order (bi fastest within a GROUP_M-row band, bj slowest). This shortens each B-panel's reuse distance to <= GROUP_M, cutting B's HBM re-reads by a factor of GROUP_M while A still streams once. Kernel: take n_groups as an arg (no longer gridDim.x), add GROUP_M=8, wrap the per-tile body in a persistent `for (tile = blockIdx.x; ...)` loop, and re-init the mbarriers per tile (fresh phase-0 bookkeeping) so barrier parity isn't carried across iterations. setmaxnreg stays a one-time per-warpgroup action before the loop. Per-tile compute is byte-identical to Phase 7. Host: query CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, launch (num_ctas = min(SMs, total_tiles), 1, 1), and pass n_groups. TMA descriptors, padding, and readback are unchanged. GROUP_M is a tuning knob (8/16/32) to sweep on the H100. Code-only on the dev box; validate on the server. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

Phase 8 shortened each B-panel's reuse distance via rasterization; this shares the remaining HBM read of a B-panel across CTAs. CLUSTER (=2) CTAs along M form a thread-block cluster: each computes a different M-tile (its own n256 accumulator) but receives the same B-panel via one multicast HBM read (cp.async.bulk.tensor.2d...multicast::cluster), cutting B's HBM traffic by a further factor of CLUSTER on top of Phase 8's GROUP_M. (A second per-CTA accumulator can't fit the 256-reg budget, so cross-M B-reuse must be cross-CTA, i.e. clusters — not single-CTA SMEM reuse.) Kernel: __cluster_dims__(CLUSTER,1,1); the schedule walks M-super-rows of CLUSTER tiles (bi = sbi*CLUSTER + rank, shared bj). A is loaded per-CTA; B is multicast by rank 0 with an all-ranks mask into every member's sB and counted against every member's full barrier. The empty barrier is cluster-wide (init count CLUSTER; each consumer arrives on every member via mapa). Pipeline barriers are initialized once and flow continuously across tiles (single qidx/p) rather than re-init per tile, which would race the cross-CTA arrivals. One barrier.cluster after init. Mirrors pranjalssh/fast.cu matmul_9. Host: pad m_tiles to a multiple of CLUSTER (every rank gets a valid M-tile), round the persistent grid down to a whole number of clusters. CLUSTER must match the kernel constant (like STAGES). GROUP_M and CLUSTER are tuning knobs to sweep on the H100. Multicast/cluster ordering is the highest-risk, hardware-unverifiable part — see HANDOFF.md. Code-only on the dev box; validate on the server. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

Diagnosis on H200: the kernel was L2-refill-bandwidth bound, not compute bound. Microbenchmarks put the single-warpgroup wgmma.b1 ceiling at ~12,468 TOPS while sustained L2->SMEM refill tops out at ~8 TB/s (full tensor rate needs ~14.6). Each CTA computed a 64x256 tile, so it reloaded 8KB A + 32KB B per k-chunk -- B was 80% of the traffic because the tile was 4x wider than tall, and the tensor core starved waiting for it. Fix: register-block the output. Each CTA now computes a TM x NB block as MSTRIPS m64n128 wgmma strips (wgmma_n128, acc[MSTRIPS][64]) that all reuse one loaded B sub-tile, so a single L2->SMEM read of B feeds every strip. Winning config MSTRIPS=3 (192x128 block), NB=128, STAGES=4, GROUP_M=16, CLUSTER=2 cuts refill bytes/MAC by 33%. Also validates the previously-unvalidated Phase 8-9 batch on hardware (bit-exact) and lands the STAGES 3->4 / GROUP_M 8->16 tuning. Kernel-only, bit-exact (32768 cube): 8,585 -> 9,344 TOPS (+8.9%); 16384 7,636 -> 8,301; now ~75% of the tensor-core ceiling. Constants MSTRIPS in the kernel and TILE_M(=64*MSTRIPS)/NG in src/lib.rs must match. NB must be a multiple of 128 (NG even) or the C-store TMA box breaks 16B alignment; MSTRIPS=4 would exceed the 255-reg/thread cap. Adds examples/tune.rs, a fast no-CPU-check sweep harness. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

The per-tile epilogue drained its C store inline (cp.async.bulk.wait_group 0 between two __syncthreads), stalling the whole CTA for the store latency every output tile -- a bubble that grows as a fraction of tile time at smaller K. Double-buffer sC and defer the wait: each tile packs into sC[titer&1], issues its store, then does cp.async.bulk.wait_group.read 1, which blocks only until every store but the newest has finished *reading* its SMEM source. So tile T's store drains during tile T+1's compute, and the buffer freed (two tiles back) is safe to reuse. A final wait_group 0 after the persistent loop drains the last store before the CTA exits. Kernel-only, bit-exact: 32768 9,344 -> 9,426; 16384 8,301 -> 8,420; 8192 6,425 -> 6,591; 4096 3,835 -> 3,948 (smaller K gains more, as the epilogue is a larger share of tile time there). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

wgmma.fence.sync.aligned is a warpgroup-wide (128-thread) sync; the consumer was issuing two per k-chunk (one before the wgmma group, one after the wait), i.e. 64 warpgroup syncs per output tile. But the fence is only needed to order non-wgmma writes to the accumulators before a wgmma reads them -- in steady state the accumulators are touched only by wgmma, so one fence before the whole K-loop (ordering the accumulator zeroing) suffices, matching CUTLASS's warpgroup_arrive-once pattern. The per-chunk wgmma.wait_group 0 still makes results readable for the epilogue. Kernel-only, bit-exact: 32768 9,426 -> 9,605; 16384 8,420 -> 8,618; 8192 6,591 -> 6,744. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

Launch (occupancy × SM-count) CTAs instead of hard-coding SM-count, so the persistent grid exactly fills the machine for whatever the resource footprint allows. No change at the shipping config (occ=1/SM), but it's the correct, self-adapting launch and it drove the 2-CTA/SM experiment. That experiment is a documented dead end (H200, 2026-07-07): 2 CTAs/SM needs the compiled register count <=128/thread (2*256*128 = the 64K reg file), but the resident accumulator that gives the kernel its arithmetic intensity is 192 regs/thread at MSTRIPS=3. The only way to reach occ=2 is MSTRIPS=1 (64-reg acc), which collapses AI and drops 16384 from ~8,600 to ~5,500 TOPS. High AI (big accumulator) and high occupancy compete for the same register file and AI wins -- so 1 CTA/SM with the largest accumulator under the 255-reg cap is optimal. This confirms on-device that the ~10-13k TOPS bandwidth wall cannot be moved by raising occupancy. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

coderabbitai · 2026-07-17T21:06:30Z

Important

Review skipped

Draft detected.

Please check the settings in the CodeRabbit UI or the .coderabbit.yaml file in this repository. To trigger a single review, invoke the @coderabbitai review command.

⚙️ Run configuration

Configuration used: Organization UI

Review profile: ASSERTIVE

Plan: Pro

Run ID: 0c3ec4e5-0f82-497d-a3a0-72c9b246e97d

You can disable this status message by setting the reviews.review_status to false in the CodeRabbit configuration file.

Use the checkbox below for a quick retry:

🔍 Trigger review

✨ Finishing Touches

🧪 Generate unit tests (beta)

Create PR with unit tests

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

_{Comment @coderabbitai help to get the list of available commands.}

Wire the Hopper GPU backend into `fp` behind an optional `gpu` feature and keep the workspace CI green: `--workspace` builds `fp-cuda`, whose build.rs emits a stub PTX when nvcc is absent (every ubuntu-latest runner) instead of panicking. fp-cuda's library API is fp-agnostic (raw limbs), breaking the dependency cycle so `fp` can dispatch large p=2 products to the device. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

- build.rs: fall back to the stub PTX only when nvcc is genuinely absent (io::ErrorKind::NotFound); fail fast on other spawn errors (permissions, broken exec) instead of silently building the stub. - flake.nix: move the CUDA toolkit out of the default dev shell into a dedicated `devShells.gpu` (`nix develop .#gpu`), matching the sibling GPU PR, so contributors/CI don't fetch the multi-GB unfree CUDA closure for the opt-in backend. Drop the obsolete cuda-oxide/libclang/bindgen env (the crate uses cudarc, which dlopens libcuda). - README/Cargo.toml/cuda.rs: call the Cargo feature `gpu` consistently; note the kernel is sm_90a-only (architecture-specific, not forward-compatible — no Blackwell/sm_100). - cuda_dispatch test: add non-tile-aligned shapes (2049x2051x2053, 3000x2112x4097) to exercise edge masks, partial limbs, and raster tails. - bench examples: drop obsolete pre-optimization diagnosis prose; label the bench_shapes L2 size as an assumption rather than a device fact. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

Design for a BLAS3 (GEMM-based) F₂ row reduction and the CPU blocked prototype (Phase 1), with proptest regression seeds. Establishes the algorithm and the CPU oracle the device port stays bit-exact against. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

CPU-side benchmarking that shows the blocked prototype is not GEMM-bound and the M4RI/blas3 ratio collapses with n — motivating the GPU port. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

Replace the O(R^2 n) row-op back-substitution with a blocked version in the same X*U GEMM shape as the forward trailing update, walking pivot blocks right-to-left: reduce each block of pivots to RREF among itself with a few full-width row ops, then clear its pivot columns from all rows above via one GEMM. Same asymptotic work, but now as data-parallel GEMMs (and cache-friendly), so the back-substitution is BLAS3 and scales with cores instead of being a serial drag that grew with n. Proptested against row_reduce (RREF + pivots) across the existing shape grid, rank-deficient generators, and the deterministic sweep. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01ESAMSPB8MQL4SSFxRJSEyR

Rewrite the panel pivot search and forward elimination -- and the back-substitution's X column-gather -- to read/write raw limbs directly instead of going through entry()/add_basis_element()/per-row split_at_mut. The pivot bit test becomes (data[row*stride + q/64] >> (q%64)) & 1, and the elimination snapshots the pivot row's panel limbs once and XORs them into below rows in a tight limb loop. Same algorithm, same results (proptest green), but the O(m*n) panel scan sheds the FpSlice/FieldElement machinery that dominated it. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01ESAMSPB8MQL4SSFxRJSEyR

DeviceMatrix (upload/download, in-place device residency) and the trailing update gemm_xor_into, built on the Hopper wgmma.b1 kernel. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

The one genuinely-new kernel (BLAS3-GPU-HANDOFF §4, design §5): forward factorization of one 64-bit column panel, the only column-indexed region of the reduction. A single CTA sweeps the 64 bit positions with a __syncthreads between them; per bit a find-first reduction picks the pivot (the lone column op) and a row-parallel masked XOR clears it from the rows below, recording the multiplier bit into L. Forward-only, matching CPU Step A. Rows are addressed through a virtual perm (design §4.3): a pivot is promoted by a perm swap, so the matrix bytes never move; L is indexed by original row id, needing no swap. Only (pr, pivcols) come back to host. - lib.rs: GpuContext loads the kernel; identity_perm, panel_factor (returns (pr, pivcols)), download_u32. - examples/panel_factor_demo: validates against a CPU transliteration of the kernel bit-for-bit — reduced panel limb, perm order, L, pr, pivcols — across full-rank and rank-deficient panels (free-column branch), partial panels (n<64), higher limbs, mid-sweep r>0, and m=100000. All match on the H200. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

forward_reduce runs the whole blocked forward elimination over one persistent device buffer (design §4): sweep 64-bit panels, and per panel factor it (panel_factor), promote the pivot rows' deferred trailing, drop the pivots from L, and apply the trailing update M[:, c+b:] ^= L·U as one wgmma GEMM. No host round-trip inside the sweep — only (pr, pivcols) come back per panel. Three small driver kernels added to matmul_b1.cu: promote_pivots (triangular trailing solve among a panel's pivot rows, one CTA), zero_pivot_l (exclude pivots from the GEMM), gather_rows (build the contiguous U operand from the scattered pivot rows via perm). examples/forward_reduce_demo validates against the ground-truth reducer: the forward pass is elementary row ops so it preserves row space, hence row_reduce(device_M) == row_reduce(original); plus rank, pivot columns, and that the non-pivot rows (perm[r..]) are zeroed. All cases pass on the H200 — full-rank, rank-deficient, multi-panel, and wide. Back-substitution to full RREF is Phase 5b. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

Completes the device-resident reduction (BLAS3-GPU-HANDOFF §5 decisive gate). back_substitute turns the echelon form into full RREF, blocked right-to-left over ≤64-pivot blocks (design §4.6): reduce each block among itself, then clear its pivot columns from every row above via one X·U wgmma GEMM, scattered back through perm. row_reduce_dev = forward_reduce + back_substitute — the whole reduction over one persistent buffer, no host round-trip beyond the per-panel pivot read-backs. Three back-sub kernels added: block_reduce_rref (block-internal reduction, one CTA), gather_cols (build X = above-rows at the block's pivot columns), xor_into_perm (scatter the GEMM result through perm). examples/row_reduce_demo validates row_reduce_dev == fp::Matrix::row_reduce bit-for-bit — the RREF matrix, the pivot list, and rank — and agreement with row_reduce_blas3, across full-rank/rank-deficient/wide/tall shapes on the H200. Correctness Definition-of-Done (#1) met. Remaining: dispatch threshold in fp (Phase 7), active-row compaction / perf (Phase 6). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

row_reduce now tries the device-resident reduction for p=2 matrices above the FP_CUDA_THRESHOLD (min(rows,cols) ≥ 2048, same knob as the mul dispatch), and falls back to the CPU M4RI path when the GPU is unavailable or below threshold. blas::cuda::try_row_reduce uploads once, runs row_reduce_dev, and materializes the canonical RREF (pivot rows at top in column order, zeros below, pivots set) — bit-identical to the CPU path. tests/cuda_dispatch: gpu_row_reduce_matches_cpu checks the dispatched row_reduce == row_reduce_blas3 (RREF, rank, pivots) on full-rank and rank-deficient matrices above threshold; FP_CUDA_DEBUG confirms the GPU path fires (218 kernel launches). Default (non-cuda) build unaffected. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

Device-vs-CPU scaling harness over 2^n half-rank matrices out to 2^18, reporting effective binary-op throughput (Tbop/s = 2*m*n*R/t). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

Profiling (Pass 0) showed the single-CTA panel_factor kernel — one SM of 132 sweeping 64 sequential bit-steps — dominated the forward pass at ~76% of GPU time, NOT the trailing GEMM as the plan assumed (GEMM was ~11%). This pivots Phase 6 to attack the single-CTA kernels first. panel_factor_coop does the identical math grid-parallel: each bit-step's find-first and masked-XOR spread across the whole grid, with a self-contained atomic sense-counting grid barrier between the 64 steps (no cooperative_groups / cudadevrt dependency, so it compiles under nvcc -ptx). Launched cooperatively so all CTAs are co-resident (required by the spin barrier). The pivot row/word is broadcast via global scratch and the perm swap is barrier-separated from both the pre-swap read and the XOR, so there are no races (3 barriers per pivot-bit). Default on; FP_CUDA_NO_COOP=1 falls back to the single-CTA kernel. Also adds the Pass 0 profiling harness: examples/probe_pr.rs (flat-K microbench, which confirmed Loss 1 — kernel time flat pr=64..1024 while effective Tbop/s climbs 16×) and FP_CUDA_PROF per-phase GPU timers in forward_reduce / back_substitute (printed by row_reduce_dev). Results (half-rank square, bit-exact vs CPU row_reduce): n=32768: device 2.688s -> 0.787s (3.4x), vs-M4RI 3.9x -> 13.3x n=65536: device 11.89s -> 4.04s (2.9x) panel_factor share 76% -> 21%; the trailing GEMM (K-padding, Loss 1) is now the top cost at ~33%, promote_pivots + block_reduce_rref (still single-CTA) ~33%. Gates: forward_reduce_demo, row_reduce_demo, cuda_dispatch all bit-exact; clippy clean. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

Two single-CTA back-of-the-pack kernels the Pass-5 profile exposed (promote ~18%, block_reduce ~15% of GPU time). promote_pivots: the deferred-trailing promotion is embarrassingly parallel over columns — each trailing limb runs its own sequential k-loop and never touches another column — so it needs no __syncthreads at all. Rewrote it grid-strided over limbs and launch a full grid (was one CTA). 132ms -> 56ms at n=2^15. block_reduce_rref: kept single-CTA (it is entangled with the sequential right-to-left back-substitution and cannot be batched across blocks), but cut its barrier count from ~bp² (a __syncthreads per (k,j) pair) to ~2·bp by gathering all of pivot k's clear-conditions into shared memory once, then doing a limb-parallel XOR that loads rowk[c] a single time and reuses it across the ≤64 rows. (A first flattened-(j,limb) attempt regressed to 233ms via emulated 64-bit division; the limb-outer form avoids it.) 110ms -> 95ms. End-to-end (half-rank, bit-exact vs CPU): n=32768 0.787s -> 0.698s (3.85x vs the 2.688s pre-Pass-5 baseline); n=65536 4.04s -> 3.64s. The trailing GEMM (Loss 1, K-padding) is now the clear top cost at ~38% + ~9% back-sub. Gates: row_reduce_demo, forward_reduce_demo, cuda_dispatch bit-exact; clippy clean. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

…idth Attacks Loss 1 (K-padding): the trailing GEMM pads the contraction dimension to TILE_K=1024, so a 64-pivot panel wastes 15/16 of the tensor-core issue. Widening the panel to b=64·bl columns raises pr toward b and reclaims that waste. panel_factor_coop now factors a bl-limb-wide panel natively: the masked XOR clears each pivot from the below rows across ALL bl panel limbs inline (the intra-panel Schur update — CPU blas3.rs Step A, lines 159-161), so a later sub-column's find-first sees fully reduced bits and ONE wide trailing GEMM (K=pr≤b) fixes up the columns beyond the panel. Multipliers go into a bl-limb L indexed by the global pivot index; the pivot's bl panel limbs are broadcast via g_pivword. bl=1 reproduces the single-limb kernel exactly. matmul_b1_dev gains a _strided variant (pack_a takes a source row stride distinct from the K-limb count) so the wide L — stored with stride bl but only ceil(pr/64) limbs occupied — feeds the GEMM without a trim copy. The width is a genuine tradeoff: wider panels halve the GEMM but the O(pr²)/panel promotion replay is O(bl) total and the inline XOR is O(bl), so past an optimum they lose. Measured optimum grows with the trailing width (bl ≈ stride/256): bl=2 @ n=2¹⁵, 4 @ 2¹⁶, 8 @ 2¹⁷ (bl=16 regressed everywhere). Default is adaptive adaptive_bl(stride)=clamp(stride/256,1,16); override with FP_CUDA_BL. Results (half-rank, bit-exact vs CPU), vs the original pre-Phase-6 baseline: n=32768: 2.688s -> 0.64-0.84s (~3-4x), vs-M4RI 12.5x n=65536: 11.89s -> 2.70s (4.4x), vs-BLAS3 24.3x n=131072: 74.90s -> 12.81s (5.8x), 176 Tbop/s (was ~50) Speedup grows with size. Balanced profile now (panel 27 / gemm 21 / promote 20 / block_reduce 17 / bs.gemm 11 %). Gates: forward_reduce_demo, row_reduce_demo bit-exact at bl=1,2,4; cuda_dispatch pass; clippy clean. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

…buffers The device GEMM's a_int / bt / c_dev and the gather outputs (u_buf, x_buf) are fully written before they are read — pack_a/pack_b write every element (padding as explicit zeros), the GEMM stores the whole C tile grid (overwrite, not accumulate), and the gather kernels fill every output limb. Their per-call alloc_zeros therefore paid for a memset that the following kernel immediately clobbers — and c_dev is m_padded × n_padded_lim, multi-GB at large n, so that memset is real HBM bandwidth wasted every trailing update. Switched those five to unsafe `alloc` (uninitialized; uses malloc_async where available). The multiplier L and the cooperative barrier scratch stay zeroed — they are OR-accumulated / must start at 0. Bit-exact (matmul_b1_dev_demo, gemm_xor_into_demo, forward_reduce_demo, row_reduce_demo). Modest end-to-end (n=32768 half-rank ~55.9 Tbop/s); the win grows with n as the C memset dominates. Full persistent-buffer reuse (avoiding the cudaMalloc itself) is left as a follow-up. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

The GPU row-reduce path shared the matmul dispatch threshold (min-dim 2048), but a full reduction is many dependent panel steps, not one GEMM, so its CPU crossover is later. Measured on an H200 (random square, device incl. upload/download vs multi-threaded M4RI row_reduce): the GPU first wins at n≈8192 (n=4096 is still ~1.7× slower), so dispatching at 2048 spent 2048–4096 losing to the CPU. Added DEFAULT_RR_THRESHOLD = 8192 (env FP_CUDA_RR_THRESHOLD), used only by try_row_reduce; the matmul path keeps its own 2048 (FP_CUDA_THRESHOLD). The cuda_dispatch row-reduce test forces the threshold down to 2048 so its small fast shapes still exercise the GPU path. cuda_dispatch passes; non-cuda `cargo build -p fp` and `cargo clippy -p fp --features cuda` clean. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

A below row that is entirely zero across the remaining columns can never pivot and carries no multiplier, so it is permanently dead. mark_live flags such rows in [r, m_active); compact_perm stable-partitions perm so live rows precede dead ones and returns the shrunk m_active, which panel_factor_coop now takes as its row bound (find-first / masked-XOR skip the dead tail). Pivot rows [0,r) are never touched, so back-substitution and RREF materialization are unaffected. Re-scanned every COMPACT_PERIOD=4 panels to amortize the mark+host-partition cost (a few hundred KB of perm/flags per call). Default on for the coop path; disable with FP_CUDA_NO_COMPACT. The bigger win rides on the early-exit it enables: once compaction finds no live below rows (m_active <= r) no later column can hold a pivot, so the forward sweep stops immediately instead of grinding through empty panels (each of which still paid 64 grid-barrier bit-steps). Effect: dense half-rank ~1% (dead rows appear late there); low-rank / rectangular inputs, which the general row_reduce must handle, benefit sharply — a rank-1024 65536×16384 reduce drops 0.198s -> 0.141s (29%). Bit-exact throughout (row_reduce_demo, forward_reduce_demo, cuda_dispatch); clippy clean. Adds examples/probe_lowrank.rs. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

Re-profiling at the larger adaptive bl (n=2^16, bl=4) showed the single-CTA block_reduce_rref had become ~22% of GPU time — it reduces each ≤64-pivot block to RREF on one SM of 132, and its full-width row XORs grow as r·stride. block_reduce_coop does the identical math grid-parallel: gather pivot k's clear-conditions into a global buffer, atomic grid barrier, XOR rowk into the flagged block rows flattened over (row × limb) across the whole grid, barrier, next k. Reuses the validated grid_sync (cooperative launch, grid-uniform k-loop). The per-block cooperative launch + grid barriers only pay once block work (≈ bp·stride) is large, so it is gated on stride >= 1024 (n >= 65536); below that the single-CTA kernel wins. Measured (H200, half-rank): neutral at n=2^15, +6% at 2^16 (2.67s -> 2.51s), +18% at 2^17 (13.1s -> 10.7s, 210 Tbop/s). FP_CUDA_NO_COOP forces single-CTA. Bit-exact (row_reduce_demo, forward_reduce_demo, cuda_dispatch); clippy clean. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

Splitting the GEMM phase showed the kernel — not packing — dominates it, and the back-sub X·U GEMM ran at K=bp_eff=64, padded to TILE_K=1024: 16× of its tensor-core issue was multiply-by-zero (bs.gemm ~16% of GPU time). block_reduce_coop's cost is ~bp-independent (its compute and its 2·r grid barriers both scale with the pivot count r, not the block width), so on the cooperative path the back-sub block can be widened for free. Raised bp from 64 to 512 (K padded 1024 → only ~2× waste, and ~8× fewer GEMM launches). The single-CTA fallback keeps bp=64 (its shared cond[] is sized 64); override with FP_CUDA_BP. Measured (half-rank, H200): n=131072 10.78s -> 7.52s (30%, 300 Tbop/s); bp=1024 is marginally better (7.40s) with diminishing returns. Bit-exact vs CPU BLAS3 at n=65536 (32.7× faster, 2.07s vs 67.6s) and all demos; clippy + cuda_dispatch clean. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

After the back-sub GEMM was de-padded, the single-CTA promote_pivots became the dominant forward cost (~29% at n=2^16): it parallelizes only over trailing limbs (a few CTAs) while doing an O(pr²) serial replay per column. promote_coop reformulates it right-looking: process pivots i = 0..pr in order, and once pivot i is final add it to every later pivot k>i carrying its multiplier (M[k] ^= M[i] across the trailing) — so pivot i is fully promoted by the time we reach it. Sequential in i with a grid barrier between steps (reusing the validated grid_sync), but each step's XOR is flattened over (later-pivot × limb) across the whole grid instead of a handful of CTAs. Gated on stride >= 1024 (the same amortization threshold as block_reduce_coop); FP_CUDA_NO_COOP forces the single-CTA kernel. Measured (half-rank, H200): n=65536 2.07s -> 1.95s; n=131072 7.40s -> 6.19s (364 Tbop/s). Bit-exact vs CPU BLAS3 at n=65536 (32×) and all demos; clippy + cuda_dispatch clean. Cumulative vs the original pre-Phase-6 baseline: n=131072 74.9s -> 6.19s (12×). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

promote_coop is ~bl-independent (its compute and 2·pr·(stride/bl) grid barriers both scale with the pivot count, not bl), so once it is active (stride ≥ 1024) the forward panel can be widened until the trailing GEMM stops K-padding (bl=16 ⇒ K=1024 exactly) at no promote cost. Below that threshold the single-CTA promote is O(bl) and narrow panels still win. adaptive_bl now uses stride/128 on the coop path and stride/256 below it — matching the measured optima (half-rank): bl=2 @ n=2^15, 8 @ 2^16, 16 @ 2^17. End-to-end (half-rank): n=32768 0.63s -> 0.59s, n=65536 1.95s -> 1.78s, n=131072 6.19s -> 5.92s (380 Tbop/s). Bit-exact (row_reduce_demo); clippy clean. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

The cooperative kernels are barrier-bound. promote_coop's clear-condition reads L, which its XOR never mutates, so the condition can be tested inline in the XOR instead of gathered under a separate grid barrier — dropping barrier [A] and its cond buffer, one barrier per pivot instead of two. Measured (half-rank, H200): n=65536 1.79s -> 1.64s (41× vs CPU BLAS3, bit-exact); n=131072 5.93s -> 5.57s (404 Tbop/s). Confirms the barrier count, not compute, gates these kernels. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

The cooperative kernels are barrier-bound and the grid barrier's spin was the cost: atomicAdd(barrier, 0) is a read-modify-write that acquires the counter's cache line exclusively every iteration, serializing the ~1000 spinning CTAs on a single line (~µs per barrier). Spinning on a plain volatile load instead leaves the line shared; arrival stays a single atomicAdd per CTA. Measured (half-rank, H200, bit-exact): n=65536 1.64s -> 1.52s; n=131072 5.57s -> 5.31s (424 Tbop/s, 4.9% of the wgmma.b1 roofline). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

block_reduce_coop is bp-independent in wall time (compute and its 2·r barriers both scale with the pivot count, not the block width), so the back-sub block can be K=1024 — an exact TILE_K multiple, eliminating the last of the back-sub GEMM's K-padding. Measured (half-rank): n=65536 44.5× vs CPU BLAS3 (1.52s, bit-exact); n=131072 5.32s -> 5.20s (433 Tbop/s, 5.0% of the wgmma.b1 roofline). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

…factor Validates the "row-reduce a 1024-wide strip" idea: factor each bl-wide macro panel as bl_micro-wide micro sub-panels, apply each micro's pivots to the rest of the macro via an intra-macro GEMM (so the next micro's find-first sees reduced bits), then one wide far GEMM with K = all macro pivots. This shrinks the elementwise panel_factor (the intra-panel XOR, the top cost at wide bl) while keeping the far GEMM at K=1024. Bit-exact vs CPU (row_reduce_demo; vs BLAS3 at n=2^16). Off by default (FP_CUDA_MICRO); ~6% at n=2^17 (micro=4, 461 vs 434 Tbop/s). Mechanics: a fresh per-micro L (offset 0) feeds each intra-macro GEMM, and a new l_shift_or kernel accumulates the micros into a contiguous macro L (offset 0) for the far GEMM — so promotes/GEMMs all use offset-0 operands; only l_shift_or handles the bit offset. Extracted the promote+zero+gather+GEMM+xor sequence into a shared trailing_update() helper used by both the single-wide-panel (default) and recursive paths. Two known limits, documented at the knob: (1) the intra-macro GEMMs still pad K to TILE_K=1024, capping the win — the full payoff needs a small-TILE_K GEMM variant; (2) per-sub-panel scratch is not yet reused, so at small micro + large n the async-malloc high-water mark can intermittently exhaust memory (a per-macro sync mitigates but doesn't eliminate it). Hence experimental / opt-in; the default single-wide-panel path is unchanged and bit-exact. clippy clean; cuda_dispatch + all demos pass. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

…is a net loss Add factor_panel_rec: a wide-panel (bl=64) halving recursion so the forward trailing GEMMs run at K≈4096 instead of K=1024. Gated behind FP_CUDA_DEEP (FP_CUDA_BASE sets the elementwise base width, default 16 limbs); the non-deep default path is byte-identical. This is the definitive test of the "raise K → compute-bound" thesis. Result (2^16, FP_CUDA_PROF): the far update drops 225ms→47ms (4.8× — the K-scaling is real), but it is a NET LOSS. The trailing GEMM is only ~20% of the row reduction; the other ~80% is elementwise (panel_factor 31%, back-sub block_reduce 29%, promote 15%). Each intra-panel GEMM needs a promote wrapper, so promote explodes 175ms→426ms and swamps the GEMM saving. The row reduction is elementwise/latency- bound, not compute-bound; panel-width sweeps confirm a flat optimum at bl≈12-16 (= adaptive_bl), so the current default is already near-best. Bit-exact: forward_reduce_demo, row_reduce_demo, cuda_dispatch all pass with FP_CUDA_DEEP=1. Kept behind the flag as the record: if promote is ever made cheap, deep recursion flips to a win. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

Refactor back_substitute's clear-from-above (gather X/U, GEMM, scatter-XOR) into a shared bs_clear_above helper, reused by both the outer loop (above=[0,s)) and a new recursive within-block TRSM (FP_CUDA_BS_TRSM). The recursion halves each pivot block, clears the right half's pivots from the left via a large X·U GEMM, and recurses — the BLAS3 triangular-solve shape, with NO promote (source/target rows disjoint and already reduced). gather_cols/xor_into_perm read perm through a .slice(above_start..) view, so no kernel changes. MEASURED (2^16, FP_CUDA_PROF): NEUTRAL. block_reduce is grid-barrier-bound, not work-bound — block_reduce_coop does 2 grid syncs/pivot (~2r total) and the recursion with a coop base keeps them all (block_reduce 334→296ms while bs.gemm 29→65ms — a wash). A barrier-free single-CTA base removes the syncs but is worse (485ms, one SM can't cover full width). The real fix is a shared-memory 64x64 triangular base kernel (reduce base block in-CTA, no grid sync; apply via GEMM), bounded ~1.2x since block_reduce is 29%. Kept behind the flag as correct BLAS3 scaffolding for that kernel; default path unchanged and bit-exact. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

…reduce) The back-sub within-block reduce (bs.block_reduce, ~29% at 2^16) is grid-barrier- bound: block_reduce_coop does 2 grid syncs/pivot and a grid_sync's cost scales with CTA count. Fix, in two parts that only work together: (1) recurse each block to narrow (<=64) base blocks + large X.U GEMMs (the BLAS3 TRSM shape, no promote — source/target rows disjoint & already reduced); (2) run those narrow base reduces on a SMALL grid (br_ctas cap 128) so the per-pivot barrier is cheap while the tiny XOR width is still covered. Measured: bs.block_reduce 334->115 ms (2.9x), +12% at 2^16, +6% at 2^17, bit-exact. Neither part alone helps — recursion+full-grid keeps every barrier (neutral); small-grid on the full 1024-block starves its XOR (worse, 485 ms). Now the default on the coop path; FP_CUDA_NO_BS_TRSM restores the elementwise full-block reduce, FP_CUDA_BS_BASE / FP_CUDA_BR_CTAS tune it. bs_clear_above is shared with the outer loop and reads perm via .slice(off..). Gates: row_reduce_demo, forward_reduce_demo, cuda_dispatch all bit-exact on both default and fallback. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

panel_factor (grid_sync/pivot-bit) and promote (the forward-substitution TRSM, grid_sync/pivot) are both partly grid-barrier-bound, and a grid_sync's cost scales with CTA count. Cap their cooperative grids (panel_factor 128, promote 384 on H200) so barriers are cheap while the per-step work is still covered — same lever as the back-sub small-grid reduce. Measured: panel_factor +3%, promote +6% at 2^16; together +8% at 2^16, +5% at 2^17, on top of the TRSM win. Env overrides FP_CUDA_PF_CTAS / FP_CUDA_PROM_CTAS; both clamp to the natural grid so small matrices are unaffected. Bit-exact (row_reduce_demo, forward_reduce_demo, cuda_dispatch). Cumulative this session: 2^16 186->227 Tbop/s, 2^17 432->487. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

…xperiment promote is a forward-substitution triangular solve; built it as a blocked TRSM (promote_rec/promote_clear): recurse the pr pivots into 64-aligned halves, apply the left to the right pivot rows via one X.U GEMM where X is gathered straight from L (no M-column gather; needs the new l_limb_off arg on promote_coop so a sub-block reads the right L bits), U is the promoted left rows. Correct (a pivot row only receives contributions from earlier pivots — no double counting), bit-exact on both paths incl. cuda_dispatch. MEASURED (2^16): a WASH. promote drops 176->107 ms in isolation but end-to-end is flat-to-worse (best +1.4% at base=512). Unlike back-sub's clear-above GEMM (spans all above rows, large m), promote only touches the pr<=1024 PIVOT rows, so its GEMMs are small-m AND small-K — too small for the tensor cores, and the many small kernels/panel pipeline worse than one elementwise sweep. The promote grid-cap (+6%, already default) was the real win; promote is intrinsically a small op. Kept behind FP_CUDA_PROM_TRSM as the definitive experiment; default unchanged. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

Post-optimization cleanup now that the row-reduction tuning is settled. Removed (flag-gated experiments, all net-loss/wash, kept in earlier commits): FP_CUDA_PROF harness, FP_CUDA_DEEP/BASE deep recursion (factor_panel_rec), FP_CUDA_MICRO, FP_CUDA_PROM_TRSM/PROM_BASE (promote_rec, promote_clear), l_shift_or + kernel, FP_CUDA_DEBUG. Dropped A/B kill-switches (committed to defaults): NO_COOP, NO_COMPACT, NO_BS_TRSM. Kept tuning knobs: BL, PF_CTAS, PROM_CTAS, BR_CTAS, BP, BS_BASE, GEMM_CTAS. RR dispatch threshold re-validated at 8192. Net -520 lines in lib.rs. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01B1XztNMycTtSC9SQtpNYjR

JoeyBF and others added 14 commits July 17, 2026 16:44

JoeyBF force-pushed the claude/blas3-row-reduction-lgjz3y branch from 52f4503 to 7ee600c Compare July 17, 2026 21:44

JoeyBF force-pushed the claude/blas3-row-reduction-lgjz3y branch from 7ee600c to a58f096 Compare July 17, 2026 22:01

JoeyBF and others added 12 commits July 17, 2026 23:28

JoeyBF and others added 19 commits July 17, 2026 23:29

JoeyBF force-pushed the claude/blas3-row-reduction-lgjz3y branch from a58f096 to 8b721f1 Compare July 18, 2026 03:30

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

fp-cuda: device-resident F₂ row reduction (RREF) on Hopper#274

fp-cuda: device-resident F₂ row reduction (RREF) on Hopper#274
JoeyBF wants to merge 46 commits into
SpectralSequences:masterfrom
JoeyBF:claude/blas3-row-reduction-lgjz3y

JoeyBF commented Jul 17, 2026

Uh oh!

coderabbitai Bot commented Jul 17, 2026 •

edited

Loading

Review skipped

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Uh oh!

Conversation

JoeyBF commented Jul 17, 2026

What's here

Testing (H200)

Uh oh!

coderabbitai Bot commented Jul 17, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Review skipped

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

coderabbitai Bot commented Jul 17, 2026 •

edited

Loading