From Abandoned Locks to Targeted Interventions: The Phase 11 Optimization Campaign for Filecoin's Groth16 Proving Pipeline

Introduction

In the high-stakes world of Filecoin storage verification, every second of proof generation time translates directly into operational cost. The SUPRASEAL_C2 Groth16 proof generation pipeline — the computational engine behind Filecoin's Proof-of-Replication (PoRep) protocol — is a marvel of cross-language engineering, spanning Go task orchestration, Rust constraint systems, C++ thread pools, and CUDA GPU kernels. But with a peak memory footprint of approximately 200 GiB per proof cycle, the pipeline pushes modern server hardware to its limits. When multiple proofs run concurrently, DDR5 memory bandwidth becomes the critical bottleneck, and the question shifts from "how fast can we generate one proof?" to "how many proofs can we generate per hour without thrashing the memory subsystem?"

This article chronicles a pivotal sub-session in the ongoing optimization of this pipeline. It is a story of a failed design, a methodical retreat, a data-driven diagnosis, and the emergence of a new, evidence-based optimization plan. The session begins with the abandonment of a flawed Phase 10 two-lock GPU interlock design, proceeds through comprehensive benchmarking and waterfall timing analysis, and culminates in the design of Phase 11 — a set of three targeted interventions to reduce DDR5 memory bandwidth contention. Along the way, we see the systematic methodology that transforms raw performance data into actionable architectural decisions.

The Phase 10 Post-Mortem: Why the Two-Lock Design Failed

The investigation leading to this point was documented across ten detailed articles [1][2][3][4][5][6][7][8][9][10], which collectively map the memory-bandwidth landscape of the CUZK proving pipeline. These articles trace the systematic reconnaissance of the codebase — from initial glob searches [2] through targeted grep-driven discovery of each memory-heavy operation [7], culminating in a comprehensive catalog of six hot operations with precise file paths, line numbers, and data size estimates [10].

The session opens with a reckoning. Phase 10 had attempted to improve upon Phase 9's single-lock GPU interlock by introducing a two-lock design — splitting the GPU mutex into separate locks for different phases of GPU work, theoretically allowing more overlap between CPU-side preprocessing and GPU kernel execution. But the design ran into fundamental architectural obstacles that no amount of clever coding could overcome.

The first problem was VRAM capacity. The two-lock design required pre-staging buffers from multiple workers simultaneously on the GPU. With 16 GB of VRAM on a typical NVIDIA GPU (e.g., an RTX 4090 or A6000), there simply wasn't enough space to hold the a/b/c evaluation vectors, SRS point arrays, and intermediate MSM results for two concurrent workers. The pre-staged buffers alone — approximately 12.5 GiB of pinned host memory per worker — would exceed VRAM before any actual computation began.

The second problem was more fundamental: CUDA memory management APIs are device-global. Calls like cudaDeviceSynchronize() and cudaMemPoolTrimTo() operate on the entire GPU device, not on individual contexts or streams within it. This means that splitting the lock cannot create independent scheduling domains — any synchronization operation on one "lock's" work affects all other work on the same device. The lock split was architecturally meaningless because the underlying hardware doesn't support the abstraction it tried to create.

The team's response was disciplined and pragmatic: revert to Phase 9's proven single-lock design. No ego, no sunk-cost fallacy — just a clean rollback and a commitment to understand the problem better before attempting another fix. This decision set the stage for the systematic investigation that followed.

Benchmarking Across the Concurrency Landscape

With the codebase restored to the Phase 9 baseline, the assistant launched a comprehensive benchmarking campaign. The goal was to understand how the system behaves under varying levels of concurrency — specifically, how the number of concurrent proving operations (controlled by the c parameter, representing the number of Curio tasks) and the number of parallel jobs per task (the j parameter) affect throughput and latency.

