Phase 9 PCIe Transfer Optimization: From 14.2% Breakthrough to the 41-Second Paradox

Introduction

The cuzk SNARK proving engine had been through eight optimization phases by the time it reached a critical plateau. Phase 8 had achieved perfect GPU-boundedness at 37.4 seconds per proof, with GPU scheduling efficiency at 100.0%. But careful analysis of GPU power and utilization dips revealed that the GPU was not actually busy 100% of the time at the Streaming Multiprocessor level — it was stalling on PCIe data transfers. Phase 9 was designed to eliminate those stalls, and it succeeded beyond expectations in single-worker mode, delivering a 14.2% throughput improvement. But when the full dual-worker production benchmark ran, the result was a paradox: the GPU kernel times were dramatically improved, yet the end-to-end throughput had regressed.

This article tells the story of Phase 9 in its entirety — from the comprehensive handoff document that synthesized everything learned across eight prior phases, through the implementation of two ambitious CUDA-level changes, through a grueling debugging session that uncovered a subtle CUDA memory pool asymmetry bug, through the validation of a 14.2% improvement, and finally through the dual-worker benchmark that revealed PCIe bandwidth contention as the next frontier. It is a story about the iterative nature of systems optimization, where each layer of improvement reveals the next constraint, and where the infrastructure around benchmarking is as important as the kernel code being optimized.

The Foundation: Eight Phases of Optimization

To appreciate what Phase 9 accomplished, one must understand the project it served. The Curio node is a Filecoin storage provider implementation that must periodically generate Proofs of Replication (PoRep) to prove it is storing data correctly. These proofs rely on Groth16 zk-SNARKs, which are computationally expensive to generate — requiring hundreds of gigabytes of polynomial data, multi-gigabyte Structured Reference String (SRS) files, and complex sequences of Number Theoretic Transforms (NTTs) and Multi-Scalar Multiplications (MSMs) on GPUs.

The optimization journey began with Phases 0–5, which established the baseline proving pipeline and identified the ~200 GiB peak memory footprint. Phase 6 introduced waterfall timeline instrumentation to measure exactly where GPU time was spent. Phase 7 implemented a per-partition dispatch architecture, where each partition of the proof is processed independently. Phase 8 introduced a dual-GPU-worker interlock system, narrowing a C++ static mutex to cover only the CUDA kernel region, allowing two GPU workers to share a single device more efficiently.

The Phase 8 benchmark results showed the system was perfectly GPU-bound at 37.4 seconds per proof, with GPU efficiency at 100.0% at the scheduling level. But the user had observed GPU power and utilization dips correlating with PCIe traffic bursts, suggesting that the GPU was not actually busy 100% of the time at the SM level — it was stalling on data transfers. This observation became the foundation for Phase 9.

The Handoff Document: A Blueprint for Action

Message [msg 2368] stands as one of the most remarkable artifacts in the entire conversation. It is a 2,000+ word comprehensive status document that the assistant wrote before beginning Phase 9 implementation. The message serves multiple purposes simultaneously: it is a retrospective analysis of everything learned in Phases 0–8, a forward-looking plan for Phase 9, a context management tool for the assistant's own working memory, and a handoff document that any future reader could use to understand the project state.

The message's "Discoveries" section is particularly rich. It includes a detailed PCIe transfer inventory table that accounts for every byte transferred during proof generation — 23.6 GiB total per partition — classified by pinning status and synchronization pattern. This table is the foundation of the optimization strategy:

| Phase | Transfer | Size | Pinned? | Problem | |---|---|---|---|---| | NTT | a/b/c polynomials | 6 GiB | No (Rust Vec) | Staged through 32 MB bounce buffer, half bandwidth | | H MSM | H SRS points (8 batches) | 6 GiB | Yes | Per-batch gpu.sync() — 8 hard stalls | | Batch additions | L/A/B_G1/B_G2 SRS | 7.7 GiB | Yes | OK — already double-buffered | | Tail MSMs | L/A/B bases + scalars | 3.9 GiB | No (std::vector) | Non-pinned + per-batch sync stalls |

From this inventory, the assistant identified two high-impact changes. Change 1 (Tier 1) targeted the 6 GiB non-pinned a/b/c polynomial upload — the largest single transfer with the worst bandwidth penalty. Change 2 (Tier 3) targeted the 8 per-batch sync stalls in the H MSM — the most frequent stall events. The assistant explicitly chose not to fix the tail MSM transfers in this phase, a pragmatic engineering decision to prioritize highest-impact, lowest-risk changes first.

