Probing the Interconnect: Measuring GPU-to-GPU Bandwidth in the GLM-5-NVFP4 Performance Investigation

Introduction

In the high-stakes world of large language model inference, every millisecond counts. When a team deploying the massive GLM-5-NVFP4 model across eight NVIDIA RTX PRO 6000 Blackwell GPUs discovered that their actual single-stream throughput of 10.36 tokens per second was a mere 3.4% of the theoretical maximum of 309 tok/s, a systematic manhunt began. The gap was so vast that it demanded a comprehensive parallel audit of every layer of the system stack — from UEFI firmware settings and Linux kernel parameters to NVIDIA driver versions, PCIe configuration, CPU power management, and GPU memory topology. In the span of a single round, the assistant dispatched ten parallel exploration agents, each tasked with investigating a different potential bottleneck. Message 1268 represents one of those agents: the NCCL/P2P bandwidth probe, a critical piece of the puzzle designed to determine whether GPU-to-GPU communication was strangling the model's performance.

The Message: A Task to Measure GPU Interconnect Speed

The message itself is a task dispatch — a structured prompt sent to a subagent that will run its own multi-round conversation to completion before returning results. The full text reads:

[assistant] [task] {"description":"Check NCCL/P2P bandwidth","prompt":"SSH into the LXC container (ssh root@10.1.230.174) and measure actual GPU-to-GPU communication bandwidth. The sglang server is currently running (don't kill it), but we can measure P2P bandwidth with existing tools:\n\n1. Check if NCCL tests are...}

The task result that came back revealed a crucial finding:

With explicit stream synchronization, ALL pairs show ~50 GB/s same-NUMA and ~37 GB/s cross-NUMA. The earlier ~20 TB/s numbers were measurement artifacts from lazy synchronization. These are the real numbers.

This result — approximately 50 GB/s within the same NUMA node and 37 GB/s across NUMA boundaries — would become a key data point in the investigation, ruling out one major hypothesis and redirecting attention toward the actual bottleneck: the FP4 GEMM kernel overhead.

Why This Message Was Written: The Reasoning and Motivation

The context for this message is essential. The team had just completed a theoretical maximum performance calculation that showed the Blackwell GPUs were capable of approximately 309 tokens per second in single-stream inference. The actual measured throughput was 10.36 tok/s — a staggering 30x shortfall. This kind of gap cannot be explained by a single cause; it typically arises from a cascade of misconfigurations, suboptimal kernel choices, and architectural bottlenecks.

The user's instruction in message 1259 was explicit: "Start explore agents to 1. learn about potential issues observed in configurations for UEFI/Linux kernel (runtime knobs as well as kernel version)/LXC settings/Nvidia, Cuda driver versions/etc. and have those agents check the state of the machine we're working with. Use 10 parallel agents."

The assistant responded by designing a comprehensive diagnostic sweep. The ten agents covered: UEFI/BIOS settings, LXC container configuration, NVIDIA driver version and CUDA setup, kernel runtime parameters, kernel version compatibility, GPU memory configuration, PCIe MaxReadReq size, power and thermal management, NCCL/P2P bandwidth, and likely a tenth agent not shown in the immediate context. This was a methodical, divide-and-conquer approach to system debugging — each agent focused on a narrow slice of the stack, gathering empirical data to either confirm or eliminate potential causes.

The NCCL/P2P bandwidth probe (message 1268) was motivated by a specific concern: in a model deployed across eight GPUs using tensor parallelism (TP) and pipeline parallelism (PP), the communication overhead between GPUs can become a dominant factor. If GPU-to-GPU bandwidth was degraded — due to PCIe generation mismatches, NUMA imbalances, or driver issues — it could explain a significant portion of the throughput gap. The GLM-5-NVFP4 model, with its Mixture-of-Experts (MoE) architecture, is particularly sensitive to inter-GPU communication because expert routing requires all-to-all communication patterns across the GPU fabric.

How Decisions Were Made

Several design decisions are visible in the task prompt. First, the assistant chose to measure bandwidth using NCCL tests — the NVIDIA Collective Communications Library — which is the standard tool for benchmarking GPU interconnect performance in ML workloads. This was the right choice because NCCL is what frameworks like SGLang, vLLM, and PyTorch use internally for tensor parallelism communication.

Second, the assistant explicitly chose not to kill the running SGLang server. This was a pragmatic decision: restarting the server would take 80+ seconds (as seen in earlier warmup times), and the measurement tool (NCCL tests) could run alongside the server without conflict. This preserved the ability to immediately resume benchmarking after the diagnostic sweep.

Third, the task prompt shows a fallback strategy: "Check if NCCL tests are available, or compile/run a simple P2P bandwidth test." The assistant anticipated that NCCL benchmarks might not be pre-installed on the container and planned for an alternative measurement approach. This reflects robust engineering thinking — always have a Plan B.

The task also implicitly decided to measure bandwidth with explicit stream synchronization. This turned out to be critical: the earlier measurements (which showed absurd ~20 TB/s) were artifacts of lazy CUDA synchronization, where the benchmark reported completion before data had actually been transferred. The subagent discovered this and corrected the methodology, producing trustworthy numbers.

Assumptions Made

The message operates on several assumptions, most of which were reasonable:

  1. That GPU-to-GPU bandwidth could be a bottleneck. Given the 30x performance gap, every layer of the system was suspect. Communication bandwidth was a legitimate candidate, especially for an MoE model spread across eight GPUs.
  2. That NCCL tests would be available or could be easily compiled. The container had a full CUDA toolkit installed (version 13.1), so NCCL tests were likely present or could be built from source.
  3. That measuring bandwidth wouldn't interfere with the running server. NCCL benchmarks allocate GPU memory and launch kernels, which could potentially disrupt the SGLang server's memory pools. The assistant assumed the impact would be minimal — a reasonable gamble given that the server was idle during the diagnostic window.
  4. That the measurement would be straightforward. The discovery of the lazy synchronization artifact shows this assumption was naive. Without explicit synchronization, the benchmark reported throughput numbers that were physically impossible (20 TB/s on PCIe Gen5 x16, which caps at ~64 GB/s). The subagent had to debug this measurement issue before producing reliable results.
  5. That the LXC container had direct GPU access for P2P. The container configuration (as revealed by another agent) passed through all eight GPUs with cgroup devices allow entries. The assumption that P2P access was functional was validated by the successful measurements.

Mistakes and Incorrect Assumptions

The most significant mistake revealed by this message is not in the task itself but in the earlier measurement methodology. The subagent's report mentions "The earlier ~20 TB/s numbers were measurement artifacts from lazy synchronization." This implies that someone (perhaps the team earlier in the session, or a default NCCL test configuration) had run P2P bandwidth tests without proper CUDA stream synchronization and obtained wildly inflated numbers. This is a classic CUDA benchmarking pitfall: without explicit cudaStreamSynchronize() calls, asynchronous GPU operations may not have completed when the timer stops, leading to measurements that reflect launch latency rather than actual data transfer time.

This mistake could have led to incorrect conclusions — if someone had seen 20 TB/s, they might have concluded "communication is not the bottleneck" based on faulty data. The subagent's correction was therefore essential to the integrity of the entire diagnostic effort.

Another subtle issue: the task prompt says "The sglang server is currently running (don't kill it), but we can measure P2P bandwidth with existing tools." This assumes that running NCCL benchmarks alongside a live server is safe. In practice, NCCL tests allocate significant GPU memory for buffers, which could cause the server's KV cache to be evicted or cause OOM errors. The fact that the measurement succeeded without incident doesn't mean the assumption was universally safe — it was a calculated risk.

Input Knowledge Required

To fully understand this message, one needs knowledge spanning several domains:

Output Knowledge Created

The message produced several important pieces of knowledge:

  1. Measured P2P bandwidth: ~50 GB/s same-NUMA, ~37 GB/s cross-NUMA for all 8 GPU pairs. These numbers are consistent with PCIe Gen5 x16 theoretical limits (~64 GB/s) minus protocol overhead, indicating the PCIe fabric was functioning correctly.
  2. Confirmation of NUMA topology impact: The ~25% bandwidth penalty for cross-NUMA transfers (50 vs 37 GB/s) confirmed that the two-socket EPYC topology was affecting communication, but not catastrophically.
  3. Elimination of communication as primary bottleneck: With 50 GB/s available bandwidth, the 95ms per-token decode time could not be explained by communication latency alone. This redirected the investigation toward compute-bound bottlenecks — specifically the FP4 GEMM kernels, MoE routing, and attention mechanisms.
  4. Methodological insight: The discovery that lazy synchronization produced 20 TB/s artifacts is a valuable lesson for anyone benchmarking GPU P2P bandwidth. It underscores the importance of explicit synchronization in CUDA performance measurement.
  5. Validation of the hardware stack: The PCIe topology, NVIDIA driver, and LXC GPU passthrough were all functioning correctly for P2P communication. This ruled out several classes of configuration errors.

The Thinking Process Visible in the Message

The task prompt reveals a careful, methodical thought process. The assistant considered:

Conclusion

Message 1268 is a small but critical piece of a much larger diagnostic operation. It represents the methodical, hypothesis-driven approach to debugging a complex ML inference system. By measuring GPU-to-GPU bandwidth and finding it healthy (~50 GB/s same-NUMA, ~37 GB/s cross-NUMA), the team was able to eliminate communication as the primary bottleneck and focus on the true culprit: the FP4 GEMM kernel efficiency on Blackwell's SM120 architecture. The message also demonstrates the power of the parallel agent architecture — while this agent was probing P2P bandwidth, nine others were simultaneously investigating different layers of the system, all contributing to a comprehensive system health picture. In the end, the answer wasn't in the interconnects; it was in the kernels.