Skip to content

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

Draft
JoeyBF wants to merge 47 commits into
SpectralSequences:masterfrom
JoeyBF:claude/blas3-row-reduction-lgjz3y
Draft

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

Conversation

@JoeyBF

@JoeyBF JoeyBF commented Jul 17, 2026

Copy link
Copy Markdown
Collaborator

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_dispatchgpu_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

@coderabbitai

coderabbitai Bot commented Jul 17, 2026

Copy link
Copy Markdown

Important

Draft PR not reviewed

Draft PRs are not automatically reviewed by default.

  • Trigger a manual review

To automatically review draft PRs, update your CodeRabbit configuration:

reviews:
  auto_review:
    drafts: true
✨ 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.

@JoeyBF
JoeyBF force-pushed the claude/blas3-row-reduction-lgjz3y branch 3 times, most recently from a58f096 to 8b721f1 Compare July 18, 2026 03:30
JoeyBF added a commit to JoeyBF/sseq that referenced this pull request Jul 23, 2026
Replaces the count-based launch cap (NASSAU_GPU_CONCURRENCY=8 exclusive
sections) with two decoupled controls:

- NASSAU_GPU_MEM_BUDGET_MB (default 4096): admission weighted by a launch's
  output bytes, so dozens of small low-stem launches run concurrently again
  (the count cap throttled exactly the region that never had a memory problem)
  while the frontier stays bounded to ~budget/block-size in flight.
- NASSAU_GPU_STREAMS (default 8): fixed CUDA stream slots, round-robin and
  SHARED (small launches serialize on a stream rather than demanding an
  exclusive one), so stream/pool count is bounded independently of concurrency.

Master device uploads are now prefix-only with doubling: a launch ships
max(need, 2*uploaded) entries, not the whole master, so frontier launches (which
append new high-degree R each t) no longer re-ship gigabytes of untouched tail.