The handoff document also includes a VRAM budget analysis, a detailed description of the Phase 8 architecture, benchmark data, git commit history, a prioritized TODO list, and references to every file that needs modification. It is, in essence, a complete engineering design document produced by an AI assistant as a natural part of the coding workflow.

The Implementation: Two Changes, Two Streams

The implementation of Phase 9 proceeded through a series of carefully orchestrated edits. Change 1 required modifying the GPU mutex region in groth16_cuda.cu to move the polynomial upload out of the critical path. The assistant introduced a pre-staging phase that runs before the mutex acquisition: host memory pages are pinned with cudaHostRegister, device-side buffers (d_a at 4 GiB and d_bc at 8 GiB) are allocated, and cudaMemcpyAsync transfers are issued on a dedicated upload stream with CUDA event-based synchronization. Inside the mutex, the code waits for the upload events to complete before launching the NTT kernel, but the actual transfer latency has been removed from the critical section.

The implementation required adding a new function, execute_ntt_msm_h_prestaged(), that accepts pre-allocated device pointers rather than allocating them internally. This function was carefully stitched into the existing code path, with fallback logic that preserves the original behavior if pre-staging fails.

Change 2 targeted the Pippenger MSM in pippenger.cuh. The existing code issued a hard gpu.sync() after each batch to ensure that device-to-host (DtoH) transfers of bucket results completed before the CPU collect() function read them. The assistant introduced double-buffered host result buffers (res_buf[2] and ones_buf[2]) and deferred the sync() call to the next iteration. This allowed GPU compute for batch N to overlap with the DtoH transfer of batch N-1's results, eliminating the per-batch stall.

Both changes required deep knowledge of CUDA programming models — stream semantics, event synchronization, pinned memory, and the subtle distinction between synchronous and asynchronous memory transfers. The assistant demonstrated this knowledge throughout the implementation, making precise edits to C++ CUDA code that spanned multiple files and functions.

The OOM Crisis: When Optimizations Collide with Reality

The initial dual-worker benchmark with gpu_workers_per_device=2 and concurrency=3 failed catastrophically. The first partition completed in a blistering 1,588 milliseconds — a 57.6% improvement over the Phase 8 baseline of 3,746 ms — but every subsequent partition crashed with out-of-memory (OOM) errors. Even the fallback path, which avoided pre-staging entirely, also failed. The GPU was left with 10 GiB of leaked memory after the failed runs.

Message [msg 2440] captures the assistant's response to this crisis. It is a masterclass in systems-level debugging, proceeding through a series of increasingly refined hypotheses:

Hypothesis 1: The pre-staging allocation is too large. The assistant computed exact memory sizes: d_a at 4 GiB, d_bc at 8 GiB, totaling 12 GiB on a 16 GiB GPU. But tracing the lg2 computation confirmed the domain size was correct — 2^27 = 134,217,728 elements for the PoRep circuit. The allocation size was expected.

Hypothesis 2: The cleanup order is wrong. The assistant traced the C++ code and discovered that d_bc was freed after the mutex release. When Worker 1 acquired the mutex, Worker 0's 8 GiB d_bc was still allocated. Worker 1 then tried to allocate its own 12 GiB, totaling 20 GiB on a 16 GiB GPU. The fix was to free d_bc before releasing the mutex, and even earlier — inside the per-GPU worker thread, right after the NTT phase completed.

Hypothesis 3: Host memory registration conflicts between workers. Both workers, processing different partitions of the same proof, attempted to pin the same host memory pages with cudaHostRegister. The second call returned cudaErrorHostMemoryAlreadyRegistered (error 712), causing Worker 1 to skip pre-staging and fall back to the original path. But the fallback path also failed, indicating a deeper issue.

Hypothesis 4: CUDA memory pool asymmetry. This was the breakthrough. The gpu.Dmalloc and gpu.Dfree functions in the supraseal-c2 codebase use cudaMallocAsync and cudaFreeAsync, which operate on CUDA memory pools. These pools cache freed memory and do not release it back to the synchronous cudaMalloc pool used by dev_ptr_t and the pre-staging allocator. When the first partition's MSM freed its ~385 MiB of buckets via cudaFreeAsync, that memory went into the async pool — invisible to the next partition's cudaMalloc call. The second partition saw a GPU that appeared to have ~14 GiB free (as reported by nvidia-smi) but the synchronous allocator could only see what remained in its own pool, which might be far less.

This diagnosis, detailed in [msg 2447], is a remarkable piece of reasoning. The assistant had to pierce the abstraction layer of CUDA's memory management API, understanding that cudaMallocAsync and cudaMalloc draw from separate pools, and that memory freed via the async path is not returned to the synchronous pool. This is not a bug in the application code — it is a fundamental property of the hardware abstraction layer that the code was unknowingly relying on.

