From Validation to Optimization: Phase 4 of the CUZK Pipeline Begins

Introduction

The investigation into the SUPRASEAL_C2 Groth16 proof generation pipeline for Filecoin PoRep had reached a pivotal moment. Phase 3 had just concluded with a resounding validation: cross-sector batching worked, delivering a 1.42× throughput improvement on real 32 GiB PoRep data with an RTX 5070 Ti. The architecture—BatchCollector, synth_batch_full, synthesize_porep_c2_multi, and split_batched_proofs—had been proven correct across four systematic tests covering timeout flush, batch-of-2 synthesis amortization, 3-proof overflow handling, and WinningPoSt bypass. The results were documented, committed, and the path forward was clear.

But validation is not optimization. The Phase 3 architecture proved that batching could work, but it did not make the pipeline fast. The baseline single-proof time stood at 89 seconds—54.7 seconds for synthesis and 34 seconds for GPU compute. The peak memory footprint remained at ~203 GiB for a single proof, ballooning to ~360 GiB for batch-of-2. The system worked, but it was far from efficient.

This chunk of the conversation marks the beginning of Phase 4: Compute-Level Optimizations. It is a transition from proving what is possible to engineering how to make it efficient. The work proceeds along two parallel tracks: a main session that implements the highest-impact micro-optimizations from the proposal document, and a subagent session that performs deep reconnaissance into the GPU kernel code—specifically the sppark MSM library—to understand the computational primitives that dominate proof generation time. Together, these tracks represent the first concrete steps toward transforming a working proof-of-concept into a production-grade proving pipeline.

The Dependency Chain: Understanding What Can Be Modified

Before any optimization could be implemented, the team needed to understand the dependency chain. The SUPRASEAL_C2 pipeline is not a monolithic codebase; it is a layered stack of dependencies, each with its own build system and release cycle. The critical insight was that two key components—bellpepper-core (the LC Indexer, responsible for the rank-1 constraint system transformation) and supraseal-c2 (the C++/CUDA layer that orchestrates GPU proof generation)—both come from crates.io as published Rust packages. Modifying them required local forks.

The team created extern/bellpepper-core/ and extern/supraseal-c2/ directories within the project workspace, then used Cargo's [patch.crates-io] mechanism to redirect the dependency resolution to these local copies. This is a standard Rust technique for working with upstream dependencies that need modification, but its application here reveals a deliberate architectural decision: rather than maintaining separate forks with their own versioning and merge overhead, the team chose to patch in-place, keeping the modifications tightly coupled to the cuzk project's build.

This decision had immediate consequences. With the patched dependencies in place, the team could modify the LC Indexer's allocation strategy, the proving assignment's memory management, and the GPU kernel launch parameters—all without waiting for upstream maintainers to accept changes. The cost was the responsibility of keeping the forks synchronized with upstream releases, but for the rapid iteration of Phase 4, this was a acceptable trade-off.

Four Optimizations, One Benchmark

The optimization proposal (c2-optimization-proposal-4.md) had identified several high-impact changes. The team selected four for the initial implementation wave, each targeting a different layer of the pipeline:

A1: SmallVec for the LC Indexer

The LC Indexer (bellpepper-core) is responsible for transforming the rank-1 constraint system into a linear combination representation suitable for the prover. During this transformation, it creates millions of small vectors—each typically containing 1–4 (usize, Scalar) pairs representing variable assignments. The original code used Vec<(usize, Scalar)>, which means every single one of these small vectors incurred a heap allocation. With approximately 780 million such allocations per partition, the cumulative overhead was substantial.

The fix was elegant: replace Vec<(usize, Scalar)> with SmallVec<[(usize, Scalar); 4]>. SmallVec is a Rust crate that provides a vector-like container with a fixed-size inline buffer. When the vector's length is within the inline capacity (4 elements in this case), no heap allocation occurs. Since the vast majority of LC indexer vectors have 1–4 elements, this change eliminates nearly all of the 780 million heap allocations per partition. The memory savings come not from reducing total bytes allocated, but from eliminating the per-allocation overhead—the malloc bookkeeping, the cache-line pollution, and the deallocation cost.

A2: Pre-sizing the ProvingAssignment

The ProvingAssignment is the data structure that accumulates the a, b, and c witness vectors during synthesis. As constraints are evaluated, new assignments are appended, and the underlying storage grows. In the original implementation, the ProvingAssignment used Rust's standard Vec, which doubles in capacity whenever it runs out of space. For a structure that ultimately holds ~328 GiB of data per partition (as the benchmark later revealed), this doubling strategy causes repeated reallocations and copies of massive memory regions.

