The b_g2_msm Investigation: When a 25-Second CPU Bottleneck Met a Pipeline That Already Solved It

Introduction

In the high-stakes world of Filecoin proof-of-replication (PoRep) proving, every millisecond counts. The Groth16 proof generation pipeline in supraseal-c2 is a massively parallel system that coordinates CPU and GPU resources to generate zk-SNARK proofs for storage verification. But like any complex system, it harbored a bottleneck so stark that it demanded immediate attention: a single CPU-side computation called b_g2_msm — a multi-scalar multiplication over BLS12-381 G2 points — consuming approximately 25 seconds out of a 27-second total proving time [1]. All other multi-scalar multiplications (h, l, a, b_g1) already ran efficiently on the GPU, completing in roughly 2 seconds combined. The asymmetry was tantalizing: if b_g2_msm could be moved to the GPU, the entire CPU contention problem would vanish.

What followed was a systematic, multi-phase investigation spanning over twenty messages, multiple subagent sessions, and deep dives into CUDA template metaprogramming, Rust pipeline orchestration, and GPU memory architecture. The investigation's ultimate finding was both surprising and elegant: the problem had already been solved by an architectural change elsewhere in the system — the partitioned pipeline — which reduced b_g2_msm from ~25 seconds to ~0.4 seconds by ensuring num_circuits=1 per GPU call. This article synthesizes that investigation, tracing the path from initial reconnaissance through the critical insight to the final feasibility assessment, and examines what this case reveals about systems-level optimization in complex codebases.

The Problem: A 25-Second CPU Bottleneck in a 27-Second Pipeline

The investigation began with a stark performance datum [1]. The GPU proving pipeline for Filecoin PoRep C2 proofs took approximately 27 seconds total. Within that window, roughly 25 seconds were consumed by a single CPU-side operation: b_g2_msm, a Pippenger multi-scalar multiplication over BLS12-381 G2 points. The other MSMs — h, l, a, and b_g1 — all ran on the GPU via the msm_t template class from the sppark library, completing in about 2 seconds combined. The GPU sat idle for most of the pipeline's duration while the CPU labored on this one computation.

The root cause lay in how the pipeline handled multiple circuits. When num_circuits > 1 (the batch-all path with 10 circuits), each circuit ran a single-threaded Pippenger MSM in parallel via a thread pool. Ten single-threaded MSMs contending for CPU resources resulted in approximately 25 seconds of wall-clock time. When num_circuits == 1, a single circuit could use the full thread pool for a parallelized Pippenger implementation, completing in approximately 0.4 seconds [20]. This 62.5× difference in throughput per circuit was the key to everything that followed.

The user's initial request [1] laid out six investigative threads: read the b_g2_msm code in groth16_cuda.cu, check for existing G2 CUDA kernels, assess VRAM feasibility, search for existing GPU G2 MSM implementations in other libraries, evaluate the alternative of using the partitioned pipeline (which might already reduce b_g2_msm via num_circuits=1), and verify the partitioned pipeline path. This structured approach set the stage for a methodical investigation.

Phase 1: Systematic Codebase Reconnaissance

The assistant began with broad reconnaissance [2]. It read the main CUDA source file groth16_cuda.cu, globbed for .cuh and .cu files across the project tree, and ran grep searches for fp2, Fp2, FP2, g2_msm, and related terms. The results were promising: 170 matches for Fp2-related patterns, including a proof-of-concept directory sppark/poc/msm-cuda/ containing an msm_fp2_correctness test and a multi_scalar_mult_fp2_arkworks function [3].

A critical early finding was that the user's assumption about function names was incorrect. The grep for msm_cuda returned no matches — the GPU MSM infrastructure was not built around standalone msm_cuda_* functions but rather a type-parameterized template class called msm_t<bucket_t> defined in pippenger.cuh [3]. This distinction was crucial: the path to GPU G2 MSM was not about writing new CUDA kernels but about instantiating existing templates with G2 types.

The assistant then pivoted to deeper investigation [11]. It read the msm_t template class definition in pippenger.cuh, discovering that it was parameterized on bucket_t, point_t, affine_t, and scalar_t — making it theoretically compatible with G2 types. It examined pippenger_inf.cu in the proof-of-concept directory, which included bls12-381-fp2.hpp and contained a mult_pippenger_fp2_inf function calling the same mult_pippenger<bucket_fp2_t> template. It read the Fp2 field arithmetic header bls12-381-fp2.hpp, confirming that GPU device functions for Fp2 arithmetic existed [3].

