The Benchmark That Failed Twice: Diagnosing GPU Memory Contention in Phase 9 PCIe Optimization

Introduction

In the course of optimizing the cuzk SNARK proving engine for Filecoin's PoRep protocol, the assistant reached a critical juncture: the Phase 9 PCIe Transfer Optimization had been implemented, rebuilt, and was ready for its first real test under the intended dual-worker configuration. The message at <msg id=2425> captures this moment — a benchmark run that would reveal whether the carefully engineered fix for GPU memory contention had succeeded, or whether deeper issues lurked beneath the surface. The answer, delivered in stark red FAILED markers, was the latter. But as with all debugging sessions, the failure itself was the most informative result.

Context: The Phase 9 Optimization

The Phase 9 optimization targeted two root causes of GPU idle gaps identified in the Phase 8 baseline. The first change (Tier 1) addressed the 6 GiB of non-pinned a/b/c polynomial uploads that were blocking the GPU pipeline. By pinning host memory with cudaHostRegister, allocating device buffers, and issuing async cudaMemcpyAsync transfers on a dedicated stream with event-based synchronization, the assistant aimed to move these uploads out of the critical path. The second change (Tier 3) eliminated per-batch hard sync stalls in the Pippenger MSM by introducing double-buffered host result buffers and deferring the sync() call to the next iteration, allowing GPU compute to overlap with DtoH transfers.

The initial implementation placed the pre-staging (allocation and async upload) outside the GPU mutex, allowing one worker to upload its a/b/c polynomials while the other worker held the GPU for kernel execution. This design assumed that the two workers' memory footprints would fit within the 16 GiB VRAM budget: the worker running kernels needed ~7.5 GiB peak, while the worker pre-staging needed ~6 GiB, totaling 13.5 GiB — theoretically comfortable.

The First Failure and the First Fix

The first benchmark run at <msg id=2417> revealed the flaw in this assumption. All five proofs failed with OOM errors. The daemon logs at <msg id=2418> showed prestage_setup=fallback err=2 (where err=2 is cudaErrorMemoryAllocation), followed by panics from the fallback path also failing with OOM. The root cause was clear: both workers tried to pre-stage simultaneously. With gpu_workers_per_device=2, worker 0 allocated 6 GiB for pre-staging, then worker 1 also tried to allocate 6 GiB — that's 12 GiB just for pre-staging on a 16 GiB GPU, leaving nothing for the rest of the pipeline.

