diff --git a/.github/workflows/gpu-tests.yml b/.github/workflows/gpu-tests.yml
index ddcce0ee3..c1fc18aa6 100644
--- a/.github/workflows/gpu-tests.yml
+++ b/.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:
diff --git a/Makefile b/Makefile
index fa80a77fe..3e4a88ecb 100644
--- a/Makefile
+++ b/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:
diff --git a/crypto/math-cuda/kernels/constraint_interp.cu b/crypto/math-cuda/kernels/constraint_interp.cu
index 4c4caf076..535a09fb5 100644
--- a/crypto/math-cuda/kernels/constraint_interp.cu
+++ b/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];
+    }
+}
diff --git a/crypto/math-cuda/src/constraint_interp.rs b/crypto/math-cuda/src/constraint_interp.rs
index 315e6eea4..0c2a620ad 100644
--- a/crypto/math-cuda/src/constraint_interp.rs
+++ b/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,
+    })
+}
diff --git a/crypto/math-cuda/src/device.rs b/crypto/math-cuda/src/device.rs
index ba63b4817..3a2f1db2a 100644
--- a/crypto/math-cuda/src/device.rs
+++ b/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,
diff --git a/crypto/math-cuda/tests/comp_h_to_slabs.rs b/crypto/math-cuda/tests/comp_h_to_slabs.rs
new file mode 100644
index 000000000..0bce949d5
--- /dev/null
+++ b/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);
+    }
+}
diff --git a/crypto/stark/src/gpu_lde.rs b/crypto/stark/src/gpu_lde.rs
index 23366d67f..8782c6923 100644
--- a/crypto/stark/src/gpu_lde.rs
+++ b/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>(
diff --git a/crypto/stark/src/prover.rs b/crypto/stark/src/prover.rs
index d31ea09a2..5078ce290 100644
--- a/crypto/stark/src/prover.rs
+++ b/crypto/stark/src/prover.rs
@@ -1090,16 +1090,21 @@ pub trait IsStarkProver<
         //  - The composition path needs a uniform zerofier with ≥1 group. An
         //    empty constraint set makes `all(end_exemptions == 0)` vacuously
         //    true here but `is_uniform()` false downstream (0 groups).
-        //  - The device-resident R2 path exists only for the d=2 quotient
-        //    decomposition, checked below once `n` is in hand.
+        //  - Device-only is entered only for the d=2 quotient decomposition,
+        //    checked below once `n` is in hand. A d=1 table also has a device R2
+        //    path, but the gate below excludes it, so it stays device-additive.
         if !air.has_aux_trace() || air.constraints_meta().is_empty() {
             return false;
         }
         let n = domain.interpolation_domain_size;
-        // The device-resident R2 path only exists for the d=2 quotient
-        // decomposition; any other part count skips it entirely and needs the
-        // host evaluator, which device-only would leave without data until the
-        // R2 downgrade recovered it.
+        // Only the d=2 quotient decomposition has a device-resident R2 path that
+        // can serve every downstream consumer from the handle alone. A d=1 table
+        // does have a device R2 path, but it always drains its single part to host
+        // (the query-0 composition canary reads it), so it gains nothing from
+        // dropping the host trace and this gate keeps it device-additive. Any other
+        // part count has no device R2 path at all and needs the host evaluator,
+        // which device-only would leave without data until the R2 downgrade
+        // recovered it.
         if air.composition_poly_degree_bound(n) / n != 2 {
             return false;
         }
@@ -1538,6 +1543,54 @@ pub trait IsStarkProver<
         }
     }
 