Phase 2: The Critical Insight — The Partitioned Pipeline

The investigation took a decisive turn when the assistant examined the partitioned pipeline code in pipeline.rs [21]. The user had hypothesized that the partitioned pipeline — which processes one partition at a time rather than batching all 10 — might already use num_circuits=1 per GPU call, making the b_g2_msm bottleneck disappear. The assistant read the architecture documentation and found the smoking gun: a comment at line 1749 of pipeline.rs stating "With num_circuits=1 per GPU call, the b_g2_msm use..." [5].

The assistant traced the dispatch logic in prove_porep_c2_slotted() [21]:

Phase 3: The GPU Template Infrastructure

With the partitioned pipeline insight in hand, the assistant nevertheless continued to investigate the GPU G2 MSM feasibility, recognizing that the batch-all path and future use cases might still benefit from GPU acceleration [6][7][12].

The central question was whether the msm_t template class in pippenger.cuh could be instantiated with G2 types. The template signature was:

template<class bucket_t, class point_t, class affine_t, class scalar_t,
         class affine_h = class affine_t::mem_t,
         class bucket_h = class bucket_t::mem_t>
class msm_t { ... };

The assistant discovered that the GPU Pippenger kernel used bucket_t::degree to determine thread layout [7][13]. The degree constant controlled how many field elements each coordinate occupied — degree = 1 for G1 (a single Fp element per coordinate) and degree = 2 for G2 (an Fp2 extension field with two Fp elements per coordinate). The kernel code at line 129 of pippenger.cuh read:

#define ACCUMULATE_NTHREADS (bucket_t::degree == 1 ? 384 : 256)

And at lines 161-163:

const uint32_t degree = bucket_t::degree;
const uint32_t warp_sz = WARP_SZ / degree;
const uint32_t lane_id = laneid / degree;

The assistant traced the degree definition to xyzz_t.hpp line 25: static const unsigned int degree = field_t::degree [13]. This confirmed that the degree was inherited from the field type — fp_t for G1 (degree=1) and fp2_t for G2 (degree=2). The GPU kernel was already designed to handle extension field arithmetic through this compile-time parameterization.

Further evidence came from the batch addition kernel. The assistant found that batch_addition&lt;bucket_fp2_t&gt; was explicitly instantiated as a GPU __global__ kernel in groth16_split_msm.cu, and execute_batch_addition&lt;bucket_fp2_t&gt; was called from the main CUDA file at line 639 [9]. This proved that the GPU infrastructure could handle G2 point arithmetic at the kernel level.

Phase 4: The Missing Template Instantiations

Despite the promising infrastructure, the assistant discovered a critical gap. The main compilation unit groth16_cuda.cu defined #define SPPARK_DONT_INSTANTIATE_TEMPLATES at line 29, which suppressed automatic template instantiation of the GPU Pippenger kernels (accumulate, integrate, breakdown) from pippenger.cuh [14]. The assistant searched for explicit instantiations of these kernels for bucket_fp2_t and found none [4][5].

The grep for accumulate&lt;bucket_fp2|integrate&lt;bucket_fp2|breakdown.*fp2 returned "No files found" [4]. This meant that while the template infrastructure was degree-aware and type-parameterized, the actual GPU machine code for G2 MSM was simply not compiled into the groth16 prover binary. The batch_addition&lt;bucket_fp2_t&gt; kernel was instantiated (proving that G2 GPU code could be compiled), but the full Pippenger pipeline was missing.

This finding transformed the feasibility question from "can we do GPU G2 MSM?" to "how much work is it to add the missing template instantiations?" The answer, as the assistant would later estimate, was approximately 20 lines of explicit template instantiations plus wiring — about 2-5 days of engineering effort [20].

Phase 5: The Smoking Gun — Proof-of-Concept Exists