The Memory-Aware Allocator: A Pivot in Design

The user's intervention at [msg 2445] was concise but crucial: "Immediate OOM - perhaps we want to memory-manage the early copies such that they only copy as much memory as we can actually allocate? With some 512MiB buffer left spare for other host processes." This suggestion reframed the problem. Instead of trying to force the full pre-staging allocation to work, the solution was to make the pre-staging adaptive.

Message [msg 2449] captures the assistant's redesign of the pre-staging allocation. The message lists four key insights:

  1. Use cudaMemGetInfo to check actual free VRAM — query the CUDA runtime for available memory at the moment of allocation, rather than assuming a fixed amount.
  2. Calculate precise memory requirements — d_a requires domain_size * 32 bytes (4 GiB for a 2^27 domain), d_bc requires domain_size * (1 + lot_of_memory) * 32 bytes (8 GiB), plus a 512 MiB safety margin.
  3. Implement a three-tier fallback strategy — if both d_a and d_bc fit, pre-stage both (optimal); if only d_a fits, pre-stage just d_a (partial optimization); if neither fits, fall back entirely to the original non-pre-staged path.
  4. Release the CUDA memory pool cache before checking — call cudaDeviceSynchronize() and trim the memory pool before the pre-staging check, ensuring the free memory query reflects what is actually available to synchronous cudaMalloc calls. The assistant also moved the pre-staging allocation inside the GPU mutex to serialize it across workers, preventing both workers from trying to allocate 12 GiB simultaneously. And it freed d_bc immediately after the NTT phase (before the mutex release) to minimize peak memory pressure.

The Validation: 14.2% Throughput Improvement

With the memory-aware allocator in place, the assistant ran the single-worker benchmark. Message [msg 2457] captures the moment of validation:

All 3 proofs COMPLETED! And the prove time is 32.1s average vs the Phase 8 baseline of 37.4s — that's a 14.2% improvement!

The detailed GPU timing logs confirmed the kernel-level gains:

| Metric | Phase 8 | Phase 9 | Delta | |---|---|---|---| | ntt_msm_h_ms | ~2430 ms | ~690 ms | -71.6% | | batch_add_ms | ~640 ms | ~650 ms | ~same | | tail_msm_ms | ~125 ms | ~82 ms | -34.4% | | gpu_total_ms | ~3746 ms | ~1450–1900 ms | -50–61% |

The NTT+MSM phase — the dominant GPU kernel — had been cut by over 70%. The tail MSM showed a 34% improvement from the deferred-sync double-buffering change. The batch_add metric was unchanged, as expected, since that phase was not a target of this optimization. Overall GPU time per partition was slashed by half to 61%, depending on the partition.

The memory-aware allocator logged its operation: prestage_vram free_mib=13884 usable_mib=13372 — detecting 13.9 GiB free, reserving a 512 MiB safety margin, and reporting 13.4 GiB usable. The prestage_setup=ok status confirmed no fallback was needed.

The Daemon That Wouldn't Die

With single-worker validation complete, the assistant needed to run the full production benchmark with gpu_workers_per_device=2 and concurrency=3 — the configuration that would actually be deployed. But transitioning from single-worker to dual-worker benchmarking turned into a multi-layered saga of daemon lifecycle management.

The assistant's first attempt to restart the daemon failed silently. When the assistant examined the log in message [msg 2459], the log showed a gpu_worker completing partition 1 with gpu_ms=2888 — a timing characteristic of the single-worker configuration. The daemon was still running the old config. The restart had not actually killed the process.

Message [msg 2460] captures the moment of recognition: "Wait — the daemon was already processing from the gw=1 run. Let me kill it properly and restart." The assistant issued a corrected restart sequence with three critical improvements: a pgrep cuzk-daemon verification step after the pkill -9 to confirm the process was actually terminated; an increased sleep duration from 2 to 3 seconds, acknowledging that GPU state cleanup might require more time; and explicit deletion of the old log file with rm -f before starting the new daemon, ensuring no stale data could contaminate the readiness check.

But even the corrected restart sequence was not enough. In message [msg 2461], the assistant ran grep "ready" /tmp/cuzk-phase9-daemon.log && echo "READY" || echo "NOT READY" and received a positive result — a timestamp showing the daemon was ready. But the match was from the previous daemon instance. The log file had not been properly cleared, and the new daemon had failed to start because port 9820 was still bound by the lingering old process.

This is a classic systems debugging pitfall: a file-based readiness check is only meaningful if you know the file's provenance. The assistant had designed a check that confirmed "a daemon was ready at some point" but interpreted it as "the daemon is ready now." The gap between these two interpretations could have invalidated the entire benchmark run.

