The Bottleneck That Wouldn't Stay Still: How Phase 10's Two-Lock Architecture Collided with CUDA's Device-Global Reality

Introduction

In the relentless pursuit of faster Groth16 proof generation for Filecoin's Proof-of-Replication (PoRep) protocol, the cuzk optimization campaign had reached a critical inflection point. Phase 9's PCIe transfer optimization had just achieved a stunning 71.6% reduction in NTT MSM kernel time and a 14.2% throughput improvement in single-worker mode. Yet the system's steady-state throughput stubbornly plateaued at approximately 41 seconds per proof, and dual-worker benchmarks showed outright regression. Something fundamental had changed: the bottleneck had shifted again.

This segment of the optimization journey captures the moment when the team discovered that the GPU was no longer the limiting factor. Through a series of increasingly sophisticated benchmarks, timing instrumentations, and awk-powered log analyses, they revealed that the true bottleneck had migrated to a far more stubborn constraint: CPU memory bandwidth contention on an 8-channel DDR5 system. This discovery led directly to the design of Phase 10's two-lock architecture—an ambitious attempt to overlap CPU and GPU work through finer-grained locking—and ultimately to the humbling realization that CUDA's device-global synchronization model imposes limits that no software abstraction can circumvent.

The story that unfolds across this segment is not one of failure but of deepening understanding. Each iteration revealed new constraints, each measurement exposed a new layer of the system's behavior, and each design iteration incorporated the lessons of the previous one. The two-lock architecture, though it did not work as intended on the first attempt, set the stage for the refined design that would eventually emerge—one that acknowledges the device-global nature of CUDA memory management and works within it, rather than fighting against it.

The Commit That Checkpointed Progress

The segment opens with the assistant committing the Phase 9 PCIe optimization code to git—a process that proved surprisingly involved. The assistant first checked git status and discovered that while groth16_cuda.cu was modified, the other Phase 9 files (groth16_ntt_h.cu and the Pippenger MSM changes in the sppark submodule) were untracked. This triggered a multi-step process: committing the submodule changes separately, then staging everything together, only to discover that the sppark submodule path was gitignored, requiring a force-add with -f.

The commit message itself was a carefully crafted summary of Phase 9's achievements, including detailed performance numbers: the 71.6% reduction in ntt_msm_h_ms (from 2430ms to 690ms), the 61.3% reduction in gpu_total_ms, and the explicit acknowledgment that "gw=2 shows regression (41.0s) due to cudaDeviceSynchronize + pool trim serialization—needs further investigation." This commit was not merely a log entry; it was a knowledge artifact that crystallized the current state of understanding, capturing what was accomplished, by how much, and what remained unsolved.

The user's observation at this point was remarkably prescient. Noting "jumpy and inconsistent gpu use," the user hypothesized that the GPU had become so fast that synthesis herding was creating visible gaps in utilization. The suggestion was to run benchmarks with higher concurrency (15–30 synthesis workers) and more proofs to let the pipeline stabilize. This diagnostic insight would prove to be the key that unlocked the true bottleneck.

The Diagnostic Benchmark Campaign

Following the user's guidance, the assistant launched a systematic benchmark sweep at concurrency levels of 15, 20, and 30 synthesis workers. The results were initially confusing: at c=15, prove times ranged from 37–45 seconds—worse than the small-run Phase 9 result of 32.1 seconds. Queue times were growing, indicating that proofs were piling up faster than they could be processed. At c=30, the system crashed with an Out-of-Memory error, revealing the limits of the 754 GiB system memory under extreme load.

The assistant turned to the TIMELINE instrumentation built into the cuzk daemon—structured event logs that record every GPU_START, GPU_END, SYNTH_START, and SYNTH_END event with precise timestamps. Using a series of grep, sed, sort, and awk pipelines, the assistant extracted per-partition GPU timing data and computed statistical distributions. The results were eye-opening:

Tracing the Phantom 1.9 Seconds

The assistant then added fine-grained timing instrumentation to the pre-staging path, measuring each phase of the per-partition workflow with microsecond precision. The results told a stark story:

The User's Provocative Proposal

It was at this moment that the user sent a message that would redirect the entire optimization effort. The proposal was deceptively simple: split the single GPU mutex into two independent locks—one for compute (GPU kernel execution) and one for memory management (VRAM allocation and pre-staging uploads)—and increase the number of GPU workers from 2 to 3 or 4.

The reasoning was elegant. With the existing single-mutex design, Phase 9's pre-staging operations (allocating 12 GiB of VRAM and uploading data) had been added inside the same lock that protected GPU kernel execution. This meant that while one worker held the mutex to pre-stage memory, no other worker could run kernels. The two-lock design would allow Worker A to run GPU kernels under compute_mtx while Worker B allocated and uploaded the next partition's data under mem_mtx, and Worker C did CPU-only preprocessing with no lock at all.

