fix(math-cuda): the launch-size bound was 256x too strict, so the LogUp aux build never ran on the device - #969
Merged
Merged
Conversation
`logup::cfg`, `batch_inverse_ext3_dev` and the inverse-denominator launch refused any one-dimensional launch above u32::MAX / BLOCK_SIZE = 16,777,215 elements, returning a synthesised `CUDA_ERROR_INVALID_VALUE` before touching the driver. The hazard they guard against is real but sits 256 times higher: the launch builds its grid as `(total as u32).div_ceil(BLOCK_SIZE)`, so the cast wraps only past u32::MAX elements; the grid at u32::MAX is 16,777,216 blocks, far inside gridDim.x's 2^31 - 1; and every kernel behind the three sites (logup.cu, inverse.cu) computes its element index as the 64-bit `blockIdx.x * (uint64_t)blockDim.x + threadIdx.x`, so the index arithmetic is exact up to the bound. The old bound applied the cast's limit to the block count instead of the element count. What it cost, read from the code and confirmed by the first loud abort: every LogUp aux build with more than 16,777,215 interaction-rows never launched. At 2^22 rows that is any table with five or more interactions, so the base layer's CPU, MEMW and LOCAL_TO_GLOBAL aux builds, and the wrap's LFM_BLAKE3 at 661M, all fell to the CPU build through `.ok()?` with no message, while the pipeline was described as device-resident. The same bound sent every R3/R4 batch inverse past 16.7M elements, which is a 2^23 LDE at three evaluation points, to the host as well. The abort that exposed it was LOCAL_TO_GLOBAL at 2^22 rows with 6 interactions, 6 x 2^22 = 25,165,824, refused, and reported as a device error once the aux build's admission stopped swallowing the result. One bound, `launch_total_fits(total) = total <= u32::MAX`, with the reasoning in its doc, at all three sites. Pure tests pin the shapes that were refused (6 x 2^22, 24 x 2^22, 1,261 x 2^21) as expressible, the largest expressible launch as u32::MAX with a 2^24-block grid, and one past it as refused. Beyond u32::MAX the refusal stands and is now loud at the admitted sites.
MauroToscano
added a commit
that referenced
this pull request
Sep 8, 2026
The fused per-table task opens with the LogUp aux build and continues into the aux commit and rounds 2 to 4. The VRAM gate that admits tasks against each other sized only the second phase, so the aux build's device transient, four fingerprint buffers alive during the batch inverse (`device_set::aux_build_device_set`), was invisible to it. With the launch-size bound fixed (#969) those builds run on the card for the first time, and ten 2^22 epochs of the real block on a 32,607 MiB card peaked at 31,896 MiB with builds overlapping under the gate. At that branch alone an overlap that exceeds the card still falls back to the CPU build silently; once the build is admitted the same overlap is a loud abort, so the gate must know the term. `device_set::fused_task_peak_bytes` is now what the gate spends: the larger of the rounds' set and the aux build's transient added to the R1 resident main LDE and snapshot it runs against. The two phases do not overlap -- the build's fingerprints are freed before the aux LDE is allocated -- so the peak is the larger, not the sum. Interactions are bounded by twice the aux width (one committed pair per column plus the one or two absorbed), since the AIR trait exposes only the layout here; the build's own admission sizes the exact descriptor at the call site. A table with no aux trace runs no build and keeps no snapshot, so its peak is its rounds' set exactly. For a 2^22 table at blowup 2 the build phase overtakes the rounds' phase at roughly six aux columns and grows about 0.5 GiB per column after that, which is why the wide-LogUp tables #969 unlocked are the ones this term binds on. Pre-registered for the stacked base layer (box A, CLI, 10 epochs x 2^22): VRAM peak at most 28,600 MiB, zero aborts, wall 65 to 72 s, host peak 33.4 +/- 0.5 GiB, verify OK. The correctness gate is the verified output, not proof bytes -- the grinding nonce makes bytes non-reproducible here.
MauroToscano
added a commit
that referenced
this pull request
Sep 8, 2026
The fused per-table task opens with the LogUp aux build and continues into the aux commit and rounds 2 to 4. The VRAM gate that admits tasks against each other sized only the second phase, so the aux build's device transient, four fingerprint buffers alive during the batch inverse (`device_set::aux_build_device_set`), was invisible to it. With the launch-size bound fixed (#969) those builds run on the card for the first time, and ten 2^22 epochs of the real block on a 32,607 MiB card peaked at 31,896 MiB with builds overlapping under the gate. At that branch alone an overlap that exceeds the card still falls back to the CPU build silently; once the build is admitted the same overlap is a loud abort, so the gate must know the term. `device_set::fused_task_peak_bytes` is now what the gate spends: the larger of the rounds' set and the aux build's transient added to the R1 resident main LDE and snapshot it runs against. The two phases do not overlap -- the build's fingerprints are freed before the aux LDE is allocated -- so the peak is the larger, not the sum. Interactions are bounded by twice the aux width (one committed pair per column plus the one or two absorbed), since the AIR trait exposes only the layout here; the build's own admission sizes the exact descriptor at the call site. A table with no aux trace runs no build and keeps no snapshot, so its peak is its rounds' set exactly. For a 2^22 table at blowup 2 the build phase overtakes the rounds' phase at SEVEN aux columns, and each column past that opens the margin by another 503,316,480 B (0.469 GiB) -- the build grows 805,306,368 B per aux column against the rounds' 301,989,888 B. Both sides carry the same resident main LDE and snapshot, so the crossover does not move with the main width. A test pins it, because it is what says which tables this term binds on: the wide-LogUp ones #969 unlocked. Pre-registered for the stacked base layer (box A, CLI, 10 epochs x 2^22): VRAM peak at most 28,600 MiB, zero aborts, wall 65 to 72 s, host peak 33.4 +/- 0.5 GiB, verify OK. The correctness gate is the verified output, not proof bytes -- the grinding nonce makes bytes non-reproducible here.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Correctness gate
Base layer, 10 × 2^22 epochs of block 25368371 through the CLI on a 5090: the proof host-verifies to a byte-identical public output against the pre-campaign record — both print
Output: 5d442ba039edfbb5edfb01fe024ebd6a4ca3c93b16f5448b483267dcb7d2906a8c852cfec692ac16db0935c14b963f7e3492813b4d61b62b0db1066954da0988588aandVerification succeeded!. Since the aux columns are committed, a device build that diverged from the CPU one would move the roots; they did not move.⚠ Proof BYTE COUNT is deliberately not the gate here and differs (record 907,507,216, this run 909,077,592). Two reasons: proof bytes are non-reproducible run to run in this project by construction, because the grinding nonce is drawn non-deterministically while roots are deterministic; and the record is at
8064a8efonorigin/main, so everything the campaign has merged sits between it and this branch. Output equality is the claim; byte equality would not have been evidence either way. The LogUp aux columns are committed, so if the device build had diverged from the CPU build the Merkle roots would have moved and the bytes could not match; that is a sharper argument than any parity test, and the parity tests (logup_gpuHost vs Dev vs CPU reference) pass as well.The finding
The GPU-resident LogUp aux build has never run for any table above 16,777,215 interaction·rows. At 2^22 rows that is any table with five or more bus interactions — CPU, MEMW, LOCAL_TO_GLOBAL among them — and the wrap's LFM_BLAKE3 at 1,261 × 2^19 = 661M is the 13 s CPU LogUp stretch visible in the q=41 host-memory curve. The same bound sent every R3/R4 batch inverse above 16.7M elements (a 2^23 LDE at three evaluation points) to the host.
Two layers of
.ok()?hid it:try_build_aux_resident_gpufell totry_build_term_columns_gpu, which tripped the same bound and fell to the CPU build, with no message. The bound itself,total > u32::MAX / BLOCK_SIZE, is 256× stricter than the truncation it guards: the launch builds its grid as(total as u32).div_ceil(256), so the cast wraps only pastu32::MAXelements; the grid atu32::MAXis 16,777,216 blocks against gridDim.x's 2^31 − 1; and every kernel behind the three sites (logup.cu,inverse.cu) indexes with the 64-bitblockIdx.x * (uint64_t)blockDim.x + threadIdx.x. The refusal wore a synthesisedCUDA_ERROR_INVALID_VALUE, so when the aux build's admission stopped swallowing the result (pt/device-fit) it surfaced as a device error on LOCAL_TO_GLOBAL at 2^22 rows with 6 interactions, 6 × 2^22 = 25,165,824.The fix
One bound,
launch_total_fits(total) = total <= u32::MAX, with the reasoning in its doc, at the three sites (logup::cfg,batch_inverse_ext3_dev, the inverse-denominator launch). Beyondu32::MAXthe refusal stands and is now loud at the admitted sites.Pure tests: 6 × 2^22, 24 × 2^22 and 1,261 × 2^21 expressible;
u32::MAXyields a 2^24-block grid;u32::MAX + 1refused.Measured (box A, RTX 5090, one run per cell)
The base layer is proved through the CLI, which has always installed jemalloc, so nothing in this measurement changes because of #966 (which gave the
cargo test --libharness the same allocator); the record row and this row are measured under one allocator and are directly comparable.5d442ba0…988588a[gpu] ABORT/ synthesisedCUDA_ERROR_INVALID_VALUEmath-cuda's own suites green. Pre-registered before the run: completes, host-verifies, wall 70–82 s (held: 64.66 s, under the optimistic end), host flat (held: 33.3 → 33.4 GiB); VRAM up into the 30,000–32,000 MiB band because the aux-build transient (4× the fingerprint buffer, ~8 GiB for CPU at 2^22) was not in the fused-phase admission estimates and several builds can overlap (held: 31,896); host down 0–3 GiB (did not hold: flat at +0.1); wall not slower, 70–82 s (held, and beyond the optimistic end: 64.7 s). Moving the aux builds onto the card, which they had never used, makes the base layer a fifth faster with the proof verifying. One run per cell.
The overlap, and where it is fixed
The VRAM peak sat at the card's ceiling: with the builds on the card, several overlapped under a gate whose estimates did not include them. At this branch alone an overlap that exceeds the card still falls back to the CPU build silently through the remaining
.ok()?, so this run cannot say whether one did. The fix is not here:pt/device-fit's second commit sizes the fused task on the larger of its two phases so builds that will not fit together serialise, and that branch's admission turns any remaining overflow into a loud abort. Merge order: this branch, thenpt/panic-shutdown(so an abort fails the prove instead of hanging it), thenpt/device-fit.