The Moment the Two-Lock Dream Died: Diagnosing a Fundamental GPU Interlock Flaw

Introduction

In the high-stakes world of Filecoin proof generation, every second counts. The cuzk SNARK proving engine had been through nine optimization phases, each squeezing more throughput out of the Groth16 pipeline. Phase 10 promised a breakthrough: a two-lock GPU interlock design that would allow multiple workers to overlap memory operations with kernel execution, theoretically eliminating the serialization that had bottlenecked the system. But when the benchmarks came back, the dream shattered. This message — message 2661 in the conversation — captures the precise moment of diagnosis, where the assistant discovers why the two-lock design is fundamentally flawed and traces the root cause to an OOM crash that reveals a deeper truth about CUDA memory management.

The Message in Full

The assistant writes:

There's the problem: OOM crash. The first worker successfully pre-staged (1.2s upload) and is holding compute_mtx with ~12 GB allocated. Then another worker enters compute_mtx (since the first one released mem_mtx), tries to allocate for its GPU kernels, and hits OOM because the first worker's buffers are still on the GPU.

>

Wait — that shouldn't happen because compute_mtx serializes the GPU kernel phase. Let me look more carefully at the sequence:

>

1. Worker 0: acquires mem_mtx, prestages OK (alloc 12 GB + upload), releases mem_mtx, then waits for compute_mtx 2. Worker 1: acquires mem_mtx, tries cudaMalloc → fails (Worker 0's buffers still on GPU), falls back, releases mem_mtx, waits for compute_mtx 3. Worker 0: acquires compute_mtx, runs kernels 4. Worker 1: acquires compute_mtx after Worker 0 releases → runs cudaDeviceSynchronize + pool trim → then runs kernels

>

But the crash is cudaMalloc inside gpu_t.cuh:331 — that's deep inside the kernel pipeline, not in our pre-staging code. This means even with cudaDeviceSynchronize + pool trim at the start of the compute_mtx fallback, there's still not enough VRAM when two workers' pre-staged buffers overlap.

>

Actually looking again: Worker 0 pre-staged 12 GB, then acquired compute_mtx and started kernels. But Worker 0's pre-staged d_a_prestaged + d_bc_prestaged are still allocated during the kernels — they're consumed during kernel execution but aren't freed until after the per_gpu threads join. Meanwhile, Worker 2 might also try to pre-stage...

>

The fundamental issue: pre-staging allocates ~12 GB that persists through the entire compute phase. When the next worker enters compute_mtx (after the first releases), the first worker's pre-staged buffers are freed, but the kernel pipeline itself also allocates memory internally (MSM temp, NTT workspace, etc.), and those are freed asynchronously.

>

Let me check more of the timing to understand the exact sequence: [bash] grep -E "(compute_wait|prestage_setup|fallback_sync|gpu_total|TIMELINE)" /home/theuser/cuzk-p10-daemon2.log | head -40

The Context: Nine Phases of Optimization

To understand why this message matters, we need to appreciate the journey that led here. The cuzk proving engine had been through eight previous optimization phases, each documented in design proposals and implemented with careful benchmarking. Phase 9 had achieved a 14.2% throughput improvement through PCIe transfer optimization, but dual-worker mode still showed performance gaps. The bottleneck had shifted from GPU utilization to what appeared to be lock contention on the CPU side.

The Phase 10 two-lock design was an ambitious architectural change. The original single-lock design used one mutex to serialize all GPU access — both memory allocation (pre-staging NTT domain data) and kernel execution (MSM, NTT, and other compute kernels). This meant that while one worker was uploading data to VRAM, another worker that wanted to run kernels had to wait, even though the GPU could theoretically overlap these operations. The two-lock split aimed to separate "memory phase" from "compute phase," allowing one worker to pre-stage data for its next partition while another worker was still running kernels on the current partition.

The design was documented, the code was written, and the first build succeeded. The assistant had tested a single proof (c=1, j=1) with 3 GPU workers and found it slow but correct. The real test was always going to be concurrent proofs, where the overlap benefits would manifest.

The Crash: What the Logs Revealed

When the assistant ran the first concurrent benchmark with 3 concurrent proofs and 3 GPU workers (c=3, j=3), the results were catastrophic. One proof failed outright, and the two that completed took 77.2 seconds per proof — nearly double Phase 9's 41.3 seconds. Something had gone fundamentally wrong.

The daemon log told the story. The assistant grepped for error messages and found:

thread 'tokio-runtime-worker' panicked at .../lib.rs:153:9:
cudaMalloc(&d_ptr, n * sizeof(T))@sppark-0.1.14/sppark/util/gpu_t.cuh:331 failed: "out of memory"

An OOM crash, deep inside the sppark library's GPU utility code. This wasn't in the pre-staging path — it was in the main kernel pipeline, where the actual MSM and NTT computations happen. The crash occurred because the GPU's 16 GB VRAM was exhausted.

The Reasoning: Tracing the Failure Sequence

The assistant's thinking in this message is a masterclass in diagnostic reasoning. It walks through the exact sequence of events that leads to the crash, correcting its own assumptions along the way.

Step 1: Initial hypothesis. The assistant first identifies the obvious problem: Worker 0 pre-stages ~12 GB, holds compute_mtx, and Worker 1 tries to allocate for kernels and hits OOM. But this seems contradictory — compute_mtx should serialize the kernel phase, preventing concurrent allocations.

Step 2: Deeper analysis. The assistant reconstructs the actual sequence:

  1. Worker 0 acquires mem_mtx, pre-stages 12 GB (allocates buffers + uploads data), releases mem_mtx, then waits for compute_mtx.
  2. Worker 1 acquires mem_mtx, tries to pre-stage, but cudaMalloc fails because Worker 0's buffers are still on GPU. Worker 1 falls back (skips pre-staging), releases mem_mtx, waits for compute_mtx.
  3. Worker 0 acquires compute_mtx, runs kernels.
  4. Worker 1 acquires compute_mtx after Worker 0 releases, runs cudaDeviceSynchronize + pool trim, then runs kernels. But the crash isn't in step 2 (pre-staging fallback) — it's in step 3 or 4, inside the kernel pipeline itself. This means even after the fallback path runs cudaDeviceSynchronize and cudaMemPoolTrimTo, there's still not enough VRAM. Step 3: The fundamental flaw. The assistant realizes the real issue: Worker 0's pre-staged buffers (d_a_prestaged and d_bc_prestaged) are ~12 GB and persist through the entire compute phase. They're consumed during kernel execution but aren't freed until after the per-GPU threads join. Meanwhile, the kernel pipeline itself allocates additional memory internally — MSM temporary buffers, NTT workspace, and other intermediate data. When Worker 1 enters compute_mtx and tries to run its own kernels, Worker 0's pre-staged buffers are still on the GPU (because Worker 0 hasn't finished its compute phase yet), AND Worker 0's kernel allocations may not have been fully freed (because CUDA memory management APIs like cudaDeviceSynchronize and pool trimming are asynchronous and device-global).

