perf(prover): device-resident rounds 2-4 and fused NTT for GPU continuations

[ColoCarletti]

Second GPU optimization round on top of #863: ethrex 10tx continuations on RTX 5090 go from 12.01s to 9.75s (-18.8%, 8 ABBA pairs, sd 1.3%). Proofs are unchanged and cross-verify against the previous prover.

• R2 composition parts, their Merkle tree, R3 OOD evals and R4 openings now read the resident device buffer instead of round-tripping ~200MB per table through the host.
• FRI keeps each layer's evaluations on device; the query phase gathers only the ~219 opened positions per layer.
• The multiplicity subset tree of preprocessed tables stays resident (no per-epoch node download and host rebuild).
• Kernels: first 8 NTT levels fused in shared memory on the row-major path, small Merkle levels built in one launch, batch-inverse total inverted on device, big trace uploads staged through pinned memory.
• LogUp aux build no longer pays a host transpose the GPU path discards.

[ColoCarletti]

/bench-gpu

[github-actions]

GPU Benchmark (ABBA) — ac3dee9e36 vs main (14 pairs)

_{RTX 5090 · Vast.ai datacenter @ $0.8013888888888888/hr · prover/cuda · ethrex 100tx continuations · drift-free A/B/B/A}

=== ABBA paired result  (improvement: - = PR faster) ===
  pairs: 14   mean A (PR): 60.956s   mean B (base): 74.924s

  [parametric] paired-t   mean -18.64%   sd 0.72%   se 0.19%
               95% CI: [-19.06%, -18.22%]   (t df=13 = 2.16)
  [robust]     median -18.48%   Wilcoxon W+=0 W-=105  p(exact)=0.0001  (z=-3.26)

  --- server stability (this run; compare across servers) ---
  run-to-run jitter:    A CV 1.02%   B CV 0.62%        (lower = steadier)
  within-session drift: -0.15% over the run, 1st->2nd half -0.13%
    (jitter -> Tier-1 cached gate floor; drift -> whether the cached baseline can be trusted)

  VERDICT: REAL IMPROVEMENT - PR faster by ~18.64% (t-CI and Wilcoxon agree)

  raw pairs: /tmp/abba_run/pairs.csv

_{- = PR faster. Trust the verdict when paired-t and Wilcoxon agree.}

[ColoCarletti]

/bench-gpu

[ColoCarletti]

/ai-review

[github-actions]

Codex Code Review

• Medium — inverse.rs: GPU inversion now silently maps a zero batch product to zero, whereas the prior host path returned Err. Any zero denominator makes every “inverse” zero, allowing the prover to continue with invalid LogUp/barycentric values instead of falling back or aborting. Preserve zero detection in the device path and return an error before consuming the result.

[github-actions]

AI Review

PR #875 · 16 changed files

Findings

Status	Sev	Location	Finding	Found by
confirmed	medium	crypto/math-cuda/kernels/inverse.cu:284	Silent zero handling in batch inverse Fermat inversion kernel	nemotron openrouter/nvidia/nemotron-3-ultra-550b-a55b glm openrouter/z-ai/glm-5.2
confirmed	low	crypto/crypto/src/merkle_tree/merkle.rs:177	is_root_only stole from_root's doc comment (doc misplacement)	glm openrouter/z-ai/glm-5.2
confirmed	low	crypto/crypto/src/merkle_tree/merkle.rs:269	get_proof_by_pos lacks root-only guard	nemotron openrouter/nvidia/nemotron-3-ultra-550b-a55b
confirmed	low	crypto/math-cuda/src/device.rs:425	htod_via API footgun: potential self-deadlock when holding PendingD2H	nemotron openrouter/nvidia/nemotron-3-ultra-550b-a55b

Status column reflects the verdict from the verifier: deepseek-verifier (openrouter/deepseek/deepseek-v4-pro).

AI-002: Silent zero handling in batch inverse Fermat inversion kernel

• Status: confirmed
• Severity: medium
• Location: crypto/math-cuda/kernels/inverse.cu:284
• Found by: nemotron:openrouter/nvidia/nemotron-3-ultra-550b-a55b, glm:openrouter/z-ai/glm-5.2
• Verified by: deepseek-verifier:openrouter/deepseek/deepseek-v4-pro
• Rejected by: -

Claim

The invert_total_ext3 kernel returns zero for zero input instead of signaling an error, which can mask invalid denominators in batch inverse.

