The Moment the Two-Lock Design Broke: Diagnosing CUDA Device-Global Synchronization Conflicts

In the high-stakes world of Filecoin proof generation, every millisecond counts. The SUPRASEAL_C2 Groth16 proving pipeline, responsible for producing Filecoin's Proof-of-Replication (PoRep) proofs, consumes roughly 200 GiB of peak memory and must coordinate GPU kernels, CPU synthesis, and PCIe transfers across multiple concurrent workers. After nine optimization phases had progressively squeezed throughput from the system, Phase 10 aimed to introduce a clever two-lock architecture: split the single GPU mutex into a mem_mtx for short-lived VRAM allocation and pre-staging (~18ms), and a compute_mtx for the actual GPU kernel execution. The theory was elegant — while one worker ran GPU kernels under compute_mtx, another worker could simultaneously allocate VRAM and upload data under mem_mtx, hiding the CPU-side memory management overhead behind GPU computation.

But theory and hardware reality collided spectacularly. Message <msg id=2618> captures the exact moment that collision became visible — a single bash command revealing the raw timing data that would ultimately force a fundamental rethinking of the locking strategy.

The Message: Raw Data from the Trenches

The message itself is deceptively simple. After a second failed correctness test of the Phase 10 implementation, the assistant writes:

Still failing. Let me check the log:

>

``bash grep -E "OOM|out of memory|panic|prestage_setup|free_mib|ERROR" /home/theuser/cuzk-p10-daemon.log | tail -20 ``

>

`` 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 CUZK_TIMING... ``

Six lines of timing instrumentation. That is all. Yet these six lines contain the entire story of why the two-lock design failed, and they illuminate a fundamental truth about GPU programming: CUDA devices are not easily partitioned.

Context: The Road to Phase 10

To understand why this message matters, one must understand the journey that led to it. The optimization pipeline had been running for many sessions, each phase peeling away a bottleneck:

• Phase 7 introduced per-partition dispatch, restructuring how partitions were assigned to GPU workers.
• Phase 8 implemented a dual-GPU-worker interlock, narrowing a C++ static mutex that was causing GPU utilization gaps.
• Phase 9 optimized PCIe transfers by pre-staging VRAM allocations, achieving a 14.2% throughput improvement in single-worker mode — but revealing PCIe bandwidth contention in dual-worker mode. By the end of Phase 9, the bottleneck had shifted dramatically. Detailed timing analysis showed that the CPU critical path — specifically prep_msm (1.9s) and b_g2_msm (0.48s) — now dominated the per-partition wall time at ~2.4s, leaving the GPU idle for ~600ms per partition. At high concurrency, the 10 synthesis workers competed with the CPU MSM operations for 8-channel DDR5 memory bandwidth, inflating CPU times by 2–12×. The user proposed a two-lock design to overlap CPU and GPU work. The assistant analyzed deadlock risks and designed a protocol using mem_mtx for VRAM allocation/upload and compute_mtx for kernel execution. The key insight was that d_a and d_bc (the GPU-side circuit data) are freed synchronously inside compute_mtx, so the next worker could safely allocate in mem_mtx without needing a cooperative free slot. The design document (c2-optimization-proposal-10.md) was written, and the implementation began.

The First Failure and the Misdiagnosis

The first implementation attempt ([msg 2611]) failed immediately with an OOM error. The assistant's initial diagnosis ([msg 2613]) focused on a plausible culprit: 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 the stream was never synchronized after the last cudaFreeAsync.

The fix seemed straightforward: add cudaDeviceSynchronize() back into the mem_mtx region. The assistant noted it was "only ~1ms when there's nothing to sync" and applied the edit ([msg 2613]).

This was a reasonable hypothesis, but it was wrong. The assumption that a device-wide synchronization would be harmless assumed that the previous worker had already released compute_mtx and finished its kernels. In practice, the very overlap the two-lock design was trying to achieve meant that another worker's kernels could still be running when mem_mtx was acquired.## Reading the Timing Data: A Story in Six Lines

The six lines of output in message <msg id=2618> tell a precise story. Let us read them carefully:

Line 1: The first worker succeeds.

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

Worker 0 acquires mem_mtx. It sees 13.8 GiB free on the GPU — plenty of room for the 12 GiB needed (4 GiB for d_a, 8 GiB for d_bc). Pre-staging succeeds. The alloc_ms=1291 (1.3 seconds for cudaMalloc) is already suspicious — that is an unusually long allocation time, suggesting the CUDA driver is doing internal defragmentation or memory pool reorganization. The sync_ms=0 confirms no device-wide synchronization was needed. Total pre-staging setup: 1.3 seconds. Worker 0 then releases mem_mtx and proceeds to acquire compute_mtx to run GPU kernels.

Lines 2–3: Subsequent workers find no room.

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

Worker 1 acquires mem_mtx and sees only 1.5 GiB free. Worker 2 acquires mem_mtx and sees 1.4 GiB free. Both correctly detect insufficient VRAM and take the skip_vram fallback path. This is expected behavior — Worker 0's 12 GiB allocation is still live because Worker 0 is still running kernels under compute_mtx. The two-lock design intended for this to happen: Worker 1 would skip pre-staging, release mem_mtx, then wait for compute_mtx. Once Worker 0 finished and freed its VRAM, Worker 1 could allocate normally in the fallback path.

But the timing data reveals the hidden problem. The sync_ms=0 and trim_ms=0 on the first worker's successful pre-stage show that no device-wide synchronization was performed. The cudaDeviceSynchronize that the assistant added in the previous fix attempt ([msg 2613]) was placed inside mem_mtx, but the first worker's sync_ms=0 indicates it didn't trigger — or rather, the sync happened but had nothing to wait for because no other worker was running kernels yet. The problem would manifest on subsequent workers.

The Hidden Serialization

The critical insight that emerges from these six lines is not explicitly stated — it must be inferred from what is missing. Notice that the log shows prestage_setup=skip_vram for Workers 1 and 2, but there is no corresponding prestage_setup=ok or prestage_setup=fallback_ok after Worker 0 finishes. This means the fallback path inside compute_mtx is also failing — or at least, it is not succeeding cleanly.

The assistant's subsequent analysis ([msg 2619]) reveals the full horror: 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 — it hasn't finished yet. Worker 1 is stuck in mem_mtx waiting for a device-wide sync that cannot complete until Worker 0 releases compute_mtx, which Worker 0 cannot do until it finishes its kernels, which are running on the same device that Worker 1 is trying to synchronize.

This is a deadlock of a particularly subtle kind — not a resource deadlock (no circular wait on mutexes), but a synchronization deadlock where a device-global operation (cudaDeviceSynchronize) blocks progress on a mutex-protected region, preventing the very completion it is waiting for. The two locks become effectively serialized: mem_mtx cannot release until the device sync completes, the device sync cannot complete until compute_mtx releases, and compute_mtx cannot release until the kernels finish, which they would do faster if another worker could pre-stage VRAM in parallel — which is blocked by mem_mtx.

The Assumptions That Failed

The two-lock design rested on several assumptions that the timing data proved false:

Assumption 1: Memory management operations are device-local. The design assumed that cudaMalloc, cudaMemPoolTrimTo, and VRAM queries could proceed independently of kernel execution on the same device. In reality, cudaDeviceSynchronize is a device-global barrier. cudaMemPoolTrimTo may also trigger implicit synchronization depending on the driver version and pool state. There is no way to say "trim only the memory that belongs to my stream" — the CUDA memory model treats the device as a single resource.

Assumption 2: VRAM allocation is fast and predictable. The design assumed that cudaMalloc(12 GiB) would complete in milliseconds. The alloc_ms=1291 (1.3 seconds) on the first worker's allocation shattered this assumption. When 12 GiB must be allocated on a device with 24 GiB total, the CUDA driver may need to defragment, evict cached allocations from the pool, or synchronize pending operations — all of which take time and may block the calling thread.

Assumption 3: The fallback path is safe. The design assumed that when pre-staging failed (skip_vram), the fallback path inside compute_mtx would allocate memory normally. But the fallback path calls gpu.Dmalloc(d_a_sz) which uses cudaMallocAsync + sync(), and then dev_ptr_t<fr_t> d_b(...) which uses synchronous cudaMalloc(8 GiB). If the CUDA memory pool still holds cached allocations from the previous worker's async frees (which were never trimmed because cudaMemPoolTrimTo was called inside mem_mtx while the previous worker's kernels were still running), the synchronous cudaMalloc can fail with OOM.

Assumption 4: Three workers can share 24 GiB of VRAM. The design set gpu_workers_per_device=3, assuming that with careful timing, three workers could share a single GPU's 24 GiB of VRAM. Each worker needs 12 GiB (4 GiB for d_a, 8 GiB for d_bc). Three workers need 36 GiB total — 50% more than available. The design relied on the hope that only one worker would hold VRAM at any given time, but the locking protocol could not enforce this because VRAM was freed inside compute_mtx while mem_mtx was already released.## The Thinking Process: From Data to Diagnosis

What makes message <msg id=2618> particularly instructive is what happens after it. The assistant does not simply stare at the six lines and give up. Instead, the message serves as a diagnostic pivot point. The raw timing data triggers a chain of reasoning that leads to the root cause.

The assistant's first observation is that the first worker's pre-stage succeeded (prestage_setup=ok), while subsequent workers hit skip_vram. This confirms the expected behavior: Worker 0's 12 GiB allocation is still live. But the assistant then notices something deeper: the alloc_ms=1291 on the first worker. A 1.3-second cudaMalloc is pathological. The CUDA driver is struggling.

The assistant then reconstructs the full sequence of events across all three workers ([msg 2619]):

1. Worker 0 acquires mem_mtx, allocates 12 GiB (1.3s), releases mem_mtx.
2. Worker 1 acquires mem_mtx, sees 1.5 GiB free, skip_vram, releases mem_mtx.
3. Worker 2 acquires mem_mtx, sees 1.4 GiB free, skip_vram, releases mem_mtx.
4. Worker 0 acquires compute_mtx, starts kernels.
5. Workers 1 and 2 wait for compute_mtx.
6. Worker 0 releases compute_mtx.
7. Worker 1 enters compute_mtx, tries fallback path — OOM. The critical realization: cudaDeviceSynchronize inside mem_mtx blocks while Worker 0 holds compute_mtx and runs kernels. Worker 1 holds mem_mtx and is blocked on cudaDeviceSynchronize, while Worker 2 is waiting for mem_mtx. The device-wide sync creates a chain of dependencies that serializes everything. The assistant correctly identifies the root cause: "the real issue is: cudaDeviceSynchronize inside mem_mtx blocks while Worker 0 has compute_mtx running kernels!" This is the moment of insight — the recognition that CUDA's device-global synchronization model defeats the attempt to partition the GPU into independent memory-management and compute domains.

Input Knowledge Required

To understand this message, the reader needs:

1. CUDA memory model: Understanding that cudaDeviceSynchronize blocks until all pending operations on the device complete, not just operations from a specific stream or context. This is a device-global barrier.
2. CUDA memory pool semantics: Knowing that cudaMallocAsync/cudaFreeAsync use a stream-ordered memory pool, and that cudaMemPoolTrimTo may not reclaim memory that is still referenced by pending operations on any stream.
3. The Groth16 proving pipeline structure: Understanding that each partition requires ~12 GiB of VRAM (4 GiB for d_a, 8 GiB for d_bc), and that the proving process involves pre-staging (alloc + upload), kernel execution (NTT, MSM, batch_add), and cleanup (free).
4. The Phase 10 two-lock design: Knowing that mem_mtx was intended for short VRAM operations (~18ms) and compute_mtx for GPU kernel execution (~2s), and that the design assumed these could overlap.
5. The cuzk architecture: Understanding that multiple GPU workers (set by gpu_workers_per_device) compete for a single GPU's VRAM, and that the system uses a partition-based dispatch model where each partition is processed by one worker.

Output Knowledge Created

This message creates several pieces of critical knowledge:

1. Empirical confirmation that device-global synchronization defeats lock splitting: The timing data proves that cudaDeviceSynchronize inside mem_mtx blocks until the other worker's kernels (running under compute_mtx) complete, destroying the intended overlap.
2. The pathological allocation time: The alloc_ms=1291 measurement reveals that cudaMalloc(12 GiB) can take 1.3 seconds — far longer than the ~18ms the design assumed for pre-staging. This alone undermines the premise that mem_mtx would be a "short hold."
3. The VRAM contention pattern: The progression from free_mib=13869 (Worker 0) to free_mib=1501 (Worker 1) to free_mib=1463 (Worker 2) shows exactly how VRAM is consumed and not released between workers.
4. The fallback path failure mode: The skip_vram entries show that the fallback path is being triggered, but the subsequent OOM (visible in the earlier message's FAILED result) proves that the fallback path itself is not safe when the pool hasn't been properly trimmed.

The Deeper Lesson: Hardware Constraints Override Software Abstraction

The Phase 10 two-lock failure is a textbook example of a phenomenon that every GPU programmer encounters eventually: hardware constraints override software abstraction. No matter how cleanly you partition your code into mutex-protected regions, the underlying hardware is a single device with global state. cudaDeviceSynchronize does not care about your lock boundaries. cudaMalloc does not respect your thread boundaries. The CUDA driver's memory manager is a single system that must coordinate all allocations and deallocations across all streams and contexts.

The assistant's ultimate solution — removing the device-wide synchronization from mem_mtx entirely and relying on the already-proven fallback path inside compute_mtx — acknowledges this reality. The two-lock design is not abandoned, but its scope is narrowed. The mem_mtx region becomes a best-effort optimization: try to pre-stage if VRAM is available, but do not block on device-wide operations. If pre-staging fails, the fallback path inside compute_mtx handles allocation, accepting the serialization cost.

This is the essence of engineering under real-world constraints: the optimal design on paper may be defeated by hardware behavior that only reveals itself under measurement. The six lines of timing data in message <msg id=2618> are not just debugging output — they are the voice of the hardware saying "this is how I actually work, now design for that."