+    /// Decompose the resident composition `H` into device-resident parts per the
+    /// AIR's part count: the trivial d=1 de-interleave (`H` is the single part on
+    /// the LDE coset) or the d=2 quotient split H₀/H₁. Both keep the parts
+    /// device-resident — the commit tree and the R4 openings read `handle.m`, while
+    /// R3 and R4 DEEP read the host part Vec's length (see
+    /// [`crate::gpu_lde::try_comp_h_to_slabs_dev`] for the invariant that ties the
+    /// two together). `None` → the caller falls back to the host path. Shared by the
+    /// R2 producer and the `xcheck` mirror so the two cannot drift. `want_host` gates
+    /// the d=2 host drain only — d=1 tables are never device-only, so they always
+    /// keep their host part.
+    #[cfg(feature = "cuda")]
+    fn decompose_comp_h_dev(
+        number_of_parts: usize,
+        h_dev: &math_cuda::constraint_interp::GpuCompH,
+        domain: &Domain<Field>,
+        twiddles: &LdeTwiddles<Field>,
+        want_host: bool,
+    ) -> Option<(
+        Vec<Vec<FieldElement<FieldExtension>>>,
+        math_cuda::lde::GpuLdeExt3,
+    )> {
+        if number_of_parts == 1 {
+            // d=1 is never device-only (`device_only_for`'s degree gate admits only
+            // d=2), so the single part is always kept on host — `want_host` must
+            // hold, and the d=1 helper ignores it by design.
+            debug_assert!(
+                want_host,
+                "d=1 composition parts are never device-only; want_host must hold"
+            );
+            // The d=1 helper trusts `h_dev.num_rows` as the LDE size; the d=2 arm
+            // gets an incidental domain check via `weights.len() == n`. Pin the
+            // same invariant here so a domain/`H` size mismatch can't slip through.
+            debug_assert_eq!(
+                h_dev.num_rows,
+                domain.interpolation_domain_size * domain.blowup_factor,
+                "d=1 H row count must equal the LDE domain size"
+            );
+            crate::gpu_lde::try_comp_h_to_slabs_dev::<Field, FieldExtension>(h_dev)
+        } else {
+            crate::gpu_lde::try_decompose_extend_d2_dev::<Field, FieldExtension>(
+                h_dev,
+                twiddles.inv_2x(domain),
+                &twiddles.composition(domain).weights,
+                want_host,
+            )
+        }
+    }
+
     /// Algebraically decompose H(x) = H₀(x²) + x·H₁(x²) on the LDE coset, then
     /// extend each half to the full LDE domain. This replaces the expensive
     /// iFFT(2N) + break_in_parts + FFT(2N)×2 pipeline with:
@@ -1671,7 +1724,8 @@ pub trait IsStarkProver<
         #[cfg(feature = "cuda")]
         let mut downloaded_h: Option<Vec<FieldElement<FieldExtension>>> = None;
         #[cfg(feature = "cuda")]