Evidence

The kernel computes a^(p^3-2) via square-and-multiply. If a == 0, all multiplications yield 0 and the result is silently 0. The host-side invert_ext3_host returns an error for zero input, but the device path (used for n==1 and total inversion) does not. Comment at inverse.rs:98-101 acknowledges: 'a zero total maps silently to zero instead of Err. Unreachable with honest inputs... callers must not rely on a zero-total error.'

Suggested fix

Either keep the host-side invert_ext3_host for the total (preserving the Err→CPU-fallback behavior on degenerate inputs), or have the kernel/host driver detect a zero total (e.g. D2H the 3 u64 total and check) and return Err so the dispatcher falls back to CPU rather than shipping zero inverses.

AI-006: `is_root_only` stole `from_root`'s doc comment (doc misplacement)

• Status: confirmed
• Severity: low
• Location: crypto/crypto/src/merkle_tree/merkle.rs:177
• Found by: glm:openrouter/z-ai/glm-5.2
• Verified by: deepseek-verifier:openrouter/deepseek/deepseek-v4-pro
• Rejected by: -

Claim

The new is_root_only method was inserted between from_root's existing doc comment and the from_root function itself. As a result, the doc comment (lines 171-176, originally describing from_root: 'Create a root only Merkle tree placeholder... must NOT be called on it.') now attaches to is_root_only, and from_root (line 188) is left with no doc comment. is_root_only's rendered docs read awkwardly, mixing from_root's description with the appended 'True when this tree carries only its root...' line.

Evidence

Lines 171-176 show the doc comment block immediately preceding pub fn is_root_only (line 177), and pub fn from_root (line 188) follows with no preceding /// doc lines. Rust attaches a doc comment to the immediately following item, so from_root lost its documentation.

Suggested fix

Move is_root_only (with its own dedicated doc comment) to a location that does not split from_root from its doc, e.g. place it after from_root or give from_root its doc back and write a fresh doc block for is_root_only.

AI-007: get_proof_by_pos lacks root-only guard

• Status: confirmed
• Severity: low
• Location: crypto/crypto/src/merkle_tree/merkle.rs:269
• Found by: nemotron:openrouter/nvidia/nemotron-3-ultra-550b-a55b
• Verified by: deepseek-verifier:openrouter/deepseek/deepseek-v4-pro
• Rejected by: -

Claim

The get_proof_by_pos method does not check is_root_only() and will panic or return an error if called on a root-only tree.

Evidence

Root-only trees (created via from_root) have empty nodes vec. get_proof_by_pos calls build_merkle_path which calls node_get on indices that don't exist. The new is_root_only() method (lines 177-186) correctly identifies these trees, but get_proof_by_pos doesn't use it. Current code paths avoid this via defensive checks (e.g., prover.rs:2652 asserts !main_commit.tree.is_root_only()), but the method itself is unsafe to call.

Suggested fix

Add an early check in get_proof_by_pos: if self.is_root_only() { return None; } or return a specific error variant.

AI-008: htod_via API footgun: potential self-deadlock when holding PendingD2H

• Status: confirmed
• Severity: low
• Location: crypto/math-cuda/src/device.rs:425
• Found by: nemotron:openrouter/nvidia/nemotron-3-ultra-550b-a55b
• Verified by: deepseek-verifier:openrouter/deepseek/deepseek-v4-pro
• Rejected by: -

Claim

The htod_via function locks the per-worker staging slot mutex. If called from a thread that already holds a PendingD2H guard on the same slot, it will deadlock.

Evidence

Comment at lines 415-418: 'Corollary: never call htod_via/async_dtoh_via on the same slot from the thread holding a live PendingD2H — the non-reentrant slot mutex self-deadlocks.' The PendingD2H struct holds a MutexGuard<'a, PinnedStaging> (line 349). Current usage in lde.rs:470 calls htod_via without holding PendingD2H, so it's safe today, but the API is fragile.

Suggested fix

Either make the staging slot reentrant (e.g., use parking_lot::ReentrantMutex), or restructure to avoid the footgun by separating H2D and D2H staging slots per worker.

Reviewer Lanes

