perf(gpu): run DECODE (num_parts=1) DEEP/FRI on device

.github/workflows/gpu-tests.yml

@@ -3,8 +3,9 @@ name: GPU Tests (merge queue)
 # Run the GPU test suite (which CPU CI can't, since GitHub runners have no GPU) on a rented
 # Vast.ai RTX 5090 when a PR is in the merge queue, and block the merge if it fails.
 # Groups (see scripts/gpu_test.sh): math-cuda kernel parity, cuda_path_integration (GPU proof
-# verifies), cuda_fallback (CPU fallback verifies), the prover/stark/crypto/ecsm suite on the
-# GPU path, and the comprehensive all-instructions prove. Orchestration runs on a GitHub-hosted
+# verifies), cuda_d1_path (the num_parts==1 device DEEP/FRI path), cuda_fallback (CPU fallback
+# verifies), the prover/stark/crypto/ecsm suite on the GPU path, and the comprehensive
+# all-instructions prove. Orchestration runs on a GitHub-hosted
 # runner; all GPU work happens on the rented box (provisioned by the template onstart). The box
 # is ALWAYS destroyed at the end.
 #
@@ -55,7 +56,7 @@ jobs:
     # Skip on PRs (reports as Skipped = required check satisfied, no GPU rental); run for
     # real on merge_group and manual dispatch.
     if: github.event_name != 'pull_request'
-    # Provisioning + cuda builds + 5 test groups; the prover suite (single-threaded, real
+    # Provisioning + cuda builds + 6 test groups; the prover suite (single-threaded, real
     # ELF proves) dominates. Generous ceiling; teardown still always destroys the box.
     timeout-minutes: 240
     steps:

Makefile

@@ -3,7 +3,7 @@ compile-programs compile-recursion-elfs clean-asm clean-rust clean-bench clean-s
 clean-recursion-elfs clean test test-asm \
 test-rust test-ethrex test-ethrex-offline test-executor test-syscalls test-flamegraph flamegraph-prover test-profile-recursion test-profile-recursion-single test-profile-recursion-multi \
 test-profile-recursion-block recursion-profile-block-input \
-test-fast test-prover test-prover-all test-prover-debug test-disk-spill test-math-cuda test-cuda-integration test-cuda-fallback \
+test-fast test-prover test-prover-all test-prover-debug test-disk-spill test-math-cuda test-cuda-integration test-cuda-d1 test-cuda-fallback \
 test-prover-cuda test-prover-comprehensive-cuda \
 bench-math-cuda bench-prover bench-prover-cuda build check clippy fmt lint regen-ethrex-fixtures \
 update-ethrex-fixture-checksums check-ethrex-fixture-checksums ethrex-real-block-fixture \
@@ -574,7 +574,7 @@ test-disk-spill:
 GPU_TEST_TIMEOUT := timeout -k 30 2700
 
 # math-cuda kernel tests (requires NVIDIA GPU + nvcc). Group 1 of gpu_test.sh,
-# so a hang here also costs Groups 2-5: they run after it, sequentially.
+# so a hang here also costs Groups 2-6: they run after it, sequentially.
 test-math-cuda:
 	$(GPU_TEST_TIMEOUT) cargo test -p math-cuda --release
 
@@ -586,6 +586,33 @@ test-cuda-integration:
 	$(GPU_TEST_TIMEOUT) cargo test -p lambda-vm-prover --release --features cuda \
 	    --test cuda_path_integration -- --ignored --nocapture --test-threads=1
 