-        if number_of_parts == 2 && !crate::gpu_lde::gpu_force_downgrade() {
+        if (number_of_parts == 1 || number_of_parts == 2) && !crate::gpu_lde::gpu_force_downgrade()
+        {
             // Serializing this window across tables (device constraint eval +
             // decompose, where H is born) empirically eliminates a transient
             // whole-buffer H corruption seen under concurrent R2 windows on
@@ -1690,11 +1744,15 @@ pub trait IsStarkProver<
                 boundary_coefficients,
                 &round_1_result.rap_challenges,
             ) {
-                match crate::gpu_lde::try_decompose_extend_d2_dev::<Field, FieldExtension>(
+                let want_host = !round_1_result.lde_trace.host_trace_empty();
+                // num_parts==1 de-interleaves `H` (the single part); num_parts==2
+                // runs the degree-2 quotient split. Both keep the parts resident.
+                match Self::decompose_comp_h_dev(
+                    number_of_parts,
                     &h_dev,
-                    twiddles.inv_2x(domain),
-                    &twiddles.composition(domain).weights,
-                    !round_1_result.lde_trace.host_trace_empty(),
+                    domain,
+                    twiddles,
+                    want_host,
                 ) {
                     Some((parts, handle)) => {
                         gpu_composition_parts = Some(handle);
@@ -1709,7 +1767,13 @@ pub trait IsStarkProver<
         }
         #[cfg(feature = "cuda")]
         if let Some(h) = downloaded_h.take() {
-            precomputed_parts = Some(Self::decompose_and_extend_d2(&h, domain, twiddles));
+            // num_parts==1: the downloaded `H` IS the single part (no host
+            // decompose); num_parts==2: run the host degree-2 split + extend.
+            precomputed_parts = Some(if number_of_parts == 1 {
+                vec![h]
+            } else {
+                Self::decompose_and_extend_d2(&h, domain, twiddles)
+            });
         }
         #[cfg(not(feature = "cuda"))]
         let precomputed_parts: Option<Vec<Vec<FieldElement<FieldExtension>>>> = None;
@@ -4194,7 +4258,14 @@ pub trait IsStarkProver<
             boundary_coefficients,
             &round_1_result.rap_challenges,
         );
-        let host_parts = Self::decompose_and_extend_d2(&host_h, domain, twiddles);
+        // num_parts==1: `H` IS the single part (no host decompose); num_parts==2:
+        // the degree-2 split. Mirrors the R2 producer so the compare is apples-to-apples.
+        let number_of_parts = air.composition_poly_degree_bound(trace_length) / trace_length;
+        let host_parts = if number_of_parts == 1 {
+            vec![host_h]
+        } else {
+            Self::decompose_and_extend_d2(&host_h, domain, twiddles)
+        };
         let device_parts: Option<Vec<Vec<FieldElement<FieldExtension>>>> = if round_2_result
             .lde_composition_poly_evaluations
             .first()
@@ -4267,13 +4338,8 @@ pub trait IsStarkProver<
                 &round_1_result.rap_challenges,
             )
             .and_then(|h_dev| {
-                crate::gpu_lde::try_decompose_extend_d2_dev::<Field, FieldExtension>(
-                    &h_dev,
-                    twiddles.inv_2x(domain),
-                    &twiddles.composition(domain).weights,
-                    true,
-                )
-                .map(|(parts, _handle)| parts)
+                Self::decompose_comp_h_dev(number_of_parts, &h_dev, domain, twiddles, true)
+                    .map(|(parts, _handle)| parts)
             });
         let rerun_verdict = match &rerun {
             None => "device rerun declined".to_string(),
diff --git a/crypto/stark/tests/gpu_constraint_interp.rs b/crypto/stark/tests/gpu_constraint_interp.rs
index 625795244..eef21953a 100644
--- a/crypto/stark/tests/gpu_constraint_interp.rs
+++ b/crypto/stark/tests/gpu_constraint_interp.rs
@@ -115,6 +115,37 @@ fn all_ops_program() -> ConstraintProgram<Gl, Ext> {
     b.finish(1) // 1 base root, 2 ext roots
 }
 
+/// DECODE-shaped program: a preprocessed LogUp-only table declares
+/// `EmptyConstraints` (no base transition roots) — only the framework's aux
+/// LogUp ext roots. Mirrors that shape (`num_base == 0`) so the composition
+/// kernel is exercised on a program with zero base-dim roots, the case the
+/// DECODE `num_parts == 1` device path relies on.
+fn decode_shaped_program() -> ConstraintProgram<Gl, Ext> {
+    let mut b = IrBuilder::<Gl, Ext>::new();
+
+    // Root 0 (ext): main(0,0)·challenge(0) + alpha_pow(1)·aux(0,0) − table_offset.
+    let m0 = b.main(0, 0);
+    let ch = b.challenge(0);
+    let ap = b.alpha_power(1);
+    let a0 = b.aux(0, 0);
+    let off = b.table_offset();
+    let t1 = b.mul(m0, ch); // base × ext → ext (auto-embed)
+    let t2 = b.mul(ap, a0); // ext × ext
+    let s = b.add(t1, t2);
+    let r0 = b.sub(s, off);
+    b.emit(0, r0);
+
+    // Root 1 (ext): aux(1,0) − aux(0,0) + const_ext (next-row aux read).
+    let a0n = b.aux(1, 0);
+    let a0c = b.aux(0, 0);
+    let ce = b.const_ext(ext3(5, 4, 3));
+    let d = b.sub(a0n, a0c);
+    let r1 = b.add(d, ce);
+    b.emit(1, r1);
+
+    b.finish(0) // 0 base roots, 2 ext roots — the EmptyConstraints (LogUp-only) shape
+}
+
 /// Derive the trace/uniform footprint the program actually touches, so the
 /// harness works for any program (synthetic or real): #main cols, #aux cols,
 /// #rap challenges, #alpha powers, and the max frame offset.
@@ -519,6 +550,24 @@ fn check_composition(prog: &ConstraintProgram<Gl, Ext>, label: &str, seed: u64)
             enc(&h_cpu)
         );
     }
+
+    // evaluate_dev parity: the device-resident `H` (keep=true) that the
+    // num_parts==1 slab path consumes must equal the host-drained `H`
+    // (keep=false) bit-for-bit — same kernel, only the D2H differs. Confirms the
+    // composition path engages AND agrees on a device-resident `H`, including
+    // the empty-base (LogUp-only) program shape.
+    let dev = match try_eval_composition_gpu(
+        prog, &main, &aux, &rap, &alpha, &offset, NEXT_STEP, NUM_ROWS, &inputs, true,
+    ) {
+        Some(stark::constraint_ir::gpu_interp::GpuComposition::Dev(h)) => h,
+        _ => panic!("[{label}] GPU composition Dev (keep) path must engage"),
+    };
+    let dev_raw = math_cuda::constraint_interp::download_comp_h(&dev)
+        .unwrap_or_else(|e| panic!("[{label}] download_comp_h failed: {e:?}"));
+    assert_eq!(
+        dev_raw, gpu,
+        "[{label}] evaluate_dev (keep=true) H != host-drained H, seed {seed:#x}"
+    );
 }
 
 #[test]