The assistant's diagnosis in message [msg 2462] was precise: "That's the old log. The daemon didn't start because the port is still in use." The corrective command — pkill -9 -f cuzk-daemon; sleep 5; ss -tlnp | grep 9820 — added a port verification step using ss to explicitly check whether anything was listening on port 9820 before proceeding. This was the missing piece: rather than inferring daemon health from log output, the assistant now verified the underlying condition directly.

Message [msg 2463] represents the successful culmination of the daemon lifecycle saga. The log was explicitly deleted before starting. The grep pattern included "error" and "panic" alongside "ready" — a defensive measure informed by the earlier OOM panics that had corrupted CUDA contexts. The daemon responded with PID 748530 and the message "cuzk-daemon ready, serving on 0.0.0.0:9820." Every element of the bash command — the log deletion, the config selection, the sleep duration, the grep pattern — encoded a lesson learned from a prior failure.

The 41-Second Paradox

With the daemon running cleanly, the assistant launched the production benchmark in message [msg 2464]:

/home/theuser/curio/extern/cuzk/target/release/cuzk-bench batch -t porep \
  --c1 /data/32gbench/c1.json -c 5 -j 3 2>&1 | tee /tmp/cuzk-phase9-gw2-bench.log

Five proofs at concurrency 3, with gpu_workers_per_device=2 — the exact configuration intended for production deployment. The results were deeply counterintuitive:

[1/5] COMPLETED — 60.0s (prove=47245 ms, queue=258 ms)
[2/5] COMPLETED — 103.7s (prove=60315 ms, queue=708 ms)
[3/5] COMPLETED — 133.5s (prove=50167 ms, queue=1155 ms)
[4/5] COMPLETED — 113.0s (prove=60350 ms, queue=250 ms)
[5/5] COMPLETED — 101.2s (prove=50468 ms, ...)

The system throughput worked out to 41.0 seconds per proof. This was worse than the Phase 8 baseline of 37.4 s/proof, and significantly worse than the Phase 9 single-worker result of 32.1 s/proof.

This was a stunning reversal. The kernel-level optimizations were working beautifully — the GPU timing logs showed gpu_total_ms around 1,616 ms per partition, ntt_msm_h_ms between 588–855 ms, and tail_msm_ms around 90 ms. The individual GPU operations were dramatically faster. Yet the end-to-end throughput had regressed.

The Diagnosis: PCIe Bandwidth Contention

The assistant's response in message [msg 2465] was immediate and methodical: "All 5 proofs completed. But the throughput is 41.0s/proof with gw=2. Let me check the detailed timing to understand what's happening."

The grep command extracted the key metrics:

The Silence That Speaks Volumes

At this critical juncture, message [msg 2466] arrives — a user message containing only empty <conversation_data></conversation_data> tags. No instructions, no questions, no redirection, no praise, no criticism. Just silence.

This is not a passive or accidental silence. It is an active, deliberate choice to stay out of the way. The user recognizes that the assistant is already doing exactly what needs to be done: investigating the regression systematically. The assistant has already formulated the right question ("why is gw=2 slower than gw=1?") and has already taken the first investigative step (extracting timing logs). The user's intervention would only slow things down.

This empty message represents a remarkable level of trust between user and assistant — trust in competence, in judgment, in methodology, and in prioritization. It is the kind of trust that only develops after many rounds of successful collaboration. The user has seen the assistant navigate OOM errors, diagnose CUDA memory pool issues, implement memory-aware allocation, and deliver consistent improvements. At this point, the user knows that the best thing they can do is get out of the way.

The Broader Significance

The 41.0 s/proof result, while disappointing relative to the 32.1 s single-worker result, is actually a success in diagnostic terms. It reveals that the Phase 9 kernel optimizations are working correctly (the GPU times are excellent) and that the next bottleneck is architectural rather than algorithmic. The path forward might involve PCIe bandwidth management (e.g., staggering transfers, using multiple PCIe lanes, or CPU-side compression), daemon scheduling improvements, or a fundamental rethinking of how GPU workers share the PCIe bus.

This pattern — each optimization layer revealing the next constraint — is characteristic of systems engineering at the highest level. Phase 6 revealed GPU idle gaps. Phase 7 revealed mutex contention. Phase 8 revealed PCIe transfer stalls. Phase 9 eliminated those stalls at the kernel level, only to discover that PCIe bandwidth contention under concurrency is the next frontier. The bottleneck has shifted from GPU compute to data movement, and the optimization strategy must evolve accordingly.