The Critical Insight: Device-Global Synchronization

This message reveals a deep truth about CUDA memory management that the Phase 10 design had overlooked. The APIs that the fallback path relies on — cudaDeviceSynchronize and cudaMemPoolTrimTo — are device-global operations. They affect all contexts on the GPU, not just the calling worker's allocations. This means:

  1. cudaDeviceSynchronize blocks the entire device, not just one stream or worker. Calling it inside compute_mtx serializes all GPU activity, defeating the purpose of splitting the lock.
  2. cudaMemPoolTrimTo trims memory pools device-wide, but it can't free memory that's still in use by active kernels or pinned allocations from other workers.
  3. Asynchronous deallocation means that when Worker 0 frees its pre-staged buffers after its compute phase, the memory isn't immediately available — CUDA's async pool may hold onto it for reuse within the same stream. The two-lock design assumed that splitting mem_mtx from compute_mtx would allow safe overlap. But the reality is that CUDA's memory model doesn't support this kind of fine-grained interleaving on a single GPU. The VRAM is a shared resource, and all workers on the same device compete for it. With 16 GB total and each worker needing ~12 GB for pre-staged buffers plus additional kernel workspace, two workers cannot coexist.

Assumptions and Their Failure

The Phase 10 design rested on several assumptions that this message exposes as incorrect:

Assumption 1: Memory allocation and kernel execution can be safely overlapped. The design assumed that splitting the lock would allow Worker 1 to pre-stage data while Worker 0 ran kernels. In practice, pre-staging allocates ~12 GB that persists through the entire compute phase, making overlap impossible without exceeding VRAM.

Assumption 2: Fallback to device sync + pool trim is sufficient. The design included a fallback path for when pre-staging fails: synchronize the device and trim memory pools, then proceed without pre-staging. But the crash occurred inside the kernel pipeline itself, meaning even the fallback couldn't reclaim enough memory.

Assumption 3: The two-lock split would improve throughput. The design assumed that reducing lock contention would improve GPU utilization. But the OOM crash and the resulting 77-second proof times showed the opposite — the design was not just suboptimal, it was broken.

Assumption 4: CUDA memory management APIs are fine-grained enough for this approach. The assistant's analysis reveals that cudaDeviceSynchronize and pool trimming are device-global operations that don't provide the isolation the design needed.