The benchmarks spanned a wide range: from c=5 j=5 (light concurrency) through c=20 j=15 (heavy concurrency). For each configuration, the assistant collected throughput data (proofs per hour), per-proof latency, and — critically — the detailed TIMELINE events logged by the daemon. These TIMELINE events record the start and end times of each major pipeline phase: synthesis, prep_msm, GPU MSM, b_g2_msm, split_vectors allocation, cudaHostRegister/cudaHostUnregister, and async_dealloc.

The results painted a clear picture. At low concurrency, throughput scaled almost linearly — adding more workers produced proportionally more proofs per hour. But as concurrency increased, the scaling curve flattened. At c=20 j=15, GPU utilization reached 90.8%, yet throughput plateaued at approximately 38 seconds per proof. The GPU was busy, but the system was not producing proofs faster. Something was limiting the pipeline's ability to convert GPU utilization into proof throughput.

The Waterfall Timing Analysis: Identifying the Bottleneck

To understand the plateau, the assistant performed a detailed waterfall timing analysis. By extracting the TIMELINE events from the daemon logs and aligning them by timestamp, the assistant could visualize how different pipeline phases overlapped across concurrent workers. This analysis revealed the root cause: DDR5 memory bandwidth contention.

The key insight came from comparing the timing of two specific phases: synthesis (the CPU-bound PCE MatVec and witness generation) and prep_msm (the CPU preprocessing that prepares data for the GPU). Under low concurrency, these phases ran sequentially or with minimal overlap. But under high concurrency, they began to overlap significantly — and when they overlapped, both phases inflated in duration.

Synthesis, which normally completes in approximately 8-10 seconds for a single circuit, stretched to 14-16 seconds under heavy concurrency. Prep_msm, normally a 3-4 second phase, inflated to 6-8 seconds. The sum of the inflations was roughly equal to the overlap time, suggesting that the two phases were competing for the same shared resource: DDR5 memory bandwidth.

This diagnosis was consistent with the hardware architecture. A modern AMD Zen4 server with 96 cores has eight DDR5 memory channels, providing approximately 400-500 GiB/s of aggregate bandwidth. The PCE MatVec during synthesis reads approximately 49 GiB of data (722 million non-zero entries × 68 bytes per entry), while prep_msm reads and writes approximately 24 GiB (12 GiB read + 12 GiB write). When both phases run simultaneously on different cores, they collectively demand more bandwidth than the memory controller can supply, causing each phase to slow down as threads stall waiting for data.

The waterfall analysis also revealed a secondary contention point: the b_g2_msm phase, a CPU-only G2 multi-scalar multiplication that runs after GPU work completes. This phase overlaps with the next proof's GPU kernel execution by design (the GPU lock is released before the prep_msm_thread is joined). But under high concurrency, the G2 MSM's memory traffic — approximately 4.9 GiB of G2 point reads — competes with the next proof's synthesis, creating a three-way contention pattern.

Designing Phase 11: Three Interventions

With the diagnosis confirmed — DDR5 memory bandwidth contention between synthesis, prep_msm, and b_g2_msm — the assistant designed Phase 11, a set of three targeted interventions to reduce contention. Each intervention addresses a specific mechanism identified in the waterfall analysis.

Intervention 1: Bounding async_dealloc to a Single Thread

The first intervention targets a subtle but significant source of memory bandwidth waste: TLB shootdown storms caused by concurrent munmap() calls. The async_dealloc thread frees approximately 37 GiB of C++ data structures and 130 GiB of Rust ProvingAssignment vectors per 10-circuit batch. When multiple workers' async_dealloc threads run concurrently, their munmap() calls trigger inter-processor interrupts (TLB shootdowns) on all cores, stalling memory accesses across the entire system.

The fix is straightforward: serialize the async_deallocation work onto a single thread. Instead of each worker spawning its own detached std::thread for destruction, a single dedicated deallocation thread processes a queue of destructors. This eliminates concurrent TLB shootdowns and reduces the peak deallocation latency from "all workers contend" to "one worker at a time."

The expected improvement is modest but meaningful: approximately 1-3% throughput gain by eliminating the TLB shootdown overhead that currently inflates synthesis and prep_msm times during concurrent deallocation.