The user's message ended with a crucial warning: "needs to be ordered correctly to not deadleck." This single sentence would drive an entire phase of analysis.

The Resource Deadlock Analysis

The assistant's response to the user's proposal was methodical and thorough. The assistant articulated the three-role model that would become the conceptual foundation of Phase 10:

Worker A: holds compute_mtx, wants mem_mtx to free d_bc (8 GiB)
Worker B: holds mem_mtx, wants to alloc 12 GiB but VRAM is full
         (waiting for Worker A to free d_bc first)
         Then wants compute_mtx (which A holds)

This was, as the assistant correctly identified, "a classic resource deadlock—not a lock-ordering deadlock, but a resource capacity deadlock." The distinction was crucial. A lock-ordering deadlock occurs when threads acquire locks in inconsistent orders; that can be fixed by enforcing a global lock hierarchy. A resource deadlock arises from a finite pool of a resource (here, VRAM) combined with the inability to release it without holding a lock that someone else needs. No amount of lock reordering can fix it—the fundamental problem is that there isn't enough VRAM for two workers' buffers to coexist.

The assistant explored a cooperative handoff design using a shared to_free slot, but ultimately realized that because d_a and d_bc are freed synchronously via cudaFree inside compute_mtx, the VRAM is genuinely available the moment compute_mtx is released. No cooperative free slot is needed.

The Simplified Two-Lock Design

The final design that emerged from this analysis was remarkably clean. The assistant traced through the entire VRAM lifecycle, mapping each allocation and deallocation to specific lines of code, and arrived at a simplified flow:

[no lock]     cudaHostRegister + prep_msm_thread starts
lock(mem_mtx)
  pool_trim + cudaMemGetInfo
  cudaMalloc 12 GiB + async upload + events
unlock(mem_mtx)
lock(compute_mtx)
  stream.wait(events), NTT kernels
  cudaFree(d_bc)
  H-MSM (uses d_a)
  gpu_ptr_t(d_a) destructor → cudaFree
  batch_add + tail_msm
unlock(compute_mtx)
[no lock]     destroy events + stream
              cudaHostUnregister
              prep_msm_thread.join()
              Epilogue

No nested locks, no cooperative free slot, no deadlock risk. The two locks were acquired and released in strict sequence: mem_mtx first, then compute_mtx. Because d_a and d_bc were freed synchronously inside compute_mtx, the VRAM was genuinely available when the next worker acquired mem_mtx.

The assistant documented this design in c2-optimization-proposal-10.md, complete with a risk assessment table, testing plan, and expected outcomes. The predicted improvement was a 30–38% throughput gain in isolation, reducing per-partition wall time from 3.7 seconds to 1.8–2.0 seconds.

Implementation: From Design to Code

The implementation unfolded across a rapid sequence of edits. The assistant introduced the gpu_locks struct and updated the FFI functions create_gpu_mutex and destroy_gpu_mutex to allocate and free the new struct on the C++ heap. The Rust side remained unaware of the change—the opaque void* pointer was preserved, maintaining the FFI contract.

The assistant then performed the critical restructure of generate_groth16_proofs_c. The message contains a numbered inventory of the existing code structure—seven items mapping the control flow from cudaHostRegister through lock acquisition, pre-staging, kernel execution, join/cleanup, unlock, and post-processing—mapped onto a target structure of four items with two lock scopes. This methodical approach, building a cognitive map before cutting code, reflects disciplined engineering practice.

The assistant then wrapped the per-GPU thread creation and join section with compute_mtx, reading the source code to find precise boundaries. At one point, the assistant paused to verify pointer lifecycle correctness—tracing the d_a_prestaged pointer from allocation through RAII wrapper to destruction—and read the full restructured section for consistency.

A grep-based audit confirmed no stale references to the old gpu_lock or mtx_ptr variables remained. The assistant declared "No leftover gpu_lock or mtx_ptr references. Now let me build," marking the transition from implementation to build verification.

The Build Error That Revealed an FFI Subtlety

The first build attempt failed. The assistant ran a targeted diagnostic command—grep -E "error|warning:" | head -20—to extract the error from the noisy build output. The error was an nvcc type conversion failure:

gpu_locks* locks = gpu_mtx ? static_cast<gpu_locks*>(gpu_mtx) : &fallback_locks;

The function parameter was still declared as std::mutex* gpu_mtx, but the new code was trying to static_cast it to gpu_locks*. The fix was to change the parameter to void*, matching the opaque pointer convention used by the Rust FFI. This was a subtle but important lesson: when changing the internal structure behind an opaque pointer, the C++ function signature must be updated to match the opaque type, not the old concrete type.

