When Hardware Overrides Abstraction: The Phase 10 Two-Lock Design and the Discovery of CUDA's Device-Global Trap

Introduction

In the long arc of optimizing the SUPRASEAL_C2 Groth16 proof generation pipeline for Filecoin's Proof-of-Replication (PoRep), Phase 10 stands as a watershed moment — not because it succeeded, but because it failed in a way that revealed a fundamental truth about GPU programming. The chunk of work spanning messages 2588 through 2627 documents a complete cycle: design, implementation, debugging, and the sobering discovery that software abstractions cannot always override hardware realities. This article synthesizes that journey, tracing how an elegant two-lock architecture collided with CUDA's device-global synchronization model, and what that collision teaches us about the limits of lock-based concurrency on shared GPU hardware.

The Bottleneck That Motivated Phase 10

The optimization campaign leading into Phase 10 had been remarkably successful. Phase 9's PCIe transfer optimization had achieved a 14.2% throughput improvement in single-worker mode, but dual-worker benchmarks revealed that the bottleneck had shifted yet again. Detailed timing instrumentation showed that the CPU critical path — specifically prep_msm at ~1.9 seconds and b_g2_msm at ~0.48 seconds — now dominated the per-partition wall time at approximately 2.4 seconds. The GPU was idle for roughly 600 milliseconds per partition, waiting for the CPU thread to finish memory-bound computations. At high concurrency, the ten synthesis workers competed with CPU MSM operations for 8-channel DDR5 memory bandwidth, inflating CPU times by 2 to 12 times [4][6].

The bottleneck had migrated from PCIe transfers and GPU kernel execution to CPU memory bandwidth contention. The existing architecture used a single std::mutex per GPU that serialized all GPU access — a worker held this mutex from the moment it started allocating VRAM until its last kernel finished. This meant no overlap was possible between workers. The user proposed a two-lock design to address this, and the assistant began the work of designing and implementing what would become Phase 10 [1][4].

The Two-Lock Design: An Elegant Abstraction

The Phase 10 design, documented in c2-optimization-proposal-10.md at <msg id=2587>, proposed splitting the single mutex into a gpu_locks struct containing two independent mutexes: mem_mtx for VRAM allocation and pre-staging uploads (expected to take ~18ms), and compute_mtx for GPU kernel execution (~1.8s). The intended flow was:

[no lock]      cudaHostRegister + prep_msm_thread starts
lock(mem_mtx)  pool_trim + cudaMemGetInfo + cudaMalloc 12 GiB + async upload
unlock(mem_mtx)
lock(compute_mtx)  per_gpu threads: NTT + batch_add + tail_msm + free d_bc + d_a
unlock(compute_mtx)
[no lock]      Destroy events + stream + cudaHostUnregister + prep_msm_thread.join()

The key insight that made the design viable was that d_a and d_bc — the two large device buffers totaling approximately 12 GiB — were freed synchronously via cudaFree inside compute_mtx, meaning the VRAM was genuinely available the moment compute_mtx was released. No cooperative free slot mechanism was needed. The design increased gpu_workers_per_device from 2 to 3, creating a pipeline where one worker could be in its mem_mtx phase (allocating and uploading), another in its compute_mtx phase (running kernels), and a third doing CPU-only work outside any lock [1][4][6].

The expected per-partition wall time reduction was from ~3.7 seconds to ~1.8–2.0 seconds, a potential 30–38% throughput improvement. The assistant methodically decomposed the implementation into tasks tracked via the todo tool: write the design spec, implement the gpu_locks struct, restructure generate_groth16_proofs_c with the two lock regions, and update the engine configuration [1].

Implementation: From Design to Code

The implementation unfolded across a rapid sequence of edits. At <msg id=2589>, 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 [2][3].

At <msg id=2591>, the assistant 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 [4].

The assistant then wrapped the per-GPU thread creation and join section with compute_mtx at <msg id=2593> and <msg id=2595>, reading the source code to find precise boundaries. At <msg id=2598>, 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 [7][8][11].

A grep-based audit at <msg id=2601> 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" at <msg id=2602>, marking the transition from implementation to build verification [15].

The Build Error That Revealed an FFI Subtlety

The first build attempt failed. At <msg id=2605>, 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 at line 161 of groth16_cuda.cu:

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 [18][19][20].

After this correction, the build succeeded at <msg id=2608>. The assistant wrote a configuration file with gpu_workers_per_device = 3 and launched the correctness test at <msg id=2609> [21][22][23].

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 [24][25].

At <msg id=2613>, 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 [26].

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 [26].

The Timing Logs That Told the Truth

At <msg id=2618>, 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 [31][38].

The critical realization came at <msg id=2619>: 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 [31][32][38].

The Fundamental Discovery: Device-Global Operations Defeat Lock Isolation

The debugging that followed at <msg id=2625> and <msg id=2626> 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 [38][39].

The assistant's proposed fix at <msg id=2626> 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 [39][40].

What the Chunk 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 at <msg id=2618> 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 chunk, 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.

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 chunk spanning messages 2588 through 2627 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.