The Two-Lock Temptation: Designing Phase 9b's GPU Lock Split for SUPRASEAL_C2
Introduction
In the relentless pursuit of throughput for Filecoin's Proof-of-Replication (PoRep) proving pipeline, optimization often follows a familiar rhythm: identify a bottleneck, eliminate it, then watch the next bottleneck emerge. By the time the conversation reaches message 2584, this rhythm has played out over nine distinct phases. Phase 9 had just shifted the bottleneck from PCIe transfer bandwidth to CPU memory bandwidth contention, revealing that the GPU was spending ~600ms per partition idle, waiting for CPU-side MSM preprocessing to finish. The natural response—split the single GPU mutex into two locks to allow better overlap of CPU and GPU work—seemed straightforward. Message 2584 presents that implementation plan in meticulous detail, laying out the code changes, risk assessment, testing strategy, and expected outcomes for what the assistant calls "Phase 9b."
This article examines message 2584 as a pivotal moment in the optimization journey: a carefully reasoned design document that, despite its thoroughness, would encounter fundamental hardware constraints that no amount of software abstraction could circumvent. The message represents the convergence of deep systems knowledge, lock-design theory, and CUDA memory management expertise, all brought to bear on a problem that had shifted from "how do we make the GPU faster" to "how do we keep the GPU fed."
Context: The Optimization Journey So Far
To understand why message 2584 was written, one must appreciate the trajectory that led to it. The SUPRASEAL_C2 pipeline generates Groth16 proofs for Filecoin PoRep, a computationally intensive operation that consumes ~200 GiB of peak memory and involves a complex chain from Curio's Go task layer through Rust FFI into C++/CUDA kernels. Each phase of optimization had progressively narrowed the bottleneck:
- Phase 7 introduced per-partition dispatch architecture, enabling parallel partition processing.
- Phase 8 implemented a dual-GPU-worker interlock, spawning two worker threads per GPU device to overlap CPU preprocessing with GPU kernel execution. This achieved 13–17% throughput improvement.
- Phase 9 optimized PCIe transfers by pre-staging VRAM allocations and uploading data asynchronously. In single-worker mode, this yielded 14.2% throughput improvement. However, in dual-worker mode, PCIe bandwidth contention between the two workers negated the benefit, returning throughput to roughly Phase 8 levels. The critical diagnostic came from fine-grained timing instrumentation added during Phase 9's analysis. The assistant and user discovered that GPU kernel time had dropped dramatically to ~1.8s per partition, but the CPU critical path—specifically
prep_msm(1.9s) andb_g2_msm(0.48s)—now dominated the per-partition wall time at ~2.4s. The GPU was idle for ~600ms per partition, waiting for the CPU thread to finish its MSM preprocessing. Moreover, at high concurrency, the ten synthesis workers competed with the CPU MSM operations for the 8-channel DDR5 memory bandwidth, inflating CPU times by 2–12×. The bottleneck had moved from PCIe transfers and GPU kernel execution to CPU memory bandwidth contention. This diagnosis set the stage for the two-lock design proposed in message 2584.
The Reasoning Behind the Two-Lock Design
The fundamental insight driving message 2584 is that the existing single-mutex design serializes two activities that could, in principle, overlap: VRAM management (allocation, upload, pool trimming) and GPU kernel execution. With a single std::mutex guarding the entire GPU-critical section, one worker holds the lock from pre-staging through kernel execution through cleanup, preventing any other worker from starting its VRAM setup even though the GPU is idle during the CPU-bound portions.
The assistant's reasoning, visible in the preceding messages (2573–2583), went through several iterations. The initial proposal from the user suggested three to four workers with two locks: a compute lock for GPU kernel execution and a memory-management lock for pre-staging allocation, upload, and cleanup. The assistant immediately identified a critical deadlock scenario:
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 is not a lock-ordering deadlock but a resource capacity deadlock—the 16 GiB GPU VRAM cannot simultaneously hold two workers' 12 GiB pre-staged buffers. The assistant correctly diagnosed that "the alloc and free MUST be serialized—you can't have Worker B pre-stage while Worker A's buffers are still live."
The breakthrough came when the assistant traced the exact VRAM lifecycle (task at message 2574) and realized that d_a and d_bc—the two largest allocations totaling 12 GiB—are freed synchronously via cudaFree inside the compute lock region. This means that by the time compute_mtx is released, that 12 GiB is genuinely available. No cooperative free slot or handoff protocol is needed. The design simplifies dramatically:
Since d_a and d_bc are freed synchronously (cudaFree) inside compute_mtx, by the time compute_mtx is released, that 12 GiB is genuinely free. We DON'T need a cooperative free slot.
This realization transformed the design from a complex nested-lock protocol with handoff slots into a clean two-phase approach: mem_mtx for VRAM management, compute_mtx for kernel execution, with no nesting and no deadlock risk.
The Implementation Plan in Detail
Message 2584 presents the implementation plan with the precision of an engineering specification. It is organized into six sections: a summary, changes by file, risk assessment, testing plan, and expected outcomes. Each section reflects deep knowledge of the codebase and the CUDA programming model.
The Core Structural Change
The plan replaces the single std::mutex* gpu_mtx with a struct containing two mutexes:
struct gpu_locks {
std::mutex compute_mtx; // Held during GPU kernel execution
std::mutex mem_mtx; // Held during VRAM alloc/free
};
The FFI boundary remains unchanged—create_gpu_mutex still returns void*, and destroy_gpu_mutex still takes void*. The Rust side, which orchestrates worker spawning and passes the opaque pointer, requires zero modifications. This is a deliberate design choice to minimize the blast radius of the change.
The Restructured Lock Regions
The plan specifies three distinct phases per worker, replacing the current monolithic lock region:
- MEM LOCK REGION (
mem_mtx): Pool trim,cudaMemGetInfo,cudaMallocford_aandd_bc(12 GiB), async upload, event creation. The key change is the removal ofcudaDeviceSynchronize—the assistant assumes that pool trim alone suffices because the previous worker's stream operations were all synchronized before it releasedcompute_mtx. - COMPUTE LOCK REGION (
compute_mtx): Per-GPU thread execution including NTT kernels, batch addition, tail MSM.d_bcis freed early viacudaFree(line 815 in the source), andd_ais freed by itsgpu_ptr_tdestructor. All internal MSM allocations are freed before the scope exits. - NO LOCK: Event and stream destruction (tiny objects, no VRAM impact),
cudaHostUnregister,prep_msm_thread.join(), and epilogue point arithmetic. This structure allows Worker B to acquiremem_mtxand start its VRAM allocation while Worker A is still running kernels undercompute_mtx. The overlap is not perfect—Worker B must wait forcompute_mtxto be released before it can run its own kernels—but it hides the CPU-side overhead ofprep_msmandb_g2_msmbehind the GPU kernel execution of the previous worker.
Assumptions Embedded in the Design
Message 2584 makes several assumptions that are critical to the plan's viability:
Assumption 1: Pool trim without cudaDeviceSynchronize is sufficient. The assistant writes: "Pool trim alone should suffice because by the time compute_mtx was released by the previous worker, all stream ops were synced." This assumes that cudaMemPoolTrimTo can reclaim memory from stream-ordered frees that were issued in a stream that has since been synchronized. In CUDA's memory pool model, this is generally true—pool trim is designed to release cached memory back to the OS. However, the assumption would prove incorrect in practice, as the subsequent chunk analysis reveals.
Assumption 2: Three workers will not saturate DDR5 bandwidth more than two. The risk assessment acknowledges this as a "Medium" likelihood concern but offers only monitoring as a mitigation. The user's earlier diagnosis had shown that at high concurrency (c=15–30), the ten synthesis workers competed with CPU MSM operations for memory bandwidth, inflating CPU times by 2–12×. Adding a third GPU worker would increase CPU-side work per partition, potentially worsening this contention.
Assumption 3: The fallback path is a sufficient safety net. The plan states: "If cudaMalloc fails (not enough VRAM), fall back to the original non-prestaged path inside the compute_mtx region." This assumes that the fallback path, which allocates VRAM lazily during kernel execution, is always available and does not introduce its own performance degradation. In practice, the fallback path would serialize with other workers' compute regions, potentially causing cascading delays.
Assumption 4: The barrier.wait() for prep_msm_thread synchronization remains correct. The plan notes that the barrier wait (line 840) stays inside compute_lock because the per-GPU thread runs under that lock. This assumes that the barrier protocol—where the prep_msm_thread signals completion of split vector computation—does not introduce ordering dependencies between the two locks. This is likely correct, but it adds a subtle coupling between the lock regions.
The Thinking Process Visible in the Message
Message 2584 is notable for what it reveals about the assistant's thinking process, even though it is presented as a finished plan rather than a stream of reasoning. The structure itself tells a story:
The plan begins with a summary that frames the goal in concrete terms: "hide the 0.5s b_g2_msm + epilogue overhead, reducing per-partition wall time from ~2.4s (CPU-limited) to ~1.8s (GPU kernel time)." This is not vague optimization—it is a specific, measurable target derived from the timing instrumentation of Phase 9.
The file-by-file breakdown shows systematic thinking about the entire codebase, not just the CUDA kernel file. The assistant checked the Rust FFI layer (lib.rs), the bellperson prover wrapper (supraseal.rs), and the engine configuration (engine.rs, config.rs). This cross-layer awareness is essential for a change that touches the lock protocol, because any mismatch between the C++ struct layout and the Rust opaque pointer handling would cause undefined behavior.
The risk assessment is particularly revealing. Each risk is assigned a likelihood and a mitigation, showing that the assistant is not blindly optimistic but has considered failure modes. The "Pool-cached memory not reclaimable without cudaDeviceSynchronize" risk is rated "Medium" likelihood—a prescient assessment, as this would become the root cause of the implementation's failure in the subsequent chunk.
The testing plan shows an understanding of the benchmark infrastructure: correctness first (-c 1 -j 1), then a sweep of concurrency levels (c=3 j=3, c=10 j=5, c=15 j=10), then comparison of per-partition timing metrics against the Phase 9 baseline. This mirrors the methodology used throughout the optimization journey.
The expected outcomes table is perhaps the most ambitious part of the plan. It projects a 30–38% throughput improvement in isolation and 15–27% at high concurrency. These numbers are derived from the timing analysis: if per-partition wall time drops from ~3.7s to ~1.8–2.0s, throughput should improve proportionally. The caveat—"The isolation number depends on whether prep_msm (1.9s) or compute (1.8s) is the bottleneck"—shows awareness that the actual improvement depends on which phase dominates after the change.
Input Knowledge Required
To fully understand message 2584, one needs knowledge spanning multiple domains:
CUDA Memory Management: The distinction between cudaFree (synchronous, immediate) and cudaFreeAsync (stream-ordered, pool-cached) is central to the design. The assistant knows that d_a and d_bc are freed via cudaFree, making their VRAM immediately available, while internal MSM allocations use cudaFreeAsync and require pool trim to reclaim. This understanding of CUDA's memory pool model is non-trivial.
Lock Design Patterns: The assistant recognizes the difference between a lock-ordering deadlock (avoided by consistent acquisition order) and a resource deadlock (caused by capacity constraints). The initial design with cooperative free slots shows familiarity with producer-consumer handoff patterns.
The Codebase Architecture: The plan references specific line numbers (line 641, lines 650-777, line 815, line 840) and specific variable names (d_a, d_bc, gpu_ptr_t, dev_ptr_t). This level of familiarity comes from the extensive tracing and code reading done in preceding messages.
Benchmark Methodology: The testing plan references specific benchmark commands (cuzk-bench batch --type porep --c1 /data/32gbench/c1.json) and concurrency configurations (c=3 j=3). This reflects the established testing infrastructure built over the optimization journey.
VRAM Capacity Constraints: The design is fundamentally shaped by the 16 GiB GPU VRAM limit and the ~13.8 GiB peak usage during H-MSM. The assistant knows that only ~2 GiB headroom exists, which is why pre-staging cannot overlap with the previous worker's kernel execution.
Output Knowledge Created
Message 2584 creates several forms of output knowledge:
A Concrete Implementation Plan: The message serves as a specification that could be handed to a developer for implementation. It identifies every file that needs to change, the nature of each change, and the rationale behind it. This is the kind of documentation that enables parallel work or future maintenance.
A Risk Registry: By explicitly enumerating risks and mitigations, the plan creates a shared understanding of what could go wrong. This is valuable for debugging—if the implementation fails, the team knows where to look first.
A Baseline for Comparison: The expected outcomes table provides quantitative targets. If the implementation achieves 30–38% throughput improvement, the design is validated. If it achieves less, the discrepancy reveals additional bottlenecks. If it achieves more, the model was too conservative.
A Testing Protocol: The testing plan establishes a repeatable methodology for evaluating the change. This is important for regression testing and for comparing future optimization phases against this baseline.
A Design Rationale: The message captures why certain decisions were made (no cooperative free slot needed, no cudaDeviceSynchronize in mem region, barrier stays inside compute lock). This rationale is essential for anyone who later needs to modify or extend the design.
What the Plan Got Wrong
Despite its thoroughness, message 2584 contains a critical incorrect assumption that would be revealed in the subsequent implementation and debugging (Chunk 1 of Segment 27). The assumption that cudaMemPoolTrimTo without cudaDeviceSynchronize is sufficient to reclaim pool-cached memory proved incorrect. In practice, cudaDeviceSynchronize and cudaMemPoolTrimTo are device-global operations that block while another worker holds compute_mtx and runs kernels. This effectively serialized the two locks, destroying the intended overlap.
The chunk analysis states: "The root cause was that cudaDeviceSynchronize and cudaMemPoolTrimTo inside mem_mtx are device-global operations that blocked while another worker held compute_mtx and ran kernels. This effectively serialized the two locks, destroying the intended overlap and forcing all subsequent partitions to use the slow fallback path."
The assistant had assumed that pool trim was a lightweight operation that could run concurrently with kernel execution on the same device. In reality, CUDA's memory pool management operations require device-wide synchronization, meaning that mem_mtx and compute_mtx could not actually run in parallel. The two-lock design degenerated into single-lock serialization, and the OOM failures from the fallback path caused performance to regress to 102s per proof—worse than the baseline.
This is a profound lesson in the limits of software abstraction over hardware. The assistant designed a clean locking protocol at the software level, but CUDA's device-global synchronization semantics overrode that abstraction. The hardware constraint—that memory management on a single CUDA device cannot be fully isolated from compute operations—made the two-lock design infeasible on a single GPU.
The Broader Significance
Message 2584 represents a pivotal moment in the optimization journey because it is the point where the team attempted to solve a CPU-side bottleneck (memory bandwidth contention) with a GPU-side architectural change (lock splitting). The subsequent failure revealed that the true bottleneck was not lock granularity but the fundamental coupling between CPU and GPU work on a shared memory bus.
The lesson extends beyond this specific codebase: when optimizing heterogeneous compute pipelines, the bottleneck often shifts between CPU, GPU, PCIe, and memory bandwidth in ways that resist clean software abstractions. The two-lock design was elegant in theory but violated a hardware constraint that no amount of software sophistication could work around.
The message also demonstrates the value of thorough design documentation. Because the plan was written down with explicit assumptions and risk assessments, the failure could be precisely attributed to a specific assumption (pool trim without device synchronize). Without this documentation, the debugging process would have been far more difficult.
Conclusion
Message 2584 is a masterclass in systems-level optimization design: it traces the VRAM lifecycle, identifies a clean lock separation, enumerates risks, specifies code changes across five files, and projects quantitative outcomes. It is the product of extensive analysis across multiple preceding messages, including code reading, task execution, and interactive Q&A with the user.
Yet the message also illustrates the fundamental unpredictability of optimization work. The most carefully reasoned design can fail when it encounters hardware constraints that were not fully accounted for. The assumption that cudaMemPoolTrimTo could run concurrently with kernel execution on the same device was reasonable from a software perspective but incorrect from a CUDA hardware perspective.
The ultimate outcome—a redesign of the locking strategy in subsequent phases—does not diminish the value of message 2584. On the contrary, the plan's explicit documentation of assumptions and risks made it possible to diagnose the failure precisely and iterate toward a working solution. In optimization engineering, knowing exactly why a design failed is often as valuable as knowing why it succeeded.