The optimization added a new_with_capacity constructor that allows the caller to pre-allocate the exact capacity needed. This eliminates the reallocation copies entirely—a single allocation at the start, sized to fit the entire partition's witness data. The potential savings were estimated at ~32 GiB of avoided reallocation copies per partition.

A4: Parallelizing B_G2 CPU MSMs

The Groth16 proof requires a multi-scalar multiplication (MSM) on the G2 curve for the B component. In the original code, this MSM was performed sequentially on the CPU—a single thread iterating over the scalars and points. The optimization replaced this sequential loop with groth16_pool.par_map, distributing the work across all available CPU cores using the existing thread pool.

This is a classic embarrassingly parallel workload: the MSM decomposes into independent bucket accumulations that can be computed in parallel and then combined. The optimization required minimal code changes—essentially wrapping the loop body in a parallel iterator—but the impact on wall-clock time was expected to be proportional to the number of cores available.

B1: Pinning a, b, c Vectors with cudaHostRegister

The a, b, and c witness vectors are produced by the CPU during synthesis and must be transferred to the GPU for the MSM computation. In the original implementation, these vectors were allocated as standard pageable memory. When the GPU accessed them via direct memory access (DMA), each page fault triggered a copy from pageable to pinned memory—a slow, serialized operation that can dominate transfer time for large buffers.

The fix added cudaHostRegister and cudaHostUnregister calls around the Rust-provided arrays, pinning the memory pages so that the GPU can access them directly without page fault overhead. This is a standard CUDA optimization for host-device transfers, but its application here required careful integration with Rust's memory management: the Rust Vec that owns the data must remain alive and unmoved while the GPU accesses the pinned region.

D4: Per-MSM Window Tuning

The MSM algorithm uses a windowing method where scalars are split into fixed-size windows, and points are accumulated into buckets corresponding to each window's value. The window size is a critical parameter: larger windows reduce the number of bucket accumulations but increase the number of buckets (and thus memory usage and final reduction cost). The optimal window size depends on the number of points being multiplied.

The original code used a single msm_t instance with uniform window parameters for all MSM operations. The optimization split this into three separate instances, each tuned for the specific popcount (number of non-zero scalars) of the L, A, and B_G1 MSMs. Since these MSMs have different sizes and sparsity patterns, custom window sizes improve both compute efficiency and memory bandwidth utilization.

The Regression: When Theory Meets Practice

With the four optimizations implemented, the team ran an end-to-end single-proof benchmark on the RTX 5070 Ti with real 32 GiB PoRep data. The result was sobering: 106 seconds total, compared to the 89-second baseline from Phase 3. A 19% regression.

The breakdown revealed the culprits:

The Response: Revert and Instrument

The team's response was immediate and disciplined. The A2 hint usage was reverted in the synthesis call sites—the new_with_capacity API was kept available (it might be useful with a different allocation strategy), but the synthesis code fell back to the original Vec-based growth. The B1 optimization was kept in place, but the team recognized that its overhead needed to be understood and potentially mitigated through batched pinning or alternative transfer strategies.

More importantly, the team added detailed phase-level timing instrumentation to the CUDA code using std::chrono. The instrumentation measures seven distinct phases:

  1. prep_msm: Preparation of MSM inputs on the GPU
  2. NTT+MSM_H: Number-theoretic transform and MSM on the H component
  3. batch_addition: The batched point addition kernel (the core of the MSM bucket accumulation)
  4. tail MSMs: The final reduction MSMs that combine bucket results
  5. B_G2: The G2 MSM for the B component
  6. proof assembly: Combining the MSM results into the final Groth16 proof
  7. pin/unpin: The overhead of cudaHostRegister and cudaHostUnregister calls This instrumentation is the foundation for the next round of optimization. Instead of guessing which changes help or hurt, the team can now run precise A/B tests, measuring each phase's contribution to the total time. If a proposed optimization speeds up batch_addition by 5% but slows prep_msm by 10%, the instrumentation will reveal the net effect.

The Subagent Reconnaissance: Exploring the Sppark MSM Code

While the main session implemented and benchmarked the Phase 4 optimizations, a subagent session was dispatched to perform deep reconnaissance on the sppark MSM code. The subagent's mission was to find and read two critical artifacts: the batch_addition.cuh GPU kernel file and the max_num_circuits constant definition.

This reconnaissance was not academic. The batch_addition.cuh file contains the cooperative GPU kernel that performs batched point additions during MSM—the computational heart of the proving system's most expensive phase. Understanding its bitmap scan pattern, its use of cooperative grid synchronization, and its __launch_bounds__ configuration is essential for any optimization targeting the GPU MSM. The max_num_circuits constant controls how many concurrent circuits the system can process, directly impacting memory usage and the feasibility of batching multiple sectors.

The search journey itself is a case study in codebase navigation. The first attempt used glob patterns (**/sppark/**/batch_addition.cuh and **/batch_addition.cuh), but both returned nothing. The file simply was not where the glob engine expected it. A grep for max_num_circuits found 18 matches, but only in Markdown design documents—not in source code. The assistant had found references about the constant but not the constant itself.

The pivot came in message 2, when the assistant shifted from targeted globbing to exploratory discovery. Instead of searching for the file, it searched for the directory structure: **/sppark/** returned nothing, confirming that the sppark directory was not at the expected location. But **/supra_seal/**/*.cuh and **/supraseal/**/*.cuh both returned results, including groth16_srs.cuh—the file that the documentation had identified as containing max_num_circuits. A bash find for directories matching *sppark* revealed only Rust build fingerprint directories, not the source tree.

The breakthrough came in message 3, when the assistant used find on the specific path extern/supraseal/deps/sppark and located batch_addition.cuh at extern/supraseal/deps/sppark/msm/batch_addition.cuh. A broader find confirmed this was the only copy. The directory listing revealed the full structure of the sppark library: subdirectories for ec/, ff/, go/, msm/, ntt/, poc/, polynomial/, rust/, and util/.

Message 4 delivered the payoff: the contents of batch_addition.cuh and both copies of groth16_srs.cuh. The batch addition kernel revealed its structure—a CUDA header with cooperative groups support, including cuda.h, cooperative_groups.h, and vector. The SRS headers confirmed static const int max_num_circuits = 10; at line 62, identical in both supraseal and supra_seal copies.

Message 5 presented the complete findings, extracting the four specific code sections the user had requested: the bitmap scan loop (lines 63–75), the point load (line 86), the cooperative grid sync (line 107), and the __launch_bounds__ annotation (lines 119–123). Each section was accompanied by engineering analysis explaining its significance—the warp-level shuffle pattern, the memory access coalescing, the grid synchronization constraint, and the register allocation trade-off.

The Convergence: What the Reconnaissance Revealed

The subagent's findings have direct implications for the Phase 4 optimization work:

The bitmap scan pattern in batch_addition.cuh uses __shfl_sync to broadcast bitmap words across warp lanes, with each lane claiming the j-th set bit. This is efficient for dense bitmaps but may suffer from load imbalance when bitmaps are sparse—a concern for the tail MSMs where few points remain.

The cooperative grid sync at line 107 means the kernel must be launched with cudaLaunchCooperativeKernel, which imposes constraints on grid dimensions and GPU architecture support. Any optimization that changes the grid size (e.g., to use fewer blocks for smaller MSMs) must respect this constraint.

The max_num_circuits = 10 constant is the hard limit that the Cross-Sector Batching proposal must overcome. Bumping it to 30 or higher would enable larger batch sizes, but the get_num_threads() function shows that increasing max_num_circuits also reduces the CPU threads available for SRS loading—a trade-off that becomes irrelevant if the SRS is persistently loaded (as proposed in the Persistent Prover Daemon).

The dual-directory structure (both supraseal and supra_seal containing identical copies of groth16_srs.cuh) means that any modification to max_num_circuits must be applied to both copies, or the build system must be updated to use a single canonical source.

Looking Ahead: The Instrumented Path

The Phase 4 work is not complete—it has barely begun. The regression from the initial benchmark was a valuable lesson: not all optimizations are beneficial in isolation, and the interaction between memory allocation strategies and operating system page management can dominate performance. The detailed timing instrumentation now in place provides the feedback loop needed to iterate toward genuine improvement.

The subagent's reconnaissance has mapped the critical GPU kernel code, revealing the bitmap scan strategy, the cooperative synchronization pattern, and the hardcoded circuit limit. These findings will inform the next wave of optimizations: perhaps a revised pinning strategy that batches cudaHostRegister calls, or a modified bitmap scan that handles sparsity more efficiently, or a refactored max_num_circuits that supports dynamic batch sizes.

The overarching theme of this chunk is the shift from validation to optimization—from proving that an architecture works to engineering it to be efficient. The Phase 3 validation showed that cross-sector batching could achieve 1.42× throughput improvement. Phase 4 aims to make that improvement reliable, predictable, and cost-effective. The regression was a setback, but the instrumentation and reconnaissance that followed it have laid the groundwork for the next breakthrough.