The daemon lifecycle saga that preceded the benchmark — the stale process, the deceptive grep, the port verification — is equally instructive. It demonstrates that the infrastructure around the benchmark is as important as the kernel code being optimized. A failed daemon restart, a stale log file, or an insufficient sleep duration can invalidate an entire benchmark run or, worse, produce misleading results. The assistant's methodical refinement of the restart procedure — from a naive kill-and-hope approach to a verified, multi-step protocol — is a microcosm of the engineering discipline that characterizes the entire optimization campaign.

The Thinking Process: What This Segment Reveals

Across the messages in this segment, a consistent thinking process emerges. The assistant operates in a structured, iterative cycle: analyze the current state, design a targeted change, implement it with precision, test it, diagnose failures, and iterate. Each message builds on the previous one, creating a chain of reasoning that is remarkably transparent.

The handoff document ([msg 2368]) reveals the assistant's ability to synthesize vast amounts of information into a coherent plan. The PCIe transfer inventory table, the VRAM budget calculation, the prioritized TODO list — these are not just documentation; they are reasoning artifacts that show the assistant working through the optimization problem systematically.

The debugging messages ([msg 2440] through [msg 2447]) reveal a methodical approach to failure diagnosis. The assistant does not guess at the cause — it traces through every allocation and deallocation path, verifies reference counting, checks destructor ordering, and considers edge cases like CUDA context corruption after errors. Each hypothesis is tested against the evidence, and the assistant maintains multiple hypotheses simultaneously, letting the evidence narrow the field.

The redesign message ([msg 2449]) reveals the assistant's ability to pivot from debugging to architectural design. The four insights are not arbitrary; they form a logical progression: check memory, compute needs, decide allocation strategy, ensure accurate measurement. The assistant then grounds the plan in the actual codebase by reading the relevant file before making edits.

The validation messages ([msg 2457], [msg 2458]) reveal the assistant's discipline in verification. The headline number (14.2% improvement) is celebrated, but the assistant immediately digs into the detailed timing logs to confirm the kernel-level gains and check for unexpected regressions. The unchanged batch_add metric is noted as confirmation that the optimization was correctly scoped.

The dual-worker benchmark messages ([msg 2464], [msg 2465]) reveal the assistant's resilience in the face of counterintuitive results. The 41-second paradox could have been demoralizing — weeks of work producing a result that was worse than the baseline. But the assistant's response was not frustration or denial; it was immediate, methodical investigation. The question shifted from "did the optimization work?" to "what is the next bottleneck?"

Knowledge Inputs and Outputs

This segment draws on an extraordinary range of knowledge domains: CUDA programming models (streams, events, pinned memory, memory pools), GPU architecture (PCIe bandwidth, SM utilization, copy engine independence), SNARK proving (Groth16, NTT, MSM, Pippenger's algorithm), C++ memory management (RAII, smart pointers, lambda capture semantics), concurrent programming (mutexes, barriers, worker threads), and systems performance analysis (benchmark methodology, bottleneck identification, instrumentation).

The segment produces equally valuable outputs: a validated 14.2% throughput improvement, a memory-aware allocator design that can be reused in future phases, a documented CUDA memory pool interaction pattern that explains phantom OOM failures, and a methodology for debugging concurrent GPU systems that future engineers can follow.

Conclusion

Phase 9 of the cuzk SNARK proving engine optimization represents a significant milestone in the project's trajectory. The 14.2% throughput improvement from 37.4 to 32.1 seconds per proof in single-worker mode is meaningful, but the real value of this phase lies in what it revealed about the system. The CUDA memory pool asymmetry bug, the cleanup ordering issue, the host registration conflicts, the daemon lifecycle pitfalls, and the PCIe bandwidth contention under concurrency — each of these discoveries deepens the team's understanding of how GPU memory management works in a concurrent, multi-worker environment.

The segment also demonstrates a particular style of engineering work that is characteristic of the best systems optimization: measure everything, identify the bottleneck, trace it to specific code paths, design targeted mitigations, implement with precision, test under realistic conditions, diagnose failures systematically, and iterate. The assistant's ability to produce a comprehensive handoff document, implement complex CUDA changes, debug subtle memory pool interactions, validate results with rigorous benchmarking, and pivot immediately when confronted with counterintuitive results is a testament to the power of structured, methodical reasoning in software engineering.

The story does not end here. The dual-worker benchmark revealed that PCIe bandwidth contention is the next bottleneck, and the assistant has already pivoted to analyzing detailed timing logs to understand it. Phase 10 awaits. But Phase 9 has delivered its breakthrough, and the path forward is clearer than ever.