NCCL Deep Dive: From P2P/CUMEM to torchrun Init

Introduction

When you run torchrun --nproc_per_node=8 train.py and PyTorch magically synchronizes gradients across 8 GPUs, the heavy lifting is done by NCCL (NVIDIA Collective Communications Library). But what actually happens under the hood? How does NCCL discover which GPUs can talk to each other? What is P2P/CUMEM? And how does torchrun bootstrap all of this?

This post traces the entire path — from the hardware interconnects between GPUs, through NCCL’s topology detection and transport selection, to the torchrun launch sequence that wires everything together.

What Is NCCL

NCCL (pronounced “nickel”) is NVIDIA’s library for multi-GPU and multi-node collective communication. It provides operations like:

Operation	What It Does	When Used
AllReduce	Sum (or max/min) a tensor across all GPUs, result on all GPUs	Gradient synchronization in DDP
AllGather	Gather tensors from all GPUs, full result on all GPUs	FSDP parameter gathering
ReduceScatter	Reduce + scatter (each GPU gets a shard of the result)	FSDP gradient reduction
Broadcast	Send tensor from one GPU to all others	Parameter initialization
Send/Recv	Point-to-point transfer between two GPUs	Pipeline parallelism

NCCL is not a general-purpose communication library like MPI. It is purpose-built for GPU-to-GPU data movement, and it understands GPU hardware topology (NVLink, NVSwitch, PCIe, InfiniBand) to pick the fastest transport for each pair of GPUs.

Where NCCL Sits in the Stack

User Training Code (PyTorch)
       │
       ▼
torch.distributed  (Python API: init_process_group, all_reduce, ...)
       │
       ▼
ProcessGroupNCCL   (C++ backend in PyTorch)
       │
       ▼