@@ -527,3 +576,19 @@ fn gpu_composition_matches_cpu_oracle_all_ops() {
         check_composition(&all_ops_program(), "ALL_OPS_COMP", seed);
     }
 }
+
+/// num_parts==1 de-risk: the DECODE-shaped (empty-base, LogUp-only) program must
+/// evaluate on the GPU composition kernel, match the CPU oracle, and produce a
+/// device-resident `H` bit-identical to the host-drained one.
+///
+/// This closes the num_parts==1 device path at the unit level — the `H` the slab
+/// de-interleave consumes. The end-to-end counterpart (de-interleave -> commit ->
+/// OOD -> DEEP -> FRI -> openings, then verify) is `prover/tests/cuda_d1_path.rs`,
+/// which needs a lowered `LAMBDA_VM_GPU_LDE_THRESHOLD` because no fixture crosses
+/// the default for a d=1 table; `make test-cuda-d1` runs it.
+#[test]
+fn gpu_composition_matches_cpu_oracle_decode_shaped() {
+    for seed in [0x0123_4567_89AB_CDEF, 0xDEAD_BEEF_CAFE_F00D, 7] {
+        check_composition(&decode_shaped_program(), "DECODE_SHAPED_COMP", seed);
+    }
+}
diff --git a/prover/tests/cuda_d1_path.rs b/prover/tests/cuda_d1_path.rs
new file mode 100644
index 000000000..da449ee47
--- /dev/null
+++ b/prover/tests/cuda_d1_path.rs
@@ -0,0 +1,80 @@
+//! End-to-end coverage for the num_parts==1 (DECODE) device DEEP/FRI path.
+//!
+//! No fixture crosses the default GPU LDE threshold for a num_parts==1 table, so
+//! this binary lowers `LAMBDA_VM_GPU_LDE_THRESHOLD` (via `make test-cuda-d1`)
+//! until DECODE engages and the whole d=1 wiring — de-interleave -> R2 commit ->
+//! R3 OOD -> R4 DEEP -> FRI -> openings — runs end to end, validated by the
+//! release query-0 composition canary and the final verify.
+//!
+//! Fixture and threshold are one choice. There are exactly two d=1 tables (a d=1
+//! table is one with a single bus interaction): DECODE, sized from the guest's
+//! instruction count, and KECCAK_RC, fixed at `NUM_ROWS = 32` => LDE 64. DECODE's
+//! ROM comes from the ELF and not from cycles, so every `fib_iterative_*` variant
+//! is 13 executable words => 16 rows => LDE 32 — below KECCAK_RC's 64, which means
+//! no threshold isolates DECODE with a fib fixture. `all_instructions_64` is 66
+//! executable words => 128 rows => LDE 256, so at threshold 128 DECODE engages and
+//! KECCAK_RC declines: a nonzero counter uniquely attributes to DECODE.
+//!
+//! Its own binary (not another test in `cuda_path_integration.rs`) on purpose:
+//! `gpu_lde_threshold()` caches the env in a `OnceLock` on first read, so the
+//! lowered value must be the one the process sees before any prove — which only
+//! holds if this is the sole test in the process.
+//!
+//! `#[ignore]`'d so the no-GPU CI path skips it. Single test thread: the dispatch
+//! counters it asserts on are process-global.
+#![cfg(feature = "cuda")]
+
+use lambda_vm_prover::test_utils::asm_elf_bytes;
+use lambda_vm_prover::{prove, verify};
+use stark::gpu_lde::{gpu_comp_h_slabs_calls, reset_all_gpu_call_counters};
+
+/// The fixture whose DECODE ROM crosses the lowered threshold: 66 executable
+/// words -> 128 rows -> LDE 256.
+const FIXTURE: &str = "all_instructions_64";
+/// DECODE's LDE for [`FIXTURE`], and the LDE of the only other d=1 table. The
+/// threshold must fall between them so the counter attributes to DECODE alone.
+const DECODE_LDE: usize = 256;
+const KECCAK_RC_LDE: usize = 64;
+
+/// With the LDE threshold lowered so the DECODE (num_parts==1) table engages, the
+/// device de-interleave path (`gpu_comp_h_slabs_calls`) must fire and the proof —
+/// whose DECODE DEEP/FRI now ran on device — must still verify. Guards a silent
+/// CPU fallback (counter == 0) and a bad-layout regression (fires but the proof
+/// fails verification); the in-prove release query-0 canary guards the
+/// composition-row gather on top.
+#[test]
+#[ignore = "requires GPU + a lowered LAMBDA_VM_GPU_LDE_THRESHOLD; run via `make test-cuda-d1`"]
+fn gpu_num_parts_1_decode_path_fires_and_verifies() {
+    // Pin the window rather than just "below the default": a threshold anywhere
+    // outside (KECCAK_RC_LDE, DECODE_LDE] silently measures the wrong table (or no
+    // table), which is exactly the failure this constant pair exists to prevent.
+    let thr: usize = std::env::var("LAMBDA_VM_GPU_LDE_THRESHOLD")
+        .ok()
+        .and_then(|s| s.parse().ok())
+        .unwrap_or(0);
+    assert!(
+        thr > KECCAK_RC_LDE && thr <= DECODE_LDE,
+        "run via `make test-cuda-d1`: LAMBDA_VM_GPU_LDE_THRESHOLD must land in \
+         ({KECCAK_RC_LDE}, {DECODE_LDE}] so {FIXTURE}'s DECODE (LDE {DECODE_LDE}) engages the \
+         device path while KECCAK_RC (LDE {KECCAK_RC_LDE}) declines; got {thr}"
+    );
+
+    let elf = asm_elf_bytes(FIXTURE);
+    // Warm-up amortises PTX load + pool warm-up so the measured prove reflects
+    // steady state (mirrors cuda_path_integration.rs).
+    let _ = prove(&elf).expect("warm-up prove");
+    reset_all_gpu_call_counters();
+
+    let proof = prove(&elf).expect("prove");
+
+    assert!(
+        gpu_comp_h_slabs_calls() > 0,
+        "num_parts==1 device de-interleave path did not fire: DECODE (the only d=1 table \
+         above the threshold for {FIXTURE}) did not take it, so the d=1 DEEP/FRI wiring \
+         was not exercised"
+    );
+    assert!(
+        verify(&proof, &elf).expect("verify"),
+        "num_parts==1 device DEEP/FRI proof failed verification"
+    );
+}
diff --git a/prover/tests/cuda_path_integration.rs b/prover/tests/cuda_path_integration.rs
index b8e540a3b..dd841d7b7 100644
--- a/prover/tests/cuda_path_integration.rs
+++ b/prover/tests/cuda_path_integration.rs
@@ -76,11 +76,14 @@ fn gpu_path_fires_end_to_end() {
     // path.
     assert!(gpu_bary_calls() > 0, "R3 GPU barycentric did not fire");
 
-    // R2 GPU composition-poly LDE. Fires via one of two paths depending on the
+    // R2 GPU composition-poly LDE. Fires via one of three paths depending on the
     // AIR's `number_of_parts`: the fused two-halves quotient decomposition for
     // the common degree-2 case (`== 2`, counted by `gpu_extend_halves_calls`),
-    // or the batched parts LDE for `> 2` (counted by `gpu_parts_lde_calls`).
-    // fib_iterative_1M only exercises the degree-2 path, so assert on either.
+    // the batched parts LDE for `> 2` (counted by `gpu_parts_lde_calls`), or the
+    // d=1 de-interleave (`== 1`, counted by `gpu_comp_h_slabs_calls` — covered
+    // separately by `cuda_d1_path.rs`, since no d=1 table here crosses the
+    // default LDE threshold). fib_iterative_1M only exercises the degree-2 path,
+    // so assert on either of the two counted here.
     assert!(
         gpu_extend_halves_calls() + gpu_parts_lde_calls() > 0,
         "R2 GPU composition LDE did not fire (neither two-halves d2 nor parts>2 path)"
diff --git a/scripts/gpu_test.sh b/scripts/gpu_test.sh
index 1c5458a67..7d40f9f67 100755
--- a/scripts/gpu_test.sh
+++ b/scripts/gpu_test.sh
@@ -3,11 +3,12 @@
 # gpu_test.sh — run the CUDA-only test groups on a GPU box.
 #
 # Exercises the CUDA path, which CPU CI can't (GitHub runners have no GPU):
-#   1. math-cuda kernel parity         (make test-math-cuda)
-#   2. end-to-end GPU dispatch + proof  (make test-cuda-integration)
-#   3. GPU error-path / CPU fallback    (make test-cuda-fallback)
-#   4. prover/stark/crypto/ecsm suite   (make test-prover-cuda) — CPU CI's prover tests on GPU
-#   5. comprehensive all-instructions   (make test-prover-comprehensive-cuda)
+#   1. math-cuda kernel parity          (make test-math-cuda)
+#   2. end-to-end GPU dispatch + proof   (make test-cuda-integration)
+#   3. num_parts==1 (DECODE) device path (make test-cuda-d1)
+#   4. GPU error-path / CPU fallback     (make test-cuda-fallback)
+#   5. prover/stark/crypto/ecsm suite    (make test-prover-cuda) — CPU CI's prover tests on GPU
+#   6. comprehensive all-instructions    (make test-prover-comprehensive-cuda)
 #
 # Runs on the rented Vast box from the gpu-tests.yml merge-queue workflow. All groups
 # run even if one fails (so the log shows every failure); the script exits non-zero if ANY
@@ -41,7 +42,7 @@ nvidia-smi --query-gpu=name,driver_version,compute_cap --format=csv,noheader
 
 # --- Build the guest ELFs the tests prove ---------------------------------------
 # math-cuda parity needs none; cuda_path_integration / cuda_fallback prove an asm ELF; the
-# prover suite (Groups 4 & 5) proves asm AND rust guests. Build both up front.
+# prover suite (Groups 5 & 6) proves asm AND rust guests. Build both up front.
 log "compiling guest programs (asm + rust)"
 make compile-programs-asm
 make compile-programs-rust
@@ -57,9 +58,10 @@ run() {  # $1 = make target
 }
 run test-math-cuda                  # Group 1: kernel parity
 run test-cuda-integration           # Group 2: end-to-end GPU dispatch + proof verifies
-run test-cuda-fallback              # Group 3: GPU error -> CPU fallback still verifies
-run test-prover-cuda                # Group 4: prover/stark/crypto/ecsm suite on the GPU path
-run test-prover-comprehensive-cuda  # Group 5: comprehensive all-instructions prove on GPU
+run test-cuda-d1                    # Group 3: num_parts==1 (DECODE) device DEEP/FRI + verify
+run test-cuda-fallback              # Group 4: GPU error -> CPU fallback still verifies
+run test-prover-cuda                # Group 5: prover/stark/crypto/ecsm suite on the GPU path
+run test-prover-comprehensive-cuda  # Group 6: comprehensive all-instructions prove on GPU
 
 if [ "$fail" -ne 0 ]; then
     log "FAILED — one or more GPU test groups failed"