The 18 Optimizations: A Compute-Level Deep Dive into SUPRASEAL_C2 Groth16 Proof Generation

Introduction

In the previous segment of this coding session, the assistant had completed an architectural analysis of the SUPRASEAL_C2 Groth16 proof generation pipeline for Filecoin's Proof-of-Replication (PoRep), producing three optimization proposals focused on memory reduction and throughput architecture: Sequential Partition Synthesis (Proposal 1), Persistent Prover Daemon (Proposal 2), and Cross-Sector Batching (Proposal 3). These proposals addressed high-level structural inefficiencies—how the pipeline was organized, where memory was wasted, and how processes were orchestrated.

But architectural optimization, no matter how elegant, can only go so far. Once the structure of computation is fixed, the next layer of performance lies in the details: the cycles spent in each kernel, the bytes moved across each bus, the allocations made in each hot loop. This segment captures the systematic, bottom-up investigation of those details—a journey that began with a meta-synthesis of everything known, descended through five parallel deep dives into GPU kernels and CPU hotpaths, triaged follow-up investigations on three critical questions, and ultimately produced a document identifying 18 specific micro-optimizations with an estimated combined speedup of 30-43%.

The segment concludes with a pivot that opens an entirely new paradigm: the exploitation of constraint structure itself—specifically, the fact that ~99% of the auxiliary assignment consists of SHA-256 boolean circuits—for precomputation or mathematical batching transpositions. This question moves the optimization frontier from generic compute improvements to domain-specific circuit exploitation, potentially unlocking gains far beyond what micro-optimizations can achieve.

The Meta-Synthesis: Consolidating the Known

The segment opens at message 26 with a comprehensive meta-synthesis document that represents the assistant's consolidation of everything discovered across the earlier architectural phase and the initial compute-level investigations [1]. This message is not a typical conversational turn; it is a structured reference document that catalogs the full architecture, memory budget, circuit characteristics, GPU pipeline bottlenecks, CPU synthesis performance, and transfer patterns.

The meta-synthesis serves as the bridge between the architectural phase and the compute-level phase that follows. It is the moment when the assistant pauses to ask: What do we actually know now? The answer is organized into a layered taxonomy of inefficiency—architecture, memory budget, circuit characteristics, GPU pipeline, CPU synthesis, and transfer patterns—each representing a distinct layer of abstraction in the optimization problem.

Architecture and Memory Budget

The full call chain is traced from Go through CGO, Rust FFI, bellperson, and ultimately into C++ CUDA kernels ([msg 26]). Each proof runs in an isolated child process spawned by ffiselect, with SRS loaded from scratch every time (~47 GiB pinned memory, 30-90s deserialization). The memory budget for a 32 GiB sector with 10 partitions is staggering: ~120 GiB for the a, b, c proving assignment vectors, ~40 GiB for auxiliary assignment, ~47 GiB for the SRS, and ~5-15 GiB for C++ temporary structures, totaling a peak of ~195-220 GiB on the host. GPU peak memory is a more modest ~12-16 GiB for NTT workspace and MSM buckets.

Circuit Characteristics

The circuit for each partition contains ~130 million constraints and ~130 million auxiliary variables, with a FFT domain of 2^27 = 134,217,728. Critically, ~99% of auxiliary assignment values are boolean (0 or 1)—dominated by SHA-256 internal bits—while only ~1% are "significant" field elements (Poseidon intermediates, column values). This extreme sparsity in the witness values would later become a key insight for both the split MSM optimization (already implemented) and the constraint structure exploitation question posed at the end of the segment.

Identified Bottlenecks

The meta-synthesis identifies nine structural bottlenecks, including: all 10 partitions synthesized in parallel before any GPU work (a hard barrier via rayon::collect()), no early free of a/b/c vectors (Rust holds them while C++ runs, with no callback mechanism), cross-circuit bitmap union for split MSM requiring all circuits upfront, per-process SRS loading, global mutex serialization in C++, sequential B_G2 MSM on CPU, no transparent huge pages (THP) for multi-GiB allocations, and no NUMA awareness. These bottlenecks form the foundation for both the earlier architectural proposals and the compute-level optimizations to come.

The Pivot: From Architecture to Micro-Architecture

