The NTT Kernel Deep Dive: From Architectural Mapping to Micro-Optimization in Filecoin's Groth16 Pipeline

Introduction

In the landscape of zero-knowledge proof systems, few operations are as computationally demanding as the Number Theoretic Transform (NTT). As the finite-field analogue of the Fast Fourier Transform, the NTT underpins polynomial multiplication in Groth16 proof generation—and in Filecoin's Proof-of-Replication (PoRep) pipeline, it accounts for a substantial fraction of a ~200 GiB peak memory footprint. This article synthesizes a deep-dive investigation into the NTT kernel internals of the SUPRASEAL_C2 proof generation pipeline, a systematic effort that moved from high-level architectural mapping through exhaustive source code reading to the identification of 18 specific micro-optimization opportunities.

The work documented across ten articles [1]–[10] represents a comprehensive methodology for GPU kernel analysis. Beginning with a meticulously structured research prompt [1], the investigation proceeded through file discovery, kernel source reading, field arithmetic analysis, and stream infrastructure tracing, culminating in a detailed optimization roadmap. This article traces that journey, examining both the technical findings and the investigative methodology that produced them.

The Investigation Framework: A Blueprint for Kernel Analysis

The investigation was launched by a user request of extraordinary precision [1]. The prompt specified six dimensions of inquiry—kernel architecture, memory access patterns, CUDA stream overlap, pointwise multiplication fusion, the lot_of_memory adaptive flag, and a sweeping survey of optimization opportunities—and demanded "detailed findings with specific line numbers." This was not a casual question; it reflected deep familiarity with both GPU architecture and the Filecoin proof pipeline.

The structure of the prompt itself embodied a sophisticated mental model of GPU optimization. The six dimensions were ordered hierarchically: architecture first because it constrains everything else, memory access patterns second because they dominate GPU performance, stream overlap third because it addresses system-level orchestration, and so on [1]. This decomposition transformed a sprawling investigation into tractable sub-questions, each targeting a specific performance dimension.

The assistant's response was equally methodical. Rather than diving immediately into kernel code, the first step was to map the terrain [3]. Three glob commands discovered the file landscape: parameter files for six elliptic curves (vesta, pallas, goldilocks, bls12_381, bls12_377, baby_bear), the main NTT dispatcher (ntt.cuh), kernel implementation files (gs_mixed_radix_wide.cu, ct_mixed_radix_wide.cu and their narrow variants), the Groth16 H-polynomial computation (groth16_ntt_h.cu), field arithmetic headers (mont_t.cuh, shfl.cuh, bls12-381.hpp), and utility headers (gpu_t.cuh). This discovery phase revealed that the NTT implementation was not a monolithic file but a distributed system of templates, parameters, and launchers spread across multiple directories [3].

The Systematic Reading Strategy

The reading order reflected a deliberate top-down investigative strategy [4]. The assistant read the dispatcher first (ntt.cuh), which provided the top-level logic: how many kernel launches occur, how steps are divided across stages, and which algorithm (Gentleman-Sande vs. Cooley-Tukey) is used for which direction. Next came the kernel implementation files (gs_mixed_radix_wide.cu and ct_mixed_radix_wide.cu), which contained the actual GPU compute logic: thread block sizes, shared memory usage, and memory access patterns.

The fourth file, groth16_ntt_h.cu, was particularly significant because it lives outside the sppark library proper—it sits in the supra_seal/c2/cuda/ directory, indicating a higher-level integration file specific to the Groth16 proving system [4]. This file bridges the gap between the generic NTT library and the specific needs of the C2 proof generation pipeline, containing the stream overlap logic and the lot_of_memory flag.

After reading the kernel code, the assistant recognized a critical gap: the kernels were written in terms of types and abstractions that had not yet been understood. What is fr_t? How large is it? How does shfl_bfly work? What is the gpu_t object that provides streams? This recognition drove the next phase of the investigation [9].

Descending into the Foundations: Field Arithmetic and Warp Primitives

The pivot to foundational analysis is captured in a single message: "Now let me look at the fr_t type, shared memory primitives, and the stream/gpu utilities to understand sizes" [9]. This seemingly simple statement represented a sophisticated methodological decision—the recognition that kernel-level analysis could not be completed without understanding the type system and arithmetic primitives upon which the kernels were built.

Reading bls12-381.hpp revealed that fr_t (the BLS12-381 scalar field element) is 32 bytes, represented in Montgomery form as 8 × uint32_t [6]. This size directly determines shared memory consumption per thread block, register pressure, and global memory bandwidth requirements. For a domain of size 2^24 (16 million elements), a single NTT pass reads and writes 768 MB of data.

The mont_t.cuh file revealed the Montgomery multiplication implementation—the core arithmetic operation of the NTT butterfly [6]. Understanding its instruction count and register usage was essential for estimating kernel occupancy and identifying compute-bound bottlenecks. The shfl.cuh file provided the warp-level shuffle primitives, and the grep for shfl_bfly across the codebase revealed 8 matches across mont_t.cuh, mont_t.hip, mont32_t.cuh, and mont32_t.hip, confirming that the shuffle butterfly was implemented for both the 384-bit (6-limb) and 256-bit (4-limb) field element representations, with separate implementations for CUDA and HIP backends [5][8].

The shfl_bfly function at line 1107 of mont_t.cuh was the critical primitive [8]. It wraps CUDA's __shfl_xor_sync intrinsic with Montgomery multiplication, implementing the NTT butterfly at the warp level. This is significantly faster than shared memory for warp-local data because it avoids the latency and bank conflicts of shared memory. The function's existence confirmed that the NTT implementation was making use of warp-level cooperation—a key optimization insight.

The Stream Infrastructure and the Flip-Flop Logic

The investigation of CUDA stream overlap required understanding the gpu_t abstraction defined in gpu_t.cuh [5]. The assistant traced how streams were created, managed, and synchronized throughout the pipeline. The groth16_ntt_h.cu file revealed a three-stream pipeline where gpu[0] handles the a array, gpu[1] handles b, and gpu[2] (when lot_of_memory is true) handles c [7]. Understanding that these were fixed stream assignments—not a flip-flop that alternates across calls—was essential for the stream overlap analysis.

The final verification before synthesis came in message 8, where the assistant paused to check the multi-stream flip-flop logic one more time [7]. This moment exemplifies a crucial principle in technical analysis: the moment of synthesis is not the end of data gathering but the point at which one knows what one still needs to know. By verifying this detail, the assistant ensured that the final analysis would be comprehensive.

Key Technical Findings

The comprehensive analysis produced in message 9 [10] documented the NTT kernel architecture with remarkable precision. The wide kernels use __launch_bounds__(768, 1), meaning 768 threads per block with a maximum of one block per multiprocessor—an unusual configuration that prioritizes per-thread register availability over occupancy. The block size varies with radix according to the formula block_size = 1 << (radix - 1), yielding sizes of 32, 64, 128, 256, and 512 for radices 6 through 10.

Shared memory usage was calculated as sizeof(fr_t) * block_size = 32 * block_size, ranging from 1 KB (radix 6) to 16 KB (radix 10). The number of kernel launches per NTT varies from 2 to 4 depending on domain size, with each launch carrying ~5-10 μs of overhead. For a typical 2^23 NTT, approximately 34 kernel launches occur, totaling 170-340 μs of launch overhead—only 1-2% of total time, confirming that launch overhead is not the primary bottleneck.

The register pressure analysis was particularly revealing. The assistant estimated 128-192 registers per thread for the NTT kernel, derived from the Montgomery multiplication implementation (approximately 46 registers per multiply, with two active multiplies plus twiddle factors). At radix-10 with 512 threads per block and 192 registers per thread, the total register demand is 98,304—far exceeding the A100's 65,536 registers per SM. This forces at most one block per SM at radix-10, confirming that register pressure is a significant occupancy constraint [10].

The shared memory bank conflict analysis identified an 8-way conflict pattern. Because fr_t is 32 bytes (8 × 4-byte words) and NVIDIA shared memory has 32 banks of 4 bytes each, consecutive threads accessing consecutive fr_t elements create a pattern where thread 0 accesses banks 0-7, thread 1 accesses banks 8-15, thread 2 accesses banks 16-23, thread 3 accesses banks 24-31, and thread 4 wraps to banks 0-7—creating an 8-way bank conflict [10].

The coalescing analysis revealed that the stride pattern in the GS kernel depends on stage - iterations. When this value is zero (final stage), consecutive threads access consecutive pairs (good coalescing). When it is positive (intermediate stages), there is a stride of 2^(stage-iterations) between adjacent thread accesses (strided, non-coalesced access). The narrow kernels already implement a coalesced load/store pattern that the wide kernels lack—a critical optimization opportunity [10].

The Optimization Roadmap

The investigation culminated in a prioritized optimization table with estimated impact and difficulty [10]:

| Priority | Optimization | Estimated Impact | Difficulty | |---|---|---|---| | HIGH | Fuse LDE_powers into NTT first/last step | 10-15% bandwidth reduction | Medium | | HIGH | Port coalesced load/store from narrow to wide kernels | 15-25% improvement on intermediate stages | Medium | | HIGH | Shared memory bank conflict mitigation (SoA layout) | 10-20% improvement on stages ≥ 6 | Medium | | MEDIUM | Fuse coeff_wise_mult/sub_mult into NTT last step | ~5% total, save 2 global memory passes | Medium | | MEDIUM | Increase shared memory for larger radix | Reduce kernel launches from 3 to 2 | High | | LOW | Double-buffer HtoD with NTT compute | Marginal, HtoD is small vs compute | Low | | N/A | Tensor cores | Not applicable to modular arithmetic | N/A |

The combined estimated speedup from these compute-level improvements is 30-43% on existing hardware [10]. This directly complements the memory and throughput gains from the three architectural proposals developed in the root session: Sequential Partition Synthesis (reducing peak memory from ~200 GiB to ~64-103 GiB), Persistent Prover Daemon (eliminating ~60s per-proof SRS loading overhead), and Cross-Sector Batching (projecting 2-3× throughput per GPU and 5-6× reduction in $/proof).

The "N/A" entry for tensor cores is particularly instructive [10]. It shows that the assistant considered a popular optimization technique, traced through the arithmetic requirements, and concluded it was infeasible—saving future investigators from pursuing a dead end. This negative knowledge is as valuable as the positive findings.

The Broader Context: Connecting Micro to Macro

This NTT kernel deep dive did not exist in isolation. It was part of a larger investigation that had already mapped the full call chain from Curio's Go task layer through Rust FFI into C++/CUDA kernels, identified nine structural bottlenecks, and produced three architectural optimization proposals [10]. The micro-optimization analysis provided the compute-level foundation needed to realize those architectural changes.

The Sequential Partition Synthesis proposal, for example, would change how partitions are processed—streaming them one-at-a-time through the GPU rather than all ten in parallel. But the NTT kernels themselves would still need to be efficient. The micro-optimizations identified in this investigation ensure that when the architectural changes are made, the kernels are operating at peak efficiency.

Similarly, the Cross-Sector Batching proposal exploits freed memory headroom to batch multiple sectors' circuits into single GPU invocations. The NTT kernel analysis reveals whether the GPU has sufficient resources (registers, shared memory, compute units) to handle the increased workload without degrading per-sector performance.

The investigation also opened a new frontier: can the known structure of the constraints—specifically the dominance of SHA-256 boolean circuits (99% of aux assignment)—be exploited for pre-computation or mathematical batching transpositions? This question moves the optimization frontier from generic compute improvements to domain-specific circuit exploitation, leveraging the deep understanding of the constraint structure already documented.

Conclusion

The NTT kernel deep dive documented across these ten articles represents a masterclass in GPU performance analysis. Beginning with a meticulously structured research prompt, the investigation proceeded through systematic file discovery, top-down source code reading, foundational type analysis, and stream infrastructure tracing, culminating in a prioritized optimization roadmap with quantified impact estimates.

The methodology exemplifies several principles of effective systems investigation: start with architecture before diving into implementation, trace the full dependency chain from application code down to arithmetic primitives, quantify everything in terms of hardware resources, compare alternative implementations to identify optimization opportunities, and prioritize findings by impact and difficulty.

The technical findings—8-way shared memory bank conflicts, strided global memory access in intermediate NTT stages, register pressure limiting occupancy at high radices, and the untapped coalescing pattern in the narrow kernels—provide a concrete roadmap for improving NTT performance by 30-43% on existing hardware. Combined with the architectural proposals from the root session, these micro-optimizations contribute to the overarching vision of a continuous, memory-efficient proving pipeline optimized for heterogeneous cloud rental markets where RAM cost dominates.

For anyone interested in GPU performance optimization, finite field arithmetic, or the engineering of proof systems, this investigation is a rich source of both technical insight and methodological guidance. It shows how expert knowledge, systematic reasoning, and careful attention to detail combine to produce actionable findings that can drive real performance improvements in a production system handling some of the most computationally demanding operations in modern cryptography.