After this correction, the build succeeded. The assistant wrote a configuration file with gpu_workers_per_device = 3 and launched the correctness test.

The Crash: OOM and Performance Regression

The correctness test failed catastrophically. The daemon log showed out-of-memory errors and severe performance regression—prove time ballooning to 102 seconds, far worse than the Phase 9 baseline of ~41 seconds.

The assistant began debugging. The initial diagnosis focused on CUDA's stream-ordered memory pool. The reasoning was that internal MSM and batch_add operations use cudaMallocAsync/cudaFreeAsync, which go to a stream-ordered pool. When cudaMemPoolTrimTo(pool, 0) was called inside mem_mtx, the async frees from the previous worker's last operations might not have been processed yet because no cudaDeviceSynchronize had been called after the last cudaFreeAsync. The fix seemed straightforward: add cudaDeviceSynchronize() back into the mem_mtx region.

This was a reasonable hypothesis, but it was incomplete. The assistant noted that cudaDeviceSynchronize would be "only ~1ms when there's nothing to sync"—an assumption that would prove disastrous because it ignored the case where another worker's kernels were still running under compute_mtx.

The Timing Logs That Told the Truth

The assistant ran a grep on the timing logs that revealed the true story in six lines:

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

The first worker successfully pre-staged VRAM, seeing 13.8 GiB free. But subsequent workers found only ~1.5 GiB free—because the first worker's 12 GiB allocation was still live on the GPU. The alloc_ms=1291 (1.3 seconds for cudaMalloc) was already pathological, far exceeding the expected ~18ms. Workers 1 and 2 correctly fell back to skip_vram, but the fallback path inside compute_mtx then ran into allocation failures.

The critical realization came next: cudaDeviceSynchronize inside mem_mtx blocks while another worker holds compute_mtx and runs kernels. Worker 1 acquires mem_mtx and calls cudaDeviceSynchronize, which blocks until Worker 0's kernels complete—but Worker 0 still holds compute_mtx. The device-wide synchronization cannot complete until the compute worker finishes, and the compute worker cannot finish faster because the memory worker is blocked. This is not a resource deadlock in the traditional sense—there is no circular wait on mutexes—but a synchronization deadlock where a device-global operation blocks progress on a mutex-protected region, preventing the very completion it is waiting for.

The Fundamental Discovery: Device-Global Operations Defeat Lock Isolation

The debugging that followed crystallized the root cause. The assistant's analysis was precise:

cudaDeviceSynchronize and cudaMemPoolTrimTo are device-global operations that interact with ALL streams, including those running under compute_mtx. You can't isolate memory management from compute on the same device.

This is the central lesson of the Phase 10 saga. The two-lock design was logically sound at the C++ abstraction level—two independent mutexes protecting two independent resources should allow concurrent access. But CUDA GPUs are not abstract resources; they are physical devices with global state. Memory management operations like cudaDeviceSynchronize, cudaMemPoolTrimTo, and even cudaMalloc have device-wide effects that cannot be confined to a single worker's context. The memory pool is device-wide. Synchronization is device-wide. VRAM is a single finite resource shared by all workers.

The assistant's proposed fix was to remove all device-global operations from mem_mtx entirely. Instead of querying free memory or trimming pools, the code would simply try cudaMalloc directly and fall back to the proven Phase 8 allocation path inside compute_mtx if it failed. This acknowledged that the two-lock design could not achieve its intended overlap—the hardware simply does not support clean separation of memory management and compute on a single device.

What the Segment Teaches Us

The Phase 10 two-lock saga is a masterclass in the tension between software abstraction and hardware reality. Several themes emerge:

1. The bottleneck migration pattern. Each optimization phase in this campaign revealed the next bottleneck. Phase 9 optimized PCIe transfers, which revealed CPU memory bandwidth contention. Phase 10 attempted to address that contention through lock splitting, which revealed CUDA device-global synchronization constraints. This pattern is universal in systems optimization: each fix shifts the bottleneck to a new component, and the next fix must be designed with the new bottleneck in mind.

2. The limits of lock-based concurrency on GPUs. Mutexes work well on CPUs because memory management and computation are genuinely independent—different CPU cores can allocate memory and run computations concurrently without interference. On a GPU, memory management and computation share the same hardware pathways, the same memory pool, and the same synchronization primitives. Operations that appear local (pool trim, memory info queries) are in fact device-global. This means that lock splitting, a standard technique for CPU concurrency, may not translate directly to GPU programming.

3. The importance of timing instrumentation. The six lines of timing data told the entire story. Without the prestage_setup, free_mib, alloc_ms, and skip_vram instrumentation, the assistant would have been guessing at the failure mode. The investment in detailed timing logging paid off by providing precise, actionable evidence.

