The Moment a Two-Lock Design Collides with Hardware Reality
In the high-stakes world of GPU-accelerated SNARK proving, every millisecond counts. The Filecoin PoRep (Proof-of-Replication) pipeline, implemented across a chain of Go, Rust, and CUDA code, must generate Groth16 proofs for 32 GiB sectors — a process that peaks at ~200 GiB of memory and demands every ounce of GPU throughput. This article examines a single message ([msg 2611]) from an optimization session targeting Phase 10 of the cuzk proving engine: the moment a carefully designed two-lock GPU interlock was first put to the test, and the revealing failure that followed.
The Message
The subject message is a straightforward correctness test invocation, executed by the AI assistant after implementing the Phase 10 two-lock architecture:
[assistant] 3 workers running (worker_id 0, 1, 2). Correctness test first:
[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:07:55.180968Z INFO cuzk_bench] loading C1 output path=/data/32gbench/c1.json
=== Batch Benchmark ===
proof type: porep
count: 1
concurrency: 1
[1/1] FAILED — 41.1s (prove=0 ms, queue=0 ms)
=== Batch Summary ===
total time: 41.1s
completed: 0
failed: 1
The message is terse — a single shell command and its output — but it carries enormous diagnostic weight. A proof that should have taken ~32 seconds (the Phase 9 baseline) not only took 41.1 seconds but failed entirely. The prove=0 ms field is particularly damning: it indicates the daemon returned an error before the proof could even be measured, meaning something went catastrophically wrong in the GPU pipeline.
Context and Motivation
To understand why this message was written, one must trace the optimization arc that led to Phase 10. The cuzk proving engine had undergone nine phases of optimization, each targeting a different bottleneck in the Groth16 proof generation pipeline. Phase 9 had achieved a breakthrough: by pre-staging NTT (Number Theoretic Transform) inputs over PCIe before GPU kernel execution, the assistant had reduced GPU kernel time from ~3.75s to ~1.82s per partition. However, this victory revealed a new bottleneck: the CPU-side critical path — prep_msm (multi-scalar multiplication preparation) at ~1.91s and b_g2_msm (G2-group MSM) at ~0.48s — now dominated the per-partition wall time at ~2.4s, leaving the GPU idle for ~600ms per partition.
The insight was that the single std::mutex protecting GPU access was too coarse. Under the Phase 8 dual-worker design, a worker held the GPU mutex for the entire duration of VRAM allocation, data upload, kernel execution, and cleanup. The CPU work (prep_msm, b_g2_msm, epilogue arithmetic) could theoretically overlap with another worker's GPU kernel execution — if only the mutex were split.
Phase 10's design, documented in c2-optimization-proposal-10.md, proposed replacing the single mutex with a gpu_locks struct containing two mutexes: mem_mtx for VRAM allocation and pre-staging upload (expected ~18ms hold time), and compute_mtx for GPU kernel execution (~1.8s hold time). By increasing gpu_workers_per_device from 2 to 3, the design hoped to pipeline three workers: while Worker A ran kernels under compute_mtx, Worker B could acquire mem_mtx, pre-stage its VRAM, and then wait for compute_mtx — hiding the CPU overhead behind GPU execution.
The expected improvement was dramatic: per-partition wall time from ~3.7s to ~1.8–2.0s, a 30–38% throughput improvement in isolation.## Assumptions Embedded in the Design
The Phase 10 design rested on several critical assumptions, and this message represents the first test of all of them simultaneously:
Assumption 1: Memory management can be isolated from compute on the same CUDA device. The design assumed that cudaMalloc, cudaMemGetInfo, and cudaMemPoolTrimTo operations inside mem_mtx would not interfere with kernel execution happening under compute_mtx on the same GPU. This assumption turned out to be fundamentally wrong: cudaDeviceSynchronize and cudaMemPoolTrimTo are device-global operations that interact with all streams, including those running kernels for another worker. As the assistant later discovered, "you cannot isolate memory operations from compute operations on the same GPU device."
Assumption 2: VRAM would be available for pre-staging when a worker acquired mem_mtx. With a 16 GiB GPU and peak VRAM usage of ~13.8 GiB during H-MSM (d_a at 4 GiB + d_bc at 8 GiB + MSM temp at ~1.8 GiB), the design assumed there would be enough headroom for a second worker to pre-stage its 12 GiB allocation. In practice, with gpu_workers_per_device=3, only the first worker ever succeeded at pre-staging; all subsequent workers found insufficient free memory and fell back to the slow path.
Assumption 3: The fallback path would be a rare exception. The design included a fallback path for when pre-staging failed, but treated it as an edge case. In reality, with 3 workers sharing 16 GiB of VRAM, the fallback became the common case — and the fallback path itself was slower than the Phase 8 baseline because it ran inside compute_mtx without proper pool trimming.
Assumption 4: The two-lock split would be transparent to the Rust FFI layer. The assistant changed the C++ function parameter from std::mutex* to void* to accommodate the new gpu_locks struct, assuming the Rust side (which passes the pointer as *mut c_void) would be unaffected. This assumption held — the build succeeded — but it masked the deeper hardware conflict.
The Failure Mode: What the Message Reveals
The message's output — FAILED — 41.1s (prove=0 ms, queue=0 ms) — is a compressed diagnostic signal. The prove=0 ms means the Rust-side prove() function returned an error immediately or never completed its timing measurement. The 41.1 seconds is the total wall time including daemon communication, queue wait, and error handling.
What the message doesn't show is the daemon-side log, which the assistant examined in the following messages ([msg 2612], [msg 2618]). Those logs revealed the true story:
- Worker 0 acquired
mem_mtx, successfully pre-staged 12 GiB of VRAM (taking ~1.3s forcudaMallocdue to internal driver defragmentation), then releasedmem_mtx. - Worker 1 acquired
mem_mtx, found only ~1.5 GiB free (Worker 0's 12 GiB still live on the GPU), and fell back withskip_vram. - Worker 2 experienced the same fate.
- Worker 0 acquired
compute_mtxand ran kernels. - When Worker 0 released
compute_mtx, Worker 1 entered the fallback path — but the fallback path'scudaMallocfor 8 GiBd_ballocation failed with an OOM because the CUDA memory pool hadn't been trimmed. The root cause was subtle:cudaDeviceSynchronizeinsidemem_mtx(added to ensure pool frees were resolved) blocked while Worker 0 heldcompute_mtxand ran kernels. This effectively serialized the two locks, destroying the intended overlap and forcing every subsequent partition through the slow fallback path.
Input Knowledge Required
To fully understand this message, one needs knowledge spanning multiple domains. On the CUDA side: the behavior of cudaMalloc vs cudaMallocAsync, the stream-ordered memory pool, the device-global nature of cudaDeviceSynchronize and cudaMemPoolTrimTo, and the VRAM budget of Groth16 proving (~13.8 GiB peak). On the system architecture side: the Phase 9 baseline performance (32.1s/proof in isolation, 41.3s at c=15 j=15), the CPU critical path breakdown (prep_msm 1.91s, b_g2_msm 0.48s), and the DDR5 memory bandwidth contention that limited high-concurrency throughput. On the software engineering side: the FFI boundary between Rust and C++, the gpu_locks struct replacing std::mutex, and the gpu_workers_per_device configuration parameter.
Output Knowledge Created
This message, though a failure report, produced critical knowledge. It validated that the two-lock design compiled and ran — the daemon started, workers spawned, and the pipeline executed — but it also revealed that the design's core premise was flawed. The output knowledge includes:
- The impossibility of device-local memory isolation: CUDA's memory management APIs are inherently device-global, making fine-grained lock splitting ineffective on single-GPU systems.
- The VRAM contention floor: With 16 GiB of VRAM and ~13.8 GiB peak usage, there is simply no room for two workers' pre-staged buffers simultaneously, making the
mem_mtxpre-staging path a dead letter for all but the first worker. - The fallback path as primary: When pre-staging always fails for workers 2 and 3, the two-lock design degenerates into a slower version of the single-lock design, because the fallback path lacks proper pool management.
- A diagnostic methodology: The assistant's subsequent debugging — examining
prestage_setup=skip_vramandfree_mibvalues in the daemon log — established a pattern for diagnosing GPU memory contention that would inform the eventual redesign.## The Thinking Process: From Failure to Diagnosis The reasoning visible in the assistant's subsequent messages reveals a systematic diagnostic process triggered by this failure. The assistant did not simply report the failure and move on; it immediately began probing the daemon log ([msg 2612]), extracting timing markers likeprestage_setup=skip_vram,free_mib=1501, andalloc_ms=1291. Each log line was a clue: prestage_setup=skip_vramfor workers 1 and 2 told the assistant that pre-staging was failing for all but the first worker.free_mib=1501(only 1.5 GiB free) confirmed VRAM exhaustion — Worker 0's 12 GiB allocation was still live.alloc_ms=1291(1.3 seconds forcudaMalloc) revealed that even the first worker's allocation was abnormally slow, suggesting CUDA driver-level defragmentation.sync_ms=616(616ms forcudaDeviceSynchronize) showed thatcudaDeviceSynchronizeinsidemem_mtxwas blocking on Worker 0's ongoing kernel execution. The assistant's thinking, visible in its commentary, moved through several hypotheses: 1. First hypothesis: "pool_trim isn't reclaiming the async pool memory because there's nocudaDeviceSynchronizebefore it." This led to addingcudaDeviceSynchronizeback intomem_mtx— which made the problem worse by serializing the locks. 2. Second hypothesis (after the 72.5s successful-but-slow run): "cudaDeviceSynchronizeinmem_mtxblocks while the other worker runs kernels, creating serialization." This was the correct diagnosis, leading to the insight thatcudaDeviceSynchronizeandcudaMemPoolTrimToare device-global operations that cannot be isolated. 3. Third hypothesis: "Just trycudaMallocdirectly — if it succeeds, great; if it fails with OOM, fall back. No need forcudaMemGetInfoor pool trim inmem_mtx." This was the pragmatic response: accept that pre-staging will almost always fail for workers 2 and 3, and focus on making the fallback path insidecompute_mtxwork correctly. The assistant's reasoning demonstrates a critical skill: the ability to distinguish between software design flaws (incorrect lock granularity) and hardware constraints (device-global synchronization, VRAM capacity). The two-lock design was logically sound as a software abstraction, but it violated a physical constraint of the CUDA programming model.
Lessons for GPU Pipeline Design
This message and its aftermath teach several enduring lessons for anyone designing high-throughput GPU pipelines:
Hardware constraints override software abstractions. No matter how cleanly you split locks at the software level, if those locks guard operations that implicitly synchronize the entire device, the split is illusory. The CUDA programming model does not support fine-grained concurrency control within a single device — memory management and compute are fundamentally entangled.
Pre-staging is a single-worker optimization. On GPUs with limited VRAM (16 GiB in this case), pre-staging buffers for multiple workers is impossible when each worker's working set is ~12 GiB. The only way to benefit from pre-staging is to ensure the pre-staging worker is the same worker that runs the kernels — which is exactly what the Phase 9 single-worker design did.
Fallback paths must be first-class citizens. The Phase 10 design treated the fallback path as an edge case, but it became the common path. Any design that relies on a "fast path" that is only available to the first worker is not a scalable design.
The value of a failing test. This message, despite reporting a failure, was more valuable than a passing test would have been. It exposed fundamental assumptions about the CUDA programming model that no amount of code review or design discussion would have revealed. The 41.1-second failure was a diagnostic goldmine.
Conclusion
Message [msg 2611] captures a pivotal moment in the optimization of a production GPU proving pipeline. It is the point at which a carefully reasoned software design — the two-lock GPU interlock — collided with the physical reality of CUDA device architecture. The failure was not due to a bug in the code (the locks were correctly implemented, the FFI boundary was clean, the build succeeded) but due to an incorrect assumption about the CUDA programming model: that memory management operations could be isolated from compute operations on the same device.
The assistant's response to this failure is a model of systematic debugging: observe the symptom (FAILED, 41.1s), examine the logs (skip_vram, free_mib, alloc_ms), form hypotheses (pool trim, DeviceSync serialization), test each hypothesis, and ultimately arrive at a correct understanding of the hardware constraint. The two-lock design was abandoned not because it was poorly implemented, but because the underlying assumption was physically impossible on a single GPU.
This message reminds us that in systems programming — especially at the intersection of CPU and GPU — the most important debugging tool is not a deeper understanding of the code, but a deeper understanding of the hardware.