The Moment of Failure: Debugging Phase 10's Two-Lock GPU Interlock

A Single Benchmark That Exposed a Fundamental CUDA Design Conflict

In the course of a deep optimization campaign targeting Filecoin's Groth16 proof generation pipeline, one message stands out as a critical inflection point: a failed correctness test that forced the team to confront a fundamental constraint of CUDA device architecture. This article examines that message in detail—a simple bash command and its output that revealed far more than a broken build.

The Message

The subject message, <msg id=2617>, is a single tool call from the AI assistant running a correctness test for Phase 10 of the cuzk SNARK proving engine optimization:

[assistant] [bash] echo "=== Phase 10 correctness test gw=3 c=1 j=1 ===" && time /home/theuser/curio/extern/cuzk/target/release/cuzk-bench batch \
  --type porep --c1 /data/32gbench/c1.json \
  -c 1 -j 1 2>&1
=== Phase 10 correctness test gw=3 c=1 j=1 ===
[2026-02-19T17:12:19.977378Z]  INFO cuzk_bench: loading C1 output path=/data/32gbench/c1.json
=== Batch Benchmark ===
proof type:  porep
count:       1
concurrency: 1

  [1/1] FAILED — 41.4s (prove=0 ms, queue=0 ms)

=== Batch Summary ===
total time:    41.4s
completed:     0
failed:        1

At first glance, this is a straightforward failure report. But in context, it represents the culmination of a complex chain of reasoning, implementation, and debugging that had unfolded over the preceding hours. To understand why this message matters, we must understand what Phase 10 was trying to achieve and why it failed.

The Context: Phase 10's Two-Lock Design

The optimization campaign had progressed through nine phases, each targeting a different bottleneck in the Groth16 proof generation pipeline. Phase 9 had successfully optimized PCIe transfers by pre-staging NTT data onto the GPU, reducing GPU kernel time from ~3.75s to ~1.82s per partition. However, this exposed a new bottleneck: the CPU-side critical path—specifically the prep_msm (~1.9s) and b_g2_msm (~0.48s) operations—now dominated the per-partition wall time, leaving the GPU idle for ~600ms per partition waiting for the CPU thread.

Phase 10 was designed to address this by introducing a two-lock architecture. The existing single gpu_lock mutex, which serialized all GPU access (both memory management and kernel execution), was to be split into two separate mutexes:

• mem_mtx: A short-duration lock for VRAM allocation and data upload (~18ms)
• compute_mtx: A longer-duration lock for GPU kernel execution (~1.8s) The theory was elegant: by releasing mem_mtx before entering compute_mtx, a second worker could begin its memory operations (allocating VRAM and uploading data) while the first worker's kernels were still running. This would effectively hide the CPU overhead by overlapping it with GPU computation, potentially reducing per-partition wall time from ~3.7s to ~1.8–2.0s—a 30–38% throughput improvement. The design was documented in c2-optimization-proposal-10.md, and the implementation involved: 1. Creating a gpu_locks struct containing two std::mutex instances 2. Changing the FFI interface from passing std::mutex* to passing void* (the opaque gpu_locks pointer) 3. Restructuring generate_groth16_proofs_c() to use mem_mtx for pre-staging and compute_mtx for kernel execution 4. Setting gpu_workers_per_device to 3 to take advantage of the overlap

The First Failure and the Misdiagnosis

Before the subject message, the assistant had already run one correctness test (<msg id=2611>) which also failed with a similar 41.1s result. The assistant examined the daemon logs and found the root cause: the first worker successfully pre-staged 12 GiB of VRAM, but subsequent workers found only ~1.5 GiB free—far below the 12 GiB needed. The pre-staging fell back to skip_vram, but the fallback path inside compute_mtx also failed with out-of-memory (OOM) errors.

The assistant's diagnosis in <msg id=2613> identified the problem as a pool trim issue: cudaMemPoolTrimTo() wasn't reclaiming async pool memory because there was no cudaDeviceSynchronize() before it. The fix seemed straightforward: add cudaDeviceSynchronize() back into the mem_mtx region. The assistant reasoned that this would only take ~1ms when there was nothing to sync.

This was the critical assumption that proved incorrect—and the subject message is the evidence of that failure.

The Subject Message: A Second Failure

The subject message shows the result of that "fix": the test still fails, with the same 41.4s wall time and prove=0 ms indicating that the proof never completed successfully. The benchmark reports completed: 0, failed: 1.

What makes this message so significant is what it doesn't show directly but what the assistant immediately deduces from it. The failure pattern is identical to the first attempt, which means the cudaDeviceSynchronize() addition didn't solve the underlying problem—it may have even made things worse by introducing serialization.

The Deeper Diagnosis That Followed

The assistant's subsequent analysis (<msg id=2618>, <msg id=2619>) revealed the true root cause, and it was far more fundamental than a missing synchronization call. The daemon logs showed:

CUZK_TIMING: prestage_vram free_mib=13869 usable_mib=13357 need_da_mib=4096 need_dbc_mib=8192 total_mib=12288
CUZK_TIMING: prestage_setup=ok domain=134217728 lot_of_memory=1 sync_ms=0 trim_ms=0 alloc_ms=1291 upload_ms=5 total_ps_ms=1297
CUZK_TIMING: prestage_vram free_mib=1501 usable_mib=989 need_da_mib=4096 need_dbc_mib=8192 total_mib=12288
CUZK_TIMING: prestage_setup=skip_vram
CUZK_TIMING: prestage_vram free_mib=1463 usable_mib=951 need_da_mib=4096 need_dbc_mib=8192 total_mib=12288
CUZK_TIMING: prestage_setup=skip_vram