The assistant then found definitive evidence that GPU G2 MSM was not just theoretically possible but had been implemented, tested, and benchmarked in the same codebase [19]. The sppark/poc/msm-cuda/ directory contained:

  1. A correctness test in tests/msm.rs with a function msm_fp2_correctness() that compared GPU G2 MSM results against arkworks (a Rust elliptic curve library).
  2. A benchmark in benches/msm.rs with a criterion_benchmark_fp2() function that measured GPU G2 MSM performance at scale.
  3. A CUDA kernel in cuda/pippenger_inf.cu that called mult_pippenger&lt;bucket_fp2_t&gt;() — the same template instantiation needed in the Groth16 prover. The proof-of-concept called the same mult_pippenger&lt;bucket_fp2_t&gt; GPU template from pippenger.cuh that the Groth16 prover would need [19]. This was not a theoretical exercise — it was a working implementation that had been verified against a reference library.

Phase 6: The VRAM Analysis

A common objection to GPU G2 MSM is that G2 points are too large to fit in GPU memory. The assistant performed a quantitative VRAM analysis [20]:

The Final Assessment

The assistant's comprehensive feasibility assessment [20] synthesized all the evidence into a nuanced recommendation. The key findings were:

  1. GPU G2 MSM is technically feasible. The template infrastructure is degree-aware, the Fp2 device functions exist, the proof-of-concept has been tested and benchmarked, and the VRAM batching works within 16 GB.
  2. The partitioned pipeline already solves the bottleneck. With num_circuits=1 per GPU call, b_g2_msm takes ~0.4 seconds on CPU — negligible compared to the ~3 seconds of GPU time per partition.
  3. The implementation effort is modest but unnecessary now. Adding the missing template instantiations would take approximately 2-5 days, but the optimization would save at most 0.4 seconds per partition in the primary deployment path.
  4. GPU G2 MSM would be valuable in specific scenarios: the batch-all path (slot_size &gt;= num_partitions), multi-sector batching, or any configuration where CPU contention for b_g2_msm becomes a bottleneck again. The recommendation was to defer GPU G2 MSM implementation until one of those scenarios arises, and to focus engineering effort on the remaining structural bottlenecks identified in the broader investigation.

Broader Significance: Optimization in Context

This investigation is a case study in the importance of understanding architectural context before pursuing optimizations [20]. A less thorough analysis might have concluded "GPU G2 MSM is feasible, let's implement it" and spent 2-5 days on an optimization that saved at most 0.4 seconds per partition in the primary deployment scenario. Instead, the assistant traced the full system — from CUDA kernel templates through Rust pipeline orchestration — and discovered that the problem had already been solved by an architectural change elsewhere.

The investigation also reveals several methodological lessons:

Trace the type system. The assistant's discovery that bucket_t::degree controlled GPU kernel thread layout was the key to understanding that the template infrastructure was already designed for extension field arithmetic [7][13]. Without tracing this type parameter, the feasibility assessment would have remained speculative.

Verify compilation, not just source code. The finding that SPPARK_DONT_INSTANTIATE_TEMPLATES suppressed the GPU kernel instantiations [14] was a reminder that what exists in source code is not always what gets compiled into the binary. The assistant checked both the template definitions and the actual instantiations.

Consider the alternative before committing to the primary approach. The assistant's decision to investigate the partitioned pipeline path in parallel with the GPU G2 MSM feasibility [5][21] was a sound engineering judgment. If the problem was already solved by a configuration change, a complex GPU implementation would be wasted effort.

Quantify before optimizing. The VRAM analysis [20] and the proportionality argument (0.4s CPU vs 3s GPU per partition) provided the quantitative basis for the deferral recommendation. Without these numbers, the recommendation would have been subjective.

Conclusion

The b_g2_msm investigation is a testament to the value of systematic, evidence-driven codebase analysis. What began as a question about moving a computation from CPU to GPU evolved into a deep exploration of CUDA template metaprogramming, Rust pipeline orchestration, GPU memory architecture, and the subtle interactions between parallelization strategies and performance characteristics. The ultimate finding — that the partitioned pipeline architecture already solved the bottleneck — was not obvious at the outset. It emerged through careful reading of source code, tracing of type parameters, verification of compilation flags, and quantitative analysis of memory budgets.

The investigation produced four key documents: a background reference mapping the full call chain with nine structural bottlenecks, and three composable optimization proposals (Sequential Partition Synthesis, Persistent Prover Daemon, and Cross-Sector Batching) [analyzer summary]. The b_g2_msm question, while resolved as a deferral rather than an implementation, contributed crucial understanding of the GPU template infrastructure and the partitioned pipeline's capabilities.

