Skip to content

milnor_gpu: fix >2^32-element master via cubecl 64-bit addressing - #17

Draft
JoeyBF wants to merge 140 commits into
rref-hpcfrom
milnor-gpu-u64-addressing
Draft

milnor_gpu: fix >2^32-element master via cubecl 64-bit addressing#17
JoeyBF wants to merge 140 commits into
rref-hpcfrom
milnor-gpu-u64-addressing

Conversation

@JoeyBF

@JoeyBF JoeyBF commented Jul 27, 2026

Copy link
Copy Markdown
Owner

Problem

The batched Milnor multiply silently corrupted results once the shared masks admissible-master crossed 2³² u16 elements (~8.6 GB, around S_2 stem ~160–180). The buffer length and per-R offsets were u32, so cubecl's default U32 address_type truncated them (len as u32) → out-of-range reads → dx != 0 (d²≠0) panics, e.g. deterministically at (180,92). (160,82) stayed just under the ceiling and was clean.

Diagnosis (instrumented): masks reached 4,452,195,742 u16 (8.29 GiB) > RESIDENT_MAX_CAP = 2³²−1; the ArrayArg length metadata also truncated (u32). Confirmed not the row-reduce, not the resident basis, not a cross-stream race — pure u32 addressing overflow.

Fix — idiomatic cubecl 0.10 64-bit addressing (no hand-rolled buffer splitting)

cubecl has AddressType {U32, U64}; a plain #[cube(launch)] defaults to U32. Two CUDA-backend quirks shaped the fix:

  1. address_type = "dynamic" miscompiles for us — it picks U32 for small blocks and then narrows the Array<u64> offset arrays on read (usize::cast_from(u64) under u32 usize), corrupting results. So the multiply is static "u64".
  2. checked-mode + U64 won't compile — the bounds clamp emits min(u64, u64), ambiguous for NVRTC. So it uses launch_unchecked (safe: every access is in-bounds by construction — uploaded need_* prefix + per-column j guards).

Changes (all in the batched multiply path):

  • multiply_batch_kernel#[cube(launch_unchecked, address_type = "u64")].
  • RInfo.cs_off/mk_off u32→u64; r_cs_offset/r_mk_offset bound as Array<u64>.
  • resident grow copies (copy_into_u16/_u32) → #[cube(launch_unchecked, address_type = "dynamic")].
  • drop RESIDENT_MAX_CAP + clamps; assert out_len ≤ u32::MAX loudly.

Validation (H200)

  • (160,82): rc=0 clean, 517 s vs 486 s baseline (+6.4% for static-u64 multiply, partly offset by unchecked dropping the clamp — no regression).
  • (180,92): ran 1500 s with zero dx panics (was a deterministic panic ~130–180 s), i.e. correct well past the 2³² masks boundary. A to-completion rc=0 run is in progress.

🤖 Generated with Claude Code

JoeyBF and others added 30 commits July 26, 2026 20:03
The batched Milnor multiply silently corrupted results once the shared
`masks` admissible-master crossed 2^32 u16 elements (~8.6 GB, around
S_2 stem ~160-180): the buffer length and the per-R offsets were u32, so
cubecl's default U32 `address_type` truncated them (`len as u32`), giving
out-of-range reads and `dx != 0` (d^2 != 0) panics at e.g. (180,92).
(160,82) stayed just under the ceiling and was clean.

Fix uses cubecl 0.10's first-class 64-bit addressing rather than
hand-rolled buffer splitting:

- `multiply_batch_kernel` -> `#[cube(launch_unchecked, address_type = "u64")]`.
  Static u64 (not "dynamic"): dynamic picks u32 `usize` for small blocks and
  then narrows the u64 offset arrays on read (`usize::cast_from(u64)` under a
  u32 address type), corrupting results. `launch_unchecked` because checked
  mode emits `min(u64, u64)`, which NVRTC rejects as an ambiguous overload;
  every access is in-bounds by construction (uploaded `need_*` prefix +
  per-column `j` guards).
- `RInfo.cs_off/mk_off` u32 -> u64; `r_cs_offset`/`r_mk_offset` bound as
  `Array<u64>`.
- Resident grow copies (`copy_into_u16`/`_u32`) ->
  `#[cube(launch_unchecked, address_type = "dynamic")]` (scalar usize offsets
  adapt without narrowing; dynamic keeps small copies on u32).
- Drop `RESIDENT_MAX_CAP` and its clamps (the u32 ceiling is gone); assert
  `out_len <= u32::MAX` loudly (the row-block splitter guarantees it).

Validated on H200: (160,82) rc=0 clean, 517s vs 486s baseline (+6.4% for the
static-u64 multiply, partly offset by unchecked dropping the bounds clamp);
(180,92) ran 1500s with zero dx panics (previously a deterministic panic
~130-180s), i.e. correct well past the 2^32 masks boundary.

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

`resolve_through_stem S_2 "" 180 92` was OOM-killed at the 500 GB cgroup limit
(~280 GB anon + ~240 GB shmem). Measurement (env `NASSAU_MEM_REPORT`, plus a
pure-CPU control at ~5-7 GB) showed the resolution's own data — differentials
and module tables — is only ~1.5 GB; the ~500 GB was entirely GPU-runtime host
memory. Two fixes cut it to ~90 GB at stem 180 (dx-clean):

1. Default `NASSAU_GPU_STREAMS` 8 -> 1. cubecl gives each CUDA stream its own
   page-locked (pinned) host pool that `memory_cleanup` never trims; ≥2 streams
   ballooned pinned host memory to 140-240 GB, single-stream holds it at ~4-6 GB.
   The payoff of multi-stream is ~nil here: cubecl's server is single-threaded,
   so extra streams buy no CPU concurrency, only GPU kernel overlap the big
   saturating multiplies barely use. Override for a dedicated large-RAM node.

2. Stop retaining the resident admissible master host-side. `RESIDENT_HOST` kept
   the full `col_sums`/`masks` (~54 GB at stem 180) forever, duplicating the
   device copy, even though after upload it is never read again (offsets come
   from `index`; growth uploads only the new tail; a capacity realloc copies the
   old *device* buffer). Now it keeps only the not-yet-uploaded tail (`*_pending`,
   ~sub-GB) plus a logical length, freeing each `R`'s data the moment it reaches
   the GPU — invariant `dev.uploaded == len - pending.len()`.

Also: bound the pinned staging on resident growth to `STAGE_CHUNK` chunks, and a
gated `NASSAU_MEM_REPORT` (differentials/modules/resident heap breakdown) for
ongoing memory work.

Validated on H200: stem 180 peak RSS ~90 GB (was 500 GB OOM), dx-clean through
the >2^32 masks region and the heavy solves. Remaining growth for higher stems
is real GPU working memory (concurrent dense output matrices) + GPU device
memory, not retained host duplicates.

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

Even with the resident host-master freed, host RSS still grew ~linearly toward
the cgroup limit at high stems — not a retained duplicate but the GPU path
materializing whole dense output matrices. The CPU walks signatures one at a
time (small working set); the GPU offload built the entire matrix at once and
held its dense readback (`num_rows × num_limbs` u32, ~12 GB regions at stem 180)
alongside the assembled matrix. Two caps bound it:

1. `reuse_full_matrix` now only builds the all-rows matrix when
   `rows × cols <= NASSAU_GPU_REUSE_MAX_WORK` (default 1e10). Above that it falls
   back to per-signature builds (each a bounded row subset, like the CPU),
   so the peak scales with the largest single signature, not the whole bidegree.

2. `get_partial_matrix_restricted` processes rows in batches sized so the dense
   readback stays under `NASSAU_GPU_MAX_READBACK_MB` (default 1024). Products are
   built in row order, so each batch is a contiguous slice with its `row` remapped
   batch-local; results XOR back into the global rows. The readback is freed
   between batches.

Both are correctness-neutral (rows are independent): verified bit-for-bit vs the
CPU under `NASSAU_GPU_VERIFY` with a 1 MB cap (many batches/build), 0 mismatches.

Effect at stem 180 (streams=1 + resident de-dup + this): host anon goes from
growing past 85 GB to a bounded ~18-24 GB, with no throughput loss — cross-
bidegree rayon concurrency keeps the GPU fed, and the default cap only splits the
few giant high-t builds into a handful of chunks. Peak RSS ~30 GB (was 500 GB
OOM). This bounds the host working set so it scales toward higher stems; a D-deep
async launch/collect pipeline (cubecl launches are async; only read_one blocks)
can hide chunk latency if a much smaller cap is ever needed.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
The resident admissible master (col_sums/masks) grows with degree and is the
stem-300 device-memory wall: ~34 GB at stem 180, extrapolating past the H200's
143 GB before stem 300. This adds an opt-in eviction that bounds it.

Policy (from the NASSAU_R_STATS probe): a *degree* threshold, not LRU. Low-degree
R's are the stable hot core (reference span 0.99 of the run); high-degree R's are
scattered-recurring (0.81) and also the biggest matrices, so a degree cap evicts
the most bytes for the fewest references and is stem-independent (the kept set
saturates). Env NASSAU_GPU_RESIDENT_MAX_DEGREE (default i32::MAX = keep all).

Design — no kernel change. Every output row's products share one operation R
(extract_restricted), so hot (deg<=cap) and cold rows are DISJOINT. multiply_batch_on_gpu
partitions products, compacts each group to its own dense row range (so total
readback stays num_rows, not 2x), runs the resident pass on hot and a Transient
pass on cold, and scatters results back. Transient builds a per-block master with
create_from_slice, freed with the launch, so the device copy never persists. Cold
admissible data is cached host-side (COLD_HOST) so the expensive recompute is
one-shot, like the resident path. Fast path (cap==MAX or all-hot) is byte-identical
to before — zero regression by default. GpuProduct gains Clone for the partition.