Lane	Model	Prompt	Status	Findings
glm	openrouter/z-ai/glm-5.2	general	success	2
kimi	openrouter/moonshotai/kimi-k2.7-code	general	success	4
minimax	minimax/MiniMax-M3	general	error: opencode failed (provider/auth/runtime error) and no findings were submitted	0
moonmath	zro/minimax-m3	general	error: opencode failed (provider/auth/runtime error) and no findings were submitted	0
nemotron	openrouter/nvidia/nemotron-3-ultra-550b-a55b	general	success	5

Verification Lanes

Lane	Model	Status	Confirmed	Rejected	Uncertain
deepseek-verifier	openrouter/deepseek/deepseek-v4-pro	success	4	5	0

Native Codex and Claude reviews run separately and post their own comments. They are not included in this structured provenance report.

Discarded candidates (5) — rejected by the verifier

• Merkle tail kernel has divergent __syncthreads causing deadlock (crypto/math-cuda/kernels/keccak.cu:415, found by kimi:openrouter/moonshotai/kimi-k2.7-code) — The kernel does NOT have early returns. The __syncthreads() at line 418 is OUTSIDE the for loop. Threads with threadIdx.x >= n_pairs have zero loop iterations (the loop condition tid < n_pairs is never true) and fall through to hit the barrier — all threads in the block converge. This is a standard CUDA grid-stride pattern: the barrier is after the loop completes for all threads, not inside a divergent branch. The finding incorrectly claims threads 'return early before the first barrier' — there is no return statement anywhere in keccak_merkle_tail's while-loop body.
• Row-major 8-level fused NTT twiddle index needs parity coverage (crypto/math-cuda/kernels/ntt.cu:448, found by kimi:openrouter/moonshotai/kimi-k2.7-code) — The finding is too speculative — its own evidence admits 'The concern is low confidence.' The twiddle math at lines 456-460 computes ggp = ((uint32_t)rb << 7) + i followed by (ggp & high_mask) + (ggp >> remaining_high_bits) and final mask by (1u << gs) - 1u. This exactly mirrors ntt_dit_8_levels_batched (lines 270-274): ggp = (blockIdx.x << 7) + i then ((ggp & high_mask) << base_step) + (ggp >> remaining_high_bits) with base_step==0, which simplifies to (ggp & high_mask) + (ggp >> remaining_high_bits). The row-major grid uses blockIdx.y as the 256-row block index and rb grid-strides over blockIdx.y, which is the correct analog of blockIdx.x in the batched column-major kernel. No evidence of incorrect twiddle indexing.
• Pinned staging H2D slot not protected against concurrent async use (crypto/math-cuda/src/device.rs:677, found by kimi:openrouter/moonshotai/kimi-k2.7-code) — htod_via (lines 640-679) holds the staging slot mutex for its entire duration: from let mut staging = slot.lock().unwrap() through staging.sync_event(). The comment at lines 77-79 of record_event explicitly states: 'slot access is serialized by its mutex, so a recorded event is always synchronized before re-recording.' No concurrent caller (including async_dtoh_via) can re-record the event because the mutex is held. The immediate sync_event() after record_event() is intentional — it drains the stream to ensure the H2D DMA has completed so src_host is safe to reuse. There is no stale-event hazard under mutex exclusion.
• Synchronous D2H in gather_ext3_at inconsistent with async pinned-staging pattern (crypto/math-cuda/src/fri.rs:286, found by nemotron:openrouter/nvidia/nemotron-3-ultra-550b-a55b) — gather_ext3_at (lines 286-287) transfers a small amount of data (3*q u64s where q is the number of FRI queries, typically a handful). Using stream.clone_dtoh + synchronize for a few hundred bytes is simpler and adds negligible latency. The pinned-staging pattern (async_dtoh_via + PendingD2H guard) introduces complexity — lock management, event recording, Drop-based safety — that provides no meaningful benefit for small transfers. This is an intentional engineering choice, not a problem. The finding itself admits 'performance impact is minimal.'
• FRI symmetric eval gather position matches CPU fallback (crypto/stark/src/gpu_lde.rs:2499, found by kimi:openrouter/moonshotai/kimi-k2.7-code) — The finding itself states 'No issue found; included only to record the verification.' The position computation ((iota >> l) ^ 1) as u32 at line 2499 correctly matches the CPU's evaluation[index ^ 1] logic. The finding confirms correctness across multiple paragraphs of analysis. There is no claim of an actual issue to verify.

Raw lane outputs, candidates, final issues, and model metrics are uploaded as workflow artifacts.