Only the first partition successfully pre-staged. All subsequent partitions fell back to skip_vram—and the fallback path inside compute_mtx was also failing. The alloc_ms=1291 (1.3 seconds for cudaMalloc) suggested the CUDA driver was struggling with internal defragmentation.

The breakthrough insight came when the assistant traced the execution flow:

1. Worker 0 acquires mem_mtx, allocates 12 GiB (prestage OK), releases mem_mtx
2. Worker 1 acquires mem_mtx, sees only 1.5 GiB free (Worker 0's 12 GiB still on GPU), skip_vram, releases mem_mtx
3. Worker 2 acquires mem_mtx, sees only 1.4 GiB free, skip_vram, releases mem_mtx
4. Worker 0 acquires compute_mtx, starts kernels
5. Worker 0 releases compute_mtx
6. Worker 1 enters compute_mtx, tries fallback path—but the cudaDeviceSynchronize() inside mem_mtx had blocked while Worker 0 was running kernels The fundamental problem: cudaDeviceSynchronize() and cudaMemPoolTrimTo() are device-global operations. They interact with ALL CUDA streams on the device, including those running under another worker's compute_mtx. This means that when Worker 1 holds mem_mtx and calls cudaDeviceSynchronize(), it blocks until Worker 0's kernels complete—effectively serializing the two locks and destroying the intended overlap.

The Assumptions That Failed

Several assumptions underlying Phase 10's design proved incorrect:

1. That CUDA memory management APIs are lightweight and local: The assumption was that cudaMalloc, cudaMemGetInfo, and cudaMemPoolTrimTo operate independently of kernel execution. In reality, these APIs can implicitly synchronize the entire device.
2. That VRAM would be available for pre-staging: With 16 GiB total VRAM and peak usage of ~13.8 GiB during H-MSM, there was never enough headroom for two workers' pre-staged buffers simultaneously. Pre-staging would almost always fail with gpu_workers_per_device ≥ 2.
3. That the two-lock split would naturally provide overlap: The design assumed that releasing mem_mtx before acquiring compute_mtx would allow another worker to start its memory phase. But device-global synchronization meant the locks were effectively serialized anyway.
4. That cudaDeviceSynchronize was a ~1ms operation: The assistant assumed the sync would be nearly free when there was nothing to synchronize. But when another worker's kernels were running, it blocked for the full kernel duration (~1.8s).

The Output Knowledge Created

Despite being a failure, this message produced invaluable knowledge:

1. CUDA device-global synchronization is a hard constraint: Memory management and compute cannot be fully isolated on a single CUDA device. Any attempt to split locks must account for implicit device-wide synchronization.
2. Pre-staging with multiple workers is infeasible on 16 GiB GPUs: The VRAM budget simply doesn't allow two concurrent allocations of 12 GiB each. The fallback path is not an optimization—it's the only viable path for workers 2+.
3. The two-lock design's real benefit is limited: Without pre-staging succeeding for multiple workers, the only gain from the lock split is hiding the ~0.5s b_g2_msm CPU operation. This may not justify the complexity.
4. A different approach is needed: The assistant concluded that the mem_mtx region should abandon cudaDeviceSynchronize and cudaMemGetInfo entirely, instead attempting cudaMalloc directly and failing fast. The fallback path inside compute_mtx already works correctly (as proven in Phase 8). The real optimization may be to let the next worker's prep_msm_thread + cudaHostRegister overlap with the current worker's GPU kernels + b_g2_msm, without attempting pre-staging at all.

The Thinking Process Revealed

The subject message reveals the assistant's debugging methodology: run a targeted test, examine the output for failure patterns, and use the timing information to reconstruct the execution flow. The prove=0 ms is a crucial signal—it indicates the failure occurred before the proof could even begin, pointing to an initialization or resource allocation issue rather than a computation error.

The assistant's reasoning in the subsequent messages shows a systematic approach to diagnosing the failure:

• First, check the daemon logs for timing information
• Second, reconstruct the interleaved execution of multiple workers
• Third, identify the device-global synchronization as the root cause
• Fourth, design a revised approach that works within the hardware constraints This is a textbook example of how performance debugging requires not just understanding your own code, but understanding the implicit behaviors of the underlying hardware and driver APIs.

Conclusion

The subject message at <msg id=2617> is deceptively simple—a single failed benchmark. But it represents the moment when a well-reasoned optimization design collided with the hard realities of CUDA device architecture. The two-lock split was elegant in theory, but it failed because it assumed that CUDA's memory management APIs were isolated from its compute APIs. In reality, operations like cudaDeviceSynchronize are device-global, and VRAM is a finite shared resource that cannot be partitioned between concurrent workers on a single GPU.

The lesson is profound: when optimizing GPU pipelines, software abstractions (like mutexes) cannot override hardware constraints (like device-global synchronization). The most valuable output of this message was not a working optimization, but a deeper understanding of the system's true boundaries—knowledge that would inform the next iteration of the design.