Validated bit-exact: NASSAU_GPU_VERIFY full S_2 stem-110 resolution at theta=10
(transient path heavily stressed) and theta=100: mismatches=0 dx=0 panics=0.

Memory vs throughput (stem 180, ExclusivePages, streams=1): control (no eviction)
51 GB / 1349 s; theta=125 11 GB; theta=100 12 GB (master 34 -> ~1 GB), dx=0. But
eviction costs 2-4x throughput even at a high theta (stem 150: control 282 s,
theta=140 613 s) because the evicted high-degree matrices are the biggest and are
re-uploaded every launch. So this is a fit-in-memory lever for stems where the
master won't fit at all (a slow completion beats an OOM), tuned to the highest
theta that fits; it is NOT a speedup. Next optimization to cut the re-upload: a
bounded LRU device-side cold cache (upload once, reuse across a bidegree's launches).

Also retained: the R-access probe (NASSAU_R_STATS/dump_r_stats) and the device/host
[MEM] report (NASSAU_MEM_REPORT) used to characterize this — both env-gated no-ops.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Foundational step toward generating admissible matrices (col_sums/masks) ON the
GPU into transient scratch instead of storing/uploading the resident master --
the direction that eliminates both the stem-300 device-memory wall and the
eviction re-upload cost (a given launch would enumerate its cold R's on-device,
never uploading them).

enumerate_admissible_ref reimplements AdmissibleMatrix::next using ONLY
flag-guarded control flow -- no break, continue, or early return -- because that
is the subset the cubecl DSL compiles cleanly (cf. multiply_pair, which tracks a
`rejected` flag rather than breaking). `found` replaces the odometer's `return
true`, `handled` replaces its `continue 'mid`. This validates the tricky
restructuring on the CPU, where it is fast to debug, before the hard-to-debug
#[cube] port -- which then becomes a mechanical transcription onto per-thread
local Arrays (state is tiny: rows = |p_part|, cols <= 32).

Test admissible_enum_ref_matches asserts bit-exact equivalence with
admissible_matrices over every real R up to degree 60 (4155 R's, all match).

Not yet ported to cubecl / not wired into the multiply -- next step.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Transcribe enumerate_admissible_ref into a cubecl #[cube] kernel that
generates each R's col_sums/masks for every admissible matrix directly
into device scratch -- the on-GPU replacement for the resident/uploaded
master. One thread per distinct R, all state in fixed-size per-thread
local Arrays; flag-based control flow (while ... && !found / handled)
since the DSL has no break/continue/return.

Backend-agnostic host driver (generic over Runtime) so the identical
kernel can run on CUDA and the cpu backend. New test
admissible_enum_gpu_matches runs it on the H200 and asserts bit-exact
output (values + per-R counts) vs the CPU reference: 1055 R's, 30385
matrices, all matching.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Wire enumerate_admissible_kernel into the Transient (evicted) cold path of
multiply_batch_block. A cold R's col_sums/masks are now GENERATED on the GPU
into transient scratch (one enumeration launch, ordered before the multiply on
the same stream) rather than built on the host and uploaded via
create_from_slice. Only the small p-parts + per-R dimensions upload; the scratch
is freed with the launch. This kills the per-launch H2D master re-upload that
made eviction a 2-4x slowdown (bench 2026-07-27), the whole point of the
in-kernel-enumeration direction.

Also replace the COLD_HOST full-array cache with COLD_COUNT, a 12-byte-per-R
(cs_len, mk_len, num_mats) shape cache. The evicted tail of the master now lives
neither on the device nor the host -- only its sizes, needed up front to lay out
the scratch offsets and the pair-count prefix sum. num_mats is counted once per
distinct R (admissible_matrices, arrays dropped) and memoized.

De-gate enumerate_admissible_kernel + ENUM_* caps + the MAX_XI_TAU import for
production. Validated bit-exact: the isolation test (1055 R's / 30385 matrices)
still passes, and NASSAU_GPU_VERIFY full S_2 stem-110 resolutions at theta=10
(nearly all R's cold -> enumeration path stressed) and theta=100 both report
mismatches=0 dx=0 panics=0 rc=0.

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

Investigation of the eviction crash + a throughput bench that together show
in-kernel enumeration cannot beat upload-based eviction:

- enumerate_admissible_kernel -> u64 addressing (was default u32); did NOT fix
  the big-block CUDA_ERROR_LAUNCH_FAILED, so the fault is elsewhere in the
  Transient wiring, not the kernel.
- admissible_enum_gpu_matches extended to degree 145, chunked per-degree:
  proves the kernel bit-exact vs the CPU reference (144903 R's, 185M matrices)
  across the full range the eviction path exercises -- so the kernel is correct.
- bench_admissible_cpu_vs_gpu (ignored): CPU admissible_matrices 1.61s vs GPU
  kernel-only 2.02s (0.8x, SLOWER) for degrees 1-130. GPU enumeration is ~3x
  slower than just transferring the same arrays (0.68s readback) it replaces:
  the odometer is sequential per R with matrices-per-R spanning 1..millions, so
  the launch bottlenecks on its few longest threads at GPU scalar speed. The
  "trade GPU integer work for PCIe bandwidth" premise is refuted -- enumeration
  worsens the eviction re-upload cost instead of curing it.

The enum wiring in multiply_batch_block's Transient path still faults at scale
and should be reverted to upload-based eviction; kept here as validated,
dormant code documenting the dead end. Default (theta=inf) path is unaffected.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
The earlier bench compared GPU enum kernel-only to the D->H READBACK (0.68s) as
a stand-in for the upload it replaces, and wrongly concluded enum is ~3x slower.
Production uses the matrices on-device (no readback), and the readback is not the
upload. Measure the actual H->D create_from_slice upload of the same host-built
arrays instead, synced via a tiny throwaway readback (upload-only):

  CPU admissible_matrices (1 core) : 1.60 s
  GPU enumerate in-kernel          : 2.02 s
  H->D upload of host-built arrays  : 1.85 s
  -> in-kernel enum is 1.09x the upload (parity), NOT 3x

So in-kernel enumeration is throughput-neutral vs uploading AND saves the host
COLD_HOST array cache (tens of GB at stem 300). The direction is viable, not
refuted; the big-block wiring crash is worth fixing.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Prototype step 1 toward no-copy master growth, to kill the realloc-doubling
transient that pushes cubecl into its memory-corruption regime (the ~2x spike
at a 32->64 GiB master grow = ~96 GiB live, the root cause of the stem-140+
dx=nonzero cubecl ManagedMemoryDescriptor corruption).

A segmented master keeps old segments and appends a new fixed-size segment on
growth -- never copies -- so device peak is live_size + one segment, flat. The
kernel selects a segment by static branch (cubecl has no array-of-buffers):
offset o -> segment o/seg_elems, local o%seg_elems (seg_read_u16, MASTER_MAX_SEG
separate Array args). seg_gather_kernel + seg_read_matches_contiguous validate
the mechanic bit-exact (8 segments, scrambled indices) before touching the
multiply hot path.

Not yet wired into multiply_batch_kernel / resident_dev_handle -- next step.

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

Replace the doubling-realloc-with-copy growth of every resident device
buffer with a single segmented, append-only, no-copy mechanism. Growth now
allocates only the new fixed-size segment(s) it needs and stage-writes the
tail into them; existing segments are never reallocated or copied, so the
device peak is `live + one_segment` instead of the ~2x transient (old+new
buffer both live) that pushed cubecl into its silent memory-corruption
regime -- the stem-140+ `dx != 0`.

One mechanism for all four resident buffers (master cs/mk, basis pp/ln):
- SegBuf + seg_grow! replace GrowBuf, resident_dev_handle!,
  basis_dev_handles!, stage_upload!, RESIDENT_REALLOC, RESIDENT_INIT_CAP
  (all deleted -- no two coexisting growth paths).
- multiply_batch_kernel binds each store as MASTER_MAX_SEG (=16) segment
  Arrays and GATHERS a thread's matrix cs/mk + its term p-part into small
  WORKING_CAP locals via seg_read_u16/seg_read_u32 (correct at any offset,
  so no layout padding), then calls the UNCHANGED pure multiply_pair with
  base 0. Only this kernel changed; multiply_pair and the test kernel are
  untouched.
- No realloc barrier: segments never change identity or get freed, so a
  reader's cloned handles stay valid across a concurrent append.
- master_seg_elems() (env NASSAU_GPU_MASTER_SEG_ELEMS, default 1<<31, < u32
  so a segment length never truncates cubecl's 32-bit metadata) => 64 GiB
  per buffer over 16 segments. Over-cap is a clean assert, not corruption.

Validated bit-exact on H200: seg_read_matches_contiguous (16-arg
primitive), multiply_batch_matches_reference across NASSAU_GPU_MASTER_SEG_ELEMS
768..4096 (multi-segment single-launch gather), and a new
multiply_batch_incremental_growth at seg_elems=8192 (cross-launch append
into the partially-filled last segment).

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Relocate every #[cfg(test)] item (the test-only cube kernels xor_f2 /
seqno_kernel / multiply_single_r_kernel / seg_gather_kernel and their host
drivers, plus enumerate_admissible_ref / _on_runtime, seg_gather_on_gpu,
check_enum_backend, enum_launch_timed) out of the production module body and
into `mod tests`, so the production path is no longer interleaved with test
scaffolding. Drops the now-redundant inner #[cfg(test)] attributes (the
module already carries it) and applies the project's nightly rustfmt (this
also formats the segmented-master rewrite from the previous commit).

Pure move; no behavior change. GPU correctness suite still green.

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

Phase 1 of reducing the per-launch allocation/copy churn that drives cubecl's
allocator into CUDA_ERROR_LAUNCH_FAILED (719) at scale (memcheck showed the
fault is on cuMemcpyHtoDAsync/cuMemAllocAsync, not our kernels). Keeps the
exclusive-pages memory mode (the ~50% footprint win) and removes the churn it
caused, instead of switching to the sub-slice pool (which reuses buffers but
~2x memory -> OOM at the 90 GB master).

- g/xi: the seqno table and (constant) xi degrees are identical every launch at
  a given degree; upload them once to shared resident handles (RESIDENT_SEQNO,
  keyed by g.len()), re-uploading only on a degree bump, instead of a
  create_from_slice each launch. Synced before publish for cross-stream reads.
- out_h: reuse a persistent per-worker (= per-stream) XOR accumulator (OUT_ACCUM
  thread-local), grown only when a larger out_len appears, instead of empty()
  + free-via-memory_cleanup every launch (the single biggest churned buffer).
  Read back only the used [0,out_len) prefix via Handle::offset_end.

memory_cleanup stays for now (trims the still-churning per-R/per-product
metadata; phase 2 makes those persistent and drops it). Bit-exact vs the prior
path: multiply_batch_matches_reference, multiply_batch_incremental_growth
(5 growing launches: out_h regrow + g/xi re-upload), multiply_single_r.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Revert the OUT_ACCUM thread-local output buffer from the previous commit: it
was per rayon worker (~100 threads), so it held ~100 persistent block-sized
buffers (~28-50 GB) and DEFEATED the existing out_h bounding — each launch's
out_h is already chunked to NASSAU_GPU_BLOCK_MB (512 MiB) and total in-flight
output is capped by the GPU_BUDGET permit (NASSAU_GPU_MEM_BUDGET_MB). Persisting
them per thread pushed device memory toward the OOM ceiling.

out_h returns to a per-launch empty() reclaimed by memory_cleanup (properly
chunked + budget-bounded). The resident g/xi seqno tables (uploaded once vs
re-uploaded per launch) are kept — a clean churn reduction. Bit-exact:
multiply_batch_matches_reference, multiply_batch_incremental_growth,
multiply_single_r.

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

When the cubecl CUDA context is poisoned by the unresolved uninit-handle bug
(CUDA_ERROR_LAUNCH_FAILED / ServerUnhealthy at the multiply readback,
milnor_gpu.rs:2074; tracel-ai/cubecl#1401), the whole context is dead — every
later launch fails, so per-call retry cannot recover. Wrap multiply_batch_on_gpu
in catch_unwind: on the first failure, set a process-wide GPU_DISABLED flag and
finish the resolution on the CPU via a new cpu_multiply_batch. The CPU output is
bit-identical to the GPU's (validated by cpu_multiply_batch_matches_gpu, including
the multi-block out_offset path a real module row uses). RREF runs on a separate
fp-cuda runtime and is intentionally not gated by this flag.

Add benches/nassau_milnor_gpu.rs: a GPU batched-multiply throughput / regression
bench (counterpart to nassau_milnor.rs) that hammers multiply_batch_on_gpu over an
output-degree sweep with no row-reduction or resolution machinery — for
`cargo bench --baseline` comparison across cubecl commits and for isolating a
multiply/allocator crash from the rest of the pipeline. Compiles to a no-op main
without the `gpu` feature.

Also pins cubecl/cubecl-common to the JoeyBF fork branch
(claude/pool-slot-map-v0.10.0) for the #1401 generational-slot-pool validation.
TEMPORARY: revert to the crates.io release before merging upstream.

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

The residual high-stem CUDA_ERROR_LAUNCH_FAILED in the GPU multiply is a cross-stream
pool-reclaim race: ~100 streams each calling client.memory_cleanup() every launch, one
stream's cleanup reclaiming a shared resident-master page still in flight on another
stream's kernel -> bad device pointer -> context poison (tracel-ai/cubecl#1401; the fork's
per-command retain_until_complete does not cover cross-stream shared-page reclaim).

Gate the cleanup call on NASSAU_GPU_CLEANUP_EVERY (default 1 = every launch, N = every Nth,
0 = never). Validated: with =0, S_2 stem-200 resolves fully clean on GPU (0 LAUNCH_FAILED,
0 dx) — the first clean GPU stem-200 — with device memory plateauing ~125 GB under cubecl's
internal pressure-triggered reclaim (which does NOT hit the race). A moderate throttle
(~16-32) is the likely stem-300 config: race-avoiding while bounding memory further.

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

An #[ignore]d GPU soak: N worker threads hammer multiply_batch_on_gpu against one
growing shared resident master, across NASSAU_GPU_STREAMS streams with per-launch
memory_cleanup on — the cross-stream access pattern the stem-200 resolution drives.
Each result is checked against the bit-identical cpu_multiply_batch oracle (up to
NASSAU_SOAK_VERIFY_MAX, to keep the CPU precompute cheap); any mid-soak context
death flips GPU_DISABLED and fails the run.

Two jobs in one:
- Correctness: catches cross-stream renumber/identity races in seconds.
- #1401 reproducer (measured): clean at max_degree<=128, but at max_degree=160 with
  NASSAU_GPU_STREAMS=48 the cubecl cross-stream pool-reclaim race fires within a ~45s
  soak (never-initialized / ServerUnhealthy cascade, gpu_disabled flips) at only
  ~28 GB host / ~22 GB GPU — the genuine timing race, NOT an OOM. That's a ~1-2 min,
  low-memory stand-in for the 40-min stem-200 crash, and the gate the coming
  single-submission-thread redesign must turn GREEN.

Repro: NASSAU_GPU_STREAMS=48 NASSAU_GPU_CLEANUP_EVERY=1 NASSAU_SOAK_MAX_DEGREE=160 \
  cargo test -p algebra --release --features gpu -- --ignored --nocapture concurrent_growth_soak

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Drop the JoeyBF/cubecl `pool-slot-map-v0.10.0` fork pin for upstream
`tracel-ai/cubecl` tag `v0.11.0-pre.1`, whose rewritten memory manager
(+1377 lines) and CUDA stream layer are the async-engine/pool subsystem
where the cross-stream reclaim race (#1401) lives.

Migrate milnor_gpu.rs to the 0.11 launch API (227 one-line replacements):
  - launchable array kernel args `&Array<T>`/`&mut Array<T>` → slices
    `&[T]`/`&mut [T]` (incl. `&mut Array<Atomic<u32>>` → `&mut [Atomic<u32>]`);
  - host-side `ArrayArg::from_raw_parts` → `BufferArg::from_raw_parts`
    (identical 2-arg signature, from `cubecl::prelude::*`);
  - local scratch `Array<T>` passed to `#[cube]` helpers now go through
    `Array::as_slice()`; non-launch helper params take `&[T]`.
Every `address_type = "u64"`/`"dynamic"` attribute and `launch_unchecked`
call is preserved verbatim — the high-stem u64-addressing correctness path
is untouched.

Effect at the harsher-than-production d=160/48-stream soak: break rate
3/5 (fork) → 1/5 (0.11-pre), independent of NASSAU_GPU_CLEANUP_EVERY, with
0 correctness mismatches across all runs. The residual fault surfaces as a
gentler CUDA_ERROR_ILLEGAL_ADDRESS at read_one (vs the old LAUNCH_FAILED
uninit-handle cascade), always caught by the CPU multiply fallback. The
race is reduced but not eliminated; the fallback remains the safety net.

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

multiply_pair computed word = row_base + (out_offset + seqno)/32 and did an
unguarded atomic XOR into out[word]. When out_offset + seqno spans past the
row's num_limbs (which nassau_gpu::get_partial_matrix_restricted already
anticipates and masks on readback — bits >= target_dim are dropped), the
device write overran: into the next row (silent corruption) or past the
buffer (CUDA_ERROR_ILLEGAL_ADDRESS). compute-sanitizer on a malloc_sync
build confirmed "Invalid __global__ atomic of size 4 bytes ... out of bounds".

Thread num_limbs through multiply_batch_kernel and skip writes with
global_bit/32 >= num_limbs — a device-side mirror of the host's existing
defensive mask. Correctness-preserving; the skipped bits are exactly the
ones the host discards. The d=160/48-stream concurrent soak that reliably
broke ~3-5/5 now runs 0/5 with 0 correctness mismatches.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
The p-part of a Milnor basis element was a `Vec<u32>`, costing a heap
allocation and a pointer chase per element. At p = 2 the internal degree of
P(R) is sum_i r_i (2^i - 1) with non-negative terms, so r_i <= deg/(2^i - 1);
sizing each field by that bound packs the whole exponent sequence into 64 bits
for every degree up to 2045. At odd primes the same bound applies divided by
q = 2(p-1), so one layout serves every prime.

`MilnorBasisElement` is now 16 bytes, `Copy`, and entirely inline. Measured
over degrees 0..=300 at p = 2, `basis_table` drops from 252 MiB in 5,036,688
allocations to 77 MiB in none.

Three things fall out of the packing:

- The packed value is a canonical key, so the hand-rolled `MilnorHashMap`
  specialization for `not(odd-primes)` is gone; a plain `HashMap` now hashes a
  single word on every path. That code also assumed a degree bound of 1536
  without enforcing it. `compute_basis` now asserts the bound up front, which
  is what lets everything downstream skip range checks.
- Trailing zeros are not represented, so the "pop trailing zeros" loops after
  building a product disappear.
- `PPartMultiplier` no longer borrows its inputs, so its lifetime parameter is
  gone, and `PPartAllocation` loses the buffer it existed to recycle.

In `ext`, `MilnorSubalgebra`'s signature test becomes one masked comparison on
the packed word instead of a loop over entries, with the mask hoisted out of
`signature_mask`'s inner loop.

Two behaviour changes worth noting:

- `basis_element_from_string("P0")` and `("Sq0")` now return the identity
  rather than `None`. P(0) is the identity, and `AdemAlgebra::try_beps_pn`
  already special-cases `x == 0` this way; the old `None` came from `vec![0]`
  and `vec![]` hashing differently, an artifact of the representation.
- `increment_p_part` now carries before incrementing. The old order
  transiently stored `max[i] + 1`, which need not fit a field whose width is
  exactly saturated by `max[i]`. The enumeration is unchanged.

The observation that every Milnor exponent sequence up to degree 512 fits in 64
bits is due to Lixiong Wu; this implementation works out the widths, finds that
the same layout holds all the way to degree 2045, and carries it through the
algebra.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XK8YCdHkCqUV9YM7D957hV
The first packing pass regressed `milnor_ppart` by up to 8% at odd primes and
mod 4, because assembling the answer went from a memcpy plus a vectorized add
to a per-entry read-modify-write through the checked `PPart::set`, and because
`PPart::get`'s range branch landed in `update`'s inner loop.

Two changes, both confined to the kernel:

- Assemble the answer in a plain `u64` and store it once. Entries are written
  in increasing index order into a value that starts at zero, so a shift and
  an `or` suffice; the range checks become debug assertions backed by
  `compute_basis`'s degree gate.
- Pad the layout tables to 16 entries so the private `PPart::entry` can mask
  its index rather than branch on it. Padded entries have width zero and so
  read as zero, which is the answer `get` would have returned anyway. The
  public `get` keeps its explicit check, since callers outside the multiplier
  index it with a q-part-derived length that is not bounded by `MAX_LEN`.

This recovers the regression (`ppart_4/a` and `ppart_3/a` back to baseline,
`ppart_4/b` -8%) and improves the Nassau regime further.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XK8YCdHkCqUV9YM7D957hV
`basis_table` held a `MilnorBasisElement` per basis element. At p = 2 with
unstable support off, that element is exactly
`from_p(ppart_table[t][i], t)` -- the p-part again, with a q-part that is
always zero and a degree that is the index. It was a redundant copy.

Deriving it on demand costs nothing now that `MilnorBasisElement` is `Copy`
and 16 bytes: `basis_element_from_index` returns by value and builds it in
registers rather than handing out a reference into a table. The multiply
family takes the element by value for the same reason.

The table is still built at odd primes, where the q-part varies within a
degree, and when unstable support is on, where the basis is re-sorted by
excess. Neither is a re-wrapping of `ppart_table`.

Measured over degrees 0..=250 at p = 2 (1,958,958 elements), RSS growth from
`compute_basis` drops 125.0 MB -> 95.0 MB, i.e. 66.9 -> 50.8 bytes per element.
Projected to degree 500 that is 5.24 GB -> 3.95 GB. Unlike the ranker, this
needs no basis renumbering and costs nothing at lookup time.

A test verifies the derivation matches what the table used to hold, for every
element, so the redundancy is asserted rather than assumed.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XK8YCdHkCqUV9YM7D957hV
`basis_element_to_index` runs once per term of every product, and is a hash map
storing an entry per basis element. The index it returns is a position in an
enumeration, so a canonical key alone cannot replace the map -- but the position
can be computed.

Let counts[i][d] be the number of exponent sequences of degree d using only
xi_1..xi_i. Splitting on whether r_i is zero gives the coin-change recurrence
counts[i][d] = counts[i-1][d] + counts[i][d - xi_i]. Ranking needs the number of
sequences with r_i > v, and substituting r_i -> r_i - (v+1) is a bijection onto
all sequences of degree d - (v+1)*xi_i, so that count is a single table lookup
rather than a sum. Walking the entries downward ranks a p-part in one lookup
each, against a table covering every degree at once, where the map it would
replace grows with the basis.

Whether that is worth it depends entirely on scale, which took some measuring to
see. Against the map, at p = 2:

    degree   per-degree map   hashmap    ranker    ratio
       120          0.10 MB    11.6us    26.7us    0.43x
       300          3.12 MB     792us    1384us    0.57x
       400         12.50 MB    4490us    5310us    0.85x
       500         37.50 MB   33118us   15865us    2.09x

A lookup probes only its own degree's map. While that fits in cache the map wins
easily: one hash round and one probe, against six to ten dependent table reads.
Once it does not -- the map is 37 MB in degree 500 -- every probe misses to DRAM
at ~33 ns, whereas the ranker's table is ~43 KB, stays in L1, and costs ~16 ns
regardless of degree. Benchmarking only up to degree 120, where the map is
0.1 MB, shows a 2x loss and hides all of this; the sweep here deliberately spans
the crossover. Tuning does not move the small-degree end: nested vs flat table,
a zero-padded prefix to drop the branch, and one- vs two-pass to break the
dependency chain were all measured, and the padded variant was worst, because
doubling the table pushed it out of L1.

So the two suit opposite ends of the range, and the ranker is on the right side
of the end where the algebra's memory is the problem worth solving: replacing the
map there is 3.3 GB smaller and 2x faster.

It stays off by default and unwired even when enabled, because it numbers the
basis in colex order rather than the order compute_ppart emits, which would
invalidate saved resolutions. That order is rankable in principle, but its
natural recursion has depth equal to the sum of the entries, which is worse than
hashing. The unstable path, which re-sorts each degree by excess, is not modelled
either.

Tests verify the table reproduces the algebra's own p-part counts and that the
rank is a bijection onto 0..dim in every degree, at p = 2 and p = 3, plus one
pinning down that it really does disagree with the current basis order.

Also measured and rejected: an `unrank` recovering the p-part at a given index,
which would let basis_element_from_index drop ppart_table entirely. It ran ~15x
slower than the array read it would replace, at every degree, with none of the
crossover above -- ppart_table is 8 bytes per element against ~43 for the map, so
it stays cache-resident. The bit-packing that makes rank worth having is the same
thing that makes unrank not. The likelier route, if it is ever revisited, is
enumerating the basis in index order, which is O(1) amortised and matches how
callers actually walk it.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XK8YCdHkCqUV9YM7D957hV
Three paths computed with unvalidated input before checking it against the
packing bounds, so the intermediate arithmetic went wrong first. All three are
reachable from public, non-panicking entry points.

- `basis_element_from_string("P^s_t")` indexed the xi-degree table with `t`,
  which has exactly `MAX_LEN` entries, so `t = MAX_LEN` was out of bounds. `p^s`
  and the degree product could also overflow. Now `t` is bounded by the table
  itself and both are computed with checked arithmetic.
- `try_beps_pn` computed `q * x + e` before bounding `x`, which overflows for a
  large `x`. The bound moves above the computation.
- `MilnorSubalgebra::packed_signature` assumed the profile was no longer than
  `PPart::MAX_LEN` and that each signature entry fit its field. Neither holds:
  `SubalgebraIterator` grows a profile without limit and `from_bytes` reads
  whatever length a file gives. Out of range, `PPart::shift` returns 64 and the
  shift overflowed; an oversized entry silently spilled into the neighbouring
  field, which could select unrelated basis elements. It now returns `None` for
  a signature no element can have, and `signature_mask` yields nothing.

`basis_element_from_string` is documented as total and `try_beps_pn` is the
non-panicking half of `beps_pn`, so these were contract violations rather than
merely untidy. Tests cover each.

The signature test checks the packed mask against the per-entry comparison it
replaced, over every element up to degree 60, for profiles that are narrower
than their fields, wider than their fields, and longer than a p-part can be.

Also adds `basis_order_at_p2_is_stable`, which pins the first nine degrees to
fixed element names. The basis order is a wire format -- saved resolutions store
coefficients by index -- so it needs a guard that does not read from
`ppart_table`, which is the thing being guarded. Verified separately that the
order is unchanged from the base commit: identical for all 4156 elements in
degrees 0..=60.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XK8YCdHkCqUV9YM7D957hV
CI lints with the nightly toolchain, where the unstable options in
`rustfmt.toml` -- `reorder_impl_items` among them -- actually take effect.
Stable rustfmt skips them with a warning, so `cargo fmt --check` passed locally
and failed in CI.

Formatting only: the constants are sorted and the blank lines between them
dropped.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XK8YCdHkCqUV9YM7D957hV
The prior guard bounded the intra-row limb (out_offset + seqno spanning past
the row) but not the row itself. compute-sanitizer on a malloc_sync build at
higher degree (soak d=200) caught a second out-of-bounds atomic on the same
out[word]: when row_base overruns, word = row_base + limb lands past the
buffer even with a valid limb. Add the complementary `word < out.len()`
bound so no atomic write can escape the allocation by either route.

Validated at the d=200/48-stream soak: 4/4 clean, 0 crashes, 0 fallbacks,
0 correctness mismatches.

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
…s time

The CPU fallback turned a hard GPU fault into a silent ~100x slowdown: a run
that had lost the GPU still reported "completed", so every measurement had to be
reconstructed by grepping stderr, and a crash at 5 seconds looked like a slow
success three hours later. GPU_DISABLED is still latched first so the soak test
can tell a context death from an ordinary assertion failure, but the panic now
propagates and the run dies at the fault.

The batch counters existed but take_batch_stats() had no callers, so the
dominant cost of a resolution was never attributed. Reporting them periodically
shows where multiply time actually goes, and the split matters because the
naive reading is wrong in two ways:

- The window called "marshal" spans GpuPermit::acquire and the gpu_lock
  acquisition, so it conflates host work with time parked on our own gates.
  Separating them shows the permit costs nothing measurable.
- The shares move by 2-3x as a run matures — early samples are small batches
  where fixed overhead dominates. At 2k launches it reads 29% prep / 47% lock /
  24% device; by 226k it is 14% / 9% / 77%. Only the mature numbers mean
  anything, and they say the multiply kernel itself is the bottleneck.

Reporting keys off the value fetch_add returns rather than a separate load: with
~100 workers a load races past exact multiples and the report can fire never
(observed: zero output over 12 minutes).

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
The per-bidegree signature loop is the tail's critical path: at stem 200 the
bidegrees use A(3), whose signature space is 1024, and the steps run
sequentially on one thread — b=(200,6) spends 1076 s that way, 1024 steps at
~551 ms median plus one outlier at 470 s. With only 3-5 bidegrees in flight at
that point in the wavefront, parallelising this loop looked like the largest
remaining win.

It is not available. Each step reads dx.entry(v) over its own signature mask and
then writes dx with rows of the *unmasked* matrix, whose support extends past
that mask. NASSAU_PROBE_SIG_INDEP=1 snapshots dx before the loop and compares
every read against it: 6703 of 38348 reads (17.5%) differ, i.e. were perturbed
by an earlier signature, starting from bidegrees as small as (14,2). At p=2 a
differing entry flips the zero test that drives the step, so solving the
signatures against the pre-loop dx and combining would compute different
answers. The loop is a genuine forward substitution and must stay ordered.

Kept as a probe (zero cost unless the variable is set) because the alternative
is rediscovering this by writing the parallel version and getting wrong answers.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Formatting only, no behaviour change; `just lint` runs `cargo fmt --all --check`.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
JoeyBF and others added 30 commits August 9, 2026 02:02
…ayon pool

Every profile in this file was taken either at theta=125 (a testbed handicap: 2.85x
slower than uncapped, and it inverts the bottleneck) or before the readback / `Arc`
term-list / `R`-memo fixes. This records one taken at theta uncapped, stem 150, after
those — i.e. the configuration we would actually ship, the 130 s run.

    multiply_batch_grouped (+closure) 17.4%    allocator                 9.7%
    memcpy + memset                   13.8%    add_masked                7.1%
    Matrix::row_reduce                11.3%    crossbeam_epoch (rayon)   5.8%
    get_partial_matrix_restricted      9.3%    resident_info             2.6%

The binary is 64.5% of cycles and `libcuda` does not reach the top 16. At correct theta
this is a CPU-bound workload that barely touches the GPU — the 1.56% occupancy and
52.67% spin-wait numbers describe the handicapped regime only. `resident_info` fell
12.93% -> 2.56%, which is the memo working as intended.

RULED OUT: shrinking the rayon pool. `crossbeam_epoch` at 5.8% is rayon's work-stealing
reclamation (confirmed by call graph — rayon_core worker threads, not our channel), and
the pool is 128 threads for a run using ~9 cores' worth of work, so idle stealers
churning epoch state looked like free money. Interleaved, 3 rounds:

    RAYON_NUM_THREADS default(128)  137, 125, 130   mean 130.7 s
    RAYON_NUM_THREADS 32            144, 144, 139   mean 142.3 s
    RAYON_NUM_THREADS 16            140, 159, 129   mean 142.7 s

9% WORSE. The overhead rides on idle workers and costs no wall time, while the wide pool
is buying parallelism elsewhere. Same lesson as R-affine sharding and the block budget:
reducing work that is not on the critical path buys nothing here.

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

`NASSAU_SPLIT_VERIFY` builds the reuse matrix over all rows and again over two row
ranges, and asserts they agree row by row.

This is the load-bearing claim for precomputing matrix rows AHEAD of a bidegree, which
is the remaining lever for capped-theta high stems: the rows coming from generators
that already exist can be computed early and concatenated with the rest. The argument
is that rows are generator-major and `FreeModule`'s basis tables are append-only
(`OnceVec`), so widening a restriction only appends and an early row block stays valid
verbatim — and a product `Sq(R)·gen` lands inside `gen`'s own block, so an old
generator can never write into a newly added column.

That is an argument from data-structure semantics. This checks it against real data:

    S_2 stem 60 / max_s 40, theta uncapped   0 mismatches
    S_2 stem 60 / max_s 40, theta=60         0 mismatches   (transient path exercised)

So the decomposition holds in both regimes, including the capped one the design is for.

WHY THIS MATTERS, recorded because it is the one lever left. At capped theta the enum
kernel runs at 1.56% occupancy in launches of 3-104 blocks, and the reason is SUPPLY:
demand-driven work exposes only ~5-10k independent `R`s because the wavefront is ~5
bidegrees wide. Precomputing blocks ahead in `t` makes that a tuning knob instead of a
ceiling — the `R` set in flight becomes "however far ahead we choose to run" — which is
what turns 3-block launches into device-filling ones. Blocks are also independent, so
the same work can spill to the ~119 idle cores via the CPU path, which needs no master
at all.

What remains unbuilt is the scheduling layer: a cache keyed by bidegree, a background
filler that computes blocks for future `t` from already-known generators, and assembly
of (cached prefix ++ fresh suffix) in `step_resolution`. The primitive it needs already
exists — `restricted_partial_matrix_maybe_gpu` takes an arbitrary row list — and is now
verified.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Builds a bidegree's full restricted differential matrix on background threads
BEFORE the bidegree runs, and hands it over when it starts.

The soundness argument was already in the code and is now used. `step_resolution_with_subalgebra`
deliberately ignores the target's degree-`t` generators, so its full matrix reads only data frozen
once `(s-1, t-1)` is committed -- strictly weaker than the condition for RUNNING `(s, t)`, which
also needs `(s, t-1)`. Everything between the two bounds is a bidegree whose matrix is fully
determined but which cannot run yet: the speculation window. `restricted_dims` makes the point
concrete -- the restriction bounds are `b.t()` and `b.t()-1`, properties of the bidegree, so no
subalgebra is needed and the matrix predates the choice of signature.

The window is widest where it pays. The bottleneck is the low-`s` rows, and for those the rows
below finished long ago, so lookahead is bounded by `NASSAU_SPECULATE_AHEAD`, not the wavefront.

Pieces:
  - `speculate`: cache keyed by bidegree, byte-capped, plus a build queue ordered by ascending
    `(t, s)`. The order matters -- a matrix that lands after its bidegree started is pure waste, so
    builders must always take what the wavefront reaches soonest; FIFO would drain the deepest
    speculation first.
  - builders run OUTSIDE the rayon pool, in a `std::thread::scope`. Their whole point is to use
    time the wavefront is not using; pool workers would take slots from the critical path.
  - `enqueue_spec` in the scheduler emits the newly opened window on each commit.
  - consumer takes from the cache, and on a shape mismatch warns and rebuilds rather than trusting
    it.

Off by default (`NASSAU_SPECULATE=0`). Verified with `NASSAU_SPECULATE_VERIFY`, which rebuilds each
cached matrix at consumption time and asserts equality: stem 40 / max_s 30, 318 hits, 0 mismatches,
alongside `NASSAU_GPU_VERIFY`.

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

Two bugs, found by measurement then by a hang.

DUPLICATION. The cache held only finished matrices, so a builder and the bidegree itself could
build the same matrix concurrently. A full build is most of a bidegree's cost, so that duplication
WAS the cost of speculation: stem 150 / theta=125 gave 442 built / 99 consumed / +18% wall at 2
threads, degrading monotonically to +44% at 8. `Slot::{InFlight,Ready,Taken}` gives every bidegree
exactly one owner: builders `claim` before building, and the consumer either takes the matrix,
waits for an in-flight build, or claims the slot so no builder starts work it is already doing.

DEADLOCK. Waiting unbounded then hung the run at stem 40, immediately. Builders call
`restricted_partial_matrix`, which is rayon-parallel, from OUTSIDE the pool: rayon injects the work
and blocks the builder until a pool worker runs it. A consumer is a pool worker, and one parked in
a condvar cannot steal -- so the wait closed a cycle. gdb showed it exactly: builders in
`in_worker_cold` -> `LockLatch::wait_and_reset`, workers parked in `take_or_claim`. The wait is now
bounded (`NASSAU_SPECULATE_WAIT_MS`, default 5s); on timeout the consumer claims the slot and
builds its own, and `publish` drops a matrix whose slot is no longer `InFlight` rather than parking
one nobody will collect.

Verified (`NASSAU_SPECULATE_VERIFY` + `NASSAU_GPU_VERIFY`, stem 40 / max_s 30): 0 mismatches,
built=411 hits=411 -- every published matrix consumed, wasted work 343 -> 9.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Sampling a stem-200 run settles where the headroom is: CPU sits at 700-900% of a possible 12800%
(~120 of 128 cores idle) while GPU 0 bursts to 100% and GPUs 1-3 idle. The run is GPU-bound with the
CPU nearly free.

That explains the stem-200 A/B: with duplication eliminated, speculation still measured only neutral
(1745 s -> 1797 s, +3%). Speculative builds went to the GPU, so they queued behind the very critical
path they were meant to relieve -- moving GPU work earlier cannot help when the GPU is the
constraint. `NASSAU_SPECULATE_CPU` sends them to `restricted_partial_matrix` instead, where a hit
costs the critical path nothing AND removes a launch from the queue it is waiting on.

Also records the regime finding that governs this whole lever: at stem 150 speculation is +14% even
with perfect dedup (built=2781, hits=2781), and at stem 200 it is +3% -- the penalty shrinks as the
wavefront narrows, which is the direction that makes larger stems the place to judge it.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Two things were wrong with the warmer, both invisible because it runs off the critical path.

DOUBLE ENUMERATION. `cold_count` runs `admissible_matrices` purely to read `num_mats`, discards the
arrays, and then `resident_info` ran the same enumeration AGAIN for every `R` that passed the
threshold. Every pinned `R` was enumerated twice. It now enumerates once, memoises the count, and
appends the arrays already in hand.

SERIAL. One thread warmed one `R` at a time while sampling a stem-200 run shows ~120 of 128 cores
idle and the GPU as the constraint. The elements of a degree are independent -- warming an `R` is a
pure function of its p-part, and the only shared state is the master's write lock, taken once per
STORED `R` -- so the degree now splits across `NASSAU_GPU_PREFETCH_THREADS` (default 16). Plain
`std::thread`, never rayon: this must not inject into the pool the wavefront uses, which is the
deadlock the speculative builders hit.

`resident_append` is factored out of `resident_info` for this, and is the same seam the batched GPU
enumeration needs: one place assigns offsets, picks the shard and mutates the master, whether the
enumeration came from a serial CPU call, many cores at once, or one big device launch.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Saturates the enum kernel by giving it work that is not tied to a multiply.

A production enum launch carries only the R's of the ONE multiply that needs them -- Rs/launch mean
1293, waves/SM 0.013 -- and the GPU submission queue is ~1.5 deep, so there is essentially never a
peer launch to merge with. That is why the two obvious fixes measured nothing: widening the grid 6.5x
bought 3.4%, and two CUDA streams bought 527s vs 528s. There was nothing to overlap.

The warmer has no such constraint. Nothing downstream waits on it, so it can hand the kernel a whole
degree at once. `prefetch_degree_gpu` groups a degree's R's by owning device and runs the shards
concurrently -- which also puts the three idle GPUs to work (sampling shows GPU 0 near 100% and
GPUs 1-3 near zero).

Two passes, and no host enumeration anywhere:
  - `enumerate_counts_gpu` runs with `emit = 0`, which compiles the stores out and leaves the
    odometer alone, filling `out_counts`. So `num_mats` -- the number that decides whether an R is
    worth keeping -- now comes from the device. Previously the threshold test alone cost a full CPU
    enumeration of every R in the degree, kept or not. Counts are memoised into COLD_COUNT for every
    R seen so the multiply path inherits them.
  - `enumerate_batch_gpu` emits the arrays for the keepers, chunked to bound device scratch
    (`NASSAU_GPU_ENUM_BATCH_MB`, default 2048); a degree's full output reaches many GB at high stems.

Off by default (`NASSAU_GPU_PREFETCH_GPU`). `NASSAU_GPU_ENUM_BATCH_VERIFY` checks every batch against
`admissible_matrices`.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Parallelising it is 14% slower and warms nothing extra. Measured at stem 150 / theta=125, PIN=200,
AHEAD=25, interleaved and counterbalanced: 1 thread 345 s / 336 s, 16 threads 396 s / 382 s -- and
all four runs reached the SAME frontier (degree 176) with the SAME enumeration load on the critical
path (Rs/launch 531).

The premise was wrong, not the implementation. The warmer is limited by `prefetch_ahead`, not by its
own throughput: it fills the window and sleeps. Extra threads cannot warm anything further ahead, so
they only contend with the wavefront. The idle-core observation that motivated it is real but
irrelevant here -- idle cores only help if the thing you parallelise is the bottleneck.

The knob stays, with the numbers recorded next to it and a note that it is only worth raising
together with `NASSAU_GPU_PREFETCH_AHEAD`, and only when `to_degree` shows the warmer falling behind.
The single-enumeration fix from the previous commit is unaffected and stays.

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

`Rs/launch` sat at exactly 531 across every warmer configuration tried -- AHEAD=25 serial, AHEAD=100
serial (warmed to degree 186), AHEAD=100 with 16 threads (degree 211, 579k R's, twice the warming),
and batched device enumeration. Warming further AHEAD never moved it, because the leftover R's were
never ones the warmer failed to reach in time: they are the ones `pin_min_mats` deliberately skips as
cheap.

Warming further DOWN the cost distribution moves it 4.5x, and that is worth real time (stem 150,
theta=125, interleaved, counterbalanced):

  pin 200 -> 418 s / 367 s, Rs/launch 531
  pin  20 -> 294 s / 317 s, Rs/launch 117
  pin   2 -> 304 s / 350 s, Rs/launch 4, a third fewer launches

20 beats 200 by 22% and beats running with no warmer at all (~384 s) by 20%. It turns over below
that: 2 drives enumeration to essentially nothing and still loses, because warming every trivial R
costs more than enumerating it.

Capped theta only -- with theta uncapped everything is resident and the knob does nothing -- which is
exactly the regime stems 250+ are forced into.

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

`NASSAU_GPU_ENUM_BATCH_MB` defaulted to 2048, which silently defeated the whole point. A pinned R
(num_mats >= 200) is often megabytes of col_sums/masks, so a 2 GB emit budget fits only a few hundred
of them. Measured consequence: the batched path ran at ~947 R's per launch -- BELOW the 531-mean
multiply-driven launches it was supposed to dwarf -- and `Rs/launch max` came out IDENTICAL (13650)
in the batched and un-batched arms, which is proof no batch was ever large. At ~30 blocks of the 3168
an H200 holds, waves/SM stayed at 0.005 exactly as before.

So the earlier "batched enumeration is 18.5% slower" result does not test batching. It compares two
small-launch configurations, one of which also pays device contention. Budget default is now 16384 MB.

Batch sizes are now LOGGED rather than inferred: `batches= Rs/batch= max=` in the [batch-stats] line,
from counters recorded at both launch sites. `Rs/launch` is dominated by the thousands of
multiply-driven launches, which is what hid this.

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

Both results were implementation artefacts, not findings. The argument they appeared to contradict --
free cores mean a bidegree does strictly less work -- is sound; the code was not using free cores.

SPECULATIVE CPU BUILDS RAN IN THE WAVEFRONT'S OWN POOL. `restricted_partial_matrix` is rayon-parallel,
so each of 16 builders injected a large parallel job into the GLOBAL pool. Those cores are not free
in that case: wavefront tasks queue behind speculative chunks and work-stealing spreads it. Measured
at stem 150 / theta=125 / pin 20: 649 s against a 310 s baseline, a 2.1x regression, while the cache
itself worked perfectly (53% hit rate, zero duplication, zero waste) -- exactly the signature of
contention rather than a bad idea. Builds now `install` into a dedicated `MaybeThreadPool`
(`NASSAU_SPECULATE_POOL`, default a quarter of the machine), which confines every nested `par_iter`
to that pool.

BATCHED ENUMERATION NEVER GOT BIG BATCHES. It swept ONE degree per launch, and a degree yields only a
few thousand R's -- split four ways, ~740 per shard, about 23 blocks of the 3168 an H200 holds. That
is the same starvation the multiply-driven launches suffer, which is why `Rs/launch max` came out
identical (2965) in the batched and un-batched arms and waves/SM never moved. It now sweeps the whole
lookahead window (`prefetch_degrees_gpu`) into one batch.

Also: the batch-size fields were computed but never printed -- the format-string edit missed a
`waves/SM=` prefix -- so batch sizes had to be inferred from `Rs/launch max`. Now printed directly as
`batches= Rs/batch= max_batch=`.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
`speculate` caches one whole matrix per bidegree, which caps how early a
builder can start: `(s, t)`'s matrix is only determined once `(s - 1, t - 1)`
is committed, giving one degree of slack per row and a ~30% hit rate.

A block is the rows coming from generators of `modules[s - 1]` in a single
degree -- the rows `Sq(R) * x` with `deg Sq(R) + deg x = t`. Two facts make a
block available far earlier than the matrix containing it:

* Its rows. A free module's basis at a fixed degree is generator-major and
  append-only, so the block's row range is frozen once `(s - 1, gen_deg)` is
  committed, not `(s - 1, t - 1)`.
* Its columns. `d(x)` lands in the radical (minimality -- already what
  licenses the existing truncation in `apply_to_basis_element_restricted`), so
  the block is supported strictly below generator degree `gen_deg` and can be
  built that narrow, then zero-extended.

So one commit opens blocks across a whole column of future bidegrees instead of
a single matrix, which is what deepens the GPU submission queue. The consumer
takes whatever landed and folds the rest into ONE coalesced build, so a miss
costs only its own rows -- no waiting, and none of `speculate`'s deadlock
hazard.

Verified at stem 30/max_s 25 under NASSAU_SPECULATE_VERIFY + NASSAU_GPU_VERIFY:
0 mismatches, 93.0% of rows served from blocks.

Two bugs the verifier caught, both fixed here:

* The consumer re-derived each block's row range at consumption time and
  disagreed with the builder (start=23 for a block whose rows sat at 22),
  misplacing every row. Blocks now carry the offset they were built against;
  the consumer never re-derives it.
* Blocks are validated for overlap and bounds before use rather than trusted,
  so a stale one degrades to a rebuild instead of corrupting the matrix.

`NASSAU_BLOCK_DIAG` keeps the attribution probe that separated the row-range
bug from the column bound (which measured clean: beyond_narrow=0 throughout).

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
At stem 200 blocks LOST (1794s vs 1728s baseline) while holding 264GB against
116GB, with row_rate=16.1% and unused=150.3GB -- the cache ended pinned at its
cap.

The cause is a gap between producer and consumer. The consumer only assembles
from blocks when the matrix is within `reuse_within_cap`; above it, it builds
per signature and never calls `take_all`. But `speculate_block` did not check
that cap (whole-matrix `speculate_build` does), so builders precomputed blocks
for bidegrees that would never collect them: wasted CPU, and the blocks stayed
in the cache until the run ended. Through `has_room()` that also starved the
bidegrees that DO collect, which is the 16.1% coverage.

Two fixes:

* `speculate_block` bails on `!reuse_within_cap`. The dims are read before the
  bidegree's frontier is frozen, so they are a lower bound -- sound in one
  direction: already over the cap now means over it for good.
* The consumer releases blocks for any bidegree taking the per-signature path,
  so whatever slips through the lower-bound test cannot leak. Counted as
  `released=` rather than inferred.

The cap check reuses the single `block_ranges` pass instead of calling
`restricted_dims`, which repeated the same `compute_basis` work per block and
cost more than speculation gained (row coverage 93% -> 55% at stem 30).

Hoisting `FreeModule::compute_basis` over the whole degree range would remove
that cost entirely but is NOT valid: a module computed through max.t() up front
can no longer have generators added back. Only the algebra is precomputed.

Verified at stem 30/max_s 25 under NASSAU_SPECULATE_VERIFY + NASSAU_GPU_VERIFY:
0 mismatches, 84.1% of rows served from blocks.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
`row_rate` only counts bidegrees that take the full-matrix reuse path, so it
cannot distinguish "speculation covered most of the work" from "speculation
covered most of the little work it was allowed to see". `[reuse-split]` splits
total multiply work (rows x cols) by whether the bidegree was within
`reuse_within_cap`.

The answer explains why block speculation won at stem 150 and did nothing at
stem 200:

  stem 150:  100.0% reachable, 0 bidegrees over cap
  stem 200:   13.1% reachable, 941 bidegrees holding 86.9% of the work

At stem 200 the 941 largest bidegrees each exceed the 1e10 cap, so they never
take the full-matrix path -- no blocks, and no one-big-launch GPU amortization
either; they fall back to per-signature builds. Blocks were competing over the
remaining 13%, which is why 60% row coverage bought no wall time (1846s vs
1801s uncapped baseline).

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Signature-axis speculation is not viable: a signature mask needs the module
basis at degree t, which needs progress[s-1] >= t-1 -- exactly the condition
for the bidegree to be runnable. Zero lookahead, queue depth ~1, which is the
starvation this design exists to escape. The generator axis is the only one
with a window, since block `g` needs only progress >= g.

So the question for the 87% of stem-200 work above `reuse_within_cap` is
whether generator-blocks can serve it without materialising the full matrix.
There is an asymmetry to exploit: blocks are stored NARROW (block `g` stops
below generator degree `g` by minimality) while the full matrix pads every row
to `next_dim`, so the block set is triangular where the matrix is rectangular.
If a signature's rows were selected straight from the blocks, peak would be
(block set + one signature) rather than (full matrix + one signature).

`[reuse-split]` now reports that ratio. Near 0.5 the idea is worth building;
near 1.0 the triangle is too shallow and relaxing the cap is the only lever.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
At stem 200, 941 bidegrees hold 86.9% of all multiply work and every one
exceeds `reuse_within_cap`, so they never assemble a full matrix and plain
blocks cannot reach them: block delivery needs every block resident at once,
which is the memory bound the cap exists to hold.

Signatures give a second axis. `signature_mask` selects, within a generator
degree, the ops matching the packed signature from `ppart_table(t - gen_deg)`
alone -- never the module's degree-`t` basis -- so a `(gen_deg, signature)`
piece is determined exactly when the generators of `gen_deg` are, keeping the
early window that makes speculation work at all. And because `signature_mask`
walks generators in layout order, a piece is a CONTIGUOUS run of the signature's
matrix, so assembly is row concatenation into a matrix preallocated at
`next_dim`: matrices are row-major, and widening one afterwards would restride
every row.

The consumer takes one signature at a time and discards it, so peak stays at one
signature however large the bidegree is. No extra multiplies: signatures
partition the operations, so this is a strict refinement of the same rows.

The builder now dispatches on the cap -- contiguous per-generator pieces below
it, signature pieces above it -- and `release` no longer marks a bidegree done,
which would have silenced signature delivery for exactly the bidegrees this
targets.

Also fixes a coverage bug: `iter_gen_offsets` yields one entry per GENERATOR,
not per degree, and the old code took only the first entry per degree, so every
other generator of a degree was silently left to the critical-path rebuild.

Verified at stem 30 with `NASSAU_GPU_REUSE_MAX_WORK=1000` forcing 92.2% of work
over the cap: 0 mismatches against fresh per-signature rebuilds, 82.0% of rows
served from pieces.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Covering every generator of a degree sounded like a free coverage win. It is
not: stem 150, theta=125, one binary, interleaved --

    first generator only    173 s   coverage 53.4%
    merged over the degree  194 s   coverage 53.3%
    one piece per generator 202 s   coverage 51.4%

Coverage is flat because a degree almost always carries one generator here, so
the extra generators are nearly empty. What changes is piece size and builder
load, and both bigger pieces and more pieces lose. Speculation is already at
its useful limit at stem 150; past it the builders just compete with the
wavefront.

The bisect also clears the signature scaffolding: with merging off, the new
binary matches the old one exactly (173 s), so none of the cost was the
(gen_deg, signature) path.

NASSAU_BLOCK_MERGE=1 restores full-degree coverage for regimes where generators
are denser.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
The first cut of signature speculation ran 5077 s at stem 200 against 3799 s for
plain blocks, and 4530 s for no speculation at all. Two defects:

* `signature_piece` rescanned every generator and its whole `ppart_table` once
  PER SIGNATURE, so a degree cost O(generators x table x signatures). Every
  operation has exactly one signature, and all signatures of a subalgebra share
  the same packed mask (it depends only on the profile), so one pass bucketing
  ops by `op.bits() & mask` does the whole degree for what a single signature
  used to cost.
* It built 2 988 096 pieces averaging 32 rows each -- far too small to repay
  allocation, locking and bookkeeping. `NASSAU_SIG_MIN_ROWS` (default 512) drops
  them. Signature mass is very uneven (one signature is often ~96% of its
  bidegree), so a threshold discards most pieces while keeping most rows.

Re-verified at stem 30 with `NASSAU_GPU_REUSE_MAX_WORK=1000` forcing 92% of work
over the cap and no threshold: 0 mismatches, 78.3% of rows served.

A huge `NASSAU_SIG_MIN_ROWS` disables the signature path exactly, so plain
blocks and signature pieces can be A/B'd within one binary.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
With no pieces available, `assemble_signature` still allocated the result
matrix, built a SECOND matrix for the missing rows, and copied row by row --
where the old code built once and returned it. That is pure overhead on every
signature speculation did not reach, and it cost 7% at stem 200: the control arm
(signature path disabled) ran 4082 s against the 3799 s the same configuration
measured before this path existed.

Return the direct build when nothing landed. Still counted as missed rows so
`row_rate` stays honest.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Stem 200, theta=125, counterbalanced within one binary (a huge
NASSAU_SIG_MIN_ROWS disables the path exactly, so both arms share a build):

    signature path off   4170 s (cold, ran first) / 3826 s
    min_rows=512         3786 s / 3775 s

Round 1 looked like a 9% win. It was not: the first stem-200 run after an idle
machine is reliably ~9% slow, and the control's round-2 number lands at 3826 s
against the treatment's 3781 s -- about 1%, i.e. neutral. `built`
(205 815 -> 207 904) and `row_rate` (28.1% -> 28.3%) say why: at 512 rows the
threshold makes the path inert.

So the affordable threshold is the inert one, and the range between "inert" and
the 5077 s of min_rows=1 is untested -- that measurement predates the
single-pass rewrite, which made low thresholds far cheaper. Sweep pending.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
The speculation pool defaulted to max_num_threads()/4, added as hygiene against
speculative par_iters stealing workers from the wavefront -- a fear that traced
back to a broken binary, not to measured contention. It capped speculation by
construction: ~20 of 128 cores in use, and block coverage plateaued near 52%
however the pieces were reshaped. Every granularity experiment (bigger blocks,
per-generator, per-signature) changed the SHAPE of speculative work while its
CAPACITY stayed pinned, which is why none of them moved the needle.

Stem 150, θ=125, control repeated last:

    builders / pool   wall            row coverage
    16 / 32           170 s  168 s    52.2%  50.6%
    24 / 64           134 s           63.6%
    32 / 96           119 s           70.8%
    48 / 120          134 s           77.0%

So the contention is real but only past ~3/4 of the machine: at 120 threads
coverage still climbs while wall time regresses. Default is now 3/4.

That is -30% on top of the win speculation already had, and -63% against the
320 s no-speculation baseline at this bidegree.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Blocks do what they were meant to at stem 200 -- bidegrees retire faster and the
run goes 4465 s -> 3708 s (-17%). But row coverage sits at 33% against 71% at
stem 150, so two thirds of the rows are still multiplied on the critical path,
and total CPU is ~13 of 128 cores. With 96 speculation threads available that
means the builders are IDLE rather than contended: they are bailing, and nothing
recorded which bail-out they take.

Two counters:

* `bails(room/done/claimed/other)` on the [blocks] line. The leading suspect is
  `has_room()` against NASSAU_SPECULATE_MAX_GB=150 while peak RSS is already
  196 GB, on matrices ~15x larger than stem 150 -- once the cache sits at its cap
  speculation simply stops, which no pool or granularity change can move.
* `[wavefront] in-flight bidegrees mean/max`, to separate "too few bidegrees are
  eligible" from "eligible ones are not being run". Speculation shortens a
  bidegree but does not change admission -- `(s,t)` still needs `(s,t-1)` and
  `(s-1,t-1)` -- so the frontier stays a slope-1 staircase over `s`.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
At stem 200, 718 171 queue items were popped and only 259 150 built -- 456 767
bailed because the bidegree was ALREADY COMPUTED. Two explanations were wrong:

* Not the memory cap. Raising NASSAU_SPECULATE_MAX_GB 150 -> 400 left coverage
  at exactly 33.2% and reported bails(room)=0. The cap was never reached.
* Not time lost discarding stale items. A `done` bail is a mutex and a bool;
  456k of them across 3935 s is nothing.

And the builders were not saturated either -- 1310% CPU of 128 cores with a
96-thread pool -- so they were neither too slow nor too busy. A block queued 20
internal degrees ahead should comfortably finish first, so something else is
wrong.

`[specqueue]` reports queue depth sampled at every successful pop, plus builder
seconds split into idle-on-empty-queue versus building. Near-zero depth with
large idle means the producer's gate (gen_deg <= min(progress[r-1],
progress[r-2])) admits far less than the nominal 24-degree window suggests, and
the lookahead is illusory. Large depth with large idle would instead mean
builders block on something CPU accounting does not show.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
Two changes to speculation queue discipline, both from the stem-200 numbers:
718 171 items popped, 259 150 built (36%), 33% of rows covered.

* Built fraction ~= coverage fraction means selection was size-NEUTRAL. Since
  only a fraction of the queue is ever built, taking the biggest first converts
  the same builder time into more coverage. So `rows` is now the secondary key.

* It is only the SECONDARY key. Blocks have deadlines -- worthless once their
  bidegree runs, equally useful any time before -- and the queue never drains,
  so near blocks expire while far ones stay valid. Imminence-first is therefore
  earliest-deadline-first, the optimal policy for meeting deadlines. Ordering by
  size INSTEAD was tried and is wrong in principle; it lets builders work far
  ahead while imminent blocks expire.

* `NASSAU_BLOCK_MIN_ROWS` (default 1024) drops tiny blocks entirely. A block that
  is not precomputed costs the consumer nothing extra -- its rows join the single
  coalesced build it was already making -- so a small block has near-zero value
  while still costing a queue slot, claim, lock, allocation and publish. With
  718k items the tail was most of the bookkeeping and almost none of the benefit.

Row count is `num_gens[g] x |ppart_table(t-g)|`, both O(1), so sizing every
candidate at enqueue costs nothing.

Stem 30 cannot discriminate the two orderings (46.8% vs 46.9% coverage) -- it is
too fast for speculation to matter. Needs a stem-150 A/B.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
`[specqueue]` reported mean depth 133 589 (max 343 227) while `builder_idle`
showed builders parked for 91% of their thread-time. Both were true and they
describe different moments: depth was sampled only at a successful POP, and a
pop only happens when there is work, so the statistic could only ever report
"deep". The producer emits a whole strip per commit, floods the queue, and
builders drain it over minutes; between bursts the queue is empty and nothing
was observing it.

A sampler thread now reads depth every 100 ms, giving a time-weighted mean and
the fraction of samples with an EMPTY queue. Both figures are printed alongside
the at-pop ones so the bias stays visible rather than being silently corrected.

The arithmetic this has to explain: ~718k items produced, ~29 ms per build, 32
builders -> ~650 s of building spread over a 3685 s run, against 481 535 bails
on already-computed bidegrees. That is consistent with bursty production the
builders cannot drain before the wavefront passes, which is what makes queue
ORDER (earliest-deadline-first, committed separately) the thing that matters.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
A mean hides the shape, and the shape is the question. Enqueues are triggered by
COMMITS, and commits get sparser as bidegrees get larger, so the queue is
expected to be front-loaded and to starve later -- exactly when the expensive
bidegrees need it most. Averaging that away is how "mean depth 133 589" and
"builders parked 91% of their thread-time" ended up in the same run.

`[qtrace]` emits every 2 s: elapsed, depth, total ROWS queued (depth alone says
nothing about whether the work is worth anything), tmin/tmax of the queued
deadlines (near work or far), and how many builders are occupied at that instant.

Together with the 100 ms time-weighted sampler this gives the full picture:
whether the queue starves, when, and what is left in it when it does.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
The first sampler took the queue mutex and iterated the heap. At stem 200 that
emitted NOTHING: with ~343k items and 32 builders contending, a sampler never
wins the lock, and iterating under it would block every builder. Depth and
queued rows are now atomics maintained on push/pop, so sampling is two loads --
which also makes 60 Hz free, and resolution matters because the burst that
revealed the behaviour lasted under 2 s.

What it shows at stem 200, clean and uncontended:

    t=61s   depth=80283  rows=48.4M  building=32  built=76400
    t=182s  depth=30987              building=32  built=139891
    t=242s  depth=0                  building=7   built=176814
    t=423s  depth=0                  building=0   built=196574

One burst saturates all 32 builders for four minutes, drains, and never refills;
`built` then creeps at ~5k/min. Enqueues are triggered by COMMITS, and once the
wavefront is grinding through a few large bidegrees the commits are sparse, so
production falls to a trickle builders consume instantly. Builders are starved
of eligible work for ~90% of the run -- which is the 6.5% busy / 33% coverage at
stem 200, against 70% busy / 70% coverage at stem 150.

Also retracted here: an earlier "speculation freezes after 6 s" reading came from
three runs racing on one machine after pkills silently failed. Clean stem-150
data shows `built` climbing steadily throughout.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
A block's lookahead is ~`t - gen_deg`, the OPERATION degree: it becomes eligible
when progress[s-1] reaches gen_deg, and its bidegree runs when progress[s-1]
reaches t-1. So blocks with gen_deg near t get almost no warning. If the ~33%
coverage ceiling at stem 200 is that rather than a tunable, misses must
concentrate in the LOW operation-degree buckets.

`NASSAU_OPDEG_HIST=1` buckets served vs rebuilt rows by operation degree.
Stem 150 confirms the mechanism -- coverage rises monotonically with lookahead:

    op_deg   0-24   32.3%
    op_deg  25-49   44.6%
    op_deg  50-74   59.5%
    op_deg  75-99   76.1%
    op_deg 100-124  77.9%
    op_deg 150-174 100.0%

The remaining question is whether stem 200's row mass has shifted into the
short-lookahead buckets, which would make the ceiling structural rather than
something capacity, memory, granularity or window size can move -- all four of
which have now been measured and rejected.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
`NASSAU_OPDEG_HIST=1` buckets served vs rebuilt rows by operation degree (a
block's lookahead). Coverage rises with lookahead at both stems, but the shape
is the same and the row mass has NOT shifted toward short lookahead:

    op degree      stem 150   stem 200
    0-24              32.3%      28.6%
    50-74             59.5%      49.9%
    100-124           77.9%      75.0%
    overall           68.7%      61.6%

The histogram runs only inside `assemble_full_restricted`, i.e. the UNDER-cap
path -- so blocks work about as well at stem 200 as at stem 150 where they
apply. The reported row_rate=25.9% is just 0.13 x 0.62 + 0.87 x 0: 87% of
stem-200 work sits in over-cap bidegrees that never assemble a full matrix and
get no coverage at all.

This corrects the previous commit's hypothesis, which guessed the ceiling was
structural lookahead. It is not; it is the reuse cap. That reopens
`(gen_deg, signature)` pieces, which exist precisely to reach over-cap work and
were called closed on a measurement using min_rows=512 -- a threshold that made
the path inert (built 205 815 -> 207 904). The range below it was never tested
after the single-pass rewrite made low thresholds ~46x cheaper.

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

"Signature pieces are neutral" was an artefact of NASSAU_SIG_MIN_ROWS=512, which
made the path inert (built moved 205 815 -> 207 904). That threshold was added
when pieces cost 5077 s and 2 988 096 of them -- BEFORE the single-pass rewrite
made a whole degree's pieces cost what one signature used to. Carrying it
forward after the fix, and then reading "inert" as "does not help", was the
error.

Stem 200, theta=125, control repeated last:

    sig_min_rows   wall     built       coverage   builder idle
    off            3836 s   246 974     32.5%      112 011 s
    128            4044 s   312 795     36.1%      118 424 s
    32             3329 s   1 278 972   43.7%        5 771 s
    off (last)     3416 s   241 482     32.3%       98 419 s

At 32 rows the starvation that dominated every stem-200 measurement disappears:
builder idle 98 000 s -> 5 771 s, coverage 32% -> 44%, 5x the blocks built. This
is the design working as intended -- an over-cap bidegree never assembles a full
matrix, but it can still be served one signature at a time, so the reuse cap is
not a coverage ceiling after all.

Wall time is -2.5% against the trailing control, but the two control arms differ
by 12%, so no speedup is demonstrated yet; only that the mechanism engages.
Sweeping lower thresholds next.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SmhK9czwJj2dZeemDhXtBf
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