PCIe and NVLink Interconnects

Reading time: ~35 min · Interview relevance: High · Target roles: ML Engineer, AI Infrastructure Engineer

PCIe 4.0 x16 gives you 32 GB/s between host and GPU. NVLink 4.0 gives you 900 GB/s between GPUs. That 28x gap is not a marketing footnote - it is the physical reason why tensor parallelism only works within a single NVLink domain, and why your DataLoader is probably the bottleneck nobody talks about.

---

The Night the Training Job Stalled at 38% GPU Utilization

It was 2:17 AM on a Tuesday. A team at a large tech company had just spun up a 16-GPU training run for a 7-billion parameter language model. The cluster was state-of-the-art: 8x A100 80GB SXM GPUs per node, dual nodes connected over InfiniBand, the entire stack blessed by the MLOps team. The job launched cleanly. The loss curves looked right. And then someone opened the GPU utilization dashboard.

38%.

Every GPU sitting at 38% utilization on a $500,000 cluster. The team lead spent the next three hours convinced something was misconfigured in NCCL. They cycled through environment variables, tried different collective backends, pinned processes to NUMA nodes. Nothing moved. GPU utilization stayed frustratingly low, as if the GPUs were waiting for something.

They were. They were waiting for data.

The dataset was 800 GB of compressed text stored on a network-attached filesystem. The DataLoader was spawning eight worker processes per GPU. Each worker was reading batches, decompressing them in CPU RAM, and then handing them to the GPU over PCIe. The PCIe bus - the physical connector between the CPU and GPU - was saturated. It could deliver at most 32 GB/s in each direction (PCIe 4.0 x16), and eight workers times two nodes times batch after batch of 2 KB tokenized sequences meant the GPUs were starved. The massive NVLink fabric connecting GPUs to each other at 900 GB/s was idle. The A100s with their 2 TB/s HBM2e memory bandwidth were idle. The whole 16-GPU cluster was bottlenecked by a bus that had not fundamentally changed topology since the 1990s.

The fix, once the team understood the actual constraint, took four hours to implement: pinned (page-locked) memory in the DataLoader, non-blocking transfers to overlap CPU preprocessing with GPU compute, and a prefetch queue that kept two batches in flight simultaneously. GPU utilization jumped to 87%. The model training time dropped by more than half.

That story - the gap between the interconnects you think matter and the ones that actually do - is what this lesson is about. PCIe and NVLink are not just bandwidth numbers. They define the physical architecture of how data moves in every AI training and inference system ever built. Understanding them is the difference between a system that runs at 38% utilization and one that runs at 87%.

---

Why This Exists

Before PCIe, expansion cards communicated with the CPU over the PCI bus - a shared, parallel bus that all cards competed for. Every device saw every other device's traffic. Bandwidth was limited by the slowest device on the bus, and as GPUs got faster, the shared bus became a crisis. A GPU that could process data at tens of gigabytes per second was waiting for a bus that delivered single-digit gigabytes per second and shared that with every other card in the system.

PCIe (Peripheral Component Interconnect Express) was the answer: a point-to-point serial bus where each device gets its own dedicated lanes to the CPU. Launched in 2004, it replaced both PCI and AGP (the graphics-specific bus). The key insight was that serial communication - one bit at a time over multiple dedicated lane pairs - could outrun parallel buses because you could run it at much higher frequencies without signal interference. PCIe scales by adding lanes: x1, x4, x8, x16, each doubling bandwidth.