Input Knowledge Required

To fully understand this message, the reader needs:

  1. CUDA memory model: Understanding of device-global vs. stream-local operations, async memory pools, and the behavior of cudaMalloc/cudaFree in multi-threaded contexts.
  2. The Phase 10 two-lock design: Knowledge that mem_mtx controls memory allocation and compute_mtx controls kernel execution, and that workers acquire mem_mtx first, then compute_mtx.
  3. The pre-staging architecture: Understanding that each worker pre-stages NTT domain data (~12 GB) before the compute phase, and that these buffers are consumed during kernel execution but freed only after threads join.
  4. The cuzk pipeline structure: Knowledge of how synthesis workers produce partitions, how GPU workers consume them, and the role of partition_workers and gpu_workers_per_device configuration.
  5. The Groth16 proof structure: Understanding that each proof requires MSM (multi-scalar multiplication), NTT (number-theoretic transform), and other GPU-intensive operations that allocate significant VRAM.

Output Knowledge Created

This message produces several critical pieces of knowledge:

  1. Root cause identification: The OOM crash is caused by VRAM exhaustion when multiple workers' pre-staged buffers overlap in the 16 GB GPU memory.
  2. Failure sequence reconstruction: A precise timeline of how the two-lock design leads to the crash, step by step.
  3. Fundamental limitation discovered: CUDA's device-global synchronization APIs cannot support the fine-grained interleaving that the two-lock design requires.
  4. Design invalidation evidence: Concrete proof that the Phase 10 approach is not just underperforming but fundamentally broken, requiring a complete redesign.
  5. New diagnostic data: The timing logs (compute_wait, prestage_setup, fallback_sync, gpu_total) provide a detailed picture of where the design fails and why.

The Thinking Process: A Window into Diagnostic Reasoning

What makes this message particularly valuable is the visible thinking process. The assistant doesn't just report a failure — it walks through its reasoning in real time, correcting itself and deepening its understanding.

The message begins with a confident diagnosis: "There's the problem: OOM crash." But then immediately questions itself: "Wait — that shouldn't happen because compute_mtx serializes the GPU kernel phase." This self-correction is the hallmark of good diagnostic reasoning. The assistant doesn't accept the surface-level explanation; it digs deeper.

The numbered sequence (1-4) shows systematic reconstruction of the event timeline. Each step builds on the previous one, creating a clear causal chain. The assistant then identifies the specific line of code where the crash occurs — gpu_t.cuh:331 — and notes that this is "deep inside the kernel pipeline, not in our pre-staging code." This distinction is crucial: it means the fallback path designed to handle pre-staging failures doesn't protect against the deeper OOM.

The final paragraph — "Actually looking again" — represents a second pass of analysis. The assistant realizes that Worker 0's pre-staged buffers persist through the entire compute phase, and that the kernel pipeline's own internal allocations compound the problem. The phrase "The fundamental issue" signals the moment of true understanding: the design is not fixable with minor tweaks; it requires a fundamentally different approach.

The Aftermath: Abandoning Phase 10

This message is the turning point. The assistant goes on to revert the Phase 10 changes, returning to the Phase 9 single-lock design. The two-lock experiment is documented as a failure, and the lessons learned inform Phase 11, which targets DDR5 memory bandwidth contention instead of GPU lock contention.

The Phase 10 post-mortem, documented in cuzk-project.md, captures the key insight: "16 GB VRAM cannot accommodate pre-staged buffers from multiple workers simultaneously, and CUDA memory management APIs (cudaDeviceSynchronize, cudaMemPoolTrimTo) are device-global, defeating the lock split's purpose." This becomes a permanent record of a design that seemed promising but failed in practice.

Conclusion

Message 2661 is a pivotal moment in the cuzk optimization journey. It captures the precise instant when a promising architectural change is revealed to be fundamentally flawed, not through abstract reasoning but through concrete evidence: an OOM crash, timing logs, and careful step-by-step reconstruction of the failure sequence. The assistant's thinking process — questioning initial assumptions, reconstructing timelines, identifying the specific code paths involved, and ultimately recognizing the fundamental limitation — exemplifies the diagnostic rigor required for low-level systems optimization.

The message also demonstrates a crucial lesson for GPU programming: CUDA's memory model is not as fine-grained as CPU programmers might expect. Device-global synchronization, async memory pools, and the sheer size of GPU buffers (12 GB per worker on a 16 GB card) create constraints that cannot be overcome with lock splitting alone. Sometimes, the only winning move is to recognize that a design cannot work and to pivot to a different approach entirely.