Intervention 2: Reducing the groth16_pool Thread Count

The second intervention addresses the most direct source of memory bandwidth contention: the groth16_pool thread pool. This C++ thread pool, controlled by the CUZK_GPU_THREADS environment variable, defaults to all available CPUs (192 threads on a 96-core Zen4 with SMT). It is used for three parallel operations: split_vectors construction, prep_msm bit packing and scalar extraction, and b_g2_msm computation.

The problem is that 192 threads running Pippenger's algorithm simultaneously generate enormous memory traffic. Each thread reads SRS point arrays and scalar vectors, and with 192 threads competing for L3 cache (32 MiB shared across all cores on Zen4), the cache becomes a bottleneck. Threads spend more time waiting for data from DDR5 than computing, and their cache misses evict data that synthesis threads (running on the same cores via rayon) need.

The proposal is to reduce gpu_threads from the default 192 to approximately 32. This may seem counterintuitive — reducing parallelism to improve performance — but it is well-supported by the waterfall data. With 32 threads, the b_g2_msm phase's memory footprint shrinks, L3 cache pressure on synthesis is reduced, and the thread pool's overhead (context switching, cache misses) drops. The expected improvement is 3-5% throughput gain, with the added benefit of freeing CPU cores for synthesis work.

Intervention 3: A Lightweight Atomic Throttle for Synthesis