In the end, the most valuable insight was not about GPU G2 MSM at all. It was about the importance of understanding the full system before optimizing any single component. The partitioned pipeline, designed to reduce peak memory, inadvertently eliminated a CPU bottleneck that would have consumed weeks of engineering effort if pursued in isolation. This is the kind of emergent property that only reveals itself through thorough, system-level investigation — and it is why the assistant's methodical approach, spanning over twenty messages of code reading, type tracing, and quantitative analysis, was not just thorough but essential.## References

[1] "The b_g2_msm Question: When a Single GPU Bottleneck Reveals a Deeper Architectural Insight" — The user's initial request laying out the six-point investigation plan.

[2] "The First Cut: A Systematic Investigation Begins into GPU G2 MSM Feasibility" — Initial reconnaissance reading groth16_cuda.cu and searching for Fp2/G2 patterns.

[3] "The Data-Gathering Pivot: How a Negative Search Result Reshaped a GPU Feasibility Investigation" — Discovery that msm_cuda doesn't exist; the GPU MSM infrastructure uses msm_t&lt;bucket_t&gt; templates.

[4] "The Missing Template Instantiation: A Negative Finding That Reshaped a GPU Optimization Strategy" — Confirmation that accumulate&lt;bucket_fp2_t&gt;, integrate&lt;bucket_fp2_t&gt;, and breakdown for fp2 are not instantiated.

[5] "The Missing G2 GPU Kernels: A Critical Finding in the Supraseal-C2 Pipeline Investigation" — Confirmation that GPU kernels are only instantiated for G1, never for G2.

[6] "The Critical Turn: Investigating GPU G2 MSM Feasibility in the Supraseal-C2 Pipeline" — Reading the msm_t template and partitioned pipeline code.

[7] "The Critical Degree: How a Single Property Determined GPU Feasibility for G2 MSM" — Discovery of the degree property controlling GPU thread layout.

[8] "The Final Detail: Verifying num_circuits=1 in the Partitioned Pipeline" — Reading the SynthesizedProof struct to confirm single-circuit dispatch.

[9] "Tracing the G2 Frontier: How One Message Revealed the GPU Infrastructure for Fp2 MSM in Supraseal-C2" — Confirmation that batch_addition&lt;bucket_fp2_t&gt; is instantiated on GPU.

[10] "The Final Verification: Confirming num_circuits=1 in the Partitioned Pipeline" — Verification of the partitioned path's single-circuit design.

[11] "The Investigative Pivot: How a Single Message Unlocks the b_g2_msm Feasibility Assessment" — The pivot from reconnaissance to targeted structural understanding.

[12] "The Critical Template Check: Assessing GPU G2 MSM Feasibility in Supraseal-C2" — Examination of the msm_t template parameters and degree-awareness.

[13] "Tracing the Degree: How One File Read Unlocked the GPU G2 MSM Feasibility Question" — Reading xyzz_t.hpp to confirm the type hierarchy for degree.

[14] "The Template Instantiation Wall: How a Single Preprocessor Flag Blocked GPU G2 MSM" — Discovery of SPPARK_DONT_INSTANTIATE_TEMPLATES suppressing kernel instantiations.

[15] "Reading the Pipeline: A Subagent's Investigation into Moving b_g2_msm to GPU" — Reading the gpu_prove function and SynthesizedProof creation path.

[16] "Reading the CPU-Side Pippenger: A Methodological Step in Assessing GPU G2 MSM Feasibility" — Reading pippenger.hpp to understand the CPU-side implementation.

[17] "Tracing the num_circuits Thread: How a Single Grep Confirmed the Foundation of Phase 7" — Grep confirming num_circuits flows through the pipeline.

[18] "Tracing the CPU-Side Pippenger: A Methodical Investigation of the b_g2_msm in Supraseal-C2" — Tracing the mult_pippenger function used for b_g2.

[19] "The Smoking Gun: How a Single Read Confirmed GPU G2 MSM Feasibility in the Supraseal Pipeline" — Discovery of the working GPU G2 MSM proof-of-concept in poc/msm-cuda.

[20] "A Deep Analysis of the b_g2_msm GPU Feasibility Assessment" — The comprehensive feasibility assessment synthesizing all findings.

[21] "The Critical Turn: Investigating Whether the Partitioned Pipeline Already Solves the b_g2_msm Problem" — Reading the partitioned pipeline dispatch logic and architecture documentation.