The Hidden Trap of nvcc's Dual-Pass Compilation: Fixing Symbol Visibility in Phase 9's PCIe Transfer Optimization

A Single Edit Confirmation, a World of Context

The subject message (msg id=2403) is deceptively simple:

[assistant] [edit] /home/theuser/curio/extern/supraseal-c2/cuda/groth16_cuda.cu Edit applied successfully.

On its surface, this is nothing more than a tool output confirmation — the system reporting that an edit to a CUDA source file was applied without error. Yet this message sits at the culmination of a multi-step debugging and fix sequence that exposed a subtle but critical assumption about how nvcc, the NVIDIA CUDA compiler, handles symbol visibility across its dual-pass compilation model. Understanding why this particular edit was necessary requires tracing backward through the assistant's reasoning, the Phase 9 PCIe Transfer Optimization being implemented, and the peculiarities of CUDA compilation that nearly derailed the entire build.

The Phase 9 Pipeline: What Was Being Built

The assistant was deep in the implementation of Phase 9 of the cuzk SNARK proving engine optimization campaign — a multi-phase effort to reduce GPU idle gaps and improve throughput for Filecoin PoRep Groth16 proof generation. Phase 9 specifically targeted two root causes of GPU utilization dips identified in the Phase 8 baseline: non-pinned host memory transfers (Tier 1) and Pippenger MSM sync stalls (Tier 3).

The Tier 1 change was the more architecturally significant of the two. It moved approximately 6 GiB of a/b/c polynomial uploads out of the GPU mutex by pinning host memory with cudaHostRegister, allocating device buffers, and issuing asynchronous cudaMemcpyAsync transfers on a dedicated stream with CUDA event-based synchronization. This allowed the data upload to overlap with another worker's GPU computation, rather than blocking inside the critical section.

To implement this, the assistant had already made several edits across two files:

1. groth16_ntt_h.cu: Added execute_ntts_prestaged() and execute_ntt_msm_h_prestaged() — variants of existing functions that skip the HtoD (host-to-device) upload and instead wait on a CUDA event signaling that data has already been transferred.
2. groth16_cuda.cu: Added the pre-staging logic before the GPU mutex acquisition — computing sizes, registering host memory, allocating device buffers, creating an upload stream and events, and issuing async transfers. Inside the per-GPU thread, the code now calls the prestaged variants instead of the original functions.

The Compilation Error: When Assumptions Meet nvcc

After implementing these changes, the assistant attempted to build the project with cargo build --release -p cuzk-daemon. The build failed with errors that ntt_msm_h::lg2 could not be resolved. This was puzzling because lg2 was a private static method of the ntt_msm_h class defined in groth16_ntt_h.cu, which was #included by groth16_cuda.cu. The assistant had already made lg2 public (in msg id=2381) to allow access from the outer file. So why would nvcc fail to find it?

The answer lies in nvcc's dual-pass compilation model. When nvcc compiles a .cu file, it performs two passes:

• Host pass: Compiles code for the CPU, where __CUDA_ARCH__ is not defined.
• Device pass: Compiles code for the GPU, where __CUDA_ARCH__ is defined. The critical detail is that nvcc parses all code in both passes, regardless of whether a particular function will generate device code. The #ifndef __CUDA_ARCH__ preprocessor guard that wrapped the ntt_msm_h class and the gib constant made them invisible during the device pass. When the assistant's new code in groth16_cuda.cu referenced ntt_msm_h::lg2() and gib, nvcc's device pass could not find these symbols because the entire class definition was hidden behind the preprocessor guard. This is a classic CUDA pitfall. The extern "C" function generate_groth16_proofs_c is a host-only function — it will never generate device code. But nvcc still parses it during the device pass, and any symbol references must be resolvable in both passes. The #ifndef __CUDA_ARCH__ guard is commonly used in CUDA libraries to exclude host-only code (like memory management, file I/O, and complex class hierarchies) from the device compilation pass, but this creates a visibility trap for any code that references those excluded symbols.

The Investigation: Tracing the Original lg2

The assistant's debugging process (visible in msg id=2398–2401) is a textbook example of systematic investigation. Rather than simply guessing at the fix, the assistant:

1. Hypothesized about why the original code compiled: "The only explanation is that ntt_msm_h::lg2 was accessible because... the code was inside a translation unit that included groth16_ntt_h.cu"
2. Tested the hypothesis: Realized that even with inclusion, the #ifndef __CUDA_ARCH__ guard would hide the class during the device pass.
3. Searched for alternatives: Used grep to find all lg2 definitions in the codebase, discovering a free-standing static __device__ __host__ constexpr uint32_t lg2(T n) in ntt/kernels.cu (msg id=2400).
4. Connected the dots: Recognized that the original code at line 607 (size_t d_a_sz = sizeof(fr_t) << (lg2(points_h.size() - 1) + 1);) was calling this free-standing lg2, not the class method. The bare unqualified call resolved to the function from ntt/kernels.cu because it was included through the NTT header chain.
5. Formulated the fix: Replace ntt_msm_h::lg2() with the bare lg2() call, and handle gib references properly.

The Edit Sequence: Three Fixes in Rapid Succession

The fix was applied across three consecutive edits (msg id=2401, 2402, 2403), each addressing a different aspect of the symbol visibility problem:

• Edit 1 (msg id=2401): Changed ntt_msm_h::lg2() to lg2() in the pre-staging code within groth16_cuda.cu. The bare lg2() resolves to the free-standing function in ntt/kernels.cu, which is declared static __device__ __host__ and is therefore visible in both nvcc passes.
• Edit 2 (msg id=2402): Fixed references to gib, the file-scope constant that was also hidden behind #ifndef __CUDA_ARCH__. Since gib is defined in groth16_ntt_h.cu inside the guard, it too was invisible during the device pass. The fix likely involved either computing the value inline or using a different constant.
• Edit 3 (msg id=2403 — the subject): Applied a further correction, possibly fixing additional gib references or adjusting related size calculations that depended on the now-inaccessible constant. The subject message (msg id=2403) is the confirmation of this third and final edit. It represents the last piece of the fix puzzle — the moment when all symbol visibility issues were resolved and the code could be expected to compile cleanly through both nvcc passes.

The Deeper Lesson: CUDA's Dual-Pass Model as a Source of Bugs

This debugging episode illuminates a fundamental challenge of CUDA programming that even experienced developers can stumble over. The nvcc compiler's dual-pass model means that every line of code in a .cu file must be valid in both the host and device compilation contexts, unless it is explicitly guarded with #ifndef __CUDA_ARCH__ or #ifdef __CUDA_ARCH__.

The assistant's initial assumption was reasonable: since groth16_cuda.cu #includes groth16_ntt_h.cu, and the ntt_msm_h class is defined in that file, the class members should be accessible. What the assistant overlooked was the preprocessor guard wrapping the entire class definition. This is an easy mistake to make — the guard is invisible when reading the code linearly, and the #include directive creates an expectation of straightforward symbol availability.

The fix itself — switching from a class-qualified method call to a free-standing function — is elegant. The free-standing lg2 in ntt/kernels.cu was already designed for cross-pass visibility (declared static __device__ __host__), making it the correct choice for code that must compile in both passes. The original codebase had already established this pattern; the assistant's new code simply needed to follow it.

Output Knowledge and Broader Implications

This message, despite its brevity, creates several forms of output knowledge:

1. A working build: The immediate output is a compilable groth16_cuda.cu that will link correctly into the cuzk-daemon binary, enabling the Phase 9 benchmarks to proceed.
2. A documented pattern: The fix reinforces the pattern that any code in .cu files that must compile in both nvcc passes should use free-standing __device__ __host__ functions rather than class methods hidden behind preprocessor guards.
3. A debugging methodology: The assistant's systematic approach — hypothesizing, searching, verifying, and applying targeted fixes — provides a template for diagnosing similar CUDA compilation issues.
4. A cautionary tale: The episode serves as a reminder that #include does not guarantee symbol visibility in CUDA, and that preprocessor guards interact with nvcc's dual-pass model in ways that can surprise even experienced developers.

Conclusion

The subject message (msg id=2403) — a simple "Edit applied successfully" — is the final stroke in a three-edit fix sequence that resolved a subtle CUDA compilation error. The error arose from a mismatch between the assistant's assumption of straightforward symbol visibility through #include and the reality of nvcc's dual-pass compilation model, where preprocessor guards like #ifndef __CUDA_ARCH__ can hide entire class definitions from the device compilation pass. The fix — switching from ntt_msm_h::lg2() to the free-standing lg2() from ntt/kernels.cu — aligned the new Phase 9 code with the existing codebase patterns and restored buildability. In the broader context of the optimization campaign, this fix was a necessary prerequisite for benchmarking Phase 9's PCIe transfer optimizations, which would ultimately yield a 14.2% throughput improvement in single-worker mode.