The user's directive at message 27 is the catalyst that transforms the investigation. The message reads: "Look further for more big ideas in compute optimization, improving everything like using more advanced avx/blas/better cpu/gpu cache use/less memory copy/paging in sequantial paths/etc. - continue this research." This is not a request for more architectural proposals; it is a mandate to descend from the architectural level to the instruction level—to examine not just what the pipeline does, but how it does it, cycle by cycle.

The assistant's response at messages 28-30 establishes a structured research plan. A todo list is created with high-priority items: review existing proposals to identify gaps, then deep-dive into NTT kernel optimization, MSM kernel optimization, CPU synthesis hotpath analysis, transfer pipeline analysis, and memory layout internals. The assistant reads all four existing documents (the background reference and Proposals 1-3) to ensure no gaps remain before launching new investigations. This methodological discipline—review before investigate—is a hallmark of the assistant's approach.

Message 30 marks the actual pivot point. The assistant declares "Good, I've reviewed everything. Now let me launch deep investigations in parallel." The choice to prioritize GPU kernel analysis over CPU synthesis is itself a decision: the assistant is implicitly betting that the GPU kernels—which handle the NTT and MSM computations—have more headroom for optimization than the CPU synthesis phase.

The Five Parallel Deep Dives

Message 31 is the data acquisition phase of the investigation. The assistant launches five parallel deep-dive research tasks, each targeting a distinct subsystem of the C2 pipeline. Each task result is a comprehensive technical document, ranging from several hundred to over a thousand words of analysis. The parallel execution pattern is a powerful demonstration of how deep systems analysis can be accelerated when the investigator can simultaneously process multiple information streams.

1. NTT Kernel Deep Dive

The Number Theoretic Transform (NTT) kernel analysis examined the sppark library's implementation for BLS12-381. The NTT uses a 3-step Gentleman-Sande decomposition for the 2^27 domain size, with mixed-radix wide kernels. Key findings included:

2. MSM Kernel Deep Dive

The Multi-Scalar Multiplication (MSM) kernel analysis examined the Pippenger implementation in pippenger.cuh. The MSM is broken into five phases: breakdown, sort, batch_add, accumulate, integrate, and collect. Critical findings included:

3. CPU Synthesis Hotpath Deep Dive

The CPU synthesis analysis examined the bellperson enforce() function, called ~130 million times per partition. Each call creates three ephemeral LinearCombination objects, each containing two Indexer instances with Vec<(usize, Scalar)> allocations. The analysis revealed:

4. Host-to-Device Transfer Pipeline Analysis

The transfer pipeline analysis studied all data movement between CPU and GPU. The key findings centered on the distinction between pageable and pinned memory:

5. Memory Layout and blst Internals

The memory layout analysis examined field element representation across the CPU (C++ and Rust) and GPU sides. The critical finding was that Rust and GPU use identical bit-level representation for Fr elements—both use Montgomery form with 4×uint64_t layout—meaning zero conversion overhead exists. This ruled out the hypothesis that Montgomery form transfer compression could save bandwidth.

The blst assembly analysis revealed the exact instruction sequences for field arithmetic: 31 MULX instructions for the multiplication core, 73 ADCX/ADOX instructions for the reduction chain, and branchless is_zero/is_one detection using vec_is_zero from vect.h. The analysis confirmed that no AVX-512 IFMA optimization is possible for the serial enforce() hotpath, as IFMA requires 8 independent multiplications simultaneously.

The Triage: Targeted Follow-Up Investigations

Messages 32-34 represent the convergence and triage phase. The assistant acknowledges that all five deep dives have returned comprehensive results and prepares to process the findings. But before moving to synthesis, three critical gaps are identified that require targeted follow-up investigations:

NUMA and Transparent Huge Pages

The deep dives revealed that the C2 pipeline operates on multi-gigabyte arrays—a, b, c vectors at ~4 GiB each, the SRS at ~47 GiB. On multi-socket machines, NUMA effects could significantly impact memory access latency. The investigation found that no THP/huge page support exists in the C2 code path, effectively ruling out this class of system-level optimization. The quantified impact was ~150 ms total TLB overhead on single-socket and ~100-300 ms NUMA penalty on multi-socket—small relative to the 250-360 second total proof time. The documentation already recommends NPS=1 (NUMA per-socket=1), which mitigates the NUMA penalty.

Streaming NTT Feasibility

Could the NTT be computed incrementally as synthesis produces constraint data? The analysis confirmed that the 3-step GS NTT for domain size 2^27 has inherent data dependencies that make true streaming difficult: Step 1 butterflies span the entire array with stride 2^26, meaning every element must be present before any NTT computation can begin. However, the investigation also identified opportunities for fusing the NTT with other computations, such as the LDE powers computation, which could eliminate redundant global memory passes.

Batch Addition Bitmap Optimization

The MSM deep dive had identified the batch addition routine as a potential bottleneck. The deep investigation of the batch_addition.cuh kernel revealed specific optimization opportunities in the bitmap scanning inner loop, including shared memory bank conflicts and suboptimal warp utilization patterns. The bitmap scan iterates over 1024 bits sequentially, but using __ffs (find first set) with Kernighan's trick could reduce the iteration count to O(popcount)—typically much smaller than 1024 for the sparse bitmaps in the split MSM.

The Synthesis: 18 Micro-Optimizations in Three Tiers

Message 35 is the critical synthesis moment. The assistant presents a structured decision document that organizes 18 distinct optimization opportunities into three impact tiers and definitively rules out seven previously plausible approaches. The tiering system is not an arbitrary ranking; it reflects a sophisticated decision framework that balances estimated impact, implementation effort, and risk.

Tier 1: High Impact (10-30% Each)

The four Tier 1 items target fundamental inefficiencies in the current implementation:

  1. Eliminate ~780 million heap allocations per partition in the CPU synthesis enforce() loop by replacing Vec<(usize, Scalar)> with SmallVec<[(usize, T); 4]>—since linear constraints typically have 1-3 terms, this eliminates nearly all heap traffic in the hottest code path. Estimated 15-30% faster synthesis time.
  2. Pin the a, b, c vectors with cudaHostRegister() to enable truly asynchronous DMA transfers. The current pageable memory means cudaMemcpyAsync blocks the CPU thread during staging, achieving only ~10-15 GiB/s instead of the theoretical 22-25 GiB/s. Estimated 0.3-0.5 seconds saved per proof.
  3. Pin the tail MSM bases to stop a pageable copy of already-pinned SRS points. A confirmed code path at groth16_cuda.cu:145-148 copies pinned data into pageable std::vector<affine_t> containers, defeating the purpose of pinning. Estimated 0.1-0.3 seconds saved per proof.
  4. Eliminate the cooperative kernel from batch_addition. The cooperative_groups::this_grid().sync() barrier forces exclusive GPU occupancy, preventing overlap between circuits. Splitting into two non-cooperative kernel launches enables pipelining. Estimated 5-15% wall-clock improvement.

Tier 2: Medium Impact (5-15% Each)

The seven Tier 2 items represent more specialized optimizations, each targeting a specific computational pattern:

  1. Fuse LDE_powers into NTT steps to eliminate redundant global memory passes (~8 separate LDE_powers kernel launches per proof).
  2. Port coalesced load/store from narrow to wide NTT kernels, improving intermediate-stage performance where strided access patterns dominate.
  3. Shared memory bank conflict mitigation for 32-byte Fr elements that span 8 banks in a 32-bank shared memory system, causing 8-way conflicts. Fix via SoA layout in shared memory or padding.
  4. Batch_addition occupancy improvement from ~12.5% to ~25% using __launch_bounds__(256, 2) to hint the compiler toward register spilling—a nuanced trade-off between register pressure and warp-level parallelism.
  5. Arena allocator for LC temporaries as an alternative or complement to SmallVec, using bumpalo thread-local arena for all LC allocations within enforce().
  6. Per-MSM window size tuning to create separate msm_t instances for each of the L, A, and B popcount distributions rather than using a single average-sized instance.

Tier 3: Lower Impact but Low Effort (1-5% Each)

The eight Tier 3 items are characterized by low implementation effort relative to their modest impact:

  1. Reuse d_b/d_c GPU allocations across circuits (currently dev_ptr_t uses synchronous cudaMalloc per circuit).
  2. Reuse msm_t object in NTT-MSM path (currently created per-circuit at groth16_ntt_h.cu:122).
  3. FFS-based bitmap scan in batch_addition replacing sequential 1024-iteration scan with __ffs+Kernighan's trick, achieving O(popcount) iterations.
  4. AVX-512 vectorized is_zero/is_one scan in prep_msm classification.
  5. Software prefetching in prep_msm for aux_assignment sequential scan.
  6. Parallelize B_G2 CPU MSMs across circuits—a one-line change from sequential loop to par_map, reducing the B_G2 phase from ~45 seconds to ~5 seconds.
  7. Prefetch affine points in batch_addition kernel to exploit 200+ cycle uadd latency.
  8. GPU-side reduction of batch_add buckets to eliminate DtoH + CPU sum_up.

The Ruled-Out Items: Knowledge Through Elimination

Perhaps the most valuable section of the synthesis is the "RULED OUT / NOT FEASIBLE" list. Seven hypotheses were actively investigated and rejected with concrete evidence:

The Delivery: From Analysis to Artifact

Messages 36-40 capture the transition from analysis to documentation. The user's single-word response at message 36—"Save"—is the critical approval gate that transforms exploratory research into actionable documentation. The assistant writes c2-optimization-proposal-4.md (590 lines) and provides a structured summary organized into three implementation waves with time estimates.

Wave 1: Quick Wins (1-2 Weeks)

Wave 1 contains six optimizations that can be implemented rapidly with minimal risk:

Wave 2: Medium Effort (2-4 Weeks)

Wave 2 tackles more complex kernel-level changes:

Wave 3: Deeper Kernel Work (4-8 Weeks)

Wave 3 contains the most technically challenging optimization:

The New Frontier: Constraint Structure Exploitation

The segment concludes with the user posing a fundamentally new question at messages 41-43: "Any optimizations possible for known constraing shapes, e.g. lots of sha256 constraints? or even more is the knowledge of constrainst structure usable in ways that with some pre-computation (compile or runtime) would allow for better performance/parallelism? Mathematical transpositions for batching possible?"

This question represents a paradigm shift. The 18 optimizations in Proposal 4 are all generic—they make the existing pipeline run faster without leveraging any knowledge of what the circuit actually contains. The user is now asking whether the structure of the constraints themselves—the fact that ~99% of auxiliary assignment consists of SHA-256 boolean circuits—can be exploited at a deeper level.

Three distinct lines of inquiry are proposed. First, precomputation: if the SHA-256 circuit is identical across many invocations, could intermediate values or partial proofs be cached and reused? Second, structural exploitation: could the constraint system be reorganized—for instance, by recognizing that SHA-256 constraints form a boolean circuit with known fan-in/fan-out patterns, permitting more efficient synthesis strategies? Third, mathematical transpositions for batching: could the algebra of the proof system itself be rearranged to enable batch processing of multiple constraints or multiple proofs simultaneously?

This question moves the optimization frontier from generic compute improvements to domain-specific circuit exploitation. It leverages the deep understanding of the constraint structure that was documented throughout the investigation—specifically, the dominance of SHA-256 boolean circuits and the deterministic nature of the constraint graph. If the user's intuition is correct, the resulting optimizations could be far more valuable than the 30-43% improvement from Proposal 4, potentially reducing proof time by factors beyond what generic micro-optimizations can achieve.

Conclusion: The Complete Arc of Optimization

This segment captures a complete arc of optimization work that is rare to see documented in such detail. It begins with a meta-synthesis of existing knowledge, pivots through a user directive into deep compute-level investigation, executes five parallel deep dives across the entire pipeline, performs targeted follow-up investigations to close critical gaps, synthesizes 18 optimization opportunities into a tiered proposal, delivers the document, and then pivots again to an entirely new paradigm of domain-specific circuit exploitation.

The arc demonstrates a fundamental principle of performance engineering: optimization is not a single pass but a layered process. Architectural optimizations (Proposals 1-3) change the structure of computation. Micro-optimizations (Proposal 4) make each individual operation faster. Domain-specific exploitation (the new direction) changes what is computed based on knowledge of the problem. Each layer builds on the previous one, and the combined effect is multiplicative, not additive.

The estimated combined speedup of 30-43% from Proposal 4, when layered on top of the 5-6x $/proof reduction from Proposals 1-3, represents a transformative improvement in the economics of Filecoin proof generation. And if the constraint structure exploitation direction proves fruitful, the gains could be even more dramatic. The journey from architecture to assembly—and now to algebra—is a testament to the power of systematic, deep investigation.

References

[1] [chunk 1.0] — From Architecture to Assembly: The Complete Arc of the SUPRASEAL_C2 Optimization Campaign