Stem 130 improved 187s -> 150s; stem 150 memory 65/37 -> 68/30 GB, verified
bit-identical (stem 80 all-builds, stem 130 all offloaded). But a slots x budget
sweep is FLAT (8/4G=150s, 16/8G=170s, 32/16G=160s, 64/32G=201s): concurrency
knobs are not the ceiling. The ceiling is Amdahl — the GPU accelerates only the
Milnor multiply (~17% of frontier wall time; row_reduce/signature_matrix/readback
dominate and are CPU/serial through cubecl's single runner thread), so the
end-to-end GPU:CPU ratio is flat ~1.13x across the 130-150 heavy bands, not
widening. Widening it would require offloading row_reduce (PR SpectralSequences#274's RREF).

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
@JoeyBF
JoeyBF force-pushed the claude/blas3-row-reduction-lgjz3y branch 2 times, most recently from 668667c to 54c07c2 Compare July 30, 2026 03:17
JoeyBF and others added 22 commits August 15, 2026 18:41
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
JoeyBF and others added 25 commits August 15, 2026 22:44
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
The device RREF launched three cooperative kernels — panel_factor_coop,
promote_coop, block_reduce_coop — synchronized by a hand-rolled grid-wide
spin barrier. cuLaunchCooperativeKernel requires all the grid's CTAs to be
co-resident, which only holds when this process owns the whole GPU. When
another CUDA runtime shares the device (cubecl's Milnor multiply in the
nassau GPU resolution), its kernels occupy SMs, the reduce's CTAs can't all
co-reside, the missing ones never reach the barrier, and the resident ones
spin forever — the intermittent stem-150 wedge (flat sm=100%).

Add an FP_CUDA_RR_COOP switch (rr_coop()). Default off: the forward pass runs
the single-CTA panel_factor one limb at a time, promotion uses the grid-
strided promote_pivots, and back-substitution uses the single-CTA
block_reduce_rref — none launched cooperatively, so the reduce composes with
concurrent GPU work. Set FP_CUDA_RR_COOP=1 to opt into the cooperative path
on a dedicated GPU (measured 2-3x faster at Nassau strides, up to ~10x on
large dense half-rank matrices).

Both paths validated bit-exact vs CPU row_reduce (row_reduce_demo) and vs the
CPU BLAS3 oracle at 2^16/2^17 (reduce_pow2_half).

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
The composable default from the previous commit fell back to single-CTA kernels
(panel_factor, block_reduce_rref) that use one SM of ~132 — 8-10x slower than the
cooperative path on large dense matrices, since the cooperative kernels' real win
is spreading each sequential bit-step across the whole grid.

Recover that all-SM parallelism WITHOUT a cooperative launch by replacing the
in-grid grid_sync with kernel-boundary (stream-ordered) synchronization, the way
cuSOLVER/cuBLAS build grid-wide multi-step algorithms:

  - panel_factor: pf_find -> pf_swap -> pf_xor per bit-step (panel_factor_streamed)
  - block_reduce: br_cond -> br_xor per pivot (block_reduce_elem streamed arm)

All per-step state stays on the device (pivot count, find-first result, pivword,
clear-conditions), so the host issues every launch without a readback and the
latency hides behind GPU work; only the final (pr, pivcols) is copied back. No
cuLaunchCooperativeKernel anywhere on the default path, so it can't deadlock
against a concurrent cubecl kernel. The streamed back-substitution also adopts the
cooperative path's wide-block TRSM (bp=1024 + X.U GEMM) since the GEMM composes.

Perf (H200 half-rank square, device-only) vs cooperative:
  2^13 s128  0.080 vs 0.066 (1.21x)   2^15 s512  0.642 vs 0.577 (1.11x)
  2^14 s256  0.195 vs 0.166 (1.17x)   2^16 s1024 2.27 vs 1.17 (1.95x)
                                       2^17 s2048 8.63 vs 4.52 (1.91x)
Within ~1.1-1.2x of cooperative through stride 512 (the Nassau regime; the wedge
matrix was stride 215) and 3-6x faster than the single-CTA fallback at large
sizes. The residual past stride 1024 is where cooperative additionally fuses
promote/block-reduce; closable with CUDA graphs if those sizes ever matter.

Both paths validated bit-exact vs CPU row_reduce (row_reduce_demo) and the CPU
BLAS3 oracle at 2^16 (reduce_pow2_half).

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Merge pf_find and pf_swap into pf_find_swap using a threadfence "last-CTA
finalize" (a grid-wide reduction, not a barrier, so no co-residency needed),
cutting the streamed forward pass from 3 to 2 launches per bit-step and removing
the serial 1-thread swap kernel. Cap the streamed grid at FP_CUDA_PF_CTAS (128)
like the cooperative kernel. Add FP_CUDA_RR_TIMING to split forward/back timing.

Diagnostic result: the streamed forward pass is work-bound (~cols^2), not
launch-bound — merging, capping, and grid size leave it unchanged — so it is the
per-step relaunch efficiency vs the persistent cooperative grid, ~2.7x at 2^17
but shrinking with size (3.0x at 2^16). Still bit-exact vs CPU row_reduce and the
BLAS3 oracle at 2^16.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
pf_step does column cc-1's clear and column cc's find+swap in one launch, since
the forward sweep alternates xor(col j) / find(col j+1) over the same below-row
range — so each thread reads its row once and sees its own XOR before scanning.
Halves the forward pass's launches (one per column instead of two).

Marginal on its own (~5%), confirming the forward-pass gap vs cooperative is the
per-step GPU relaunch cost of non-persistent kernels, not launch count — it
resists launch reduction and only closes with scale (the O(cols) term vanishing
against O(cols^2) work): ~2x total at 2^17, ~1.4x at 2^18. Bit-exact vs the CPU
BLAS3 oracle at 2^16.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
…ation

Two CUDA consumers share this GPU: the cubecl Milnor multiply and the fp-cuda
row reduction. Moving the reduction off a single global mutex onto per-thread
streams lets concurrent rayon workers overlap, but on its own it made a
stem-200 resolution die within 5-10 seconds with CUDA_ERROR_LAUNCH_FAILED,
reproducibly (3/3).

Isolating the two runtimes shows the fault needs both of them on one device:
multiply on GPU + reduction on CPU ran clean, and reduction on GPU + multiply
on CPU ran clean, while both together died every time. compute-sanitizer finds
no invalid access in either runtime — including with cubecl's allocations
forced synchronous so every buffer is tracked — so this is contention, not an
out-of-bounds bug. Giving fp-cuda its own non-primary CUDA context did not help
either (3/3 still died); the shared *device* is what matters, not the shared
context.

Overlap is also catastrophic for throughput, not just stability. The composable
reduction is a chain of thousands of tiny sequential per-column relaunches, so
sharing the device with the multiply's saturating kernels makes every launch
queue: the same reductions take 1.8-9.7 ms on an unshared GPU and 8.6-96.8 s
co-running, with nvidia-smi showing 99% SM at 10% memory utilisation (queueing,
not compute). The comment claiming this path "needs no cross-runtime exclusion"
had it backwards — being composable means it *can* overlap without deadlocking,
not that it should.

gpu_lock arbitrates: multiplies take the shared side and still overlap each
other; a large reduction takes the device exclusively for its ~10 ms. Total cost
is ~5 s of multiply pause across a whole stem-200 run. Writer preference is
required because multiplies are continuous and would starve the reduction
indefinitely.

Two properties are load-bearing and easy to get wrong:

- WHERE the shared guard is taken. Acquiring it at multiply entry deadlocks: the
  marshalling par_iter runs chunks on other workers, which steal another
  bidegree's multiply, block on the shared side behind a waiting reduction, and
  never let the original join finish. It is taken alongside the existing
  GpuPermit, past every rayon section, for exactly the reason documented there.
- Every wait is bounded, so an unforeseen cycle degrades to lost exclusivity
  rather than a hang. The bounds must exceed how long a reduction holds the
  device; at 25 ms multiplies barged back in mid-reduction and both the slowdown
  and the crashes returned.

With this, a stem-200 resolution completes on the GPU in 2h47m with 0 crashes,
where every prior attempt died. FP_CUDA_DEVICE / NASSAU_GPU_DEVICE put the two
runtimes on separate GPUs when more than one is available, which removes the
contention by construction and makes the arbitration a no-op.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
…and log it

`gpu_lock` arbitration was gated on `FP_CUDA_DEVICE == multiply_device()`, where
`multiply_device()` read `NASSAU_GPU_DEVICE` — a variable nothing in `algebra` reads any
more, left behind when the multiply became multi-GPU. It answered "device 0" however many
GPUs the multiply was actually saturating, so `FP_CUDA_DEVICE=2` on a 4-GPU node would
conclude "separate devices, no arbitration needed" while the multiply hammered device 2
too. The multiply shards across every visible device, so the test is
`FP_CUDA_DEVICE < multiply_devices()`.

This is a latent bug, NOT the cause of the theta=125 LAUNCH_FAILED: with default settings
both sides resolved to 0, so arbitration was already enabled. Verified by the log line
this adds, which exists because `[batch-stats] lock=` cannot distinguish "arbitration off"
from "arbitration on but uncontended" — and the answer turned out to be a third thing.

What the instrumentation actually exposed: `lock=0.0s` in every run *with arbitration
enabled*, because the multiply takes `gpu_lock::shared()` inside the SUBMISSION closure and
drops it when submission returns. Submission only enqueues; the kernels run long after. So
the multiply releases the guard while its saturating kernels are still executing, the
reduction then takes `exclusive()` believing the device is idle, and its thousands of tiny
sequential launches interleave with them — the exact overlap the lock exists to prevent.
Fixing that means holding the guard through the fence rather than the submit (the same
scope error as the earlier GpuPermit bug); left for its own change since it alters the hot
path's concurrency and invalidates the ~5 s/stem-200 cost estimate in the lock's docs.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
…e multiply

`gpu_lock::exclusive()` early-returned whenever `arbitration_needed()` was false, which
conflated two independent questions:

  - does the MULTIPLY have to yield?  (only when it shares the reduction's device)
  - do REDUCTIONS have to serialize?  (always)

The second is unconditional. The row reduction's GEMM is a persistent whole-device grid —
`num_ctas = occupancy x SMs`, cluster-aligned, with cluster sync and DSMEM multicast — so
two concurrent reductions each demand the entire GPU and neither can be placed. On Hopper
that surfaces as a bare CUDA_ERROR_LAUNCH_FAILED, which compute-sanitizer does not
attribute (0 invalid accesses across a whole run: it was never a memory bug).

This is why putting the row reduction on its own GPU did not help on its own:
`FP_CUDA_DEVICE=3` with the multiply on 0..2 turned arbitration off wholesale, so
reductions stopped serializing against each other and the run still failed 63 times in
300 s. Isolating the device removes multiply contention and leaves reduction-vs-reduction
contention untouched.

With the split, an isolated reduction GPU runs clean at the FULL GEMM grid — no CTA cap,
no throughput sacrificed. Measured on the theta=125 stem-200 repro that faulted at ~105 s
in every other configuration: 400 s, 0 launch failures, 0 panics.

The alternative was capping the persistent grid (`FP_CUDA_GEMM_CTAS`). The sweep on
bench_kernel_only (16384^3, idle H200) shows why that is the wrong trade: the largest cap
that survives the workload is 32, and 32 CTAs is 2107 TOPS against 8674 at full grid —
24% of peak. 64 CTAs reaches 49% but still fails. Throughput is linear in CTA count to
~128 (97% of peak), so the kernel only needs the device it is not being given.

Not fixed here: the multiply takes `shared()` inside the SUBMISSION closure and drops it
when submission returns, while its kernels are still resident — so on a SHARED device the
yield is ineffective (`lock=0.0s` in every run). Correct fix is a dedicated fp-cuda driver
thread that fences between reductions and overlaps transfers (copy engines do not consume
SMs, so uploads can pipeline against compute). Isolating the reduction GPU sidesteps it.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Both fp-cuda entry points launch persistent grids sized to fill the machine
(`num_ctas = occupancy x SMs`): the row reduction's trailing GEMM and the
standalone `try_mul`. Since PR 273 removed the thread-block cluster layer these
are ordinary grids, so two co-scheduled ones do not fail — they QUEUE. That is
the problem, not a fix for it.

The reduction is a chain of thousands of tiny relaunches, each a dependency of
the next. Every relaunch that waits behind a saturating GEMM adds its whole
queueing delay to a serial critical path, so the co-running cost is not the
usual graceful throughput split: 1.8-9.7 ms standalone becomes 8.6-96.8 s
co-running, a three-to-four order of magnitude blow-up. The GEMM meanwhile is
one big launch and barely notices. Serializing the two costs the multiply a
bounded pause (~5 s across a whole stem-200 resolution) and buys the reduction
back its standalone latency.

`gpu_lock::exclusive()` covered `row_reduce` only; `try_mul` was deliberately
lock-free ("concurrent callers do not interfere"), so even a dedicated reduction
GPU had two independent machine-filling consumers and was not actually owned by
anything. The cooperative reduction path (`FP_CUDA_RR_COOP`) fails harder than
queueing — it spins forever at a grid-wide barrier for CTAs a concurrent
`try_mul` is holding — which is why it stays off by default.

Routing both through one thread makes single-ownership structural instead of a
discipline each new call site must remember. Jobs run to completion there, and
both end in a device-to-host download, so serialization is on COMPLETION, not
submission — the distinction that matters, and the one `gpu_lock::shared()`
still gets wrong on the multiply side (taken inside the submit closure, dropped
when submission returns, which is why `[batch-stats] lock=` reads 0.0s
everywhere).

Measured on the theta=125 stem-200 repro, reduction isolated to device 3, full
GEMM grid: 200 s, 120 GPU reductions, 0 launch failures, both with and without
FP_CUDA_RR_COOP. Default shared-device config unchanged: 147 reductions, 0
failures, max_t=245 in 200 s.

Cooperative mode is now SAFE but not a win at this workload (closed=20406 vs
20306 in 200 s — noise). The reduction path is too small a share of the
resolution for its ~2x to show. It stays off by default.

Two things this does NOT do. The driver takes no `gpu_lock` guard: serialization
among fp-cuda jobs is structural, and taking the guard there deadlocked the run
(it waits for the multiply's readers while workers block on the driver), so
yielding to the multiply on a SHARED device still needs arranging without a
guard held across a blocking job. Transfers are serialized with compute, though
copy engines do not consume SMs, so a later change can pipeline the next job's
upload against the current job's kernels.

`contended_acquisition_terminates_and_writers_are_exclusive` now measures reader
overlap in its own uncontended phase: with `exclusive` unconditional, the
contended phase keeps a writer queued almost always and writer preference
correctly holds readers off, so asserting overlap there measured the scheduler
rather than the lock.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Instrument Matrix::row_reduce to emit a `fp::rr` tracing event for every
p=2 reduction with min(rows,cols) >= 1024, recording rows/cols/min and
whether the device RREF was taken (path="gpu") or it fell back to CPU
M4RI (path="cpu"). The event inherits the active nassau span so each line
carries its bidegree/signature context.

Confirms on a stem-150 run that every reduction with min >= 8192 (up to
25091x30275, incl. the heavy zero-signature base solves) dispatches to
the GPU with no fallbacks; everything below the 8192 FP_CUDA_RR_THRESHOLD
stays on CPU. tracing is added as a gpu-gated optional dep.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Rebase adaptation, same false premise as the driver-thread doc. This comment
argued that two concurrent reductions "neither can be placed, which surfaces as
CUDA_ERROR_LAUNCH_FAILED". That was true of the cluster-era GEMM, whose CTAs had
to be co-resident by construction; PR 273 removed the thread-block cluster layer,
so the GEMM is now an ordinary grid that queues rather than failing.

The conclusion is unchanged and the code is untouched — reductions must still
serialize against each other — but for the contention reason, not a placement
failure: each reduction is a chain of thousands of tiny sequential relaunches,
and running two at once makes each one's launches queue behind the other's
machine-filling GEMM.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Apply ext/CLAUDE.md's comment rules to the row-reduction work, the pass PR 273
got in 12d97b5 and 80beab2 but that this branch's older content missed.

- Every new `unsafe` block gets a SAFETY: comment. fp-cuda's lib.rs went from
  3 unsafe / 3 SAFETY to 33 / 3 as the row reduction landed; it is 33 / 33
  again. Each names what it discharges: for a `stream.alloc`, which kernel
  writes every element before anything reads it; for a launch, that the pushed
  arguments match the kernel's parameter list and that the borrowed buffers
  outlive the call. The cooperative launches also name where the grid bound
  comes from, since co-residency is what makes them able to deadlock.

- Module docs are one line again. gpu_lock.rs's 26-line header and blas3.rs's
  29-line one moved onto the items they describe -- the arbitration rationale
  onto `exclusive`, the blocked Gauss-Jordan shape onto `row_reduce_blas3`.

- Experiment logs moved out of the source and into fp-cuda/EXPERIMENTS.md,
  which is the crate's designated home for them: the multiply-vs-reduction
  contention measurements and the three alternatives that were rejected (a
  plain RwLock, a 25 ms shared yield, a private CUDA context), the grid caps
  for the three barrier-bound kernels, the panel-width optima, the row-reduce
  CPU crossover, and why a dedicated reduction GPU does not remove the lock.

- One fact, one home: the standalone-vs-co-running figure appeared five times
  across gpu_lock.rs and blas/cuda.rs. It now lives once, in EXPERIMENTS.md,
  and the code links to it.

- Constants by name, not by value. The four tuning defaults that comments
  quoted as literals are now MAX_BL, DEFAULT_PF_CTAS, DEFAULT_PROM_CTAS,
  DEFAULT_BR_CTAS and DEFAULT_BS_BASE.

- Documented the functions that had no doc comment, and rewrapped every
  comment that ran past 100 columns.

Comments only, apart from those five constants -- each a literal replaced by a
name holding the identical value, so the behaviour is unchanged.

Checks: cargo check for fp-cuda and fp with --all-targets and again with
--features gpu against a real nvcc; nightly fmt --check; clippy; the docs gate
with RUSTDOCFLAGS=-D warnings in both feature configurations; the fp test suite
including the blas3 proptests; and row_reduce_demo on an H200, 12/12 bit-exact
against the CPU reduction.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
`gpu_lock` arbitrates between the row reduction and the Milnor multiply, but
the multiply is not on master yet, so nothing in this PR takes either guard --
`shared()` and `exclusive()` are reached only from the module's own unit test.
It ships as dead public API here and belongs with its caller instead.

Removed from this PR, to land with the multiply:

- `fp::gpu_lock` in full, and its `pub mod` declaration.
- `multiply_devices()` and the `set_devices_shared` call in `blas::cuda`. The
  former existed only to work out how many GPUs the multiply spans, and read
  `NASSAU_GPU_DEVICES`, a variable owned by code that has not landed.
- The startup line reporting whether arbitration is live, which described a
  mechanism this PR no longer contains.
- The two EXPERIMENTS.md sections on device sharing and on why a dedicated
  reduction GPU does not remove the lock.

The single-owner property this PR does rely on is unaffected: it comes from the
`driver` thread, which is internal to `fp-cuda` and references nothing external.

Also fixes a comment at the reduction site that had accumulated three
contradictory layers -- "lock-free, overlapping instead of serializing", then
"take the device exclusively for the duration", then "the exclusive guard now
lives on the driver thread, which holds it for the whole job". Only the first
was true; the driver thread explicitly does *not* take the guard, and no caller
did. Replaced with what the code actually does.

Remaining mentions of the unlanded multiply are now phrased in terms of "another
CUDA runtime", which is the property that matters at each site: the cooperative
launch needs the whole grid co-resident and any other runtime occupying SMs
breaks that, and the primary context is shared with whatever else retains it.

Checks: cargo check for fp-cuda and fp with --all-targets and again with
--features gpu against a real nvcc; nightly fmt --check; clippy clean; the docs
gate in both feature configurations; row_reduce_demo on an H200, 12/12
bit-exact.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
BLAS3-ROW-REDUCTION.md and BLAS3-GPU-HANDOFF.md were the design plan and the
CPU-to-GPU handoff brief written while the port was being built. They describe
the work as a sequence of numbered phases with gates, several of which never
landed in the form written, and they will not be maintained once the code is in.
That is not durable documentation -- the crate is described by its README and
its item docs, and what was tried and rejected belongs in EXPERIMENTS.md.

The four example headers that announced themselves as phase gates now say what
they check instead of which step of a plan they unblocked, and point at
`row_reduce_demo` where they used to forward-reference a later phase. Their
banner lines lose the phase numbers too, so the output does not date itself.

The transfer-overlap note this branch added to the handoff document goes with
it. It was future work, not a description of the code.

Checks: cargo check for fp-cuda and fp with --all-targets and again with
--features gpu against a real nvcc; nightly fmt --check; the docs gate in both
feature configurations. No reference to either document remains anywhere in the
tree.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
The BLAS3 planning documents were mostly plan -- phase gates, an API sketch, and
setup instructions for the agent that did the port -- but a few sections were a
genuine experiment log that would be a real loss. Distilled into a new
crates/fp/EXPERIMENTS.md, the CPU counterpart to the fp-cuda one, and cut down
from ~200 lines of narrative to what still informs a decision:

- **The first reading was wrong, and why.** At n <= 4096 the GEMM reduction
  measured 10-16x slower than M4RI, which read as "GPU-only technique, don't
  route CPU row_reduce here". That was an artefact of small matrices. Extending
  the sweep reverses it: the ratio collapses monotonically, and M4RI degrades
  superlinearly at 32768 (a 12.9x step where n^3 predicts ~8x) once the matrix
  overflows cache. This is worth keeping precisely because it is a conclusion
  that got overturned -- the small-n table on its own is misleading.

- **What the two constant-factor fixes bought.** Blocked back-substitution and a
  limb-wise panel, with the full progression per size. Together they move the
  four-core crossover from "few x 10^5" onto n ~ 1.2e5, the size of the real
  workload -- which is why dispatch is a size threshold rather than GPU-only.

- **Panel width is not a tuning knob.** 64 through 1024 all land within noise at
  n = 2048, so DEFAULT_BLOCK_COLS is chosen for shape, not fitted. Without this
  the next reader has no way to know the constant was already swept.

- **Recursive PLE, deferred with a reason.** The sub-cubic alternative and the
  bookkeeping cost that made it not-first.

- **A methodology trap.** The first thread-scaling numbers were taken with
  `concurrent` off, so the "parallel" column measured a serial GEMM.

Also cross-linked: each experiment log now names the other, and
DEFAULT_BLOCK_COLS points at the sweep instead of asserting its shape unbacked.

Deliberately not salvaged: the phase plan, the device API sketch (the API now
exists and documents itself), the repo setup instructions, and the gotchas list,
whose surviving items are either implemented or point at documents that are also
gone.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Three things, all in the crate's comments.

**The SAFETY notes had become a wall.** Nineteen kernel launches carried a
near-identical two-line note repeating the same two obligations. The contract is
stated once now, on `cfg_1d` -- pushed arguments must match the kernel's
parameter list in order and by type, and every device buffer they reference must
outlive the launch -- with `launch_cooperative` adding that its grid has to come
from an occupancy query. Each site is one line naming its kernel, plus a clause
only where it genuinely adds an obligation: dynamic shared memory for
`panel_factor`, the occupancy-derived grid for the three cooperative launches.
Coverage is unchanged at one `SAFETY:` per `unsafe` block.

**Nineteen `design §N` citations pointed at a document this branch deleted.**
Where the citation was the whole parenthetical it is gone; where it trailed real
prose the prose stays. The `.cu` section banners lose it too.

**Two more measurements moved to EXPERIMENTS.md** -- the width gate gating the
multi-CTA block reduction (neutral at 2^15, +6% at 2^16, +18% at 2^17) and the
panel-width optimum, which `adaptive_bl` already documents.

Also deduplicated the operand-layout description, which `matmul_b1_dev` restated
in full from `matmul_b1_raw` a few hundred lines above.

What deliberately stays: the layout arithmetic on `interleave_a`/`transpose_b`
and the algorithm structure on the reduction entry points. That is the class of
thing the code cannot say for itself, and cutting it would trade real
documentation for a smaller line count.

Checks: cargo check for fp-cuda and fp with --all-targets and again with
--features gpu against a real nvcc; nightly fmt --check; the docs gate in both
feature configurations; row_reduce_demo on an H200, 12/12 bit-exact -- which
matters here because the `.cu` banners were edited.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Stripping the `design §N` citations in the previous commit left short lines
mid-paragraph in nine doc comments. Reflowed those nine to 100 columns and
nothing else -- much of the rest of the file wraps at about 80, which is
pre-existing drift and not this pass's business to renormalize.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
An audit against ext/CLAUDE.md turned up ten doc comments this PR adds whose
summary paragraph ran several lines before the body, where the rule asks for a
one-line summary, a blank, then the body. Split each so the first line stands
alone, and reflowed the paragraphs underneath -- four of which the previous
commit's citation removal had left ragged in the middle.

Two other findings from the same audit, both in code this PR introduces:

- `try_row_reduce` documented "Assumes `m.prime() == 2` (the caller has
  checked)", which is precisely the precondition-restatement the rules exclude.
  Dropped.

- The row-reduce dispatch comment in `matrix_inner` named the threshold by value
  ("below the 8192 threshold") and referred to "the active nassau span" -- code
  that is not on master. Both are now phrased in terms of what is actually here:
  `blas::cuda`'s threshold, and whatever span the caller is in. The earlier sweep
  for unlanded-code references missed this one because it matched `NASSAU_` and
  `nassau_gpu` rather than the bare word.

Not touched, and not this PR's to fix: eleven doc comments and fourteen `unsafe`
blocks in `matrix_inner.rs` that predate this branch, plus the 25-line module doc
on `blas/mod.rs`. The PR adds 47 lines to `matrix_inner.rs` and no `unsafe` at
all there.

Checks: cargo check for fp-cuda and fp with --all-targets and again with
--features gpu against a real nvcc; nightly fmt --check; the docs gate in both
feature configurations.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
@JoeyBF
JoeyBF force-pushed the claude/blas3-row-reduction-lgjz3y branch from 54c07c2 to dd3211d Compare August 26, 2026 05:49
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants