The Device-Global Trap: When a CUDA Synchronization Call Destroyed Phase 10's Two-Lock Design

Introduction

In the high-stakes world of Filecoin PoRep proof generation, every millisecond counts. The SUPRASEAL_C2 pipeline, responsible for Groth16 proofs consuming ~200 GiB of peak memory, had already undergone nine optimization phases. Each phase had shifted the bottleneck: from PCIe transfers to GPU kernel execution, from mutex contention to CPU memory bandwidth. Now, Phase 10 promised a breakthrough — a two-lock architecture that would overlap CPU-side memory management with GPU kernel execution, hiding the ~2.4s of CPU-critical-path work that had emerged as the dominant bottleneck.

Message 2624 records the moment this elegant design met reality. It is the first successful run of the Phase 10 implementation — and it is a disaster.

The Message: A Correctness Test That Passed Too Slowly

The message itself is deceptively simple — a bash command and its output:

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

The output shows:

=== Phase 10 correctness test gw=3 c=1 j=1 ===
[1/1] COMPLETED — 72.5s (prove=101919 ms, queue=257 ms)

=== Batch Summary ===
total time:    72.5s
completed:     1
failed:        0
wall time:     avg=72.5s min=72.5s max=72.5s
prove time:    avg=101.9s min=101.9s max=101.9s

On the surface, this is good news: the proof completed successfully, the OOM errors that plagued earlier attempts were resolved. But the numbers tell a different story. A prove time of 101.9 seconds is dramatically worse than the Phase 9 baseline of ~41 seconds. The two-lock design, which was supposed to improve throughput by 30–38%, had instead produced a 2.5× regression.

The Context: Phase 10's Two-Lock Design

To understand why this message is so significant, we need to trace the debugging journey that led here. Phase 10 was born from a critical insight in Phase 9 ([msg 2602]): the bottleneck had shifted from GPU kernel execution to CPU-side operations — specifically prep_msm (~1.9s) and b_g2_msm (~0.48s) — that left the GPU idle for ~600ms per partition. The proposed solution was a two-lock architecture ([msg 2603]) using:

The Debugging Journey: Three Attempts and Counting

The implementation hit immediate problems. The first build ([msg 2605]) failed with a type-casting error — nvcc rejected static_cast<gpu_locks*>(gpu_mtx) because the parameter was declared as std::mutex*. This was fixed by changing the FFI parameter to void* ([msg 2607]).

The first correctness test ([msg 2611]) failed with OOM. The daemon log revealed that only the first worker successfully pre-staged VRAM; subsequent workers found insufficient free memory because the first worker's 12 GiB allocation was still live. The assistant diagnosed the root cause ([msg 2613]): cudaMemPoolTrimTo wasn't reclaiming async pool memory because there was no cudaDeviceSynchronize before it. The fix was to add cudaDeviceSynchronize() back into the mem_mtx region.

The second test ([msg 2617]) still failed. Deeper analysis ([msg 2619]) revealed a more fundamental problem: cudaDeviceSynchronize inside mem_mtx blocks while another worker holds compute_mtx running kernels. This meant that Worker 1, holding mem_mtx, would block on cudaDeviceSynchronize waiting for Worker 0's kernels to finish — effectively serializing the two locks. The assistant added a pool trim at the start of compute_mtx for the fallback path ([msg 2620]), rebuilt ([msg 2622]), and started the daemon ([msg 2623]).

Then came message 2624 — the first successful run.## The Deeper Lesson: Hardware Constraints Override Software Abstraction

The 102-second prove time in message 2624 is not just a performance regression — it is a profound lesson about the limits of software abstraction when dealing with CUDA hardware. The two-lock design was conceptually elegant: separate the short memory-management operations from the long compute operations, and let them overlap. But this abstraction collided with three hard realities:

1. Device-Global Synchronization

cudaDeviceSynchronize() is a device-global operation. It does not care about your mutex boundaries, your lock hierarchy, or your elegant design documents. When any thread calls it, the entire CUDA device synchronizes — blocking all streams, all kernels, all memory operations. The mem_mtx region, designed to be a quick ~18ms hold, became a blocking point that waited for the other worker's entire kernel execution to finish.

2. VRAM as a Non-Sharable Resource

The two-lock design assumed that memory management and compute could be cleanly separated because they use different hardware resources. But VRAM is a single, finite pool. When Worker 0 allocates 12 GiB for its pre-staged buffers, that memory is simply not available to Worker 1 — regardless of which mutex is held. The mem_mtx region's cudaMalloc call for 12 GiB took 1.3 seconds (alloc_ms=1291), suggesting the CUDA driver was doing internal defragmentation. Meanwhile, Workers 1 and 2 sat idle.

3. The Fallback Path's Hidden Cost

When pre-staging fails (as it did for Workers 1 and 2), the code falls back to allocating inside compute_mtx. But this fallback path itself needs VRAM — and the pool hasn't been trimmed. The dev_ptr_t<fr_t> d_b(...) call does a synchronous cudaMalloc(8 GiB) that can OOM if the async pool still holds cached memory from previous allocations. The assistant's fix — adding a pool trim at the start of compute_mtx — addressed the OOM but couldn't fix the fundamental serialization.

The Thinking Process Visible in the Reasoning

The assistant's reasoning in the messages leading up to 2624 reveals a sophisticated debugging process. In [msg 2619], the assistant walks through the exact sequence of events:

  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 The critical insight comes when the assistant corrects itself mid-thought: "Wait — the real issue is: cudaDeviceSynchronize inside mem_mtx blocks while Worker 0 has compute_mtx running kernels!" This self-correction is the moment the device-global nature of CUDA synchronization becomes clear. The assistant then traces the timing numbers: alloc_ms=1291 (1.3s for cudaMalloc) and sync_ms=616 (616ms for cudaDeviceSynchronize waiting for Worker 0's kernels). These numbers confirm that the two locks are effectively serialized — the mem_mtx holder blocks on the compute_mtx holder's kernels via the implicit device-wide sync.

What This Message Creates

Message 2624 creates several forms of knowledge:

Output knowledge: The first successful (but catastrophically slow) run of Phase 10. The proof passed correctness — the cryptographic output was valid — but the performance was unacceptable. This creates a clear signal that the two-lock design, as implemented, is fundamentally flawed.

Diagnostic signal: The prove time of 101.9s versus the Phase 9 baseline of ~41s provides a quantitative measure of the regression. The assistant immediately follows up in [msg 2625] by examining the timing logs to understand why.

Validation of the hypothesis: The assistant's earlier analysis in [msg 2619] — that cudaDeviceSynchronize in mem_mtx blocks on the other worker's compute_mtx kernels — is confirmed by the 102s result. The sync_ms=616 in the timing logs directly shows the device-wide synchronization cost.

The Broader Lesson: When Abstractions Hit Hardware

The Phase 10 two-lock design is a textbook example of a software abstraction that makes perfect sense on paper but fails when it meets hardware reality. The assumption was that memory management and compute are independent operations that can proceed in parallel. But on a CUDA device with a single memory pool and device-global synchronization primitives, they are not independent.

This is a recurring theme in high-performance GPU programming. The CUDA programming model presents an abstraction of independent streams and concurrent kernels, but the hardware underneath has hard limits: a single VRAM pool, a single memory controller, device-wide synchronization points. When you push against these limits, the abstraction leaks.

The assistant's debugging process — tracing the exact sequence of lock acquisitions, measuring timing with fine-grained instrumentation, and reasoning about the CUDA driver's internal behavior — is a model for how to diagnose these abstraction leaks. The key was not to look at the code in isolation, but to understand how the CUDA runtime actually executes those operations.

Conclusion

Message 2624 is a turning point in the SUPRASEAL_C2 optimization journey. It represents the moment when a promising design — the two-lock overlap architecture — was tested against reality and found wanting. The 102-second prove time is not a failure of implementation but a failure of abstraction: the assumption that two mutexes could independently control CUDA operations on a single device.

The message also demonstrates the value of rapid iteration. Within a span of minutes, the assistant went from build errors to OOM failures to a working-but-slow proof, each time diagnosing the root cause and adjusting the approach. The next step, visible in [msg 2625], is to remove the device-wide synchronization from mem_mtx and rely on the fallback path — a pragmatic retreat from the elegant two-lock design toward something that works with the hardware rather than against it.

In the end, the most important knowledge created by this message is the recognition that on a single CUDA device, memory management and compute cannot be fully isolated. The solution, as the assistant will discover, is not to split locks differently but to remove the device-wide synchronization entirely — accepting that the two-lock design was a step too far, and that the real optimization lies elsewhere.