4. The value of methodical engineering. Throughout the segment, the assistant demonstrates disciplined practices: building a cognitive map before editing code, verifying pointer lifecycles, grepping for stale references, separating diagnosis from treatment, and formulating precise queries to extract signal from noise. These practices did not prevent the failure—no amount of careful planning could have predicted the device-global synchronization conflict without testing—but they enabled rapid diagnosis and recovery.

5. The hardware always wins. The most profound lesson is that software abstractions cannot override hardware constraints. The two-lock design was elegant, well-reasoned, and correct in its analysis of data dependencies. But CUDA's device model does not support the clean separation of memory management and compute. The assistant's ultimate solution—removing device-wide synchronization from mem_mtx and relying on the fallback path—is a pragmatic retreat that acknowledges this reality. It is not a failure of engineering; it is the essence of systems programming: learning what the hardware actually does and designing software that works with, not against, those realities.

The Deeper Lessons for Systems Optimization

Beyond the specific technical findings, this segment offers several profound lessons about systems engineering at the frontier of hardware capacity.

First, optimization is a waterfall process. Each improvement reveals the next bottleneck deeper in the system. Phase 9 made the GPU dramatically faster, which exposed CPU memory bandwidth as the new constraint. The two-lock design attempted to address this, but it in turn revealed that CUDA's device-global synchronization model imposes fundamental limits on how much parallelism can be achieved on a single GPU. The engineer's job is not to eliminate bottlenecks—that is impossible—but to identify them, understand them, and systematically push them deeper into the system.

Second, hardware constraints override software abstractions. The two-lock design was logically sound by every analytical measure. It identified the deadlock scenario, traced through the VRAM lifecycle, verified that cudaFree is synchronous, and designed a lock protocol with no nested locks and no deadlock risk. Yet it failed because the CUDA device itself is a single, shared resource with global synchronization semantics. No amount of careful lock ordering can work around the fact that cudaDeviceSynchronize affects all streams and all outstanding operations.

Third, the gap between design and reality is where the real learning happens. The commit point captured in this segment is valuable precisely because it preserves the moment before failure—the point at which a team commits to a design that seems correct by every available analytical tool. The subsequent debugging revealed the true nature of the constraint, leading to a deeper understanding of the hardware and a more robust solution. In systems optimization, failed experiments are often more informative than successful ones.

Fourth, simple tools can reveal complex truths. The assistant's use of grep, sed, sort, and awk pipelines to extract timing distributions from log files is a masterclass in pragmatic diagnostics. The 3.7-second barrier, the 3× spread between median and maximum GPU times, the 600ms idle gap—these insights came not from a sophisticated profiling tool but from decades-old Unix utilities applied with surgical precision.

Fifth, the design document is a commitment protocol. The act of writing c2-optimization-proposal-10.md was not merely documentation; it was a formal commitment to a specific hypothesis about how the system behaves. The proposal made explicit predictions about expected performance improvements (30–38%), identified risks (deadlock, VRAM exhaustion), and specified a testing plan. When the implementation failed, the design document provided a clear baseline against which to measure the failure, enabling rapid diagnosis of which assumptions were violated.

Conclusion

Phase 10's two-lock design did not achieve its ambitious goals. The intended overlap between memory management and compute was defeated by CUDA's device-global synchronization model. But the knowledge gained—that cudaDeviceSynchronize and cudaMemPoolTrimTo are device-wide operations that cannot be isolated by software locks—is a permanent addition to the team's understanding of GPU programming. The assistant's methodical approach to design, implementation, debugging, and diagnosis ensured that this lesson was learned thoroughly and documented clearly.

The segment spanning the Phase 9 commit through the Phase 10 debugging is a complete case study in the optimization lifecycle: identify a bottleneck, design a solution, implement it, test it, discover its flaws, diagnose the root cause, and adapt. The two-lock design may have failed, but the engineering process succeeded in producing deep knowledge about the system's behavior. In systems optimization, that knowledge is often more valuable than any single performance improvement.

The story is not one of failure but of deepening understanding. Each iteration revealed new constraints, each measurement exposed a new layer of the system's behavior, and each design iteration incorporated the lessons of the previous one. The bottleneck shifted from PCIe transfers to GPU kernels to CPU memory bandwidth to device-global synchronization—and with each shift, the team's understanding of the system grew deeper and more nuanced. That is the essence of engineering at the frontier of performance.

References

[1] Chunk 0: "The Bottleneck Shift: How Phase 9's PCIe Optimization Revealed the CPU Memory Wall and Sparked Phase 10's Two-Lock Architecture" — covers the benchmark campaign, bottleneck diagnosis, and two-lock design proposal.

[2] Chunk 1: "When Hardware Overrides Abstraction: The Phase 10 Two-Lock Design and the Discovery of CUDA's Device-Global Trap" — covers the implementation, OOM failure, timing log analysis, and root cause diagnosis.