The assistant's response at <msg id=2419> was a thorough analysis of the memory contention problem, weighing several approaches:

  1. A separate lightweight mutex for allocation only
  2. Using the same mutex with a try_lock pattern
  3. Moving pre-staging inside the main mutex The assistant ultimately chose option 3, reasoning that while this sacrificed the overlap benefit (uploads would now happen during the worker's own kernel time rather than overlapping with the other worker), the remaining benefits were still substantial: pinned memory for full PCIe bandwidth, truly async cudaMemcpyAsync that doesn't block the CPU thread, and event-based synchronization that lets NTT compute start as soon as the upload for that polynomial finishes. The fix was applied, the daemon was rebuilt and restarted, and the stage was set for the second benchmark run.

The Subject Message: A Second Round of Failure

The subject message at <msg id=2425> shows the assistant launching the benchmark with the fixed code:

[assistant] Daemon is ready. Now run the benchmark:
[bash] /home/theuser/curio/extern/cuzk/target/release/cuzk-bench batch -t porep --c1 /data/32gbench/c1.json -c 5 -j 3 2>&1 | tee /tmp/cuzk-phase9-bench.log

The output is devastating:

  [1/5] FAILED — 37.1s (prove=0 ms, queue=0 ms)
  [2/5] FAILED — 63.9s (prove=0 ms, queue=0 ms)
  [3/5] FAILED — 94.3s (prove=0 ms, queue=0 ms)
  [4/5] FAILED — 88.5s (prove=0 ms, queue=0 ms)
  [5/5] FAILED — 93.7s (prove=0 ms, queue=0 ms)

Every single proof failed. The prove=0 ms in each line is particularly telling — it means the proving step itself returned zero milliseconds, indicating that the benchmark harness detected a failure before any meaningful GPU work completed. The queue times are also zero, suggesting the proofs never made it past the submission stage.

The timings show a pattern: the first proof fails quickly (37.1s), while subsequent proofs take progressively longer (63.9s, 94.3s, 88.5s, 93.7s). This is characteristic of a system where the first failure corrupts state, and subsequent attempts take longer because they're waiting for timeouts or struggling with degraded resources. The daemon was still running and accepting requests, but every request resulted in a GPU panic.

The Deeper Investigation

What makes this message particularly interesting is what follows. The assistant immediately pivots to investigation at <msg id=2426>, examining the daemon logs with surgical precision:

CUZK_TIMING: prep_msm_ms=2782
CUZK_TIMING: prestage_setup=fallback err=2
CUZK_TIMING: b_g2_msm_ms=1790 num_circuits=1
CUZK_TIMING: prep_msm_ms=3472
thread 'tokio-runtime-worker' panicked at .../src/lib.rs:153:9:
cudaMalloc(&d_ptr, n * sizeof(T))@sppark-0.1.14/sppark/util/gpu_t.cuh:331 failed: "out of memory"

But then at <msg id=2427>, a more detailed log inspection reveals something surprising:

83:CUZK_TIMING: split_vectors_ms=0 setup_to_split_ms=0
85:CUZK_TIMING: prestage_setup=ok domain=134217728 lot_of_memory=1
87:CUZK_TIMING: gpu_tid=0 ntt_msm_h_ms=840

The first partition actually succeeded. prestage_setup=ok at line 85, and ntt_msm_h_ms=840 — the NTT+MSM completed in 840 ms, dramatically faster than the Phase 8 baseline of ~2430 ms. This confirmed that the optimization itself was working. But then the second partition failed.

The assistant's investigation at <msg id=2428-2430> uncovered two critical issues:

First, the domain size calculation revealed a memory problem. The log showed domain=134217728 (2^27 = 128M elements), which the assistant traced to the H SRS having domain_size + 1 = 67108865 points. The calculation lg2(points_h.size() - 1) + 1 yields lg2(67108864) + 1 = 26 + 1 = 27, meaning the pre-staging buffers were sized for 2^27 elements. At 32 bytes per element, d_a required 4 GiB and d_bc required 8 GiB — totaling 12 GiB for pre-staging alone.

Second, and more critically, the cleanup ordering was wrong. The assistant had placed the pre-staging allocation inside the mutex, but the deallocation of d_bc (the 8 GiB buffer) was happening after the mutex release. This meant that when worker 1 acquired the mutex for the next partition, worker 0's d_bc was still allocated. Worker 1 then tried to allocate its own 12 GiB on top of the 8 GiB still held by worker 0, totaling 20 GiB — well beyond the 16 GiB VRAM limit.

The Thinking Process Revealed

The assistant's reasoning at <msg id=2419> shows a careful, methodical approach to diagnosing the memory contention. The key insight was recognizing that the original design (pre-staging outside the mutex) couldn't work because both workers would try to allocate simultaneously. The assistant walked through several alternatives:

"The simplest fix: make the pre-staging allocation happen under a separate, lightweight mutex that only guards the cudaMalloc calls (not the kernel execution). Or simpler: do the device allocation inside the mutex but the upload outside."

But then the assistant realized a deeper problem:

"The fundamental issue: with 2 workers, we can't have BOTH pre-staged simultaneously. Only the worker that's about to acquire the kernel mutex should pre-stage. The other worker should wait."

This led to the decision to move pre-staging inside the mutex. However, the assistant's reasoning contained an incorrect assumption — that freeing d_bc after the mutex release was safe because it was "GPU resource cleanup." In reality, the freed memory needed to be returned to the pool before the next worker could allocate, and CUDA's memory management doesn't guarantee immediate reclamation across different allocation paths.

At <msg id=2430>, the assistant corrected this:

"I need to split the cleanup: free d_bc and GPU resources BEFORE unlocking the mutex (they use GPU resources), but unregister host pages AFTER (those are CPU-only)."

Input and Output Knowledge

To understand this message, the reader needs knowledge of: CUDA memory management and the distinction between pinned and pageable host memory; the cuzk dual-worker GPU interlock architecture developed in Phase 8; the Groth16 proof generation pipeline for Filecoin PoRep, including the a/b/c polynomial uploads and NTT/MSM compute phases; the SRS (Structured Reference String) structure and domain size calculations; and the memory accounting from earlier phases showing ~200 GiB peak system memory and ~16 GiB per-GPU VRAM.

The message creates critical output knowledge: the Phase 9 optimization works at the kernel level (840 ms NTT+MSM vs 2430 ms baseline) but fails at the system level due to memory contention between workers. This knowledge drives the subsequent fix — moving d_bc deallocation before the mutex release — and reveals a deeper principle for the architecture: GPU memory allocation and deallocation must be serialized under the same mutex that governs kernel execution, or a memory-aware allocator must be introduced.

Conclusion

The message at <msg id=2425> is a study in the value of failure. The benchmark didn't just fail — it failed in a way that exposed the precise boundary between kernel-level optimization and system-level integration. The PCIe transfer optimization itself was validated (840 ms NTT+MSM), but the dual-worker orchestration that made it practical was broken. The assistant's response — immediate investigation, log analysis, and iterative correction — demonstrates the disciplined approach required for systems programming at this scale. The fix that followed (freeing d_bc before the mutex release) was a single line movement, but the reasoning that led to it required understanding the entire memory lifecycle across two concurrent workers, a CUDA runtime's allocation behavior, and the subtle interactions between pinned memory, async transfers, and GPU mutex boundaries.