NCCL Library       (libnccl.so — NVIDIA's collective comms)
       │
       ├──► NVLink / NVSwitch   (intra-node GPU-GPU)
       ├──► PCIe                (intra-node fallback)
       ├──► InfiniBand / RoCE   (inter-node)
       └──► TCP/IP sockets      (inter-node fallback)

PyTorch’s torch.distributed is the user-facing API. Under the hood, it delegates to ProcessGroupNCCL (a C++ class), which calls NCCL’s C API (ncclAllReduce, ncclCommInitRank, etc.). NCCL then decides how to move data based on the hardware topology it detects.

GPU Interconnects: The Physical Layer

Before understanding NCCL’s transport selection, we need to understand the physical connections between GPUs.

NVLink

NVLink is NVIDIA’s high-bandwidth, low-latency GPU-to-GPU interconnect. It bypasses the CPU and PCIe entirely.

Generation	Per-Link BW	Links per GPU	Total BW per GPU
NVLink 3 (A100)	50 GB/s	12	600 GB/s
NVLink 4 (H100/H200)	50 GB/s	18	900 GB/s
NVLink 5 (B200)	100 GB/s	18	1800 GB/s

For comparison, PCIe Gen5 x16 is ~64 GB/s bidirectional. NVLink is 10-30x faster.

NVSwitch

On systems with more than 2 GPUs, NVLink alone can’t fully connect all GPUs (not enough links). NVSwitch is a chip that acts as a crossbar switch, connecting all GPUs with full bandwidth:

8-GPU Node with NVSwitch (e.g., DGX H100):

  GPU0 ──┐                    ┌── GPU4
  GPU1 ──┤   ┌────────────┐   ├── GPU5
  GPU2 ──┼───┤  NVSwitch   ├───┼── GPU6
  GPU3 ──┘   └────────────┘   └── GPU7

  Any GPU can talk to any other GPU at full NVLink bandwidth.
  No hops, no congestion — it's a full crossbar.

PCIe

When NVLink is not available (e.g., consumer GPUs, or communication between GPUs on different PCIe switches), data goes through PCIe. This requires staging through CPU memory:

GPU 0 → PCIe → CPU Memory → PCIe → GPU 1
         ~32 GB/s           ~32 GB/s

This is 10-30x slower than NVLink and adds latency from the CPU memory hop.

InfiniBand / RoCE (Multi-Node)

For multi-node communication, NCCL uses RDMA (Remote Direct Memory Access) over InfiniBand or RoCE (RDMA over Converged Ethernet):

Node 0, GPU 3 → NVLink → GPU 7 → PCIe → HCA (IB NIC) ──► IB Switch
                                                              │
Node 1, GPU 3 ← NVLink ← GPU 7 ← PCIe ← HCA (IB NIC) ◄──┘

With GPUDirect RDMA (GDR), the IB NIC can read/write GPU memory directly, bypassing CPU memory entirely:

GPU Memory → PCIe → HCA → IB Fabric → HCA → PCIe → GPU Memory
  (no CPU memory copy!)

NCCL Transport Types

When you see NCCL INFO Channel 00/0 : 0[0] -> 1[1] via P2P/CUMEM in the logs, the last part tells you which transport NCCL chose. Here’s what each means:

P2P/CUMEM (CUDA Memory — the best intra-node transport)

GPU 0                       GPU 1
┌─────────┐   NVLink/NVSwitch   ┌─────────┐
│ CUDA Mem │ ◄═══════════════► │ CUDA Mem │
│  (HBM)   │    Direct access   │  (HBM)   │
└─────────┘                     └─────────┘

P2P/CUMEM means NCCL is using CUDA peer-to-peer memory access over NVLink. GPU 0 can directly read/write GPU 1’s memory without any CPU involvement. This is the fastest intra-node transport:

• Bandwidth: Full NVLink bandwidth (900 GB/s on H200)
• Latency: ~1-2 microseconds
• CPU involvement: None
• Requirement: GPUs must be P2P-capable (NVLink connected, or on same PCIe switch with P2P enabled)

Under the hood, NCCL calls cuMemCreate() / cuMemMap() (CUDA Virtual Memory Management API, hence “CUMEM”) to create memory mappings that span GPU boundaries. This is more efficient than the older cudaDeviceEnablePeerAccess() API because it works with CUDA’s unified virtual address space.

P2P/IPC (Inter-Process Communication)

Process A (GPU 0)              Process B (GPU 1)
┌─────────────┐                ┌─────────────┐
│ CUDA Memory  │◄──IPC Handle──►│ CUDA Memory  │
│ cuMemCreate  │   (fd passing) │ cuMemMap     │
└─────────────┘                └─────────────┘

Similar to P2P/CUMEM but uses CUDA IPC handles to share memory between different OS processes. Still goes over NVLink, still fast, but has slightly more setup overhead due to the IPC handle exchange.

SHM (Shared Memory)

Process A (GPU 0)              Process B (GPU 1)
┌─────────┐                    ┌─────────┐
│ GPU Mem  │                    │ GPU Mem  │
└────┬─────┘                    └────┬─────┘
     │ cudaMemcpy D→H                │ cudaMemcpy H→D
     ▼                               ▲
┌────────────────────────────────────────┐
│          /dev/shm  (shared memory)      │
│          (CPU DRAM, mmap'd)             │
└─────────────────────────────────────────┘

When P2P is not available, NCCL falls back to staging through CPU shared memory (/dev/shm). This involves:

1. GPU 0 copies data to CPU memory (cudaMemcpyDeviceToHost)
2. Data sits in shared memory (visible to both processes)
3. GPU 1 copies data from CPU memory (cudaMemcpyHostToDevice)

This is much slower — bandwidth is limited by PCIe and CPU memory bandwidth, and latency is higher due to the two-hop path.

NET (Network — for multi-node)

Node 0, GPU 0                    Node 1, GPU 0
┌─────────┐                      ┌─────────┐
│ GPU Mem  │                      │ GPU Mem  │
└────┬─────┘                      └────┬─────┘
     │ GDR (optional)                  │ GDR (optional)
     ▼                                 ▲
┌─────────┐    IB / RoCE / TCP    ┌─────────┐
│   HCA    │ ◄═══════════════════► │   HCA    │
└─────────┘                       └─────────┘

For inter-node communication, NCCL uses a network plugin. Common options:

• IB (InfiniBand): NET/IB — RDMA, lowest latency
• Socket: NET/Socket — TCP/IP, highest compatibility
• Custom plugins (e.g., ClockworkIB as seen in our experiments)

With GPUDirect RDMA (GDR 1 in logs), the IB NIC reads directly from GPU memory over PCIe, skipping CPU memory entirely.

Transport Selection Summary

Transport	Path	Bandwidth	Latency	When Used
P2P/CUMEM	GPU↔GPU via NVLink (same process or CUMEM-capable)	900 GB/s (H200)	~1-2 us	NVLink + CUMEM support
P2P/IPC	GPU↔GPU via NVLink (cross-process, IPC handles)	900 GB/s	~2-5 us	NVLink + cross-process
SHM	GPU→CPU→GPU via /dev/shm	~30-50 GB/s	~10-50 us	No P2P, same node
NET/IB	GPU→HCA→IB→HCA→GPU	~25-50 GB/s per rail	~5-10 us	Cross-node, IB available
NET/Socket	GPU→CPU→TCP→CPU→GPU	~1-10 GB/s	~50-500 us	Cross-node, fallback

NCCL Initialization: From init_process_group to C++ Internals

When your training script calls torch.distributed.init_process_group(backend="nccl"), a complex initialization sequence unfolds. Let’s trace it from the Python API down to the NCCL C++ source code. All NCCL links point to NVIDIA/nccl, PyTorch links to pytorch/pytorch.

The entry point is ncclCommInitRankConfig() (called from PyTorch’s NCCLUtils.cpp:93), which dispatches to the async worker ncclCommInitRankFunc() at init.cc:1624. This function orchestrates all the phases below.

Step 1: Rendezvous (All Processes Find Each Other)

Before NCCL can do anything, all processes need to discover each other. This is called rendezvous. At the C++ level, bootstrapInit() (called from init.cc:1702) establishes TCP connections between all ranks using the ncclUniqueId (which encodes rank 0’s IP:port).

With init_method="env://" (the default for torchrun):

Environment variables set by torchrun:
  MASTER_ADDR = 10.0.0.1       (IP of rank 0)
  MASTER_PORT = 29500           (port for rendezvous)
  RANK        = 3               (this process's global rank)
  WORLD_SIZE  = 16              (total number of processes)
  LOCAL_RANK  = 3               (rank within this node)

PyTorch creates a TCPStore — a simple key-value store backed by a TCP server running on rank 0. All other ranks connect to MASTER_ADDR:MASTER_PORT and use this store to exchange bootstrap information.

Step 2: NCCL Communicator Creation (ncclCommInitRank)

PyTorch’s ProcessGroupNCCL calls NCCL’s ncclCommInitRank() (or the newer ncclCommInitRankConfig()). This is where the real work happens.

From the NCCL debug logs:

ncclCommInitRankConfig comm 0x61ee0b20e510 rank 0 nranks 2 cudaDev 0 nvmlDev 0
  busId 4000 commId 0xd05b1321b6675ed0 - Init START

Step 3: Topology Detection

NCCL reads the system’s GPU topology via ncclTopoGetSystem() (called from init.cc:1046 inside the massive initTransportsRank() function at init.cc:903):

ncclTopoGetSystem()                                    [topo.cc:1444]
├── ncclTopoFillGpu() for each GPU                     [xml.cc:883]
│   ├── ncclTopoGetXmlFromSys() — read sysfs           [xml.cc:574]
│   │   └── Reads: vendor, device, link_speed, link_width from /sys/
│   └── ncclTopoGetXmlFromGpu() — query NVML            [xml.cc:735]
│       ├── NVLink detection: ncclNvmlDeviceGetNvLinkCapability()  [xml.cc:767]
│       └── Max links by SM: SM60=4, SM70=6, SM80=12, SM90+=18
├── ncclTopoProcessNet() — detect NICs                  [topo.cc:1491]
└── bootstrapIntraNodeAllGather() + ncclTopoFuseXml()   [topo.cc:1542]
    └── Merge all local ranks' topology views

From the logs you see:

NCCL INFO MNNVL busId 0x4000 fabric UUID 0.0 cliqueId 0x0 state 3
NCCL INFO ncclTopoGetCpuAffinity: Affinity for GPU 0 is 0-63,128-191
NCCL INFO NVLS multicast support is available on dev 0 (NVLS_NCHANNELS 16)

This tells us:

• MNNVL (Multi-Node NVLink): checked but not active (fabric UUID 0.0)
• CPU affinity: GPU 0 is closest to CPU cores 0-63 and 128-191 (NUMA node 0)
• NVLS (NVLink SHARP): multicast support available — enables hardware-accelerated all-reduce

Step 4: P2P Capability Detection

NCCL checks whether each pair of GPUs can do peer-to-peer memory access via ncclTopoCheckP2p() (called from paths.cc:649 — ncclTopoComputePaths()):

// p2p.cc:103-122 — convert bus ID to CUDA device index
static int busIdToCudaDev(int64_t busId) {
    int ndev;
    cudaGetDeviceCount(&ndev);   // ← CUDA_VISIBLE_DEVICES affects this!
    for (int i = 0; i < ndev; i++) { ... }
}

// p2p.cc:171 — the actual P2P capability query
cudaDeviceCanAccessPeer(&p2p, cudaDev1, cudaDev2);

From the logs:

NCCL INFO Check P2P Type isAllDirectP2p 1 directMode 0 isAllCudaP2p 1

• isAllDirectP2p 1: All GPU pairs have direct P2P access (NVLink)
• isAllCudaP2p 1: All GPU pairs support CUDA P2P

This is the critical check. If CUDA_VISIBLE_DEVICES was restricted when this runs, NCCL may not detect all GPU pairs, and may fall back to SHM.

Step 5: Channel and Ring Construction

NCCL organizes GPUs into channels — independent communication paths that can run in parallel. Each channel has a ring (for ring-based all-reduce) and a tree (for tree-based all-reduce):

Ring (2 GPUs, 1 channel):     Tree (2 GPUs):
  GPU 0 → GPU 1 → GPU 0        GPU 0
                                  ↑
                                GPU 1

From the logs:

NCCL INFO 24 coll channels, 24 collnet channels, 16 nvls channels,
         32 p2p channels, 32 p2p channels per peer

More channels = more parallelism = higher bandwidth utilization. H200 with NVLink gets 24 collective channels because it has enough NVLink bandwidth to feed them all.

NCCL INFO Channel 00/24 : 0 1
NCCL INFO Channel 01/24 : 0 1
...
NCCL INFO Channel 23/24 : 0 1

Each channel is a ring: 0 → 1 → 0 (with 2 GPUs, the ring is trivial).

Step 6: Transport Selection Per Channel

For each channel, NCCL selects the transport for each link:

NCCL INFO Channel 00/0 : 0[0] -> 1[1] via P2P/CUMEM

Breaking this down:

• Channel 00/0: Channel 0, connection index 0
• 0[0]: Rank 0, GPU index 0
• 1[1]: Rank 1, GPU index 1
• via P2P/CUMEM: Using peer-to-peer with CUDA memory mapping (NVLink)

If P2P were unavailable, you’d see:

NCCL INFO Channel 00/0 : 0[0] -> 1[1] via SHM

Step 7: Proxy Threads and Completion

NCCL spawns proxy threads for network and IPC communication:

NCCL INFO [Proxy Service] Device 0 CPU core 189
NCCL INFO [Proxy Service UDS] Device 0 CPU core 29

Finally:

ncclCommInitRankConfig comm 0x61ee0b20e510 rank 0 nranks 2 cudaDev 0 nvmlDev 0
  busId 4000 commId 0xd05b1321b6675ed0 - Init COMPLETE
Init timings: total 1.90 (kernels 0.26, alloc 1.39, bootstrap 0.14,
  topo 0.02, graphs 0.00, connections 0.09)

The 1.9 seconds is dominated by memory allocation (1.39s) and kernel compilation (0.26s). Topology detection is fast (0.02s).

How torchrun Bootstraps NCCL

Now let’s trace exactly what happens when you run:

torchrun --nnodes=1 --nproc_per_node=8 --rdzv_backend=c10d \
         --rdzv_endpoint=localhost:29400 train.py

Phase 1: Process Launch

# torch/distributed/run.py → torch/distributed/launcher/api.py
# torchrun is a thin wrapper around torch.distributed.elastic.multiprocessing

class LocalElasticAgent:
    def _start_workers(self, worker_group):
        # For each local rank 0..7:
        for local_rank in range(nproc_per_node):
            env = {
                "LOCAL_RANK":      str(local_rank),
                "RANK":            str(node_rank * nproc_per_node + local_rank),
                "LOCAL_WORLD_SIZE": str(nproc_per_node),
                "WORLD_SIZE":      str(nnodes * nproc_per_node),
                "MASTER_ADDR":     master_addr,
                "MASTER_PORT":     str(master_port),
            }
            # Fork a new process with these env vars
            process = subprocess.Popen(["python", "train.py"], env=env)

Key point: torchrun does NOT set CUDA_VISIBLE_DEVICES. Each process inherits the parent’s full GPU visibility. The process uses LOCAL_RANK to select its GPU via torch.cuda.set_device(local_rank).

Phase 2: User Script Runs

# train.py
def main():
    dist.init_process_group(backend="nccl")  # ← triggers NCCL init
    local_rank = int(os.environ["LOCAL_RANK"])
    torch.cuda.set_device(local_rank)        # ← select GPU, NOT restrict visibility
    
    model = MyModel().to(local_rank)
    ddp_model = DDP(model, device_ids=[local_rank])
    # Training loop...

Phase 3: init_process_group Internals

# torch/distributed/distributed_c10d.py

def init_process_group(backend="nccl", init_method=None, ...):
    # 1. Create a Store for rendezvous
    #    With env://, creates TCPStore connecting to MASTER_ADDR:MASTER_PORT
    store = _rendezvous_handler(init_method, rank, world_size)
    
    # 2. Create the NCCL process group
    #    This calls ncclCommInitRank() internally
    pg = ProcessGroupNCCL(store, rank, world_size, timeout)
    
    # 3. Register as the default process group
    _default_pg = pg

Inside ProcessGroupNCCL (C++), the NCCL communicator is created lazily — on the first collective operation, not during init_process_group(). When the first all_reduce is called:

// torch/csrc/distributed/c10d/ProcessGroupNCCL.cpp

ncclComm_t comm;
// All ranks exchange a unique ID via the Store
ncclUniqueId ncclId;
if (rank == 0) {
    ncclGetUniqueId(&ncclId);
    store.set("ncclUniqueId", ncclId);
} else {
    ncclId = store.get("ncclUniqueId");
}
// Every rank calls this with the same ncclId
ncclCommInitRank(&comm, worldSize, ncclId, rank);
// → This triggers the full topology detection + transport selection

Phase 4: The AllReduce Hot Path

Once initialized, every loss.backward() in DDP triggers gradient synchronization. But AllReduce is not a single round-trip — it runs in 2 * (N-1) steps for N GPUs, across multiple channels in parallel.

Ring AllReduce: The Multi-Round Algorithm

The implementation is in src/device/all_reduce.h. For N GPUs on a ring, the algorithm has two phases:

Phase A: Reduce-Scatter (N-1 steps)

Each GPU holds a chunk of the data. In each step, a GPU sends its chunk to the next GPU in the ring, which reduces (sums) it with its own chunk. After N-1 steps, each GPU holds the fully reduced result for one chunk.

Example: 4 GPUs, data split into 4 chunks [A, B, C, D]

Step 0: Each GPU sends its "own" chunk to the next GPU
  GPU 0 sends chunk D → GPU 1        (directSend)

Step 1-2: Receive, reduce with local, send to next
  GPU 1 receives chunk D, reduces with local D, sends → GPU 2  (directRecvReduceDirectSend)
  ... repeat for N-2 = 2 steps ...

Step 3: Final reduce step
  GPU 0 receives fully-reduced chunk A  (directRecvReduceCopyDirectSend)

From the source code (all_reduce.h:43-65):

// step 0: push data to next GPU
prims.directSend(offset, offset, nelem);

// k-2 steps: reduce and copy to next GPU
for (int j = 2; j < nranks; ++j) {
    prims.directRecvReduceDirectSend(offset, offset, nelem);
}

// step k-1: reduce this buffer and data, which will produce the final result
prims.directRecvReduceCopyDirectSend(offset, offset, nelem, /*postOp=*/true);

Phase B: AllGather (N-1 steps)

Now each GPU has one fully-reduced chunk. In N-1 more steps, each GPU sends its chunk around the ring so every GPU ends up with all chunks.

// k-2 steps: copy to next GPU
for (int j = 1; j < nranks - 1; ++j) {
    prims.directRecvCopyDirectSend(offset, offset, nelem);
}

// Make final copy from buffer to dest
prims.directRecv(offset, nelem);

Total: 2 * (N-1) steps where N = number of GPUs. For 2 GPUs: 2 steps. For 8 GPUs: 14 steps.

Each step involves a GPU-to-GPU data transfer over NVLink. With P2P/CUMEM, these transfers are direct memory reads/writes — no CPU involvement. But at each step boundary, the sending GPU’s NCCL kernel must signal completion to the receiving GPU. This is where CPU proxy thread latency matters — if the proxy thread is delayed by even 100 microseconds at any step, the entire ring stalls for that step.

Tree AllReduce (Alternative)

NCCL also implements tree-based AllReduce (all_reduce.h:87-147) with two phases:

1. Reduce Up: Leaf GPUs send to parents, parents reduce and send up. Root gets the final sum.
2. Broadcast Down: Root sends result to children, children forward to their children.

Tree has O(log N) latency but lower bandwidth utilization than ring. NCCL automatically selects ring vs tree based on message size — ring for large messages (bandwidth-bound), tree for small messages (latency-bound).

Multi-Channel Parallelism

The data is split across multiple channels (24 in our H200 tests). Each channel runs an independent ring/tree. This is why NCCL achieves high throughput — 24 rings running in parallel, each using a subset of the NVLink bandwidth:

Channel 0: GPU 0 →chunk0→ GPU 1 →chunk0→ GPU 0  (ring step)
Channel 1: GPU 0 →chunk1→ GPU 1 →chunk1→ GPU 0  (ring step, in parallel)
...
Channel 23: GPU 0 →chunk23→ GPU 1 →chunk23→ GPU 0

The overall call chain:

DDP.backward()
  → Gradient ready for bucket N
    → ProcessGroupNCCL::allreduce(bucket_tensor)
      → ncclAllReduce(sendbuff, recvbuff, count, datatype, op, comm, stream)
        → NCCL kernel launch on CUDA stream
          → For each channel (24 in parallel):
            → Ring: 2*(N-1) steps of GPU→GPU transfer over NVLink
              → Each step: directSend / directRecvReduceDirectSend / directRecv
              → P2P/CUMEM: cuMemMap'd peer memory, zero CPU copy

How NCCL Selects the Algorithm and Protocol

NCCL doesn’t always use the same algorithm or protocol. For every collective, it evaluates a cost table of all (algorithm, protocol) combinations and picks the cheapest.

The selection happens in src/enqueue.cc:1934-1955 — topoGetAlgoInfo():

// enqueue.cc:1940-1955 — pick minimum-time (algorithm, protocol) pair
float minTime = FLT_MAX;
for (int a = 0; a < NCCL_NUM_ALGORITHMS; a++) {    // Ring, Tree, NVLS, CollNet, ...
    for (int p = 0; p < NCCL_NUM_PROTOCOLS; p++) {  // Simple, LL, LL128
        if (table[a][p] >= 0.0 && table[a][p] < minTime) {
            algorithm = a;
            protocol = p;
            minTime = table[a][p];
        }
    }
}
info->algorithm = algorithm;
info->protocol = protocol;

The cost table is populated by src/graph/tuning.cc:587-605 — ncclTopoGetAlgoTime():

// tuning.cc:587 — estimated time = latency * count + bytes / bandwidth
float bw = comm->bandwidths[coll][algorithm][protocol];
float lat = comm->latencies[coll][algorithm][protocol];
*time = lat * latCount + nBytes / (1000 * bw);

The bandwidth and latency tables are computed during initialization from the detected topology. The key question is: why does Simple win for large messages and LL for small?

Why LL Has a Bandwidth Ceiling

The LL protocol embeds a 4-byte flag into every 4-byte data word. Each 8 bytes on the NVLink wire carries only 4 bytes of useful data — 50% efficiency. This is how it avoids CPU involvement: the GPU receiver spin-polls these flags to detect new data (prims_ll.h:37-44).

This overhead is hardcoded in tuning.cc:305:

if (a == NCCL_ALGO_RING && p == NCCL_PROTO_LL) {
    busBw = std::min(llMaxBw, busBw * .5);  // 50% wire efficiency + hard cap
}

The llMaxBw caps are per-GPU-generation in tuning.cc:170-175:

.llMaxBws = {
    {39.0,  39.0,  20.4},   // Volta  (1-node / 2-node / 4-node)
    {87.7,  22.5,  19.0},   // Ampere
    {141.0, 45.0,  35.0},   // Hopper (H100/H200)
    {282.0, 90.0,  70.0},   // Blackwell
},

Simple protocol has no flag overhead — full NVLink bandwidth. The CPU proxy cost per step is small relative to the large data transfer.

Crossover Example: H200 NVLink, 24 Channels

	LL	Simple
Max bandwidth	141 GB/s (capped)	~900 GB/s (full NVLink)
Per-step HW latency	0.6 us (tuning.cc:152)	3.4 us
Base latency	14.0 us	8.4 us
CPU proxy on data path?	No	Yes

For a 100 KB gradient bucket (small DDP bucket):

• LL: 100 KB / 141 GB/s + 14.0 us = ~15 us (latency-dominated, fast)
• Simple: 100 KB / 900 GB/s + 8.4 us + proxy overhead = ~20+ us (proxy overhead dominates)
• LL wins — lower per-step latency, CPU-free data path

For a 500 MB gradient bucket (large-model LLM training):

• LL: 500 MB / 141 GB/s = 3.5 ms (bandwidth-capped, flag overhead kills throughput)
• Simple: 500 MB / 900 GB/s = 0.55 ms (6x faster, proxy cost is negligible)
• Simple wins — 6x more bandwidth, proxy overhead is <1% of transfer time

This is why large-model LLM training uses Simple protocol where the CPU proxy IS on the critical path, while small-model benchmarks use LL where it isn’t. Any CPU contention (background threads, OS scheduler jitter) would be more pronounced with Simple protocol in large-model training.

What Runs on CPU vs GPU During AllReduce

This is a critical distinction. It depends on which protocol NCCL selected above:

LL (Low Latency) Protocol — selected for smaller messages (our benchmark uses this: AllReduce_Sum_f32_RING_LL):

Component	Runs on	What it does
ncclAllReduce() call	CPU	Enqueues NCCL kernel onto CUDA stream. Returns immediately.
ncclDevKernel_AllReduce_Sum_f32_RING_LL	GPU	The entire ring algorithm. Each step writes data to peer GPU’s LL FIFO buffer via P2P/CUMEM.
Flag-based signaling	GPU	Sender writes a 4-byte flag (NCCL_LL_FLAG) into the LL FIFO line. Receiver spin-polls this flag from GPU threads (prims_ll.h:57-60). No CPU involved.
Proxy thread	CPU	NOT on the data path for LL. The proxy’s p2pSendProxyProgress (p2p.cc:817) explicitly skips non-Simple protocols: if (p != NCCL_PROTO_SIMPLE) { ... continue; }

For LL protocol with P2P/CUMEM, the data path is entirely on the GPU:

GPU kernel thread (sender)                      GPU kernel thread (receiver)
  │                                               │
  ├─ Write data to peer's LL FIFO buffer          │
  │  (via cuMemMap'd NVLink P2P address)          │
  ├─ Write LL flag to same FIFO line              │
  │                                               ├─ Spin-poll flag from GPU memory
  │                                               ├─ Flag detected → read data
  │                                               └─ Reduce with local data, advance to next step
  │                                               
  └─ (no CPU involvement at any step)

Simple Protocol — used for larger messages:

Component	Runs on	What it does
Data transfer	GPU	Kernel copies data to/from FIFO buffers
Completion tracking	CPU (proxy)	Proxy thread polls cudaEventQuery() to check if GPU transfers completed, then advances the FIFO tail pointer (p2p.cc:822-847)

For Simple protocol, the CPU proxy IS on the critical path — it must call cudaEventQuery() between each step to advance the state machine.

So which protocol was our benchmark using?

From our benchmark trace: ncclDevKernel_AllReduce_Sum_f32_RING_LL — the LL protocol. The proxy thread was NOT on the data path.

What Still Runs on CPU

Even with LL protocol running entirely on GPU, the CPU is still involved at the boundaries of each collective:

1. ncclAllReduce() enqueue (CPU): PyTorch calls ncclAllReduce() from ProcessGroupNCCL::collective(), which enqueues the NCCL kernel onto a CUDA stream. This is a CPU call that returns quickly, but any CPU preemption here delays the kernel launch.

2. CUDA stream launch latency (CPU → GPU): The CUDA driver must schedule the kernel onto the GPU. CPU preemption between the API call and driver processing adds latency.

3. DDP bucket readiness (CPU): The Reducer::autograd_hook → mark_bucket_ready → all_reduce_bucket chain runs on CPU. This determines when the AllReduce is enqueued.

4. Stream synchronization (CPU): ProcessGroupNCCL::collective() calls syncStream() to ensure input tensors are ready before launching NCCL. This involves CPU-side CUDA event operations.

In summary:

• During AllReduce (LL protocol): Pure GPU, no CPU. Proxy thread is idle.
• Before/after AllReduce: CPU is needed to enqueue the kernel, manage streams, and run DDP hooks.

Reading NCCL Debug Logs

Set NCCL_DEBUG=INFO to see everything. Here’s how to read the key lines:

# Topology detection
NCCL INFO comm 0x... rank 0 nRanks 8 nNodes 1 localRanks 8 localRank 0 MNNVL 0
#                     │        │       │          │            │           │
#                     │        │       │          │            │           └ Multi-Node NVLink?
#                     │        │       │          │            └ my local rank
#                     │        │       │          └ workers on this node
#                     │        │       └ number of nodes
#                     │        └ total workers
#                     └ my global rank

# Transport per channel
NCCL INFO Channel 00/0 : 0[0] -> 1[1] via P2P/CUMEM
#         │          │   │ │     │ │       │
#         │          │   │ │     │ │       └ transport type
#         │          │   │ │     │ └ GPU index
#         │          │   │ │     └ rank
#         │          │   │ └ GPU index
#         │          │   └ rank
#         │          └ connection index
#         └ channel number

# P2P capability check
NCCL INFO Check P2P Type isAllDirectP2p 1 directMode 0 isAllCudaP2p 1
#                        │                │              │
#                        │                │              └ all pairs support CUDA P2P?
#                        │                └ using direct mode (single-GPU optimization)?
#                        └ all pairs have direct P2P (NVLink)?

# Init timing breakdown
NCCL INFO Init timings: total 1.90 (kernels 0.26, alloc 1.39, bootstrap 0.14,
                                     topo 0.02, graphs 0.00, connections 0.09)
#                                    │            │           │
#                                    │            │           └ setting up transports
#                                    │            └ rendezvous time
#                                    └ topology detection (fast!)

Practical: What We Observed on H200

We ran the following DDP benchmark on an 8x NVIDIA H200 (143 GB HBM3e) pod with NVSwitch, using NCCL 2.28.9 and CUDA 13.0:

# ddp_example.py — run with: NCCL_DEBUG=INFO torchrun --nproc_per_node=2 ... ddp_example.py
import torch, torch.distributed as dist, torch.nn as nn
from torch.nn.parallel import DistributedDataParallel as DDP
from torch.profiler import profile
import torch.optim as optim

SIZE = 4000

class ToyModel(nn.Module):
    def __init__(self):
        super().__init__()
        self.net1 = nn.Linear(SIZE, SIZE)
        self.relu = nn.ReLU()
        self.net2 = nn.Linear(SIZE, SIZE)
        self.net3 = nn.Linear(SIZE, SIZE)

    def forward(self, x):
        return self.net3(self.relu(self.net2(self.relu(self.net1(x)))))

def demo_basic():
    dist.init_process_group("nccl")
    rank = dist.get_rank()
    model = ToyModel().to(rank)
    ddp_model = DDP(model, bucket_cap_mb=25, device_ids=[rank])
    loss_fn = nn.MSELoss()
    optimizer = optim.SGD(ddp_model.parameters(), lr=0.001)

    with profile(activities=[torch.profiler.ProfilerActivity.CPU,
                             torch.profiler.ProfilerActivity.CUDA]) as prof:
        for i in range(10):
            optimizer.zero_grad()
            outputs = ddp_model(torch.randn(1000, SIZE, device=rank))
            labels = torch.randn(1000, SIZE, device=rank)
            loss_fn(outputs, labels).backward()
            optimizer.step()
    if rank == 0:
        prof.export_chrome_trace("trace_ddp_example.json")

if __name__ == "__main__":
    demo_basic()

Running with NCCL_DEBUG=INFO torchrun --nproc_per_node=2, the NCCL logs showed:

Check P2P Type isAllDirectP2p 1 directMode 0 isAllCudaP2p 1
Channel 00/0 : 0[0] -> 1[1] via P2P/CUMEM
24 coll channels, 16 nvls channels, 32 p2p channels
Connected all rings, use ring PXN 0 GDR 1

Key results:

• P2P/CUMEM on all 24 channels (NVLink direct memory access)
• GDR 1 (GPUDirect RDMA ready for multi-node)
• NVLS (NVLink SHARP for hardware all-reduce)
• 24 collective channels (maximum parallelism)

8-GPU Large Model: Still LL, Still P2P/CUMEM

We scaled the test to 8 GPUs with a 1.2 GB model (SIZE=10000, three 10000x10000 linear layers, 100 MB DDP buckets) and 20 timed steps after warmup:

[torchrun 8GPU] model params=1200 MB, bucket=100MB
  avg 5.96 ms/step, NCCL kernel: AllReduce_Sum_f32_RING_LL

Even with 100 MB DDP buckets and 1.2 GB of parameters, NCCL still chose LL protocol on H200 NVSwitch. We even tried forcing NCCL_PROTO=^LL,LL128 (which should disable LL) — NCCL still used LL. On H200 with P2P/CUMEM intra-node, the LL bandwidth cap (141 GB/s for Hopper) is high enough that LL remains optimal even for large messages.

This means:

1. On modern NVSwitch systems, the CPU proxy is never on the intra-node collective data path
2. The LL protocol’s GPU-only flag-based signaling handles all ring steps without CPU involvement
3. Simple protocol (with CPU proxy) only activates for multi-node NET transport or with the explicit NCCL_P2P_USE_CUDA_MEMCPY=1 flag

Key Takeaways

1. NCCL is not magic — it’s a library that queries hardware topology (NVML, sysfs, CUDA P2P) and selects the fastest transport for each GPU pair.

2. P2P/CUMEM is the gold standard for intra-node communication. It means NCCL is using NVLink with CUDA virtual memory mapping — zero CPU involvement, maximum bandwidth.

3. SHM means trouble — if you see via SHM in NCCL logs for GPUs that should be NVLink-connected, something went wrong with P2P detection. Check CUDA_VISIBLE_DEVICES and nvidia-smi topo -m.

4. CUDA_VISIBLE_DEVICES affects topology detection — NCCL can only detect P2P between GPUs that are visible. If a process only sees 1 GPU, NCCL builds a 1-GPU topology graph, even if more GPUs are available.

5. torchrun keeps full GPU visibility — it does NOT restrict CUDA_VISIBLE_DEVICES. Each process sees all GPUs and uses torch.cuda.set_device(local_rank) to select one. This is why NCCL always gets the full topology.

6. Channels = parallelism — more channels means NCCL can overlap more data transfers. The number of channels depends on available bandwidth (NVLink links) and the number of GPUs.

7. Check your logs — NCCL_DEBUG=INFO is your best friend. Look for P2P/CUMEM (good), SHM (bad for NVLink systems), and GDR 1 (GPUDirect RDMA enabled for multi-node).

Useful NCCL Environment Variables

Variable	Purpose	Example
NCCL_DEBUG=INFO	Enable detailed logging	See topology, transports, timing
NCCL_DEBUG_SUBSYS=INIT,NET	Filter debug to specific subsystems	Reduce noise
NCCL_P2P_LEVEL=NVL	Force NVLink P2P even if auto-detect fails	Recover from stale topology
NCCL_P2P_DISABLE=1	Disable P2P (force SHM)	Debugging only
NCCL_SHM_DISABLE=1	Disable shared memory transport	Force NET even intra-node
NCCL_SOCKET_IFNAME=eth0	Bind NCCL to specific network interface	Avoid wrong NIC for multi-node
NCCL_IB_DISABLE=1	Disable InfiniBand, fall back to sockets	Debugging only
NCCL_NET_GDR_LEVEL=5	GPUDirect RDMA aggressiveness	Higher = more direct GPU↔NIC
NCCL_ASYNC_ERROR_HANDLING=1	Fail fast on timeout instead of hang	Production recommended
NCCL_ALGO=Ring or Tree	Force specific algorithm	Benchmarking

C++ Deep Dive: NCCL Transport Selection and P2P/CUMEM

The NCCL initialization code covered above (from ncclCommInitRank through topology detection) culminates in transport selection. Here we trace the specific code paths.

Transport Selection

src/transport.cc:23 — selectTransport():

// Transport priority order (tried in sequence):
struct ncclTransport* ncclTransports[] = {
    &p2pTransport,      // 1. P2P (NVLink direct)
    &shmTransport,      // 2. SHM (shared memory via CPU)
    &netTransport,      // 3. NET (IB/RoCE/TCP)
    &collNetTransport,  // 4. CollNet (in-network compute)
    &profilerTransport  // 5. Profiler
};

// For each transport, call canConnect():
for (int t = 0; t < NTRANSPORTS; t++) {
    if (ncclTransports[t]->canConnect(send, recv, ...)) {
        return ncclTransports[t];  // use first match
    }
}

P2P/CUMEM vs P2P/IPC Decision

src/transport/p2p.cc:410-481 — Inside P2P transport setup:

// p2p.cc:410 — Send setup
if (ncclCuMemEnable()) {
    resources->type = P2P_CUMEM;   // Modern: cuMemCreate/cuMemMap
} else {
    resources->type = P2P_IPC;     // Legacy: cudaIpcGetMemHandle
}

// p2p.cc:476-479 — Recv setup, generates the log line you see:
if (ncclCuMemEnable()) {
    resources->type = P2P_CUMEM;
    INFO(NCCL_INIT, "Channel %02d : %d[%d] -> %d[%d] via P2P/CUMEM", ...);
}

CUMEM enablement is auto-detected in src/misc/cudawrap.cc:46:

int ncclCuMemEnable() {
    int param = ncclParamCuMemEnable();  // NCCL_CUMEM_ENABLE env var
    return param >= 0 ? param : (param == -2 && ncclCuMemSupported);
    // -2 = auto-detect (default), checks cuMemCreate support + CUDA 12+
}

The CUMEM path uses CUDA Virtual Memory Management:

// p2p.cc:219-237 — Allocate shareable buffer (sender side)
cuMemCreate(&handle, size, &prop, 0);
cuMemExportToShareableHandle(&shareableHandle, handle, handleType, 0);

// p2p.cc:285-300 — Import on receiver side
cuMemImportFromShareableHandle(&handle, shareableHandle, handleType);
cuMemAddressReserve(&dptr, size, 0, 0, 0);
cuMemMap(dptr, size, 0, handle, 0);        // ← maps peer GPU memory into local addr space
cuMemSetAccess(dptr, size, &accessDesc, 1); // ← grants access

---

C++ Deep Dive: PyTorch DDP → NCCL Call Chain

Now let’s trace how PyTorch’s DDP backward pass triggers NCCL. All links point to pytorch/pytorch on GitHub.

DDP Backward: Autograd Hook → AllReduce

The journey starts when loss.backward() computes gradients. DDP registers hooks on every parameter’s gradient accumulator.

1. Hook Registration (Reducer constructor)

torch/csrc/distributed/c10d/reducer.cpp:86-240:

Reducer::Reducer(...) {
    // Lines 188-207: Register post-accumulate hooks on each parameter
    auto grad_accumulator = param.grad_fn()->next_edge(0).function;
    hooks_.emplace_back(
        grad_accumulator->add_post_hook(std::make_unique<LambdaPostHook>([=](...) {
            this->autograd_hook(variable_index);  // ← fires after grad is computed
        }))
    );
}

2. Autograd Hook Fires → Mark Variable Ready

When autograd computes a gradient, the hook fires:

// reducer.cpp:647
void Reducer::autograd_hook(size_t variable_index) {
    mark_variable_ready(variable_index);
}

3. Bucket Fills → AllReduce Triggered

DDP groups parameters into buckets (default 25 MB). When all gradients in a bucket are ready:

// reducer.cpp:1025
void Reducer::mark_bucket_ready(size_t bucket_index) {
    // All parameters in this bucket have gradients
    all_reduce_bucket(bucket);
}

// reducer.cpp:952
void Reducer::all_reduce_bucket(Bucket& bucket) {
    GradBucket grad_bucket(bucket_index, tensor, ...);
    run_comm_hook(grad_bucket);
}

4. Communication Hook → ProcessGroupNCCL

torch/csrc/distributed/c10d/default_comm_hooks.cpp:11:

c10::intrusive_ptr<c10::ivalue::Future> AllReduceCommHook::runHook(GradBucket& bucket) {
    std::vector<at::Tensor> tensors = {bucket.getBufferRef()};
    tensors[0] /= state_->getSize();    // divide by world_size (average gradient)
    return state_->allreduce(tensors)->getFuture();  // ← calls ProcessGroupNCCL
}

5. ProcessGroupNCCL::allreduce → ncclAllReduce

torch/csrc/distributed/c10d/ProcessGroupNCCL.cpp:

// Line 4489 — entry from Python/comm hook
c10::intrusive_ptr<Work> ProcessGroupNCCL::allreduce(std::vector<at::Tensor>& tensors, ...) {
    return allreduce_impl(tensors, opts);
}

// Line 4461 — creates the NCCL call lambda
c10::intrusive_ptr<Work> ProcessGroupNCCL::allreduce_impl(std::vector<at::Tensor>& tensors, ...) {
    return collective(tensors, tensors, [&](at::Tensor& input, at::Tensor& output,
                                            ncclComm_t comm, at::cuda::CUDAStream& stream) {
        // Line 4475 — THE actual NCCL call
        return ncclAllReduce(input.data_ptr(), output.data_ptr(),
                             input.numel(), ncclDataType, ncclReduceOp,
                             comm, stream.stream());
    }, ...);
}

6. The collective() Method — Lazy Init + Stream Management

ProcessGroupNCCL.cpp:3644:

c10::intrusive_ptr<Work> ProcessGroupNCCL::collective(...) {
    // Line 3673: Get or create NCCL communicator (lazy init)
    auto ncclComm = getNCCLComm(key);
    if (ncclComm == nullptr) {
        ncclComm = initNCCLComm(key, device, opType);  // first-time creation
    }

    // Line 3700: Use a SEPARATE CUDA stream for NCCL (not the default stream)
    auto ncclStream = ncclStreams_.at(key);

    // Line 3703: Sync: make NCCL stream wait for input tensor to be ready
    syncStream(device, ncclEvents_[key], ncclStream);

    // Line 3765: Get raw NCCL handle
    ncclComm_t comm = ncclComm->getNcclComm();

    // Line 3785: Call the NCCL function (allreduce lambda from above)
    fn(input, output, comm, ncclStream);

    // Line 3798: Record completion event
    work->ncclEndEvent_->record(ncclStream);

    return work;  // caller awaits this via future
}

7. First-Time Communicator Creation

ProcessGroupNCCL.cpp:2962 — initNCCLComm():

// Line 2786 — Rank 0 generates ID, broadcasts via Store
void ProcessGroupNCCL::broadcastUniqueNCCLID(ncclUniqueId* ncclID, ...) {
    if (rank == 0) {
        ncclGetUniqueId(ncclID);
        store_->set(storeKey, vec);   // rank 0 publishes ID
    } else {
        store_->get(storeKey);        // all others block until rank 0 publishes
    }
}

Then in NCCLUtils.cpp:62:

// Line 93-96 — modern non-blocking init
ncclConfig_t config = NCCL_CONFIG_INITIALIZER;
config.blocking = 0;
ncclCommInitRankConfig(&ncclComm_, numRanks, commId, rank, &config);
// → This enters the NCCL init.cc code path traced above

Complete Call Chain Summary

loss.backward()
  ↓ autograd engine computes gradients
  ↓ post-accumulate hook fires

Reducer::autograd_hook(variable_index)                      [reducer.cpp:647]
  → mark_variable_ready() → mark_bucket_ready()            [reducer.cpp:1025]
    → all_reduce_bucket()                                   [reducer.cpp:952]
      → run_comm_hook() → AllReduceCommHook::runHook()      [default_comm_hooks.cpp:11]
        → ProcessGroupNCCL::allreduce()                     [ProcessGroupNCCL.cpp:4489]
          → allreduce_impl() with ncclAllReduce lambda      [ProcessGroupNCCL.cpp:4461]
            → collective()                                  [ProcessGroupNCCL.cpp:3644]
              ├─ getNCCLComm() or initNCCLComm()            [ProcessGroupNCCL.cpp:3673]
              │   └─ ncclCommInitRankConfig()                [NCCLUtils.cpp:93]
              │       └─ NCCL init.cc → topology → transport selection
              ├─ syncStream() — wait for input tensor       [ProcessGroupNCCL.cpp:3703]
              ├─ ncclAllReduce(data, comm, ncclStream)      [ProcessGroupNCCL.cpp:3785]
              │   └─ NCCL kernel on ncclStream → NVLink P2P/CUMEM
              └─ record ncclEndEvent                        [ProcessGroupNCCL.cpp:3798]

---

The Proxy Thread: When It Matters and When It Doesn’t

When Is the Proxy Thread Created?

The proxy thread is created during ncclCommInitRank, specifically at init.cc:1340 → proxy.cc:1871:

// init.cc:1340 — called during communicator initialization
NCCLCHECKGOTO(ncclProxyCreate(comm), ret, fail);

// proxy.cc:1898 — spawns the service thread
comm->proxyState->thread = std::thread(ncclProxyService, comm->proxyState);
ncclSetThreadName(comm->proxyState->thread, "NCCL Service %2d", comm->cudaDev);

// proxy.cc:1903 — spawns the UDS service thread
comm->proxyState->threadUDS = std::thread(ncclProxyServiceUDS, comm->proxyState);

A third thread (the progress thread) is created lazily on first use, at proxy.cc:1022:

static ncclResult_t ncclProxyProgressCreate(struct ncclProxyState* proxyState) {
    struct ncclProxyProgressState* state = &proxyState->progressState;
    if (!state->thread.joinable()) {
        state->thread = std::thread(ncclProxyProgress, proxyState);
    }
}

So every GPU gets 3 proxy threads: Service, UDS Service, and Progress. They are always created, regardless of transport.

Why the Proxy Is Never on the Data Path for Intra-Node P2P

This is the key finding from our 8-GPU benchmarks. The proxy’s proxyProgress function is NULL by default for P2P transport — it’s only set when an opt-in memcpy mode is explicitly enabled.

The P2P transport struct is defined in p2p.cc:1331-1336:

struct ncclTransport p2pTransport = {
  "P2P",
  p2pCanConnect,
  //       setup,         connect,        free,  proxySharedInit, proxySetup,         proxyConnect, proxyFree,        proxyProgress, ...
  { p2pSendSetup, p2pSendConnect, p2pSendFree, NULL, p2pSendProxySetup, NULL, p2pSendProxyFree, NULL, ... },
  //                                                                                               ^^^^
  //                                                                               proxyProgress = NULL by default!
  { p2pRecvSetup, p2pRecvConnect, p2pRecvFree, NULL, p2pRecvProxySetup, NULL, p2pRecvProxyFree, NULL, ... }
};

The proxyProgress field (position 8 in ncclTransportComm, defined at transport.h:120) is NULL. It only gets set in initCeOperation() at p2p.cc:1338-1348:

static void initCeOperation() {
    static int init = 0;
    if (!init) {
        useMemcpy = ncclParamP2pUseCudaMemcpy();  // NCCL_P2P_USE_CUDA_MEMCPY env var
        if (useMemcpy) {
            p2pTransport.send.proxyProgress = p2pSendProxyProgress;  // only set if memcpy mode!
        }
        init = 1;
    }
}

The useMemcpy flag is controlled by p2p.cc:121:

NCCL_PARAM(P2pUseCudaMemcpy, "P2P_USE_CUDA_MEMCPY", 0);  // default: 0 (disabled)

Default is 0. So proxyProgress stays NULL. When proxy.cc polls for work at proxy.cc:428:

args->progress = op->connection->tcomm->proxyProgress;  // NULL for P2P!

There’s nothing to call. The proxy thread idles.

At connect time (p2p.cc:541-542), the NULL propagates:

// We must assign the proxyConn's proxyProgress property for proper checking at enqueue-time
send->proxyConn.proxyProgress = p2pTransport.send.proxyProgress;  // NULL

This means for all intra-node P2P/CUMEM connections, the proxy never has work to do. The data path is entirely GPU-driven via LL flag polling in prims_ll.h.

When DOES the Proxy Matter?

The proxy proxyProgress is set to a real function for:

• NET transport (inter-node IB/RoCE): net.cc:519 — send->proxyConn.proxyProgress = sendProxyProgress
• CollNet transport: coll_net.cc:275 — send->proxyConn.proxyProgress = sendProxyProgress
• P2P with NCCL_P2P_USE_CUDA_MEMCPY=1: opt-in memcpy mode (not default)

So the CPU proxy is on the critical path for multi-node training (where NET transport is used) but never for intra-node P2P on modern NVSwitch systems.

Summary: CPU Proxy Is Irrelevant for Intra-Node Training

On modern NVSwitch systems (H200, H100, B200), all intra-node communication uses P2P/CUMEM with LL protocol. The proxy thread is created but has proxyProgress = NULL — it never executes data-path work. The entire AllReduce runs on GPU via flag-based signaling over NVLink.

The CPU proxy only becomes relevant for:

• Multi-node training over InfiniBand/RoCE (NET transport)
• CollNet (in-network compute with SmartNICs)
• The rarely-used NCCL_P2P_USE_CUDA_MEMCPY=1 mode

References

• NCCL Documentation
• NCCL Source Code (GitHub)
• PyTorch Distributed Overview
• PyTorch ProcessGroupNCCL Source
• PyTorch Reducer Source (DDP)
• torchrun (torch.distributed.elastic)
• NVLink & NVSwitch Architecture
• GPUDirect RDMA