+# num_parts==1 (DECODE) device DEEP/FRI coverage (requires NVIDIA GPU + nvcc).
+# No fixture crosses the default LDE threshold (1<<14) for a num_parts==1 table,
+# so lower it here until DECODE engages the d=1 device path end to end.
+#
+# Threshold and fixture are one choice, because there are exactly two d=1 tables
+# (a d=1 table is one with a single bus interaction): DECODE, whose rows come from
+# the guest's instruction count, and KECCAK_RC, fixed at NUM_ROWS=32 => LDE 64.
+# DECODE's ROM is derived from the ELF, NOT from cycles, so the whole
+# fib_iterative_* family is 13 executable words (the variants differ only in the
+# `li a0, <count>` immediate) => 16 rows => LDE 32. That sits BELOW KECCAK_RC's 64,
+# so with a fib fixture no threshold isolates DECODE: <=32 engages both and
+# 33..=64 engages only KECCAK_RC.
+#
+# all_instructions_64 is 66 executable words => 128 rows => DECODE LDE 256. At 128,
+# DECODE engages with 2x margin and KECCAK_RC (64) declines, so a nonzero
+# gpu_comp_h_slabs_calls() uniquely attributes to DECODE. 128 is also ABOVE the
+# PR's original 64, so it sends strictly fewer tables onto the GPU-committed path
+# and narrows -- rather than widens -- the R4 gather_proofs_dev abort site that
+# crypto/stark/src/gpu_lde.rs warns about for lowered thresholds.
+#
+# Its own binary + a process-wide env because gpu_lde_threshold() caches the value
+# on first read (OnceLock), so it must be set before any prove in the process.
+test-cuda-d1:
+	LAMBDA_VM_GPU_LDE_THRESHOLD=128 $(GPU_TEST_TIMEOUT) cargo test -p lambda-vm-prover \
+	    --release --features cuda \
+	    --test cuda_d1_path -- --ignored --nocapture --test-threads=1
+
 # GPU error-path coverage (requires NVIDIA GPU + nvcc).
 # Forces cuda dispatch errors and asserts the CPU fallback still produces a verifying proof.
 test-cuda-fallback:

crypto/math-cuda/kernels/constraint_interp.cu

@@ -495,3 +495,21 @@ extern "C" __global__ void decompose_d2_ext3(
         out[5 * slab_stride + i] = h1.c;
     }
 }
+
+// ============================================================================
+// Degree-1 (num_parts==1) composition part: H IS the single part, already on
+// the LDE coset, so there is no decompose and no re-extension. Only de-interleave
+// the resident ext3 composition evals `h` (num_rows rows, interleaved
+// `h[row*3 + k]`) into the 3-slab layout the commit / DEEP / FRI consumers
+// expect (`out[k*num_rows + row]`).
+extern "C" __global__ void comp_h_to_slabs_ext3(
+    const uint64_t *__restrict__ h,
+    uint64_t num_rows,
+    uint64_t *__restrict__ out) {
+    for (uint64_t i = (uint64_t)blockIdx.x * blockDim.x + threadIdx.x; i < num_rows;
+         i += (uint64_t)gridDim.x * blockDim.x) {
+        out[0 * num_rows + i] = h[i * 3];
+        out[1 * num_rows + i] = h[i * 3 + 1];
+        out[2 * num_rows + i] = h[i * 3 + 2];
+    }
+}

crypto/math-cuda/src/constraint_interp.rs

@@ -502,3 +502,71 @@ pub fn decompose_d2_into_slabs(
     }
     Ok((out, stream, n))
 }
+
+/// Degree-1 (num_parts==1) composition part: `H` is already the single part on
+/// the LDE coset, so there is neither a decompose nor a re-extension — only a
+/// de-interleave of the resident interleaved ext3 evals `h` (`num_rows` rows,
+/// `h[row*3 + k]`) into the 3-slab layout the commit / DEEP / FRI consumers read
+/// (`out[(0*3 + k) * lde_size + row]`, i.e. one column of 3 slabs). Returns a
+/// device-resident [`GpuLdeExt3`] with `m = 1` and `lde_size == h.num_rows`,
+/// kept live on `h`'s stream with a recorded event so cross-stream consumers
+/// wait device-side (no host block).
+pub fn comp_h_to_slabs(h: &GpuCompH) -> Result<GpuLdeExt3> {
+    let lde_size = h.num_rows;
+    assert!(
+        lde_size.is_power_of_two() && lde_size >= 2,
+        "H row count must be a power of two"
+    );
+    let be = backend()?;
+    let stream = h.stream.clone();
+    // The kernel writes every one of the `3 * lde_size` slab u64s, so an
+    // uninitialized allocation is sound (no zero-pad tail, unlike the d=2
+    // decompose which only fills the first `n` rows).
+    let mut out = unsafe { stream.alloc::<u64>(3 * lde_size) }?;
+
+    let grid = (lde_size as u32)
+        .div_ceil(BLOCK_DIM)
+        .clamp(1, MAX_THREADS / BLOCK_DIM);
+    let cfg = LaunchConfig {
+        grid_dim: (grid, 1, 1),
+        block_dim: (BLOCK_DIM, 1, 1),
+        shared_mem_bytes: 0,
+    };
+    let num_rows_u64 = lde_size as u64;
+    unsafe {
+        stream
+            .launch_builder(&be.comp_h_to_slabs_kernel)
+            .arg(&h.buf)
+            .arg(&num_rows_u64)
+            .arg(&mut out)
+            .launch(cfg)?;
+    }
+
+    let ready = be.take_event()?;
+    ready.event().record(&stream)?;
+
+    Ok(GpuLdeExt3 {
+        buf: Arc::new(out),
+        m: 1,
+        lde_size,
+        tree: None,
+        ready: Some(Arc::new(ready)),
+    })
+}
+
+/// Parity helper: build a resident [`GpuCompH`] from interleaved ext3 evals on
+/// host (`h[row*3 + k]`, `num_rows * 3` u64), uploaded on a fresh stream. On the
+/// prove path `H` is born on device (never uploaded); this exists only so the
+/// de-interleave kernel can be exercised in isolation against a host oracle.
+pub fn comp_h_from_host_interleaved(interleaved: &[u64], num_rows: usize) -> Result<GpuCompH> {
+    assert_eq!(interleaved.len(), num_rows * 3, "interleaved ext3 length");
+    let be = backend()?;
+    let stream = be.next_stream();
+    let buf = stream.clone_htod(interleaved)?;
+    stream.synchronize()?;
+    Ok(GpuCompH {
+        buf,
+        num_rows,
+        stream,
+    })
+}

