From 47b14218964301e899da1009055bb79d9b29068d Mon Sep 17 00:00:00 2001 From: MauroFab Date: Tue, 8 Sep 2026 03:44:24 -0300 Subject: [PATCH] math-cuda: the launch-size guard refused launches 256 times too small `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. --- crypto/math-cuda/src/inverse.rs | 68 ++++++++++++++++++++++++++++----- crypto/math-cuda/src/logup.rs | 34 ++++++++++++++--- 2 files changed, 86 insertions(+), 16 deletions(-) diff --git a/crypto/math-cuda/src/inverse.rs b/crypto/math-cuda/src/inverse.rs index 1087e2ae4..030a0df02 100644 --- a/crypto/math-cuda/src/inverse.rs +++ b/crypto/math-cuda/src/inverse.rs @@ -148,6 +148,56 @@ fn launch_invert_total( Ok(()) } +/// Whether a one-dimensional launch over `total` elements is expressible. +/// +/// The launches in this crate build their grid as `(total as u32).div_ceil( +/// BLOCK_SIZE)` blocks of `BLOCK_SIZE` threads, and every kernel behind them +/// (`logup.cu`, `inverse.cu`) computes its element index as the 64-bit +/// `blockIdx.x * (uint64_t)blockDim.x + threadIdx.x`. The only thing that +/// silently breaks is therefore the cast: past `u32::MAX` elements `total as +/// u32` wraps, too few blocks launch, and a tail of the output is never +/// written. Up to `u32::MAX` the grid is at most 16,777,216 blocks, far inside +/// `gridDim.x`'s 2^31 − 1, and the index arithmetic is exact. +/// +/// Until 2026-09 this bound was `u32::MAX / BLOCK_SIZE` — 256× too strict, as +/// if the cast applied to the block count rather than the element count. Every +/// LogUp aux build with more than 16,777,215 interaction·rows (at 2^22 rows, +/// any table with five or more interactions) and every R3/R4 batch inverse +/// past that size was refused before any launch and fell to the host, and the +/// refusal wore a `CUDA_ERROR_INVALID_VALUE` that looked like the driver's. +pub(crate) const fn launch_total_fits(total: usize) -> bool { + total <= u32::MAX as usize +} + +#[cfg(test)] +mod launch_bound_tests { + use super::*; + + /// LOCAL_TO_GLOBAL at 2^22 rows with 6 interactions — the first shape the + /// loud aux-build admission surfaced — is expressible; so is the CPU table + /// at 2^22 with its full 24, and LFM_BLAKE3 at 2^21 with its 1,261. + #[test] + fn the_shapes_that_were_refused_fit() { + assert!(launch_total_fits(6 << 22)); + assert!(launch_total_fits(24 << 22)); + assert!(launch_total_fits(1261 << 21)); + assert!(launch_total_fits(u32::MAX as usize)); + } + + /// The truncation hazard is real one past `u32::MAX`, and the old bound + /// sat 256× below it. + #[test] + fn the_bound_is_the_cast_not_the_block_count() { + assert!(!launch_total_fits(u32::MAX as usize + 1)); + assert!(launch_total_fits( + u32::MAX as usize / BLOCK_SIZE as usize + 1 + )); + let blocks = (u32::MAX).div_ceil(BLOCK_SIZE); + assert_eq!(blocks, 1 << 24); + assert!((blocks as u64) < (1u64 << 31)); + } +} + /// Device-input batch inverse. Allocates and returns a fresh `CudaSlice` /// of length `3 * n` holding the inverses. Requires `n >= 1`. /// @@ -160,12 +210,11 @@ pub fn batch_inverse_ext3_dev( stream: &Arc, ) -> Result> { assert!(n >= 1, "batch_inverse_ext3_dev requires n >= 1"); - // Runtime guard (not debug_assert): a u32 grid_dim is truncated past - // u32::MAX / BLOCK_SIZE, which would silently launch too few blocks - // and leave a tail uninverted. Reachable on LDE size 2^23+ × multi- - // eval-point R4. Returning Err lets the dispatcher's Err(_) => None - // route the caller to the CPU `inplace_batch_inverse` fallback. - if n > u32::MAX as usize / BLOCK_SIZE as usize { + // Runtime guard (not debug_assert), see [`launch_total_fits`]: past the + // bound the u32 element count the launch is built from would truncate + // and leave a tail uninverted. Returning Err lets the dispatcher route + // the caller to the CPU `inplace_batch_inverse` fallback, or abort. + if !launch_total_fits(n) { return Err(cudarc::driver::DriverError( cudarc::driver::sys::CUresult::CUDA_ERROR_INVALID_VALUE, )); @@ -259,10 +308,9 @@ pub fn compute_and_invert_denoms_ext3_dev( let total = k_scalars .checked_mul(n) .expect("compute_and_invert_denoms_ext3_dev: k_scalars * n overflow"); - // See `batch_inverse_ext3_dev` for the rationale: runtime Err, not - // debug_assert, so release builds also route past the silent-truncation - // hazard via the caller's CPU fallback. - if total > u32::MAX as usize / BLOCK_SIZE as usize { + // See [`launch_total_fits`]: runtime Err, not debug_assert, so release + // builds also route past the truncation hazard. + if !launch_total_fits(total) { return Err(cudarc::driver::DriverError( cudarc::driver::sys::CUresult::CUDA_ERROR_INVALID_VALUE, )); diff --git a/crypto/math-cuda/src/logup.rs b/crypto/math-cuda/src/logup.rs index ac449e989..e340208c6 100644 --- a/crypto/math-cuda/src/logup.rs +++ b/crypto/math-cuda/src/logup.rs @@ -44,12 +44,11 @@ pub struct LogupDescriptor<'a> { } fn cfg(total: usize) -> Result { - // See `batch_inverse_ext3_dev` for the rationale: a u32 grid_dim is - // truncated past u32::MAX / BLOCK_SIZE, which would silently launch too - // few blocks and leave a tail of the (uninitialized) output unwritten. - // Runtime Err, not debug_assert, so release builds also route to the - // caller's CPU fallback. - if total > u32::MAX as usize / BLOCK_SIZE as usize { + // See `inverse::launch_total_fits` for the bound and its reasoning: past + // it the u32 element count would truncate and leave a tail of the + // (uninitialized) output unwritten. Runtime Err, not debug_assert, so + // release builds also route to the caller's fallback or abort. + if !crate::inverse::launch_total_fits(total) { return Err(cudarc::driver::DriverError( cudarc::driver::sys::CUresult::CUDA_ERROR_INVALID_VALUE, )); @@ -483,3 +482,26 @@ pub fn logup_aux_resident( table_contribution: [l_host[0], l_host[1], l_host[2]], }) } + +#[cfg(test)] +mod launch_cfg_tests { + use super::*; + + /// 6 interactions × 2^22 rows (LOCAL_TO_GLOBAL in a 2^22 epoch) launches + /// 98,304 blocks; it used to be refused as if it were 256× larger. + #[test] + fn the_first_refused_aux_build_shape_launches() { + let c = cfg(6 << 22).expect("expressible"); + assert_eq!(c.grid_dim, (98_304, 1, 1)); + assert_eq!(c.block_dim, (BLOCK_SIZE, 1, 1)); + } + + /// The largest expressible launch is exactly u32::MAX elements; one more + /// is refused with the same error the driver would have used. + #[test] + fn the_bound_is_u32_max_elements() { + let c = cfg(u32::MAX as usize).expect("expressible"); + assert_eq!(c.grid_dim.0, 1 << 24); + assert!(cfg(u32::MAX as usize + 1).is_err()); + } +}