The third intervention is the most innovative. The waterfall analysis revealed that b_g2_msm creates a brief but intense window of memory bandwidth demand — approximately 0.4 seconds during which all 192 groth16_pool threads run Pippenger simultaneously. During this window, synthesis threads (running on the same cores via rayon's work-stealing scheduler) compete for the same DDR5 channels.

The proposal is a lightweight atomic throttle flag. Before entering the b_g2_msm computation, the prep_msm_thread sets an atomic flag. Synthesis workers, at the start of each row chunk in the SpMV loop, check this flag. If set, they yield briefly (e.g., std::this_thread::yield() or a short spin_loop_pause()) before proceeding. This is not a hard lock — it is a cooperative hint that allows synthesis to slow down slightly during the bandwidth-critical window, reducing contention without introducing the complexity of a full interlock.

The expected improvement is modest — approximately 1-3% — but the implementation cost is near zero. A single atomic boolean, a few lines of code in the SpMV loop, and a yield() call. The throttle is also composable with the other two interventions, and its effect can be measured independently by toggling it on and off.

Refining the Plan: The prep_msm Discovery

The initial Phase 11 design included a more aggressive semaphore-based interlock between prep_msm and synthesis. The idea was to prevent prep_msm from running while synthesis was active, and vice versa. But further analysis of the code revealed a critical detail that rendered this approach overkill.

The prep_msm function, when called with num_circuits=1 (the common case for individual proof generation), uses par_map with a single item. Rust's par_map with one item runs on a single thread — there is no parallel execution. This means that for single-circuit proofs, prep_msm is effectively single-threaded and its memory bandwidth demand is negligible compared to the multi-threaded synthesis phase.

The full semaphore interlock was therefore unnecessary. The only bandwidth-critical overlap is between b_g2_msm (which uses the full thread pool) and synthesis (which uses rayon). The atomic throttle addresses this specific overlap without the complexity of a full interlock.

This discovery illustrates the importance of verifying assumptions against actual code behavior. The initial design assumed prep_msm was a multi-threaded bandwidth hog; the code analysis revealed it was single-threaded for the common case. By checking the implementation rather than guessing, the team avoided building an overly complex solution for a non-existent problem.

Documentation and Project Management

The Phase 11 design was not left as a conversation artifact. The assistant produced c2-optimization-proposal-11.md, a detailed design specification covering all three interventions with implementation steps, risk assessments, and expected throughput improvements. The document includes:

The Broader Significance

The Phase 11 optimization campaign represents a mature approach to performance engineering. It is not about finding a single magic optimization that doubles throughput — it is about understanding the system's bottlenecks at multiple levels, designing targeted interventions, and measuring their effects independently.

The three interventions — bounding async_dealloc, reducing the thread pool, and adding an atomic throttle — are individually modest, with expected improvements of 1-5% each. But their combined effect, projected at 3-11% throughput improvement, is meaningful for a production system where every percentage point translates into real cost savings. More importantly, the methodology that produced them — benchmark, analyze, hypothesize, verify, implement — is reusable for future optimization cycles.

The session also demonstrates the value of failing fast and learning from failure. The Phase 10 two-lock design was abandoned not because it was difficult to implement, but because it was architecturally unsound. The team recognized this, reverted cleanly, and invested the saved time in understanding the actual bottleneck. This discipline — knowing when to cut losses and pivot to a better approach — is rare in engineering and invaluable in performance work.

Conclusion

The Phase 11 optimization campaign for the SUPRASEAL_C2 Groth16 proof generation pipeline is a case study in systematic performance engineering. It begins with the abandonment of a flawed design, proceeds through comprehensive benchmarking and waterfall timing analysis, identifies DDR5 memory bandwidth contention as the root cause, and designs three targeted interventions to reduce that contention. The plan is documented, committed, and ready for implementation.

The expected outcomes — 3-11% throughput improvement through three modest code changes — are achievable and measurable. But the lasting value of this session is not the specific interventions. It is the methodology: the willingness to abandon a failed approach, the discipline to gather data before designing solutions, the precision to identify specific contention mechanisms, and the humility to verify assumptions against actual code behavior. These habits, more than any single optimization, are what separate effective performance engineering from guesswork.## References

[1] "Mapping the Memory Bandwidth Battlefield: A Deep Dive into the cuzk Proving Pipeline Investigation" — Analyzes the opening message of the subagent session, revealing the user's six-category research mandate and the strategic thinking behind the reconnaissance effort.

[2] "Mapping the Memory Bandwidth Frontier: The Initial Reconnaissance in a Groth16 Proving Pipeline Investigation" — Examines the first assistant response, a set of parallel glob searches that map the codebase terrain across three major subsystems.

[3] "Reading the Witness and Persistence Layers: A Methodical Step in Mapping the cuzk Memory Bandwidth Pipeline" — Details the reading of witness_cs.rs and disk.rs, filling gaps in the data lifecycle understanding.

[4] "The Final Piece: Completing the Memory-Bandwidth Investigation with Targeted Searches" — Covers the closing-the-loop grep for CUZK_GPU_THREADS and the WitnessCS implementation.

[5] "Systematic Codebase Reconnaissance: Reading the Remaining Key Files in the cuzk Proving Pipeline Investigation" — Analyzes the selection of six "remaining key files" including density bitmaps, recording CS, prover, pipeline, NTT kernels, and SRS management.

[6] "Reading the Source: How a Subagent Begins Mapping Memory-Bandwidth Hotspots in the cuzk Proving Pipeline" — Examines the foundational read of CSR matrix code, SpMV evaluator, PCE library, and CUDA kernel files.

[7] "Precision Targeting: How a Single Grep-Driven Message Mapped the Memory Bandwidth Hotspots in a Groth16 Proving Pipeline" — Analyzes the five-grep message that located async_dealloc, cudaHostRegister, split_vectors, b_g2_msm, and thread_pool_t.

[8] "Tracing the Tail: How One grep Uncovered the 130 GiB Deallocation Bottleneck in Filecoin's CUZK Prover" — Details the discovery of the Rust-side async_dealloc thread and its ~130 GiB memory footprint.

[9] "The Orchestration Layer: A Systematic Deep-Dive into the cuzk Proving Pipeline's Mid-Level Architecture" — Covers the pivot to reading the bellperson supraseal prover and the engine/scheduler code.

[10] "Mapping the Memory Bandwidth Battlefield: A Deep Dive into CUZK's Six Hot Operations" — The comprehensive catalog of six memory-bandwidth-heavy operations with data size estimates, parallelization analysis, and the summary table.