crypto/math-cuda/src/device.rs

@@ -243,6 +243,7 @@ pub struct Backend {
     pub constraint_interp_kernel: CudaFunction,
     pub constraint_composition_kernel: CudaFunction,
     pub decompose_d2_kernel: CudaFunction,
+    pub comp_h_to_slabs_kernel: CudaFunction,
 
     // Twiddle caches keyed by log_n.
     fwd_twiddles: Mutex<Vec<Option<Arc<CudaSlice<u64>>>>>,
@@ -474,6 +475,7 @@ impl Backend {
             constraint_composition_kernel: constraint_interp
                 .load_function("constraint_composition_kernel")?,
             decompose_d2_kernel: constraint_interp.load_function("decompose_d2_ext3")?,
+            comp_h_to_slabs_kernel: constraint_interp.load_function("comp_h_to_slabs_ext3")?,
             fwd_twiddles: Mutex::new(vec![None; max_log]),
             inv_twiddles: Mutex::new(vec![None; max_log]),
             ctx,

crypto/math-cuda/tests/comp_h_to_slabs.rs

@@ -0,0 +1,65 @@
+//! Parity for the degree-1 (num_parts==1) composition-parts de-interleave
+//! kernel (`comp_h_to_slabs_ext3`).
+//!
+//! On the prove path a table with `num_parts == 1` has `H` itself as its single
+//! composition part, already on the LDE coset. The device path keeps it resident
+//! by de-interleaving the interleaved ext3 evals `H` (`h[row*3 + k]`) into the
+//! 3-slab layout every downstream consumer (R2 commit, R3 OOD, R4 DEEP, openings)
+//! reads (`buf[(0*3 + k) * lde_size + row]`). It is a pure transpose — no
+//! arithmetic — so raw u64 equality must hold bit-for-bit.
+//!
+//! Requires a visible GPU (like the other math-cuda GPU parity tests).
+
+use math_cuda::constraint_interp::{comp_h_from_host_interleaved, comp_h_to_slabs};
+use math_cuda::device::backend;
+
+fn check(num_rows: usize, seed: u64) {
+    // Deterministic interleaved ext3 `H` (raw, possibly non-canonical limbs —
+    // the stronger test, and exactly what a real resident `H` carries).
+    let mut state = seed;
+    let mut next = || {
+        state = state
+            .wrapping_mul(6364136223846793005)
+            .wrapping_add(1442695040888963407);
+        state
+    };
+    let interleaved: Vec<u64> = (0..num_rows * 3).map(|_| next()).collect();
+
+    let h = comp_h_from_host_interleaved(&interleaved, num_rows).expect("upload H");
+    let handle = comp_h_to_slabs(&h).expect("de-interleave H into slabs");
+    assert_eq!(handle.m, 1, "num_rows={num_rows}: single part");
+    assert_eq!(handle.lde_size, num_rows, "num_rows={num_rows}: lde_size");
+    assert_eq!(
+        handle.buf.len(),
+        3 * num_rows,
+        "num_rows={num_rows}: slab buffer"
+    );
+
+    let be = backend().expect("cuda backend");
+    let stream = be.next_stream();
+    handle
+        .wait_ready_on(stream.as_ref())
+        .expect("wait on de-interleave");
+    let slab = stream
+        .clone_dtoh(handle.buf.as_ref())
+        .expect("download slabs");
+    stream.synchronize().expect("sync download");
+
+    for row in 0..num_rows {
+        for k in 0..3 {
+            let got = slab[k * num_rows + row];
+            let want = interleaved[row * 3 + k];
+            assert_eq!(
+                got, want,
+                "num_rows={num_rows} row={row} comp={k}: slab {got:#018x} vs interleaved {want:#018x}"
+            );
+        }
+    }
+}
+
+#[test]
+fn comp_h_to_slabs_parity() {
+    for log in 1..=14 {
+        check(1usize << log, 0x00C0_FFEE_0000_0000 ^ log as u64);
+    }
+}

crypto/stark/src/gpu_lde.rs

@@ -154,6 +154,7 @@ pub fn gpu_lde_calls() -> u64 {
 pub fn reset_all_gpu_call_counters() {
     GPU_LDE_CALLS.store(0, Ordering::Relaxed);
     GPU_EXTEND_HALVES_CALLS.store(0, Ordering::Relaxed);
+    GPU_COMP_H_SLABS_CALLS.store(0, Ordering::Relaxed);
     GPU_LEAF_HASH_CALLS.store(0, Ordering::Relaxed);
     GPU_MERKLE_TREE_CALLS.store(0, Ordering::Relaxed);
     GPU_PARTS_LDE_CALLS.store(0, Ordering::Relaxed);
@@ -187,6 +188,15 @@ pub fn gpu_extend_halves_calls() -> u64 {
     GPU_EXTEND_HALVES_CALLS.load(Ordering::Relaxed)
 }
 
+/// Device-resident num_parts==1 composition-parts dispatches: one per table
+/// whose single composition part (`H` itself) was de-interleaved into a slab
+/// [`math_cuda::lde::GpuLdeExt3`] on device instead of the host arm. Nonzero
+/// confirms the degree-1 device DEEP/FRI path engaged.
+pub(crate) static GPU_COMP_H_SLABS_CALLS: AtomicU64 = AtomicU64::new(0);
+pub fn gpu_comp_h_slabs_calls() -> u64 {
+    GPU_COMP_H_SLABS_CALLS.load(Ordering::Relaxed)
+}
+
 /// 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).
@@ -682,10 +692,34 @@ where
     Some((lde_h0, lde_h1))
 }
 
+/// Shared admission gate for the device composition-parts producers: the tower
+/// must be the Goldilocks/ext3 pair the kernels are written for, and the LDE must
+/// be a power of two at or above the commit threshold. Returns the validated LDE
+/// size so callers can derive from it. Kept in one place so a future condition
+/// (a VRAM check, a tower widening) cannot land on only one of the d=1/d=2 arms.
+fn dev_comp_parts_gate<F, E>(num_rows: usize) -> Option<usize>
+where
+    F: IsField + 'static,
+    E: IsField + 'static,
+{
+    if TypeId::of::<F>() != TypeId::of::<GoldilocksField>() {
+        return None;
+    }
+    if TypeId::of::<E>() != TypeId::of::<Degree3GoldilocksExtensionField>() {
+        return None;
+    }
+    if num_rows < gpu_lde_threshold() || !num_rows.is_power_of_two() {
+        return None;
+    }
+    Some(num_rows)
+}
+
 /// Fully device-resident degree-2 decomposition + half extension: takes the
 /// resident composition evals `H`, decomposes into H0/H1 on device, LDE-extends
-/// both and keeps the de-interleaved parts buffer as a `GpuLdeExt3` (commit
-/// tree, R3 OOD, R4 DEEP and openings all read the handle). With `want_host`
+/// both and keeps the de-interleaved parts buffer as a `GpuLdeExt3` (the commit
+/// tree and the R4 openings read `handle.m`; R3 and R4 DEEP read the host part
+/// Vec's length and DEEP validates the handle against it — see
+/// [`try_comp_h_to_slabs_dev`] for why the two must stay equal). With `want_host`
 /// the evaluations are also drained to host for the fallback consumers;
 /// without it (device-only) the returned part Vecs are empty placeholders.
 /// `None` → the caller downloads `H` and runs the host decompose path.
@@ -699,16 +733,7 @@ where
     F: IsField + 'static,
     E: IsField + 'static,
 {
-    if TypeId::of::<F>() != TypeId::of::<GoldilocksField>() {
-        return None;
-    }
-    if TypeId::of::<E>() != TypeId::of::<Degree3GoldilocksExtensionField>() {
-        return None;
-    }
-    let lde_size = h.num_rows;
-    if lde_size < gpu_lde_threshold() || !lde_size.is_power_of_two() {
-        return None;
-    }
+    let lde_size = dev_comp_parts_gate::<F, E>(h.num_rows)?;
     let n = lde_size / 2;
     if weights.len() != n || inv_2x.len() < n {
         return None;
@@ -771,6 +796,55 @@ where
     Some((vec![lde_h0, lde_h1], handle))
 }
 
+/// Fully device-resident num_parts==1 composition-parts path: `H` itself is the
+/// single part, already on the LDE coset, so — unlike [`try_comp_h_to_slabs_dev`]'s
+/// d=2 sibling [`try_decompose_extend_d2_dev`] — there is no decompose and no
+/// re-extension, only a de-interleave into the slab layout the downstream consumers
+/// read. No consumer needed changing for `m == 1`, but they do not agree on where
+/// the part count comes from, and the difference matters to anyone editing this:
+///
+/// - R2 commit and the R4 openings read `handle.m`.
+/// - R3's `z^P` exponent and R4 DEEP's gamma count read
+///   `lde_composition_poly_evaluations.len()` — the HOST part Vec's length. DEEP only
+///   *validates* the handle against it and declines on a mismatch.
+/// - FRI never sees the handle at all; it consumes the DEEP codeword.
+///
+/// So the invariant to preserve is `handle.m == lde_composition_poly_evaluations.len()`
+/// (`materialize_composition_parts_host` also requires it), not "the handle is
+/// authoritative".
+///
+/// The single part is always drained to host — not just because it can be
+/// (num_parts==1 tables are never device-only; `device_only_for`'s degree gate admits
+/// only d=2), but because that host part is what feeds the query-0
+/// composition-opening canary: release-active for `qi == 0` and guarded on a
+/// non-empty host part, it is the only *in-prove* check that the device m=1 gather is
+/// correct. It does not cover DEEP or FRI, which consume separate downstream buffers;
+/// those are covered by proof verification (`prover/tests/cuda_d1_path.rs`). Returning
+/// empty parts (`vec![Vec::new()]`, as the d=2 device-only arm does) would save the
+/// D2H and keep num_parts==1 — but silently disable that canary.
+/// `None` → the caller downloads `H` and uses it directly as the single host part.
+pub(crate) fn try_comp_h_to_slabs_dev<F, E>(
+    h: &math_cuda::constraint_interp::GpuCompH,
+) -> Option<(Vec<Vec<FieldElement<E>>>, math_cuda::lde::GpuLdeExt3)>
+where
+    F: IsField + 'static,
+    E: IsField + 'static,
+{
+    dev_comp_parts_gate::<F, E>(h.num_rows)?;
+
+    // The interleaved `H` download IS the single composition part on the LDE
+    // coset — same values the slab handle holds, just interleaved. Downloading
+    // first keeps the blocking D2H off the tail of the de-interleave launch; a
+    // later handle failure just re-drains in the caller's fallback (both values
+    // drop by RAII on any early return, in either order).
+    let host = vec![download_comp_h_to_field::<E>(h)?];
+
+    let handle = math_cuda::constraint_interp::comp_h_to_slabs(h).ok()?;
+    GPU_COMP_H_SLABS_CALLS.fetch_add(1, Ordering::Relaxed);
+
+    Some((host, handle))
+}
+
 /// D2H bridge for the fallback: download a resident `H` and lift it into
 /// field elements (the exact input the host decompose expects).
 pub(crate) fn download_comp_h_to_field<E: IsField + 'static>(