But PCIe connects GPUs to CPUs. It was never designed for GPU-to-GPU communication. In 2016, when NVIDIA started shipping the Pascal GPU and the first DGX-1, researchers discovered that multi-GPU training was PCIe-limited in a different way: AllReduce operations (where every GPU needs to sum its gradients with every other GPU's gradients) had to route through the CPU over PCIe. A machine with 8 GPUs would generate 8 streams of gradient data, all trying to traverse the same PCIe bus to reach the CPU, get aggregated, and come back. At 2x PCIe 3.0 x16 = 32 GB/s total, this was catastrophically slow.

NVLink was NVIDIA's solution: a direct GPU-to-GPU interconnect that bypasses the CPU entirely. First generation NVLink launched in 2016 with Pascal. Each generation since has roughly doubled bandwidth. NVLink 4.0, shipping with the H100, delivers 900 GB/s total bidirectional bandwidth across all links. A single NVLink 4.0 connection is already 10x faster than the entire PCIe 4.0 x16 connection to the CPU. NVSwitch is the fabric chip that connects all GPUs in a DGX node so every GPU can communicate with every other GPU at full NVLink bandwidth simultaneously.

Understanding both PCIe and NVLink - their limits, their appropriate use cases, and how to code around their constraints - is foundational to building AI systems that run at the hardware's actual ceiling rather than a fraction of it.

---

Historical Context

1992 - PCI: Intel introduces PCI as a shared parallel bus. 133 MB/s peak bandwidth shared among all devices.

2001 - AGP 8x: Graphics-specific bus offering 2.1 GB/s. Faster than PCI, but still limiting compared to GPU compute capabilities.

2004 - PCIe 1.0: Point-to-point serial lanes. x16 slot delivers 4 GB/s per direction (8 GB/s bidirectional). This changes everything - dedicated lanes mean no contention.

2007 - PCIe 2.0: Doubles signaling rate to 5 GT/s per lane. x16 = 8 GB/s per direction.

2010 - PCIe 3.0: 8 GT/s per lane with 128b/130b encoding. x16 = 15.75 GB/s per direction (~16 GB/s). This generation dominates data centers for over a decade.

2016 - NVLink 1.0: Ships with Pascal (P100). 160 GB/s total bidirectional, GPU-to-GPU. Also the year NVIDIA ships the first DGX-1 with 8 P100s. The era of GPU cluster thinking begins.

2017 - NVLink 2.0: Ships with Volta (V100). 300 GB/s bidirectional. NVSwitch introduced in DGX-2 (16 V100s, full any-to-any bandwidth).

2019 - PCIe 4.0: 16 GT/s per lane. x16 = 31.5 GB/s per direction (~32 GB/s). AMD Ryzen 3000 and Zen 2 data center CPUs launch with PCIe 4.0 support.

2020 - NVLink 3.0: Ships with Ampere (A100). 600 GB/s bidirectional. A100 SXM form factor uses NVSwitch, while PCIe A100 uses PCIe 4.0 - a critical distinction that changed how people buy GPUs.

2022 - PCIe 5.0: 32 GT/s per lane. x16 = 63 GB/s per direction. Intel Sapphire Rapids and AMD Genoa are first mainstream server platforms with PCIe 5.0.

2023 - NVLink 4.0: Ships with Hopper (H100). 900 GB/s bidirectional total across all links. The gap between PCIe (64 GB/s) and NVLink (900 GB/s) grows to 14x per direction and ~28x aggregate. At this point, the choice of interconnect topology is the dominant factor in multi-GPU system design.

The engineer who internalized this history - that PCIe and NVLink represent two fundamentally different communication philosophies solving two different problems - understands why model parallelism strategies changed completely between 2016 and 2023.

---

Core Concepts

PCIe: The CPU-GPU Highway

PCIe is a point-to-point, bidirectional serial interconnect. Every PCIe device has a dedicated set of lanes running directly to the CPU (technically, to the PCIe Root Complex inside the CPU package). There is no sharing and no arbitration - your GPU's lanes are yours alone.

Each lane in PCIe is actually two unidirectional serial pairs: one transmit, one receive. This means PCIe is always full-duplex - it can send and receive simultaneously. Bandwidth numbers for PCIe are usually quoted as unidirectional (one direction at a time) or bidirectional (both directions simultaneously). GPUs are almost always installed in x16 slots.

Bandwidth by generation:

Generation	Signaling Rate	Encoding	x16 Unidirectional	x16 Bidirectional
PCIe 3.0	8 GT/s/lane	128b/130b	~15.75 GB/s	~31.5 GB/s
PCIe 4.0	16 GT/s/lane	128b/130b	~31.5 GB/s	~63 GB/s
PCIe 5.0	32 GT/s/lane	128b/130b	~63 GB/s	~126 GB/s

In practice, when ML engineers talk about PCIe bandwidth for a GPU, they mean the unidirectional bandwidth for host-to-device (H2D) or device-to-host (D2H) transfers. This is what determines how fast you can load a batch of training data onto the GPU.

Transfer Latency

Raw bandwidth is not the only constraint. PCIe transfers have a minimum latency floor of roughly 10-20 microseconds even for tiny payloads. This is because every transfer involves:

1. The CPU issuing a DMA (Direct Memory Access) command
2. The DMA engine programming the PCIe transaction
3. The data traversing the PCIe bus (propagation delay)
4. The GPU signaling completion via an interrupt or polling

For a tiny 4 KB tensor, the latency overhead often exceeds the actual transfer time. This is why batching data into large transfers is critical - you want to amortize that fixed overhead across as much data as possible.

The latency-bandwidth product:

$\text{Effective transfer time} = \text{latency} + \frac{\text{data size}}{\text{bandwidth}}$

For a 1 MB tensor over PCIe 4.0:

$t = 15 \mu s + \frac{1 \text{ MB}}{16 \text{ GB/s}} \approx 15 \mu s + 62.5 \mu s = 77.5 \mu s$

For a 100 MB tensor:

$t = 15 \mu s + \frac{100 \text{ MB}}{16 \text{ GB/s}} \approx 15 \mu s + 6250 \mu s \approx 6.3 \text{ ms}$

At 100 MB, the latency is negligible. At 1 KB, it dominates. This is why ML frameworks batch tokenized sequences into large tensors rather than sending them one at a time.

---

Pinned Memory: The DMA Fast Path

When you call tensor.to('cuda') on a normal CPU tensor, something unfortunate can happen. The CPU tensor lives in pageable memory - memory that the operating system can swap to disk, move around, or reclaim at any time. The GPU's DMA engine needs a stable, guaranteed-present physical address to read from. It cannot handle page faults mid-transfer.

The solution CUDA uses for pageable memory is to first copy the data to a temporary staging buffer in pinned (page-locked) memory, and then transfer from that staging buffer over PCIe to the GPU. This means every H2D transfer of pageable memory involves:

1. CPU copies data from pageable RAM to pinned staging buffer
2. DMA engine transfers from pinned staging buffer to GPU VRAM

You are paying for two copies when you only need one.

Pinned (page-locked) memory is RAM that the OS has promised will never be swapped out and will not be moved. The physical address is stable for the lifetime of the allocation. The DMA engine can access it directly. The transfer path is:

1. DMA engine transfers directly from pinned CPU RAM to GPU VRAM

One copy. Nearly double the effective bandwidth for the PCIe link because you are not spending CPU memory bandwidth on the extra copy.

import torch

# Pageable memory - default, the slow path
pageable_tensor = torch.zeros(1024, 1024, dtype=torch.float32)
# to('cuda') will internally: copy to staging buffer, then DMA to GPU

# Pinned memory - locked in RAM, DMA can access directly
pinned_tensor = torch.zeros(1024, 1024, dtype=torch.float32).pin_memory()
# to('cuda') will: DMA directly from pinned buffer to GPU - no staging copy

The tradeoff: pinned memory is a finite, precious resource. The OS cannot reclaim it under memory pressure. Allocating too much pinned memory can starve the OS and other processes. A practical rule: allocate only what you need for active data transfers, not your entire dataset.

The DataLoader in PyTorch exposes this directly:

from torch.utils.data import DataLoader

# Without pinned memory - staging copy on every batch
loader_slow = DataLoader(dataset, batch_size=64, num_workers=4)

# With pinned memory - direct DMA, ~1.5-2x faster H2D transfers
loader_fast = DataLoader(dataset, batch_size=64, num_workers=4, pin_memory=True)

When pin_memory=True, the DataLoader allocates the output tensors in pinned memory in the worker processes. When the main process receives a batch and calls .to('cuda'), the DMA engine can go directly to GPU memory without the staging copy.

---

Non-Blocking Transfers: Overlap Is the Real Gain

Pinned memory cuts the CPU overhead. Non-blocking transfers go further: they allow the CPU to continue doing work while the PCIe transfer is in progress.

A standard tensor.to('cuda') call is synchronous: it blocks the CPU until the transfer completes. The CPU waits. The GPU might also wait. Nothing overlaps.

A non-blocking transfer (tensor.to('cuda', non_blocking=True)) returns immediately on the CPU side. The DMA engine handles the actual transfer asynchronously. This only works with pinned memory - the DMA engine needs a stable source address to read from after the CPU function returns.

The pattern this enables is double buffering (also called prefetching):

While GPU is computing on batch N:
  CPU is loading and preprocessing batch N+1
  DMA is transferring batch N+1 to GPU
When GPU finishes batch N:
  Batch N+1 is already on GPU, ready to go

This hides nearly all PCIe transfer latency behind GPU compute, as long as the compute time is longer than the transfer time (which it almost always is for large models).

---

NVLink: The GPU-to-GPU Superhighway

PCIe connects GPUs to CPUs. NVLink connects GPUs to GPUs. The bandwidth difference reflects this: CPU-GPU communication is a secondary concern in modern AI workloads, but GPU-GPU gradient synchronization during training is the critical path.

NVLink bandwidth by generation:

Generation	Ships With	Total Bandwidth	Links per GPU
NVLink 1.0	Pascal (P100)	160 GB/s	4 x 40 GB/s
NVLink 2.0	Volta (V100)	300 GB/s	6 x 50 GB/s
NVLink 3.0	Ampere (A100)	600 GB/s	12 x 50 GB/s
NVLink 4.0	Hopper (H100)	900 GB/s	18 x 50 GB/s

Each NVLink connection is bidirectional. The "900 GB/s" for H100 is the total across all 18 links in both directions combined.

For comparison, PCIe 4.0 x16 delivers 32 GB/s unidirectional, or 64 GB/s bidirectional. A single H100 NVLink 4.0 link (50 GB/s unidirectional) already outperforms the entire PCIe bus. All 18 links together are 14x faster unidirectional, 14x faster bidirectional.

Why GPU-to-GPU Bandwidth Matters: AllReduce

In data-parallel distributed training, every GPU maintains a copy of the model. After the backward pass, each GPU has computed gradients on its local mini-batch. Before the optimizer step, all GPUs need to agree on the same gradient - the average across all GPUs. This operation is called AllReduce.

AllReduce requires:

• Each GPU sends its gradients to all other GPUs (scatter phase)
• Each GPU receives all other GPUs' gradients and sums them (reduce phase)
• The result is distributed back to all GPUs (broadcast phase)

The communication volume scales with model size. For a 7B parameter model with float16 gradients:

$\text{Gradient size} = 7 \times 10^9 \times 2 \text{ bytes} = 14 \text{ GB}$

With ring-AllReduce (the standard algorithm), each GPU sends and receives $2 \times \frac{N-1}{N}$ times the gradient size, where $N$ is the number of GPUs. For 8 GPUs:

$\text{Data per GPU per AllReduce} = 2 \times \frac{7}{8} \times 14 \text{ GB} \approx 24.5 \text{ GB}$

Over NVLink 4.0 at 900 GB/s aggregate: approximately 27 ms per AllReduce step.

Over PCIe 4.0 at 32 GB/s unidirectional (routing through CPU): approximately 765 ms per AllReduce step.

That 28x difference means the difference between AllReduce taking 2.7% of your step time vs. 76.5%. On PCIe, your GPUs would spend more time communicating than computing. This is why tensor parallelism - which requires extremely frequent all-to-all communication - only works within a single NVLink domain.

---

NVSwitch: All-to-All at Full Bandwidth

NVLink point-to-point connections work for pairs of GPUs. But in an 8-GPU DGX system, you need every GPU to communicate with every other GPU simultaneously. This is where NVSwitch comes in.

NVSwitch is a dedicated ASIC (not a GPU) that acts as a non-blocking crossbar switch for NVLink. It has many NVLink ports - enough to connect to all GPUs in the system. Any GPU can send to any other GPU at full NVLink bandwidth, simultaneously with all other GPU pairs doing the same thing.

The DGX H100 contains 4 NVSwitch chips, each connecting to all 8 GPUs. This creates a fat-tree topology with full bisection bandwidth: you can split the 8 GPUs into two groups of 4, and the total bandwidth between the two groups equals the total bandwidth any single GPU could send or receive.

Full bisection bandwidth means there are no oversubscription bottlenecks. With 4 NVSwitches each connecting all 8 GPUs in a DGX H100:

$\text{Total NVLink bandwidth} = 8 \text{ GPUs} \times 900 \text{ GB/s per GPU} / 2 = 3.6 \text{ TB/s}$

This is the total data the NVSwitch fabric can move per second. For reference, PCIe 4.0 for all 8 GPUs totals $8 \times 32 = 256$ GB/s unidirectional.

The NVSwitch fabric is why DGX H100 supports tensor parallelism across all 8 GPUs. Without NVSwitch, you would need to carefully map parallelism strategies to the physical NVLink topology to avoid bandwidth contention.

---

P2P Access: GPU-A Reading GPU-B's Memory

NVLink enables something more than just fast transfers: peer-to-peer (P2P) direct access. With P2P enabled, GPU A can directly read from and write to GPU B's VRAM without involving the CPU at all. The data never touches CPU memory or the PCIe bus.

This matters for pipeline parallelism. In pipeline parallelism, the model is split across GPUs by layer groups. GPU 0 runs layers 1-10, GPU 1 runs layers 11-20, and so on. After GPU 0 finishes its forward pass, the activations need to be on GPU 1 for it to continue. With P2P over NVLink, GPU 0 can push those activations directly to GPU 1's memory at NVLink bandwidth.

Without P2P: activations go from GPU 0 VRAM, across PCIe to CPU RAM, across PCIe to GPU 1 VRAM - two PCIe traversals at 32 GB/s each.

With P2P over NVLink: activations go from GPU 0 VRAM directly to GPU 1 VRAM - one NVLink transfer at 900 GB/s total fabric bandwidth.

---

Unified Memory: The Abstraction That Hides Complexity

CUDA Unified Memory (also called managed memory) presents a single memory space that is accessible from both CPU and GPU code. You allocate it once with cudaMallocManaged() or torch.cuda.memory.malloc_managed(), and the CUDA driver handles moving pages between CPU and GPU as needed.

The mechanism is hardware page migration. When the GPU accesses a memory address that is currently in CPU RAM, a GPU page fault fires. The CUDA driver catches this fault, migrates the relevant memory page from CPU to GPU, and retries the memory access. All of this is transparent to your code.

When Unified Memory helps:

• Datasets larger than GPU VRAM where you cannot predict access patterns ahead of time
• Code that accesses the same data from CPU and GPU without a clear handoff point
• Rapid prototyping where you want to skip explicit memory management

When Unified Memory hurts:

• Training loops with predictable access patterns (every step, same batch size): page faults add latency and the driver's prediction often lags behind actual access patterns
• When you need predictable, low-latency transfers: page migration latency is variable and can spike
• Multi-GPU scenarios: unified memory with multiple GPUs requires careful attention to memory affinity

In practice, explicit tensor.to('cuda') with pinned memory and non-blocking transfers almost always outperforms unified memory for training workloads. Unified memory is most useful for inference with irregular access patterns or for debugging.

---

The Interconnect Hierarchy

The key insight in this diagram: PCIe is the vertical axis (CPU-to-GPU), NVLink/NVSwitch is the horizontal axis (GPU-to-GPU). These are parallel communication channels. While the CPU is loading the next batch over PCIe to GPU 0, GPUs 0-7 can simultaneously be running AllReduce over NVLink. The hardware supports this overlap; your code must be written to take advantage of it.

---

Transfer Direction Asymmetry

One underappreciated fact: H2D (host-to-device, CPU to GPU) and D2H (device-to-host, GPU to CPU) transfers are often asymmetric. In practice, most systems see H2D bandwidth 10-20% higher than D2H bandwidth. This is because the PCIe protocol slightly favors the direction that the CPU initiates (H2D), and GPU-initiated transfers (D2H) carry slightly more protocol overhead.

For training workloads this barely matters - you move far more data from CPU to GPU (batch data, model checkpoints to load) than from GPU to CPU (losses, small validation metrics). But for inference workloads where you might stream generated tokens back to the CPU for postprocessing, D2H throughput becomes important and you should measure it explicitly, not assume symmetry.

---

Code Examples

1. Measuring PCIe Bandwidth

import torch
import time

def measure_pcie_bandwidth(size_mb: float, direction: str = "h2d", n_trials: int = 20) -> float:
    """
    Measure PCIe bandwidth for a given transfer size.

    Args:
        size_mb: Tensor size in megabytes
        direction: 'h2d' (host-to-device) or 'd2h' (device-to-host)
        n_trials: Number of repetitions to average

    Returns:
        Bandwidth in GB/s
    """
    n_elements = int(size_mb * 1024 * 1024 / 4)  # float32 = 4 bytes

    if direction == "h2d":
        src = torch.randn(n_elements, dtype=torch.float32)  # pageable, on CPU
        dst_device = torch.device("cuda:0")
    else:
        src = torch.randn(n_elements, dtype=torch.float32, device="cuda:0")
        dst_device = torch.device("cpu")

    # Warm up - first transfer is always slower due to driver init
    _ = src.to(dst_device)
    torch.cuda.synchronize()

    start_event = torch.cuda.Event(enable_timing=True)
    end_event = torch.cuda.Event(enable_timing=True)

    start_event.record()
    for _ in range(n_trials):
        _ = src.to(dst_device)
    end_event.record()

    torch.cuda.synchronize()

    elapsed_ms = start_event.elapsed_time(end_event)
    elapsed_s = elapsed_ms / 1000.0

    total_bytes = n_elements * 4 * n_trials
    bandwidth_gbs = total_bytes / elapsed_s / 1e9

    return bandwidth_gbs


if __name__ == "__main__":
    # Test across different tensor sizes to see bandwidth vs. size curve
    sizes_mb = [0.001, 0.01, 0.1, 1.0, 10.0, 100.0, 500.0]

    print(f"{'Size (MB)':<12} {'H2D (GB/s)':<14} {'D2H (GB/s)':<14}")
    print("-" * 40)

    for size_mb in sizes_mb:
        h2d_bw = measure_pcie_bandwidth(size_mb, "h2d")
        d2h_bw = measure_pcie_bandwidth(size_mb, "d2h")
        print(f"{size_mb:<12.3f} {h2d_bw:<14.2f} {d2h_bw:<14.2f}")

    print()
    print("Note: Small sizes show low bandwidth due to fixed ~15us latency overhead.")
    print("Large sizes approach theoretical PCIe limit (~15-32 GB/s depending on gen).")

Expected output (PCIe 4.0 system):

Size (MB)    H2D (GB/s)    D2H (GB/s)
----------------------------------------
0.001        0.05          0.04
0.010        0.48          0.42
0.100        4.21          3.87
1.000        13.45         12.18
10.000       26.33         24.71
100.000      30.87         28.54
500.000      31.24         29.01

The pattern is exactly what the latency-bandwidth formula predicts: small tensors are latency-dominated (low effective bandwidth), large tensors saturate the PCIe link.

---

2. Pinned Memory vs. Pageable Memory Benchmark

import torch
import time
from typing import Tuple

def benchmark_transfer(
    tensor_size_mb: float,
    use_pinned: bool,
    n_trials: int = 50
) -> Tuple[float, float]:
    """
    Compare H2D transfer speed with pageable vs pinned memory.

    Returns:
        (mean_ms, bandwidth_gbs) tuple
    """
    n_elements = int(tensor_size_mb * 1024 * 1024 / 4)

    if use_pinned:
        cpu_tensor = torch.randn(n_elements, dtype=torch.float32).pin_memory()
    else:
        cpu_tensor = torch.randn(n_elements, dtype=torch.float32)

    # Warm up
    _ = cpu_tensor.to("cuda:0")
    torch.cuda.synchronize()

    start = torch.cuda.Event(enable_timing=True)
    end = torch.cuda.Event(enable_timing=True)

    start.record()
    for _ in range(n_trials):
        gpu_tensor = cpu_tensor.to("cuda:0")
        # Note: .to() is synchronous - it blocks until transfer completes
    end.record()
    torch.cuda.synchronize()

    elapsed_ms = start.elapsed_time(end) / n_trials
    bytes_transferred = n_elements * 4
    bandwidth_gbs = bytes_transferred / (elapsed_ms / 1000.0) / 1e9

    return elapsed_ms, bandwidth_gbs


def run_pinned_vs_pageable_benchmark():
    sizes_mb = [1.0, 10.0, 100.0, 500.0]

    print(f"{'Size':<10} {'Pageable (GB/s)':<20} {'Pinned (GB/s)':<20} {'Speedup':<10}")
    print("-" * 60)

    for size_mb in sizes_mb:
        _, pageable_bw = benchmark_transfer(size_mb, use_pinned=False)
        _, pinned_bw = benchmark_transfer(size_mb, use_pinned=True)
        speedup = pinned_bw / pageable_bw
        print(f"{size_mb:<10.0f} {pageable_bw:<20.2f} {pinned_bw:<20.2f} {speedup:<10.2f}x")


if __name__ == "__main__":
    run_pinned_vs_pageable_benchmark()

Typical output:

Size       Pageable (GB/s)      Pinned (GB/s)        Speedup
------------------------------------------------------------
1          9.14                 13.45                1.47x
10         14.83                26.33                1.78x
100        17.21                30.87                1.79x
500        17.89                31.24                1.75x

Pinned memory delivers roughly 1.5-1.8x higher bandwidth because it eliminates the internal staging copy. The speedup is consistent across sizes because the staging copy cost scales linearly with size.

---

3. Double Buffering - Overlapping Transfer With Compute

import torch
import torch.nn as nn
import time

class DoubleBufferLoader:
    """
    Implements double-buffered data loading: while the GPU computes on batch N,
    DMA is transferring batch N+1 to the GPU in the background.

    This pattern hides PCIe latency behind GPU compute time.
    """

    def __init__(self, dataset, batch_size: int, device: str = "cuda:0"):
        self.dataset = dataset
        self.batch_size = batch_size
        self.device = device
        self.n_batches = len(dataset) // batch_size

    def iterate_with_overlap(self):
        """Generator that yields batches, pre-loading the next batch non-blocking."""
        # Pre-allocate two pinned memory buffers - the double buffer
        batch_a = torch.zeros(self.batch_size, *self.dataset[0].shape).pin_memory()
        batch_b = torch.zeros(self.batch_size, *self.dataset[0].shape).pin_memory()

        # Load first batch into buffer A
        for i in range(self.batch_size):
            batch_a[i] = self.dataset[i]

        # Start non-blocking transfer of batch 0 to GPU
        gpu_current = batch_a.to(self.device, non_blocking=True)

        buffers = [batch_a, batch_b]

        for batch_idx in range(1, self.n_batches):
            # Fill the OTHER buffer on CPU (while GPU works on current)
            next_buf = buffers[batch_idx % 2]
            start = batch_idx * self.batch_size
            end = start + self.batch_size
            for i, idx in enumerate(range(start, end)):
                next_buf[i] = self.dataset[idx]

            # Start non-blocking transfer of next batch to GPU
            # This starts DMA BEFORE we yield, so transfer overlaps with GPU compute
            gpu_next = next_buf.to(self.device, non_blocking=True)

            # Yield current batch (GPU processes this while next batch is transferring)
            # The synchronization happens implicitly when GPU kernels consume gpu_current
            yield gpu_current

            # Swap: next becomes current
            gpu_current = gpu_next

        # Yield the final batch
        yield gpu_current


def benchmark_double_buffer_vs_naive(n_batches: int = 100, batch_size: int = 256):
    """Compare naive sequential loading vs double-buffered loading."""
    feature_dim = 512
    dataset = [torch.randn(feature_dim) for _ in range(n_batches * batch_size)]
    device = "cuda:0"

    # Simulate a simple model step
    model = nn.Linear(feature_dim, 256).to(device)

    def fake_compute_step(batch):
        """Simulates GPU compute work proportional to model size."""
        with torch.no_grad():
            out = model(batch)
        torch.cuda.synchronize()
        return out

    # --- Naive approach: synchronous transfer then compute ---
    torch.cuda.synchronize()
    t0 = time.perf_counter()

    for i in range(n_batches):
        # Create batch from CPU data
        batch_cpu = torch.stack(dataset[i * batch_size:(i + 1) * batch_size])
        # Blocking transfer - CPU waits for DMA to finish
        batch_gpu = batch_cpu.to(device)
        # Then compute
        _ = fake_compute_step(batch_gpu)

    torch.cuda.synchronize()
    naive_time = time.perf_counter() - t0

    # --- Double buffered approach: overlap transfer with compute ---
    loader = DoubleBufferLoader(dataset, batch_size, device)

    torch.cuda.synchronize()
    t0 = time.perf_counter()

    for batch_gpu in loader.iterate_with_overlap():
        # Transfer of NEXT batch is already in flight via non_blocking=True
        # GPU compute on CURRENT batch runs while DMA fetches NEXT batch
        _ = fake_compute_step(batch_gpu)

    torch.cuda.synchronize()
    overlap_time = time.perf_counter() - t0

    print(f"Naive sequential:        {naive_time:.3f}s ({n_batches/naive_time:.1f} batches/s)")
    print(f"Double buffered overlap: {overlap_time:.3f}s ({n_batches/overlap_time:.1f} batches/s)")
    print(f"Speedup: {naive_time/overlap_time:.2f}x")


if __name__ == "__main__":
    benchmark_double_buffer_vs_naive()

In practice, the double buffer speedup depends on the ratio of compute time to transfer time. If compute takes 50 ms and transfer takes 10 ms, you hide most of the transfer cost. If compute takes 5 ms and transfer takes 10 ms, you cannot fully hide the transfer - the GPU has to wait. This is the "roofline" for data loading: your compute must be at least as slow as your transfers for full overlap.

---

4. Checking NVLink Topology

import torch
import subprocess
import os

def check_nvlink_topology():
    """
    Query NVLink peer-to-peer connectivity between all GPUs.
    Uses both PyTorch and nvidia-smi for topology information.
    """
    n_gpus = torch.cuda.device_count()
    print(f"Number of GPUs: {n_gpus}")
    print()

    # Check CUDA compute capability (useful for feature support)
    for i in range(n_gpus):
        major, minor = torch.cuda.get_device_capability(i)
        props = torch.cuda.get_device_properties(i)
        print(f"GPU {i}: {props.name}")
        print(f"  Compute capability: {major}.{minor}")
        print(f"  Total memory: {props.total_memory / 1e9:.1f} GB")
        print(f"  Multi-processor count: {props.multi_processor_count}")
        print()

    # Check P2P access between all GPU pairs
    print("P2P (peer-to-peer) access matrix (1=enabled, 0=not available):")
    print(f"{'':>8}", end="")
    for j in range(n_gpus):
        print(f"  GPU{j}", end="")
    print()

    for i in range(n_gpus):
        print(f"GPU{i:>4}  ", end="")
        for j in range(n_gpus):
            if i == j:
                print("    - ", end="")
            else:
                # Check if P2P access is available between GPU i and GPU j
                can_access = torch.cuda.can_device_access_peer(i, j)
                print(f"    {'1' if can_access else '0'}", end="")
        print()

    print()

    # Use nvidia-smi to get NVLink status (requires nvidia-smi to be in PATH)
    try:
        result = subprocess.run(
            ["nvidia-smi", "topo", "--matrix"],
            capture_output=True,
            text=True,
            timeout=10
        )
        if result.returncode == 0:
            print("nvidia-smi topology matrix:")
            print(result.stdout)
        else:
            print("nvidia-smi topo not available (may need root or specific driver version)")
    except (subprocess.TimeoutExpired, FileNotFoundError):
        print("nvidia-smi not available in PATH")


def enable_p2p_access():
    """
    Explicitly enable peer-to-peer access between all GPU pairs.
    PyTorch's NCCL backend does this automatically, but for manual P2P
    tensor operations you may need to call this.
    """
    n_gpus = torch.cuda.device_count()

    for i in range(n_gpus):
        for j in range(n_gpus):
            if i != j and torch.cuda.can_device_access_peer(i, j):
                with torch.cuda.device(i):
                    torch.cuda.enable_peer_access(j)
                    print(f"Enabled P2P: GPU {i} -> GPU {j}")


def benchmark_p2p_transfer(src_gpu: int = 0, dst_gpu: int = 1, size_mb: float = 100.0):
    """
    Benchmark GPU-to-GPU transfer speed.
    On NVLink systems this should approach NVLink bandwidth.
    On PCIe-only systems this routes through CPU at PCIe speed.
    """
    n_elements = int(size_mb * 1024 * 1024 / 4)

    src_tensor = torch.randn(n_elements, device=f"cuda:{src_gpu}")
    torch.cuda.synchronize()

    n_trials = 50
    start = torch.cuda.Event(enable_timing=True)
    end = torch.cuda.Event(enable_timing=True)

    # Warm up
    _ = src_tensor.to(f"cuda:{dst_gpu}")
    torch.cuda.synchronize()

    start.record()
    for _ in range(n_trials):
        dst_tensor = src_tensor.to(f"cuda:{dst_gpu}")
    end.record()
    torch.cuda.synchronize()

    elapsed_ms = start.elapsed_time(end) / n_trials
    bandwidth_gbs = (n_elements * 4) / (elapsed_ms / 1000.0) / 1e9

    p2p_available = torch.cuda.can_device_access_peer(src_gpu, dst_gpu)
    path = "NVLink (P2P)" if p2p_available else "PCIe (via CPU)"

    print(f"GPU{src_gpu} -> GPU{dst_gpu} transfer ({size_mb:.0f} MB)")
    print(f"  Path: {path}")
    print(f"  Bandwidth: {bandwidth_gbs:.2f} GB/s")
    print(f"  Time: {elapsed_ms:.3f} ms")


if __name__ == "__main__":
    check_nvlink_topology()
    benchmark_p2p_transfer(0, 1, size_mb=100.0)

Expected output on a DGX A100 (NVLink 3.0):

P2P access matrix (1=enabled, 0=not available):
         GPU0  GPU1  GPU2  GPU3  GPU4  GPU5  GPU6  GPU7
GPU0      -     1     1     1     1     1     1     1
GPU1      1     -     1     1     1     1     1     1
...

GPU0 -> GPU1 transfer (100 MB)
  Path: NVLink (P2P)
  Bandwidth: 274.83 GB/s
  Time: 0.364 ms

Expected output on a PCIe-only 4-GPU server (no NVLink):

P2P access matrix:
         GPU0  GPU1  GPU2  GPU3
GPU0      -     0     0     0
...

GPU0 -> GPU1 transfer (100 MB)
  Path: PCIe (via CPU)
  Bandwidth: 12.84 GB/s
  Time: 7.789 ms

---

5. NCCL AllReduce Timing: NVLink vs. PCIe

import torch
import torch.distributed as dist
import os
import time
import subprocess
from typing import Optional

def simulate_allreduce_timing(
    model_params_billions: float,
    n_gpus: int,
    bandwidth_gbs: float,
    algorithm: str = "ring"
) -> float:
    """
    Estimate AllReduce time for a given model and interconnect.

    Uses ring-AllReduce communication volume formula:
        bytes per GPU = 2 * (N-1)/N * gradient_bytes

    Args:
        model_params_billions: Model size (e.g., 7.0 for 7B params)
        n_gpus: Number of GPUs in the AllReduce group
        bandwidth_gbs: Available interconnect bandwidth per GPU (GB/s)
        algorithm: 'ring' (standard) or 'tree'

    Returns:
        Estimated AllReduce time in milliseconds
    """
    # float16 gradients: 2 bytes per parameter
    gradient_bytes = model_params_billions * 1e9 * 2

    if algorithm == "ring":
        # Ring-AllReduce: each GPU sends/receives 2*(N-1)/N * total_bytes
        bytes_per_gpu = 2 * ((n_gpus - 1) / n_gpus) * gradient_bytes
    elif algorithm == "tree":
        # Tree-AllReduce: each GPU sends/receives 2*log2(N) * total_bytes / N
        import math
        bytes_per_gpu = 2 * math.log2(n_gpus) * gradient_bytes / n_gpus
    else:
        raise ValueError(f"Unknown algorithm: {algorithm}")

    # Time = data / bandwidth (assuming perfect pipelining)
    bandwidth_bytes_per_s = bandwidth_gbs * 1e9
    time_s = bytes_per_gpu / bandwidth_bytes_per_s
    time_ms = time_s * 1000.0

    return time_ms


def print_allreduce_comparison():
    """
    Compare AllReduce times across different interconnects and model sizes.
    Shows why interconnect choice is critical for large-scale training.
    """
    model_sizes = [0.1, 1.0, 7.0, 70.0, 175.0]  # billions of parameters
    n_gpus = 8

    # Bandwidth assumptions (per-GPU unidirectional for ring-AllReduce)
    interconnects = {
        "PCIe 3.0 x16 (16 GB/s)": 16.0,
        "PCIe 4.0 x16 (32 GB/s)": 32.0,
        "NVLink 3.0 A100 (300 GB/s effective)": 300.0,
        "NVLink 4.0 H100 (450 GB/s effective)": 450.0,
    }

    print(f"AllReduce time for ring-AllReduce, {n_gpus} GPUs, float16 gradients")
    print("=" * 90)
    print(f"{'Model Size':<15}", end="")
    for ic_name in interconnects:
        print(f"  {ic_name[:20]:<22}", end="")
    print()
    print("-" * 90)

    for model_b in model_sizes:
        print(f"{model_b:>8.1f}B    ", end="")
        for ic_name, bw in interconnects.items():
            time_ms = simulate_allreduce_timing(model_b, n_gpus, bw)
            if time_ms < 1:
                print(f"  {time_ms*1000:.1f}us          ", end="")
            elif time_ms < 1000:
                print(f"  {time_ms:.1f}ms           ", end="")
            else:
                print(f"  {time_ms/1000:.2f}s            ", end="")
        print()

    print()
    print("Reference: A typical training step for a 7B model takes ~100-500ms on H100.")
    print("PCIe 4.0 AllReduce (765ms) > step compute time = interconnect is the bottleneck.")
    print("NVLink 4.0 AllReduce (27ms) << step compute time = interconnect is NOT bottleneck.")


def run_actual_allreduce_benchmark():
    """
    Run an actual NCCL AllReduce and time it.
    Must be run in a distributed context (e.g., via torchrun).

    Usage:
        torchrun --nproc_per_node=8 this_script.py
    """
    if not dist.is_available():
        print("torch.distributed not available")
        return

    # Initialize process group
    dist.init_process_group(backend="nccl")
    rank = dist.get_rank()
    world_size = dist.get_world_size()
    device = torch.device(f"cuda:{rank}")
    torch.cuda.set_device(device)

    # Simulate 7B model gradients in float16
    n_params = 7_000_000_000
    grad_tensor = torch.randn(n_params // world_size, dtype=torch.float16, device=device)

    # Warm up
    for _ in range(3):
        dist.all_reduce(grad_tensor, op=dist.ReduceOp.SUM)
    torch.cuda.synchronize()

    # Time the AllReduce
    n_trials = 20
    start = torch.cuda.Event(enable_timing=True)
    end = torch.cuda.Event(enable_timing=True)

    start.record()
    for _ in range(n_trials):
        dist.all_reduce(grad_tensor, op=dist.ReduceOp.SUM)
    end.record()
    torch.cuda.synchronize()

    elapsed_ms = start.elapsed_time(end) / n_trials

    if rank == 0:
        total_bytes = n_params * 2  # float16
        # Ring AllReduce effective bytes per GPU
        effective_bytes = 2 * ((world_size - 1) / world_size) * total_bytes
        bandwidth_gbs = effective_bytes / (elapsed_ms / 1000.0) / 1e9

        print(f"AllReduce benchmark: {world_size} GPUs, ~7B parameter gradients (float16)")
        print(f"  AllReduce time: {elapsed_ms:.2f} ms")
        print(f"  Effective bandwidth: {bandwidth_gbs:.1f} GB/s")

        # Check if NVLink is being used
        topo_result = subprocess.run(
            ["nvidia-smi", "topo", "-m"],
            capture_output=True, text=True
        )
        if "NV" in topo_result.stdout:
            print("  Topology: NVLink detected - NCCL should use NVLink backend")
        else:
            print("  Topology: No NVLink detected - NCCL using PCIe")

    dist.destroy_process_group()


if __name__ == "__main__":
    print_allreduce_comparison()

Output of print_allreduce_comparison():

AllReduce time for ring-AllReduce, 8 GPUs, float16 gradients
==========================================================================================
Model Size       PCIe 3.0 x16 (16 GB/s)    PCIe 4.0 x16 (32 GB/s)    NVLink 3.0 A100 (300 GB/s effective)    NVLink 4.0 H100 (450 GB/s effective)
-------------------------------------------------------------------------------------------
     0.1B         26.3ms                     13.1ms                     1.4ms                  0.9ms
     1.0B         262.5ms                    131.3ms                    14.0ms                 9.3ms
     7.0B         1837.5ms                   918.8ms                    98.0ms                 65.3ms
    70.0B         18375.0ms                  9187.5ms                   980.0ms                653.3ms
   175.0B         45937.5ms                  22968.8ms                  2450.0ms               1633.3ms

The 7B row tells the story: PCIe 4.0 AllReduce takes 918 ms. NVLink 4.0 takes 65 ms. A typical H100 training step for a 7B model takes 200-400 ms of pure compute. PCIe AllReduce would be 2-4x slower than the compute step itself. NVLink AllReduce is a small fraction of compute time.

---

Production Engineering Notes

NUMA Affinity and PCIe Root Complex Placement

Modern servers have multiple CPU sockets (NUMA nodes). Each NUMA node has its own memory and its own PCIe Root Complex. GPUs are attached to specific NUMA nodes. If you run a process on CPU 0 (NUMA node 0) and it does PCIe transfers to a GPU attached to CPU 1 (NUMA node 1), the data crosses the QPI/UPI inter-socket fabric - adding latency and reducing bandwidth.

Always pin your training processes to the NUMA node that owns the GPU:

# Check which NUMA node each GPU is attached to
nvidia-smi topo --matrix

# Run a process on the correct NUMA node (GPU 0 on NUMA node 0 in this example)
numactl --cpunodebind=0 --membind=0 python train.py

# Or use PyTorch's launcher which handles this automatically
torchrun --nproc_per_node=8 train.py

DataLoader Worker Count and PCIe Saturation

The optimal number of DataLoader workers is not "as many cores as you have." It is the number that keeps the PCIe bus saturated without exceeding it.

A PCIe 4.0 bus at 32 GB/s with 64-byte cache lines processes about 500 million cache line transfers per second. For typical ML batch sizes (256 samples of 224x224x3 images = ~37 MB), you need roughly one batch every 1.2 ms to keep a fast GPU busy. With worker overhead, 4-8 workers is usually optimal. Beyond that, you add CPU overhead and context switching without gaining PCIe bandwidth because the bus is already saturated.

Profile with nvtx or torch.profiler to identify if you are PCIe-bound:

with torch.profiler.profile(
    activities=[
        torch.profiler.ProfilerActivity.CPU,
        torch.profiler.ProfilerActivity.CUDA,
    ],
    record_shapes=True,
    with_stack=True
) as prof:
    for batch_idx, (data, target) in enumerate(train_loader):
        if batch_idx >= 10:
            break
        data = data.to("cuda:0")
        output = model(data)
        loss = criterion(output, target.to("cuda:0"))
        loss.backward()
        optimizer.step()

# Look for gaps between CUDA kernels in the trace
print(prof.key_averages().table(sort_by="cuda_time_total", row_limit=20))

If you see large gaps between CUDA kernel executions that align with aten::_to_copy operations, you are PCIe-bound.

NCCL Environment Variables

NCCL uses several environment variables that can dramatically affect whether it uses NVLink or PCIe:

# Force NCCL to use P2P (NVLink) - usually auto-detected but useful for debugging
export NCCL_P2P_LEVEL=NVL

# Disable P2P to force PCIe path (useful for comparing performance)
export NCCL_P2P_DISABLE=1

# Show NCCL topology detection output
export NCCL_DEBUG=INFO
export NCCL_DEBUG_SUBSYS=GRAPH

# Force NCCL to use specific network interface for inter-node communication
export NCCL_SOCKET_IFNAME=ib0  # InfiniBand interface

# Set number of channels (parallel communication streams)
# Higher = more parallelism but more overhead. Default is usually optimal.
export NCCL_NCHANNELS_PER_NET_PEER=2

When diagnosing slow distributed training, always run with NCCL_DEBUG=INFO first. NCCL will log which topology it detected and which transport it selected. If it says "PCIe" instead of "NVLink" on a system with NVLink hardware, something is wrong with P2P configuration or IOMMU settings.

IOMMU and P2P Transfers

IOMMU (Input-Output Memory Management Unit) virtualizes memory addresses for PCIe devices. On some Linux configurations, IOMMU is enabled in a mode that blocks P2P direct GPU-to-GPU transfers because the IOMMU cannot validate the peer memory access.

Symptoms: torch.cuda.can_device_access_peer(i, j) returns False even though the system has NVLink hardware. Solution:

# Check current IOMMU mode
dmesg | grep -i iommu

# Disable IOMMU in passthrough mode (allows direct DMA between devices)
# Add to /etc/default/grub GRUB_CMDLINE_LINUX_DEFAULT:
# For Intel: intel_iommu=on iommu=pt
# For AMD:   amd_iommu=on iommu=pt

# Alternatively, use NVIDIA's ACS (Access Control Services) override
# WARNING: reduces security isolation between PCIe devices
echo 1 > /sys/bus/pci/devices/XXXX:XX:XX.X/enable_acs

This is a common gotcha in cloud instances where IOMMU is configured for VM isolation, inadvertently breaking NVLink P2P access.

Gradient Checkpointing and Interconnect Bandwidth

Gradient checkpointing trades compute for memory by recomputing activations during the backward pass instead of storing them. A less-discussed interaction: gradient checkpointing also reduces the amount of data that needs to move over interconnects.

In tensor parallelism, forward pass activations are sharded and must be communicated between GPUs. By checkpointing activations rather than storing them, you reduce the size of communication payloads during the backward pass. On PCIe systems (where interconnect is the bottleneck), gradient checkpointing can improve throughput not just by freeing memory, but by reducing communication volume.

---

Interconnect Bandwidth Summary

---

Model Parallelism and Interconnect Requirements

Different parallelism strategies have radically different communication requirements. The interconnect you have determines which strategies are viable.

Data Parallel (DDP): AllReduce once per step. Communication volume scales with model size. On PCIe, feasible for models up to ~1B parameters where the AllReduce time is a small fraction of compute time. Above that, the AllReduce takes longer than the step itself.

Tensor Parallel: Every matrix multiplication (every forward and backward pass attention and FFN operation) requires an AllReduce. For a transformer with 96 layers, this means hundreds of AllReduce operations per step. Requires NVLink bandwidth. On PCIe this is completely impractical.

Pipeline Parallel: P2P activation transfers between stages. Volume scales with batch size times activation size, not model size. Viable on PCIe but introduces pipeline bubbles (idle GPU time while waiting for activations from the previous stage). NVLink reduces bubble size.

Hybrid (Megatron-LM style): Tensor parallelism within a single NVLink domain (all GPUs in one DGX node), data or pipeline parallelism across nodes over InfiniBand. This is the standard approach for training models above 20B parameters.

---

Common Mistakes

:::danger Forgetting to Pin Memory in Production DataLoaders The PyTorch DataLoader default is pin_memory=False. In production training, this means every batch transfer involves an internal staging copy that halves your effective PCIe bandwidth. Always set pin_memory=True for GPU training. The only exception is when GPU memory is extremely constrained and you cannot afford the pinned memory reservation, but even then the tradeoff is usually worth analyzing.

# WRONG - pageable memory, staging copy on every transfer
loader = DataLoader(dataset, batch_size=256, num_workers=8)

# CORRECT - pinned memory, direct DMA to GPU
loader = DataLoader(dataset, batch_size=256, num_workers=8, pin_memory=True)

:::

:::danger Using tensor.to(device) Instead of tensor.to(device, non_blocking=True) in Training Loops Blocking transfers stall the entire training loop waiting for PCIe DMA to complete. In a tight training loop, this means your GPU compute pipeline is idle for 5-15 ms per batch. With non-blocking transfers and a prefetch structure, that time overlaps with GPU compute.

# WRONG - CPU blocks until transfer completes, GPU sits idle
for batch, labels in loader:
    batch = batch.to(device)
    labels = labels.to(device)
    output = model(batch)

# CORRECT - transfer is in flight while previous batch computes
for batch, labels in loader:
    batch = batch.to(device, non_blocking=True)
    labels = labels.to(device, non_blocking=True)
    # Synchronization happens automatically when GPU kernel consumes these tensors
    output = model(batch)

:::

:::warning Assuming Tensor Parallel Will Work Across PCIe Boundaries If your server has 4 GPUs connected only via PCIe (no NVLink), tensor parallelism is impractical for any model above toy size. The AllReduce required for every matmul would take 50-200x longer than on NVLink, making it slower than running on a single GPU. Check your hardware before designing your parallelism strategy.

# Before designing your parallelism strategy, check P2P availability
for i in range(torch.cuda.device_count()):
    for j in range(i + 1, torch.cuda.device_count()):
        p2p = torch.cuda.can_device_access_peer(i, j)
        print(f"GPU {i} <-> GPU {j}: P2P = {'NVLink' if p2p else 'PCIe only'}")

:::

:::warning Misunderstanding Bidirectional vs. Unidirectional PCIe Bandwidth PCIe bandwidth is often quoted as bidirectional (e.g., "PCIe 4.0 = 64 GB/s"). But for a single H2D or D2H transfer, you only use one direction at a time. The effective bandwidth for loading a batch onto the GPU is the unidirectional bandwidth: 32 GB/s for PCIe 4.0, not 64 GB/s. Make sure your estimates and bottleneck analysis use the correct direction. :::

:::warning Ignoring IOMMU Settings on Cloud Instances Many cloud GPU instances (especially those using SR-IOV for virtualization) have IOMMU configured in a way that disables NVLink P2P access between GPUs. You may have a server with NVLink hardware but torch.cuda.can_device_access_peer() returns False because IOMMU is blocking direct GPU-to-GPU DMA. Always verify P2P access in your actual environment, not based on hardware specs alone. Contact cloud support or check IOMMU kernel parameters if you suspect this. :::

:::danger Allocating Too Much Pinned Memory Pinned memory cannot be swapped by the OS. If you pin too much memory (e.g., the entire dataset), you starve the OS of reclaimable memory, leading to system instability, OOM kills of other processes, and potential kernel crashes. Pin only the active transfer buffers - typically 2-4 batches worth of data, not the entire dataset. The DataLoader with pin_memory=True handles this correctly by only pinning the buffers for the current and prefetched batches. :::

---

Interview Questions and Answers

Q1: PCIe 4.0 x16 is rated at 32 GB/s. You measure only 17 GB/s in your benchmark. What are five possible reasons for the gap?

A: The gap between rated and measured bandwidth can come from several sources:

1. Encoding overhead: PCIe uses 128b/130b encoding, meaning 2 bits out of every 130 are protocol overhead. The theoretical maximum is $32 \times 128/130 \approx 31.5$ GB/s, not 32 GB/s. But that still doesn't explain 17 GB/s.

2. Pageable memory staging copy: If the source tensor is in pageable (non-pinned) memory, CUDA internally copies it to a pinned staging buffer before DMA. This wastes CPU memory bandwidth and effectively halves measured throughput. Fix: use pin_memory=True.

3. Small transfer sizes: The fixed PCIe latency (~15 us) dominates for small tensors. A 1 MB tensor sees $15 us + 62 us = 77 us$ total, yielding only 13 GB/s even at full bandwidth.

4. PCIe width negotiation failure: If the x16 slot does not have x16 electrical lanes (common in some consumer motherboards that share PCIe lanes with NVMe slots), the GPU may be running at x8 or x4, halving or quartering bandwidth.

5. NUMA mismatch: The GPU is attached to CPU socket 1 (NUMA node 1) but the process is running on CPU socket 0. Data crosses the QPI/UPI inter-socket link, adding latency and reducing bandwidth by 20-40%.

---

Q2: Explain why tensor parallelism requires NVLink but data parallelism can work over PCIe.

A: The difference is communication frequency relative to compute.

In data parallelism (DDP), each GPU runs the full forward and backward pass on its local mini-batch, then performs one AllReduce per parameter per step to synchronize gradients. For a 7B parameter model, this is about 14 GB of float16 gradients transferred once per step. If each step takes 300 ms of compute, the AllReduce over NVLink takes ~27 ms (9%), and over PCIe 4.0 takes ~765 ms (well over 100% of step time). PCIe DDP for large models is slow but at least the communication-to-compute ratio is manageable for smaller models.

In tensor parallelism, the model's weight matrices are split across GPUs column-wise or row-wise. Every matrix multiplication requires an AllReduce to sum partial results. In a transformer with 96 layers and 2 matmuls per layer (attention and FFN), that is 192 AllReduce operations per forward pass, and 192 more per backward pass - roughly 400 AllReduces per step. Each AllReduce is small (just the activation tensor, not gradients), but at 400 operations, even 5 ms per AllReduce (PCIe, small tensors) would add 2 seconds of communication overhead to a 300 ms compute step. NVLink reduces this to microseconds per AllReduce, making the total communication overhead a few milliseconds rather than seconds.

---

Q3: What is the difference between pinned memory and unified memory, and when should you use each?

A: Both provide different solutions to the CPU-GPU memory transfer problem:

Pinned (page-locked) memory is CPU RAM that the OS has promised will not be swapped or relocated. The physical address is guaranteed stable, allowing the GPU's DMA engine to access it directly without CPU involvement. The key properties: fast for H2D/D2H transfers (no staging copy), limited in quantity (cannot be swapped), must be managed explicitly (allocate before the transfer loop, free after).

Unified memory (CUDA managed memory) presents a single virtual address space accessible from both CPU and GPU. Pages migrate automatically between CPU and GPU RAM as the hardware detects access patterns via page faults. Key properties: simple programming model (no explicit transfers), variable performance (page migration latency can spike), works for any data size (the OS can handle pages going to disk if needed).

Use pinned memory when: you have predictable, repeated data transfers (training data loading), transfer size is fixed and known, and latency predictability matters. The DataLoader pin_memory=True is the standard use case.

Use unified memory when: you are prototyping and don't want to manage explicit transfers, working with datasets that are accessed irregularly (cannot predict which pages will be needed when), or the model's memory access pattern doesn't have a clear CPU-write / GPU-read handoff structure.

In production training loops, pinned memory with explicit transfers almost always outperforms unified memory because the access pattern is fully predictable (batch in, compute, next batch).

---

Q4: You're building a 4-node training cluster, 8 GPUs per node. Which parallelism strategy would you use and why?

A: This is a 32-GPU system. The key constraint is the interconnect asymmetry: within each node, you have NVLink 4.0 (900 GB/s total); between nodes, you have InfiniBand HDR (200 Gb/s = 25 GB/s).

The standard answer for 32-GPU training of large models (say 13B-70B parameters) is hybrid parallelism: tensor parallelism within each 8-GPU NVLink domain, and data or pipeline parallelism across nodes over InfiniBand.

• Tensor parallelism degree = 8 (within one NVLink node): the frequent AllReduce operations required for TP stay on NVLink at 900 GB/s. TP across nodes over 25 GB/s InfiniBand would be catastrophically slow.

• Data parallelism degree = 4 (across 4 nodes): gradient AllReduce happens once per step over InfiniBand. At 25 GB/s, a 13B model's gradients (26 GB float16) take about 2 seconds over ring-AllReduce with 4 nodes. This is slow but can be partially hidden with gradient compression or async gradient updates.

• Pipeline parallelism is an alternative for the inter-node axis, especially for 70B+ models. Split the model into 4 stages (one per node), and pipeline micro-batches through the stages. This reduces inter-node communication from 26 GB (full gradients) to just activation tensors between adjacent pipeline stages.

This is essentially the Megatron-LM approach used by NVIDIA and many large-scale training teams.

---

Q5: How would you diagnose and fix a training pipeline where GPU utilization is 40% and you suspect PCIe is the bottleneck?

A: I would work through a systematic diagnosis:

Step 1: Confirm the bottleneck with profiling.

with torch.profiler.profile(
    activities=[ProfilerActivity.CPU, ProfilerActivity.CUDA]
) as prof:
    run_training_steps(model, loader, n_steps=20)

# Export to Chrome trace and look at timeline
prof.export_chrome_trace("trace.json")

In the Chrome trace (chrome://tracing), look for gaps between GPU kernels that correlate with aten::_to_copy on the CPU timeline. If you see the GPU sitting idle while CPU is active with copy operations, you are PCIe-bound.

Step 2: Measure actual PCIe bandwidth vs. theoretical. Use the bandwidth benchmark code above. If you are getting 17 GB/s on a PCIe 4.0 system, check for pageable memory issues.

Step 3: Apply fixes in order of impact:

1. pin_memory=True in DataLoader - typically 1.5-1.8x bandwidth improvement.
2. non_blocking=True on .to(device) calls - overlaps transfer with compute.
3. Increase DataLoader num_workers to 4-8 (but not so high that it creates CPU contention).
4. Implement prefetching / double buffering for the data pipeline.
5. Move preprocessing to GPU (using cuDF, DALI, or custom CUDA kernels for data augmentation) to reduce PCIe transfer volume by only transferring compressed/raw data.

Step 4: Verify fix impact.

After each change, re-run the profiler and benchmark. GPU utilization should climb toward 85-95%. If it remains below 70% even with pinned memory and non-blocking transfers, check whether the bottleneck has shifted to GPU compute (model is too small for the batch size) or CPU preprocessing (data augmentation is slow).

---

Q6: What happens during a NCCL AllReduce and why does the ring algorithm perform better than the naive approach?

A: A naive AllReduce implementation would have one GPU collect all other GPUs' tensors, sum them, and broadcast the result. This creates a hub-and-spoke bottleneck: the collecting GPU saturates its bandwidth first, and the operation time grows linearly with the number of GPUs.

Ring-AllReduce (the algorithm used by NCCL) distributes the work evenly:

1. Scatter-Reduce phase: Arrange GPUs in a ring (0 - 1 - 2 - 3 - 4 - 5 - 6 - 7 - 0). Divide each gradient tensor into N chunks (one per GPU). In N-1 steps, each GPU sends one chunk to the next GPU and receives one chunk from the previous GPU, accumulating partial sums. After N-1 steps, each GPU has the complete sum for one chunk of the gradients.

2. AllGather phase: In another N-1 steps, distribute the fully-summed chunks back to all GPUs.

Total communication per GPU: $2 \times \frac{N-1}{N} \times \text{gradient\_bytes}$

This approaches $2 \times \text{gradient\_bytes}$ as N grows. The key insight: the total bandwidth used scales with bandwidth-per-GPU, not total-GPUs. Adding more GPUs does not increase the communication bottleneck per GPU (it only slightly approaches the $2\times$ limit). This makes ring-AllReduce near-optimal in terms of bandwidth utilization.

On NVLink with NVSwitch, NCCL can do even better: the NVSwitch enables all GPUs to communicate simultaneously at full bandwidth, so the AllReduce can complete in roughly $\frac{\text{gradient\_bytes}}{\text{per-GPU NVLink bandwidth}}$ time rather than $\frac{2 \times \text{gradient\_bytes}}{\text{per-GPU NVLink bandwidth}}$.

---

Summary

PCIe and NVLink are not interchangeable - they solve different problems in AI infrastructure:

• PCIe connects CPUs to GPUs. It is the data loading highway. PCIe 4.0 x16 at 32 GB/s unidirectional is almost always the bottleneck in data-heavy pipelines. Pinned memory and non-blocking transfers are the two most important optimizations.

• NVLink connects GPUs to GPUs. It is the gradient synchronization highway. NVLink 4.0 at 900 GB/s is what makes tensor parallelism and fast distributed training possible. NVSwitch extends this to full bisection bandwidth across all GPUs in a DGX node.

• The 28x bandwidth gap between PCIe 4.0 (32 GB/s) and NVLink 4.0 (450 GB/s per direction per GPU) is the physical reason why parallelism strategies are constrained to specific topological domains.

• In production, always: enable pin_memory=True, use non_blocking=True transfers, implement double buffering for data loading, check P2P access before designing parallelism strategies, and profile before optimizing.

The engineer who internalizes these numbers and constraints will never design a system that stalls at 38% GPU utilization while sitting on a $500,000 cluster.