From a0ead23a60d74fb90d11d504b12026177b0c26b5 Mon Sep 17 00:00:00 2001 From: MauroFab Date: Fri, 3 Jul 2026 18:22:49 -0300 Subject: [PATCH 1/2] perf(cuda): release the retained main-trace snapshot after the aux build MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The trace-domain snapshot kept by the R1 main LDE (GpuLdeBase.trace_dev, mirrored into TraceTable.main_trace_dev) is consumed exactly once, by the LogUp aux build. Both Arcs previously lived to the end of the proof, holding a main-trace-sized device buffer through the aux-commit + DEEP/FRI VRAM peak — including for tables that never take the GPU aux path. Drop them right after the aux-build pass instead. --- crypto/stark/src/prover.rs | 16 ++++++++++++++++ crypto/stark/src/trace.rs | 8 ++++++++ 2 files changed, 24 insertions(+) diff --git a/crypto/stark/src/prover.rs b/crypto/stark/src/prover.rs index f47b837d1..6096ab46b 100644 --- a/crypto/stark/src/prover.rs +++ b/crypto/stark/src/prover.rs @@ -2256,6 +2256,22 @@ pub trait IsStarkProver< }) .collect(); + // The trace-domain snapshots retained by the R1 main LDE (both Arcs: + // trace.main_trace_dev and GpuLdeBase.trace_dev) have exactly one + // consumer — the aux build above. Drop them now so the main-trace-sized + // device buffers are reclaimed before the aux-commit + DEEP/FRI VRAM + // peak instead of living to the end of the proof. + #[cfg(feature = "cuda")] + { + for (_, trace, _) in air_trace_pairs.iter_mut() { + trace.clear_main_trace_dev(); + } + for handle in main_gpu_handles.iter_mut().flatten() { + handle.trace_dev = None; + handle.trace_rows = 0; + } + } + // Spill all aux trace tables to mmap before any Round 1 aux LDE work. #[cfg(feature = "disk-spill")] if storage_mode == StorageMode::Disk { diff --git a/crypto/stark/src/trace.rs b/crypto/stark/src/trace.rs index 41e1eb675..6d40425b7 100644 --- a/crypto/stark/src/trace.rs +++ b/crypto/stark/src/trace.rs @@ -205,6 +205,14 @@ where .map(|r| (r.buf.as_ref(), r.rows)) } + /// Drop the retained device-resident main trace. Its only consumer is the + /// aux build, so the prover clears it right after that pass to reclaim the + /// snapshot's VRAM before the aux-commit + DEEP/FRI peak. + #[cfg(feature = "cuda")] + pub fn clear_main_trace_dev(&mut self) { + self.main_trace_dev = None; + } + pub fn num_steps(&self) -> usize { debug_assert!(self.main_table.height.is_multiple_of(self.step_size)); self.main_table.height / self.step_size From 9f5f85b5a5ddae2072e5b6ab93f1af8d308e67fe Mon Sep 17 00:00:00 2001 From: MauroFab Date: Fri, 3 Jul 2026 18:23:12 -0300 Subject: [PATCH 2/2] Address GPU LogUp review findings: guards, test canonicalization, docs, CI MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Hardening: - cfg() returns Err past the u32 grid limit instead of silently truncating the launch (same rationale and pattern as batch_inverse_ext3_dev). - num_out_cols == 0 is a runtime Err, not a debug_assert: in release it would wrap num_out - 1 and launch the assemble kernel with a bogus column count. - Validate descriptor column indices against the table width before launch; the kernels index main[col*num_rows + row] unchecked, so a mis-authored table was silent OOB device reads where the CPU path panics cleanly. - logup_aux_resident asserts num_rows >= 1 (reads row_sum[(num_rows-1)*3..]). Tests: - Canonicalize both sides of the term-column and resident-aux parity asserts (canon3 pattern from the batch_inverse tests): the GPU pipeline computes the same field values through different op trees, so raw limbs can rarely differ in the non-canonical band. The fingerprint test stays raw on purpose — that kernel mirrors the CPU evaluator op for op. - Fix clippy findings in the cuda-gated test module. Docs: - Reattach three doc comments orphaned by items inserted under them (coset_lde_row_major_inner, u64_to_ext3_vec, logup_aux_resident) and update the inner LDE return-tuple description to the new 4-tuple. - try_expand_..._keep_dev borrows the resident buffer (D2D copy), it does not consume it; GPU_LOGUP_CALLS counts successes, not attempts. - disable_event_tracking comment now states the real cross-stream invariant (producer host-syncs before a handle escapes) instead of claiming slices are never shared across streams. - Deduplicate split_interactions (reuse the lookup.rs definition) and name build_accumulated_column_from_terms correctly in logup.cu. CI: - New job runs the CPU-runnable logup_gpu descriptor/CPU-mirror parity tests (cargo test -p stark --features cuda --lib logup_gpu); previously no CI job compiled the module at all since nothing enables the cuda feature. --- .github/workflows/pr_main.yaml | 34 ++++++++- crypto/math-cuda/kernels/logup.cu | 2 +- crypto/math-cuda/src/device.rs | 8 +- crypto/math-cuda/src/lde.rs | 24 +++--- crypto/math-cuda/src/lib.rs | 7 +- crypto/math-cuda/src/logup.rs | 49 +++++++++---- crypto/stark/src/gpu_lde.rs | 17 +++-- crypto/stark/src/logup_gpu.rs | 118 +++++++++++++++++++----------- crypto/stark/src/lookup.rs | 2 +- 9 files changed, 176 insertions(+), 85 deletions(-) diff --git a/.github/workflows/pr_main.yaml b/.github/workflows/pr_main.yaml index 7b6258179..3f80c0582 100644 --- a/.github/workflows/pr_main.yaml +++ b/.github/workflows/pr_main.yaml @@ -127,7 +127,7 @@ jobs: test: name: Test if: always() - needs: [test-executor, test-cli, test-prover, test-disk-spill] + needs: [test-executor, test-cli, test-prover, test-disk-spill, test-stark-cuda-lib] runs-on: ubuntu-latest steps: - name: Check results @@ -136,11 +136,13 @@ jobs: cli="${{ needs.test-cli.result }}" prover="${{ needs.test-prover.result }}" disk_spill="${{ needs.test-disk-spill.result }}" + stark_cuda_lib="${{ needs.test-stark-cuda-lib.result }}" echo "test-executor: $executor" echo "test-cli: $cli" echo "test-prover: $prover" echo "test-disk-spill: $disk_spill" + echo "test-stark-cuda-lib: $stark_cuda_lib" # Allow "success" or "skipped" (skipped on merge queue pushes) if [[ "$executor" != "success" && "$executor" != "skipped" ]]; then @@ -155,6 +157,9 @@ jobs: if [[ "$disk_spill" != "success" && "$disk_spill" != "skipped" ]]; then exit 1 fi + if [[ "$stark_cuda_lib" != "success" && "$stark_cuda_lib" != "skipped" ]]; then + exit 1 + fi test-disk-spill: name: Disk-spill tests @@ -215,6 +220,33 @@ jobs: run: | cargo test --release -p lambda-vm-prover --features disk-spill -- disk_spill count_table_lengths + test-stark-cuda-lib: + name: Stark cuda-feature lib tests + runs-on: ubuntu-latest + if: github.event_name != 'push' || github.actor != 'github-merge-queue[bot]' + steps: + - name: Checkout sources + uses: actions/checkout@v4 + + - name: Setup Rust Environment + uses: ./.github/actions/setup-rust + + - name: Cache cargo build artifacts + uses: Swatinem/rust-cache@v2 + with: + shared-key: "lambda-vm-stark-cuda-lib" + cache-all-crates: "true" + + # The cuda feature gates the logup_gpu module, whose descriptor / + # CPU-mirror parity tests are pure CPU (they never touch the driver). + # Without nvcc the kernels build as empty PTX stubs, so these run on a + # plain runner; GPU-dependent tests are #[ignore] and stay skipped. + # Scoped to logup_gpu: the rest of the cuda-feature lib suite dispatches + # to the GPU and needs the CUDA driver. + - name: Run stark logup_gpu parity tests (no GPU required) + run: | + cargo test --release -p stark --features cuda --lib logup_gpu + build-prover-tests: name: Build prover tests runs-on: ubuntu-latest diff --git a/crypto/math-cuda/kernels/logup.cu b/crypto/math-cuda/kernels/logup.cu index 8823a9ed2..0c01a2f46 100644 --- a/crypto/math-cuda/kernels/logup.cu +++ b/crypto/math-cuda/kernels/logup.cu @@ -115,7 +115,7 @@ extern "C" __global__ void logup_term_ext3( // Accumulated column (K4): running sum of the term columns, on device. // row_sum[i] = sum over all term columns of term[col][i] // S = inclusive prefix scan of row_sum ; L = S[n-1] ; offset = L / N -// acc[i] = S[i] - (i+1) * offset (matches build_accumulated_column) +// acc[i] = S[i] - (i+1) * offset (matches build_accumulated_column_from_terms) // Additive 3-phase Hillis-Steele scan (mirrors inverse.cu, add not mul). // =========================================================================== diff --git a/crypto/math-cuda/src/device.rs b/crypto/math-cuda/src/device.rs index 4d25184ef..b1d6ca489 100644 --- a/crypto/math-cuda/src/device.rs +++ b/crypto/math-cuda/src/device.rs @@ -265,9 +265,11 @@ impl Backend { fn init() -> Result { let ctx = CudaContext::new(0)?; // cudarc's default per-slice CudaEvent tracking adds two driver calls - // per alloc and serialises under the context lock. We never share - // slices across streams (every call scopes its own buffers and syncs - // before returning), so the tracking is pure overhead. Disable it. + // per alloc and serialises under the context lock. Slices are only + // shared across streams after the producing stream has been host- + // synchronised (e.g. the retained trace snapshot and the resident + // LogUp aux buffer; every producer syncs before its handle escapes), + // so the tracking is pure overhead. Disable it. unsafe { ctx.disable_event_tracking() }; // Retain freed device memory in the stream ordered pool for reuse. diff --git a/crypto/math-cuda/src/lde.rs b/crypto/math-cuda/src/lde.rs index 9a36196bf..30e524c2c 100644 --- a/crypto/math-cuda/src/lde.rs +++ b/crypto/math-cuda/src/lde.rs @@ -379,17 +379,6 @@ fn launch_row_to_col_major( Ok(dst) } -/// Shared row-major LDE + Keccak + Merkle pipeline for the base and ext3 paths. -/// -/// `total_cols` is the number of base-field columns in the row-major layout: -/// `m` for base, `m * 3` for ext3. Because `Fp3 = [u64; 3]`, the three ext3 -/// components are just three adjacent base-field columns, so the same row-major -/// NTT and Keccak kernels process all of them simultaneously — no de-interleave. -/// -/// Single H2D, row-major NTT, single D2H — no CPU-side extract or transpose. -/// Returns (merkle_nodes, column-major device buffer, row-major LDE Vec). The -/// buffer is transposed to column-major (as required by the downstream GPU -/// kernels DEEP/barycentric); callers wrap it in the appropriate LDE handle. /// Row-major LDE input: either a host slice (uploaded) or an already-resident /// device buffer (copied device-to-device, no PCIe upload). enum InnerInput<'a> { @@ -397,6 +386,19 @@ enum InnerInput<'a> { Dev(&'a CudaSlice), } +/// Shared row-major LDE + Keccak + Merkle pipeline for the base and ext3 paths. +/// +/// `total_cols` is the number of base-field columns in the row-major layout: +/// `m` for base, `m * 3` for ext3. Because `Fp3 = [u64; 3]`, the three ext3 +/// components are just three adjacent base-field columns, so the same row-major +/// NTT and Keccak kernels process all of them simultaneously — no de-interleave. +/// +/// Single H2D (or D2D), row-major NTT, single D2H — no CPU-side extract or +/// transpose. Returns (merkle_nodes, column-major device buffer, row-major LDE +/// Vec, optional trace-domain column-major snapshot — `Some` iff +/// `retain_trace_col_major`). The buffer is transposed to column-major (as +/// required by the downstream GPU kernels DEEP/barycentric); callers wrap it in +/// the appropriate LDE handle. #[allow(clippy::type_complexity)] fn coset_lde_row_major_inner( input: InnerInput, diff --git a/crypto/math-cuda/src/lib.rs b/crypto/math-cuda/src/lib.rs index 14eead415..abc487750 100644 --- a/crypto/math-cuda/src/lib.rs +++ b/crypto/math-cuda/src/lib.rs @@ -1,8 +1,9 @@ //! GPU backend for the lambda-vm STARK prover. //! -//! Primary entry point: [`lde::coset_lde_base`]. Everything else (`ntt`, -//! element-wise arith) is either internal to the LDE pipeline or used by the -//! parity test suite. +//! Primary entry points: [`lde::coset_lde_base`] for the LDE pipeline and +//! [`logup::logup_aux_resident`] for the device-resident LogUp aux build. +//! Everything else (`ntt`, element-wise arith) is either internal to those +//! pipelines or used by the parity test suite. pub mod barycentric; pub mod deep; diff --git a/crypto/math-cuda/src/logup.rs b/crypto/math-cuda/src/logup.rs index 4fe7e9cd3..e9d120ee1 100644 --- a/crypto/math-cuda/src/logup.rs +++ b/crypto/math-cuda/src/logup.rs @@ -17,6 +17,8 @@ use crate::Result; use crate::device::backend; use crate::inverse::batch_inverse_ext3_dev; +// Must match LOGUP_BLK in kernels/logup.cu: the block scan kernel assumes +// exactly this many threads per block for its shared-memory array. const BLOCK_SIZE: u32 = 256; /// Flat LogUp descriptor for one table (CSR arrays, canonical Goldilocks). Built @@ -41,12 +43,22 @@ pub struct LogupDescriptor<'a> { pub mult_term_col: &'a [u32], } -fn cfg(total: usize) -> LaunchConfig { - LaunchConfig { +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 { + return Err(cudarc::driver::DriverError( + cudarc::driver::sys::CUresult::CUDA_ERROR_INVALID_VALUE, + )); + } + Ok(LaunchConfig { grid_dim: ((total as u32).div_ceil(BLOCK_SIZE), 1, 1), block_dim: (BLOCK_SIZE, 1, 1), shared_mem_bytes: 0, - } + }) } /// Fingerprint kernel over a device-resident main trace. Returns the ext3 fp @@ -94,7 +106,7 @@ fn fingerprints_into_dev( .arg(&z1) .arg(&z2) .arg(&mut out) - .launch(cfg(total))?; + .launch(cfg(total)?)?; } Ok(out) } @@ -168,7 +180,7 @@ pub fn logup_term_columns( .arg(&mult_term_coef) .arg(&mult_term_col) .arg(&mut out) - .launch(cfg(total))?; + .launch(cfg(total)?)?; } if timing { stream.synchronize()?; @@ -266,9 +278,6 @@ impl PartialEq for ResidentAux { } impl Eq for ResidentAux {} -/// Full aux build on device: fingerprints → invert → term columns → accumulate -/// scan → assemble the row-major aux trace buffer, all resident. `inv_n` is -/// `1/num_rows` embedded in ext3. The stream is synchronised before return. /// Main trace input for the resident aux build: either a host column-major /// buffer to upload, or an already-resident device buffer (from the R1 main /// LDE) to read in place. The device form skips the ~3 GB main re-upload. @@ -278,6 +287,10 @@ pub enum ResidentMain<'a> { Dev(&'a CudaSlice), } +/// Full aux build on device: fingerprints → invert → term columns → accumulate +/// scan → assemble the row-major aux trace buffer, all resident. `inv_n` is +/// `1/num_rows` embedded in ext3. Requires `num_rows >= 1`. The stream is +/// synchronised before return. #[allow(clippy::too_many_arguments)] pub fn logup_aux_resident( main: ResidentMain, @@ -288,6 +301,7 @@ pub fn logup_aux_resident( inv_n: [u64; 3], stream: &Arc, ) -> Result { + assert!(num_rows >= 1, "logup_aux_resident requires num_rows >= 1"); let be = backend()?; // Per-phase timing (env LAMBDA_VM_LOGUP_TIMING): sync between phases so wall // time is attributed correctly. Off = no extra syncs, production path. @@ -326,10 +340,13 @@ pub fn logup_aux_resident( // Term columns (committed + virtual), resident, layout [col][row]. // num_out is always >= 1 (the accumulated column); num_committed = num_out - 1. - debug_assert!( - d.num_out_cols >= 1, - "logup descriptor has no output columns" - ); + // Runtime Err, not debug_assert: in release a zero would wrap num_out - 1 + // to usize::MAX and launch the assemble kernel with a bogus column count. + if d.num_out_cols == 0 { + return Err(cudarc::driver::DriverError( + cudarc::driver::sys::CUresult::CUDA_ERROR_INVALID_VALUE, + )); + } let num_out = d.num_out_cols; let mut terms = unsafe { stream.alloc::(num_out * num_rows * 3) }?; let out_col_offsets = stream.clone_htod(d.out_col_offsets)?; @@ -356,7 +373,7 @@ pub fn logup_aux_resident( .arg(&mult_term_coef) .arg(&mult_term_col) .arg(&mut terms) - .launch(cfg(num_out * num_rows))?; + .launch(cfg(num_out * num_rows)?)?; } sync_if(timing)?; let t_term = std::time::Instant::now(); @@ -370,7 +387,7 @@ pub fn logup_aux_resident( .arg(&num_out_u32) .arg(&num_rows_u32) .arg(&mut row_sum) - .launch(cfg(num_rows))?; + .launch(cfg(num_rows)?)?; } scan_add_inplace(stream, be, &mut row_sum, num_rows)?; // row_sum now holds S let (i0, i1, i2) = (inv_n[0], inv_n[1], inv_n[2]); @@ -385,7 +402,7 @@ pub fn logup_aux_resident( .arg(&i1) .arg(&i2) .arg(&mut accumulated) - .launch(cfg(num_rows))?; + .launch(cfg(num_rows)?)?; } // Assemble row-major aux buffer: committed (num_out-1) cols + accumulated. @@ -401,7 +418,7 @@ pub fn logup_aux_resident( .arg(&accumulated) .arg(&num_rows_u32) .arg(&mut aux) - .launch(cfg(num_rows))?; + .launch(cfg(num_rows)?)?; } sync_if(timing)?; let t_accum_done = std::time::Instant::now(); diff --git a/crypto/stark/src/gpu_lde.rs b/crypto/stark/src/gpu_lde.rs index 5cfdbd8a6..f5e1683c8 100644 --- a/crypto/stark/src/gpu_lde.rs +++ b/crypto/stark/src/gpu_lde.rs @@ -84,7 +84,9 @@ pub fn gpu_extend_halves_calls() -> u64 { GPU_EXTEND_HALVES_CALLS.load(Ordering::Relaxed) } -/// LogUp aux-build term-column GPU path attempts (one per table that took it). +/// Successful LogUp aux-build GPU dispatches (one per table that took either +/// the resident or the term-column path; failed attempts fall back to CPU and +/// are not counted). pub(crate) static GPU_LOGUP_CALLS: AtomicU64 = AtomicU64::new(0); pub fn gpu_logup_calls() -> u64 { GPU_LOGUP_CALLS.load(Ordering::Relaxed) @@ -1108,12 +1110,10 @@ unsafe fn ext3_slice_to_u64(col: &[FieldElement]) -> &[u64] { unsafe { from_raw_parts(ptr, len) } } -/// Convert ext3 evals (3*n u64s, interleaved) into a freshly allocated -/// `Vec>` of length `n`. Caller must have established -/// `E == Ext3`. /// Like [`try_expand_leaf_and_tree_ext3_row_major_keep`] but the aux columns are /// already resident on device (from the GPU LogUp aux build) — no host upload. -/// Consumes the resident buffer via the device-input LDE. +/// The resident buffer is only borrowed: the device-input LDE copies it +/// device-to-device into its own scratch, so `ra` stays valid afterwards. pub(crate) fn try_expand_leaf_and_tree_ext3_row_major_keep_dev( ra: &math_cuda::logup::ResidentAux, blowup_factor: usize, @@ -1150,6 +1150,10 @@ where let lde_out: Vec> = unsafe { let mut v = std::mem::ManuallyDrop::new(lde_u64); + debug_assert!( + v.len() % 3 == 0 && v.capacity() % 3 == 0, + "lde_u64 len/capacity must be a multiple of 3 for Fp3 reinterpret" + ); Vec::from_raw_parts( v.as_mut_ptr() as *mut FieldElement, v.len() / 3, @@ -1161,6 +1165,9 @@ where Some((tree, handle, lde_out)) } +/// Convert ext3 evals (3*n u64s, interleaved) into a freshly allocated +/// `Vec>` of length `n`. Caller must have established +/// `E == Ext3`. pub(crate) fn u64_to_ext3_vec(raw: &[u64]) -> Vec> where E: IsField + 'static, diff --git a/crypto/stark/src/logup_gpu.rs b/crypto/stark/src/logup_gpu.rs index 10f2d990c..bc9e88302 100644 --- a/crypto/stark/src/logup_gpu.rs +++ b/crypto/stark/src/logup_gpu.rs @@ -13,7 +13,7 @@ use std::any::TypeId; use crate::lookup::{ BusInteraction, BusValue, LOGUP_CHALLENGE_ALPHA, LinearTerm, Multiplicity, Packing, - compute_alpha_powers, + compute_alpha_powers, split_interactions, }; use math::field::element::FieldElement; use math::field::extensions_goldilocks::Degree3GoldilocksExtensionField; @@ -49,17 +49,6 @@ fn neg_canonical(x: u64) -> u64 { if x == 0 { 0 } else { GOLDILOCKS_P - x } } -/// Committed-pair / absorbed split, mirroring `lookup::split_interactions`. -fn split_interactions(num: usize) -> (usize, usize) { - if num <= 2 { - (0, num) - } else if num % 2 == 1 { - ((num - 1) / 2, 1) - } else { - ((num - 2) / 2, 2) - } -} - /// Encode a multiplicity as a signed linear form `const + Σ coef·col` (sign /// baked in: negated for receivers so `term = m'·recip` needs no extra sign). /// Mirrors `Multiplicity::evaluate_with` for every variant. @@ -153,6 +142,21 @@ pub struct FingerprintDescriptor { } impl FingerprintDescriptor { + /// Panic if any fingerprint or multiplicity term references a main column + /// index `>= num_cols`. The kernels index `main[col*num_rows + row]` + /// unchecked (they are never told the column count), so a mis-authored + /// table would otherwise be a silent out-of-bounds device read; this makes + /// it fail loudly, like the CPU path's slice indexing. O(#terms), run once + /// per table build. + fn assert_columns_in_bounds(&self, num_cols: usize) { + for &col in self.term_col.iter().chain(self.mult_term_col.iter()) { + assert!( + (col as usize) < num_cols, + "logup descriptor references main column {col} but the table has {num_cols} columns" + ); + } + } + fn push_element(&mut self, alpha_idx: u32, const_val: u64, terms: &[(u64, u32)]) { self.elem_alpha_idx.push(alpha_idx); self.elem_const.push(const_val); @@ -371,6 +375,7 @@ where // main trace -> column-major u64. SAFETY: F == Goldilocks (repr(u64)). let num_cols = main_cols.len(); + desc.assert_columns_in_bounds(num_cols); let mut main_flat = vec![0u64; num_cols * trace_len]; for (c, col) in main_cols.iter().enumerate() { for (r, e) in col.iter().enumerate() { @@ -438,6 +443,7 @@ where } let num_cols = main_cols.len(); + desc.assert_columns_in_bounds(num_cols); // Reuse the resident main trace from the R1 main LDE (column-major // `[col*trace_len + row]`, same column order as `main_cols`) when it matches // this table exactly; otherwise flatten + upload the host columns. The @@ -555,6 +561,7 @@ mod tests { use crate::lookup::{PackingShifts, compute_alpha_powers}; use math::field::extensions_goldilocks::Degree3GoldilocksExtensionField; use math::field::goldilocks::GoldilocksField; + use math::field::traits::IsPrimeField; type F = GoldilocksField; type E = Degree3GoldilocksExtensionField; @@ -663,6 +670,17 @@ mod tests { [*v[0].value(), *v[1].value(), *v[2].value()] } + // Reduce raw limbs to canonical form before comparing. The GPU pipeline + // computes the same field values as the CPU reference but through different + // op trees (tree-scan batch inverse, closed-form accumulate), and Goldilocks + // representatives in [p, 2^64) are legal, so raw limbs may rarely differ + // even when the values match. Same pattern as math-cuda's batch_inverse + // parity tests. The fingerprint test deliberately compares raw limbs: the + // kernel mirrors the CPU evaluator op for op, so bit-identity holds there. + fn canon(a: &[u64]) -> Vec { + a.iter().map(F::canonical).collect() + } + // GPU fingerprint kernel vs the CPU evaluator, byte for byte (full ext3 // alpha/z so mul_base is exercised). Runs on the GPU box. #[test] @@ -712,6 +730,7 @@ mod tests { // CPU reference, layout [(k*num_rows + row)*3 + limb]. let mut cpu = vec![0u64; interactions.len() * num_rows * 3]; + #[allow(clippy::needless_range_loop)] // main is column-major: main[c][row] for k in 0..interactions.len() { for row in 0..num_rows { let fp = eval_fingerprint::(&desc, k, |c| &main[c][row], &alpha_powers, &z); @@ -773,7 +792,7 @@ mod tests { for (k, it) in ints.iter().enumerate() { let m = it.multiplicity.evaluate_at_row(main, row); let t = &m * &fps[k * num_rows + row]; - acc = acc + if it.is_sender { t } else { -t }; + acc += if it.is_sender { t } else { -t }; } *slot = acc; } @@ -839,6 +858,7 @@ mod tests { // Reciprocals of every interaction's fingerprint, laid out [k*num_rows+row]. let mut recips: Vec> = Vec::with_capacity(interactions.len() * num_rows); + #[allow(clippy::needless_range_loop)] // main is column-major: main[c][row] for k in 0..interactions.len() { for row in 0..num_rows { recips.push(eval_fingerprint::( @@ -895,31 +915,32 @@ mod tests { // CPU reference term columns: 2 committed pairs + virtual (last 1). let mut cpu = vec![0u64; desc.num_out_cols * num_rows * 3]; - let mut ref_cols: Vec>> = Vec::new(); - ref_cols.push(reference_term_column( - &[&interactions[0], &interactions[1]], - &main, - num_rows, - &alpha_powers, - &z, - &shifts, - )); - ref_cols.push(reference_term_column( - &[&interactions[2], &interactions[3]], - &main, - num_rows, - &alpha_powers, - &z, - &shifts, - )); - ref_cols.push(reference_term_column( - &[&interactions[4]], - &main, - num_rows, - &alpha_powers, - &z, - &shifts, - )); + let ref_cols: Vec>> = vec![ + reference_term_column( + &[&interactions[0], &interactions[1]], + &main, + num_rows, + &alpha_powers, + &z, + &shifts, + ), + reference_term_column( + &[&interactions[2], &interactions[3]], + &main, + num_rows, + &alpha_powers, + &z, + &shifts, + ), + reference_term_column( + &[&interactions[4]], + &main, + num_rows, + &alpha_powers, + &z, + &shifts, + ), + ]; for (col, rc) in ref_cols.iter().enumerate() { for (row, v) in rc.iter().enumerate() { let o = (col * num_rows + row) * 3; @@ -944,7 +965,11 @@ mod tests { .unwrap(); assert_eq!(desc.num_out_cols, 3); - assert_eq!(gpu, cpu, "GPU term columns mismatch CPU reference"); + assert_eq!( + canon(&gpu), + canon(&cpu), + "GPU term columns mismatch CPU reference" + ); } // Reference accumulated column, mirroring build_accumulated_column_from_terms. @@ -968,7 +993,7 @@ mod tests { rs = &rs + &c[row]; } acc = &acc + &rs - &offset; - out.push(acc.clone()); + out.push(acc); } (out, total) } @@ -1068,11 +1093,15 @@ mod tests { stream.synchronize().unwrap(); assert_eq!( - ra.table_contribution, - limbs(&total), + canon(&ra.table_contribution), + canon(&limbs(&total)), "table_contribution L mismatch" ); - assert_eq!(gpu, expected, "resident aux buffer mismatch CPU reference"); + assert_eq!( + canon(&gpu), + canon(&expected), + "resident aux buffer mismatch CPU reference" + ); // Resident-main path: upload main col-major to device, then build via // ResidentMain::Dev (no host upload). Must be byte-identical to Host. @@ -1095,7 +1124,8 @@ mod tests { "resident-main L mismatch vs host-upload path" ); assert_eq!( - gpu_dev, expected, + canon(&gpu_dev), + canon(&expected), "resident-main aux buffer mismatch CPU reference" ); } diff --git a/crypto/stark/src/lookup.rs b/crypto/stark/src/lookup.rs index 0c5e5cc22..4273f29a7 100644 --- a/crypto/stark/src/lookup.rs +++ b/crypto/stark/src/lookup.rs @@ -114,7 +114,7 @@ const LOGUP_CHUNK_SIZE: usize = 1024; /// Returns `(num_committed_pairs, absorbed_count)` where: /// - Committed pairs get dedicated auxiliary term columns (2 interactions per column) /// - Absorbed interactions (1 or 2) are folded into the accumulated constraint -fn split_interactions(num_interactions: usize) -> (usize, usize) { +pub(crate) fn split_interactions(num_interactions: usize) -> (usize, usize) { if num_interactions <= 2 { (0, num_interactions) } else if num_interactions % 2 == 1 {