Learn › GPU Operations & NCCL Troubleshooting › NCCL Fundamentals
NCCL_DEBUG - a hands-on Linux lab on a real virtual machine.
A training job finishes cleanly and runs slow. Switch on the NCCL debug log, find the transport selection line, reproduce the loss with NCCL_P2P_DISABLE=1, and read the wiring on two real RTX 4000 Ada cards.
It is 02:14 on a Tuesday in the us-sea region. Node gpu-node-07 carries two NVIDIA RTX 4000 Ada cards. It has run the same training job every night for three weeks.
Tonight the job finished. Exit code 0. No warnings, nothing in the kernel log, both cards healthy. It just took longer than it used to.
So the on-call engineer ran the standard bandwidth benchmark. Last week that benchmark reported 22.08 GB/s. Tonight it reports 17.95 GB/s.
Same box. Same command. Same two cards. No error either time.
Four gigabytes per second went missing, and every tool on the machine still says the box is fine. Work Order SEA-GPU-0729 is open. The only witness is a log that ships switched off.
This is a Practice Zone. The terminals below are a safe sandbox, and each one checks only the single command it is teaching. The real machine comes at the end, under Ready to practice.
A GPU is the card that does the training math. When a job uses two of them, the cards must constantly swap results with each other. That swapping is the job of NCCL, the NVIDIA Collective Communications Library.
A collective is one communication move that every card takes part in at the same time. The common one is all-reduce: every card hands over a pile of numbers, the piles are added together, and every card gets the total back.
Each participating GPU is a rank. This node has two ranks, rank 0 and rank 1.
Here is the part that matters tonight. NCCL does not have one way to move those numbers. It has several, and it calls each one a transport. At startup it inspects the machine, picks a transport, and starts moving data.
It does not ask you. It does not warn you when it picks a slow one. The job still finishes, still returns exit code 0, still reports no errors. The choice is written down in exactly one place: the debug log.
NVIDIA open-sourced the first version of NCCL in 2015. The name is pronounced "nickel".
That log has shipped with a typo in it for years. When NCCL cannot find the InfiniBand library, it prints Count not open libibverbs.so. Somebody typed Count instead of Could, it shipped, and it is still there. You will see it in a moment. It is not an error and it is not your fault.
>>> Standard error. That is why every command in this lesson carries 2>&1, which means "send standard error to the same place standard output is going". Leave it off and you capture the benchmark table while the entire debug log scrolls past and disappears.
You switch the log on with an environment variable: a setting written in front of a command, which applies to that one run and nothing else. Two of them do the work here.
NCCL_DEBUG=INFO turns the log on. Without it, NCCL says nothing at all.
NCCL_DEBUG_SUBSYS=INIT,GRAPH,TUNING aims it. NCCL sorts its messages into subsystems, and this picks three: how it started up (INIT), what map it built of the machine (GRAPH), and how it tuned itself (TUNING). Without the filter you also get every bootstrap and socket message, and the useful lines drown.
The benchmark is all_reduce_perf, the standard NCCL test program. -b 128M -e 128M runs one message size, 128 megabytes. -f 2 sets the size step factor. -g 2 says use two GPUs.
head -12 keeps the first twelve lines, because the full log runs to hundreds.
Before you run it, commit to an answer. Does turning the log on change the benchmark result?
prompt: ops@gpu-node-07:~$ answer: NCCL_DEBUG=INFO NCCL_DEBUG_SUBSYS=INIT,GRAPH,TUNING all_reduce_perf -b 128M -e 128M -f 2 -g 2 2>&1 | head -12 output:
#
localhost:24133:24133 [0] NCCL INFO ncclOsDlopen(libnccl-env.so) failed: libnccl-env.so: cannot open shared object file: No such file or directory localhost:24133:24133 [0] NCCL INFO ENV/Plugin: Could not find: libnccl-env.so localhost:24133:24133 [0] NCCL INFO cudaDriverVersion 13030 localhost:24133:24133 [0] NCCL INFO NCCL version 2.30.7+cuda13.3 localhost:24133:24133 [0] NCCL INFO NCCL git version HEAD a5187fa2e hint: Put both settings in front of the command, then all_reduce_perf with -b 128M -e 128M -f 2 -g 2, then 2>&1 piped into head -12.
The benchmark result did not change. The log is free.
Now learn to read a line, because you will read thousands of them. Take localhost:24133:24133 [0] NCCL INFO cudaDriverVersion 13030 apart:
localhost is the node name as NCCL sees it24133 is the process id24133 is the thread id inside that process[0] is which GPU this message is aboutNCCL INFO is the message levelThe numbers change on every run. The words do not. When you compare two logs, compare words.
Three facts already landed. The driver reports CUDA version 13030. The library is NCCL 2.30.7 built against CUDA 13.3. And both cards are present, an RTX 4000 Ada at PCI address 0000:00:02 and another at 0000:00:03.
That full log is several hundred lines for a single run. One line in it names the transport, and finding it is the whole skill.
grep prints only the lines that contain a pattern. Every transport line in NCCL contains the word via, so grep via finds them no matter which transport was chosen. That is the point: you are not looking for a transport you already guessed, you are asking the log which one it picked.
Swap head -12 for grep via and run the same benchmark again.
prompt: ops@gpu-node-07:~$ answer: NCCL_DEBUG=INFO NCCL_DEBUG_SUBSYS=INIT,GRAPH,TUNING all_reduce_perf -b 128M -e 128M -f 2 -g 2 2>&1 | grep via output: localhost:24133:24160 [1] NCCL INFO Channel 00/0 : 1[1] -> 0[0] via P2P/direct pointer localhost:24133:24161 [0] NCCL INFO Channel 00/0 : 0[0] -> 1[1] via P2P/direct pointer localhost:24133:24160 [1] NCCL INFO Channel 01/0 : 1[1] -> 0[0] via P2P/direct pointer localhost:24133:24161 [0] NCCL INFO Channel 01/0 : 0[0] -> 1[1] via P2P/direct pointer localhost:24133:24160 [1] NCCL INFO Channel 02/0 : 1[1] -> 0[0] via P2P/direct pointer localhost:24133:24161 [0] NCCL INFO Channel 02/0 : 0[0] -> 1[1] via P2P/direct pointer localhost:24133:24160 [1] NCCL INFO Channel 03/0 : 1[1] -> 0[0] via P2P/direct pointer localhost:24133:24161 [0] NCCL INFO Channel 03/0 : 0[0] -> 1[1] via P2P/direct pointer hint: Same command as before. Replace the head filter with grep and the pattern via.
Eight lines, and they are the same sentence eight times. Read one:
Channel 00/0 : 0[0] -> 1[1] via P2P/direct pointer
Channel 00/0 is the channel number. 0[0] -> 1[1] is the direction, rank 0 sending to rank 1. Four channels, two directions each, gives eight lines.
Then the payoff: via P2P/direct pointer. P2P means peer to peer. GPU 0 writes straight into GPU 1's memory across the PCIe link between them. Nothing is copied into host memory on the way. This is the fastest road this machine has, and this string is the proof that the data used it.
Learn this string by sight. On a healthy two-GPU node it is what you expect to see, and noticing that it is missing is the entire diagnostic.
You could wait for the next slow night to study a degraded run. Better: cause one on purpose, on a machine where nothing depends on the answer.
NCCL_P2P_DISABLE=1 tells NCCL it may not use peer to peer. The hardware is untouched. Only NCCL's permission changes.
This one setting is worth knowing for another reason. It is a real configuration that shows up in production, set by a container runtime, a scheduler default, or a colleague who copied it from a fix for a different bug two years ago.
>>> It runs to completion and says nothing. NCCL treats a missing transport as a routing question, not a failure: it falls back to the next road down and keeps going. The result validation still passes, the exit code is still 0, and the only visible difference is a number on a dashboard that nobody is watching at 02:14.
Run the same benchmark with peer to peer switched off, and widen the sweep while you are at it.
-b 8 -e 256M -f 2 runs every message size from 8 bytes to 256 megabytes, doubling each time. You want the whole bandwidth curve now, not one point on it.
Filter for via exactly as before. Same question, different machine state.
prompt: ops@gpu-node-07:~$ answer: NCCL_P2P_DISABLE=1 NCCL_DEBUG=INFO NCCL_DEBUG_SUBSYS=INIT,GRAPH,TUNING all_reduce_perf -b 8 -e 256M -f 2 -g 2 2>&1 | grep via output: localhost:24162:24189 [1] NCCL INFO Channel 00 : 1[1] -> 0[0] via SHM/direct localhost:24162:24190 [0] NCCL INFO Channel 00 : 0[0] -> 1[1] via SHM/direct localhost:24162:24189 [1] NCCL INFO Channel 01 : 1[1] -> 0[0] via SHM/direct localhost:24162:24190 [0] NCCL INFO Channel 01 : 0[0] -> 1[1] via SHM/direct localhost:24162:24189 [1] NCCL INFO Channel 02 : 1[1] -> 0[0] via SHM/direct localhost:24162:24190 [0] NCCL INFO Channel 02 : 0[0] -> 1[1] via SHM/direct localhost:24162:24189 [1] NCCL INFO Channel 03 : 1[1] -> 0[0] via SHM/direct localhost:24162:24190 [0] NCCL INFO Channel 03 : 0[0] -> 1[1] via SHM/direct hint: Put NCCL_P2P_DISABLE=1 in front of the settings you already use, sweep with -b 8 -e 256M -f 2 -g 2, and filter for via again.
Two things changed, and both are the diagnosis.
P2P/direct pointer became SHM/direct. SHM is shared memory. The data no longer crosses between the cards directly. GPU 0 copies into host memory, and GPU 1 copies out of it. Every byte now makes two trips instead of one.
Channel 00/0 became Channel 00. The trailing /0 is gone. That is a formatting difference, not a bug, and it matters to you because a grep pattern built around Channel 00/0 will silently match nothing on a degraded node. Match on via, not on the channel format.
Now the cost. Here are the last rows of both sweeps, healthy first:
# size count type redop root time algbw busbw #wrong time algbw busbw #wrong
# (B) (elements) (us) (GB/s) (GB/s) (us) (GB/s) (GB/s)
134217728 33554432 float sum -1 6139.29 21.86 21.86 0 6135.19 21.88 21.88 0
268435456 67108864 float sum -1 12157.1 22.08 22.08 0 12137.1 22.12 22.12 0
# Out of bounds values : 0 OK
# Avg bus bandwidth : 8.49873And the same sizes with peer to peer switched off:
134217728 33554432 float sum -1 7735.91 17.35 17.35 0 7737.74 17.35 17.35 0
268435456 67108864 float sum -1 14956.0 17.95 17.95 0 15043.9 17.84 17.84 0
# Out of bounds values : 0 OK
# Avg bus bandwidth : 6.57446busbw is bus bandwidth, the useful number: how fast the wire between the cards actually moved data. 22.08 GB/s becomes 17.95 GB/s at 256 megabytes. Both runs report Out of bounds values : 0 OK, which means both produced correct math.
That is the missing four gigabytes per second from Work Order SEA-GPU-0729, reproduced on demand, with the string that explains it.
Peer to peer and shared memory are the first two roads. There is a third, and a switch that closes the second one: NCCL_SHM_DISABLE=1.
Close both and NCCL still refuses to fail. It falls all the way back to the network stack and sends the data out through the network interface, to itself, on the same machine. That road is called NET/Socket.
Click each road below, then flip the switch and watch which one the data is allowed to take.
{ "height": 440, "caption": "One pair of GPUs, three possible transports. Click any road, then flip the switch. The third road stays dark until both faster roads are closed.", "nodes": [ { "id": "gpu0", "label": "GPU 0 (rank 0)", "kind": "host", "x": 0, "y": 200, "detail": "NVIDIA RTX 4000 Ada Generation, 20475 MiB, driver 610.43.02, PCI address 0000:00:02. In the log it is rank 0, written 0[0]." }, { "id": "p2p", "label": "P2P/direct pointer", "kind": "nic", "x": 420, "y": 40, "detail": "Peer to peer across PCIe. GPU 0 writes straight into GPU 1 memory, no host copy. Measured 22.08 GB/s bus bandwidth at 256 MB. Log line: Channel 00/0 : 0[0] -> 1[1] via P2P/direct pointer." }, { "id": "shm", "label": "SHM/direct", "kind": "sensor", "x": 420, "y": 200, "detail": "Shared memory through host RAM. Every byte is copied out and back in. Measured 17.95 GB/s at 256 MB, average 6.57446. Log line: Channel 00 : 0[0] -> 1[1] via SHM/direct." }, { "id": "net", "label": "NET/Socket/0", "kind": "net", "x": 420, "y": 360, "detail": "The network stack, out through eth0 and back to the same machine. Measured 2.73 GB/s at 64 MB, average 0.86666, and only 2 channels instead of 4. Log line: Channel 00/0 : 0[0] -> 1[1] [send] via NET/Socket/0." }, { "id": "gpu1", "label": "GPU 1 (rank 1)", "kind": "host", "x": 820, "y": 200, "detail": "The second RTX 4000 Ada, 20475 MiB, same driver 610.43.02, PCI address 0000:00:03. In the log it is rank 1, written 1[1]. No NVLink bridge joins these two cards: the link between them is PHB, PCIe through the host bridge." } ], "edges": [ { "from": "gpu0", "to": "p2p", "label": "default", "kind": "inband" }, { "from": "p2p", "to": "gpu1", "kind": "inband" }, { "from": "gpu0", "to": "shm", "label": "P2P closed", "kind": "mgmt" }, { "from": "shm", "to": "gpu1", "kind": "mgmt" }, { "from": "gpu0", "to": "net", "label": "both closed", "kind": "oob" }, { "from": "net", "to": "gpu1", "kind": "oob" } ], "toggle": { "label": "Peer to peer:", "on": "Allowed (default)", "off": "NCCL_P2P_DISABLE=1", "dimOn": ["shm", "net", "e:gpu0-shm", "e:shm-gpu1", "e:gpu0-net", "e:net-gpu1"], "dimOff": ["p2p", "net", "e:gpu0-p2p", "e:p2p-gpu1", "e:gpu0-net", "e:net-gpu1"] } }
Notice what never changes: the two GPUs. Same cards, same job, same correct answer at the end. Only the road changes, and only the log records which one was taken.
You can now do something most people who run these jobs cannot. You can name the road the data took, and prove it from the machine.
>>> Out through the network interface and back. NCCL picks NET/Socket/0 on eth0, and the job still completes with correct results. It also drops from 4 channels to 2 and collapses to 2.73 GB/s at 64 MB. A working job, at roughly one eighth of the bandwidth, with no error anywhere.
The GRAPH subsystem you switched on earlier prints something worth more than the transport line: the map NCCL built of this machine, before it moved any data.
These lines appear near the top of the same log:
localhost:24133:24154 [1] NCCL INFO === System : maxBw 24.0 totalBw 24.0 ===
localhost:24133:24154 [1] NCCL INFO CPU/0-ffffffffffffff (1/2/-1)
localhost:24133:24154 [1] NCCL INFO + PCI[5000.0] - NIC/0-0
localhost:24133:24154 [1] NCCL INFO + NET[1.2] - NET/0-0 (0/0/0/1.250000)
localhost:24133:24154 [1] NCCL INFO + PCI[24.0] - DEV/0-20 (10de27b210de181b)
localhost:24133:24154 [1] NCCL INFO + LOC[5000.0] - GPU/0-20
localhost:24133:24154 [1] NCCL INFO + PCI[24.0] - DEV/0-30 (10de27b210de181b)
localhost:24133:24154 [1] NCCL INFO + LOC[5000.0] - GPU/0-30
localhost:24133:24154 [1] NCCL INFO ==========================================
localhost:24133:24154 [1] NCCL INFO GPU/0-20 :GPU/0-20 (0/5000.0/LOC) GPU/0-30 (4/24.0/PHB) CPU/0-ffffffffffffff (2/24.0/PHB) NET/0-0 (4/1.2/PHB) DEV/0-20 (1/5000.0/LOC) DEV/0-30 (3/24.0/PHB)Read the last line as a row of distances from GPU 0. To itself it is LOC, local, effectively free. To the other GPU it is (4/24.0/PHB): NCCL estimates 24.0 GB/s over a link type called PHB. To the network it is 1.2, twenty times worse.
PHB means PCIe Host Bridge. The two cards reach each other through the PCIe complex on the CPU. NCCL predicted 24.0 GB/s across it. You measured 22.08. The model was right.
One more line from the same log settles the hardware question: NVLS multicast support is not available on dev 1.
Ask the driver directly and it agrees.
prompt: ops@gpu-node-07:~$ answer: nvidia-smi topo -m output: GPU0 GPU1 CPU Affinity NUMA Affinity GPU NUMA ID GPU0 X PHB 0-7 0 N/A GPU1 PHB X 0-7 0 N/A
Legend:
X = Self SYS = Connection traversing PCIe as well as the SMP interconnect between NUMA nodes (e.g., QPI/UPI) NODE = Connection traversing PCIe as well as the interconnect between PCIe Host Bridges within a NUMA node PHB = Connection traversing PCIe as well as a PCIe Host Bridge (typically the CPU) PXB = Connection traversing multiple PCIe bridges (without traversing the PCIe Host Bridge) PIX = Connection traversing at most a single PCIe bridge NV# = Connection traversing a bonded set of # NVLinks hint: The nvidia-smi subcommand for wiring is topo, and the matrix view is -m.
The matrix says PHB where the two cards meet, and the legend spells out what that means. It does not say NV1 or NV2 or any NV#, so there is no NVLink on this machine. Peer to peer here runs over PCIe, not over an NVLink bridge.
That is worth stating plainly, because most NCCL writing assumes NVLink. The same tool also reports a peer-to-peer capability matrix, and on this node it reads:
GPU0 GPU1
GPU0 X OK
GPU1 OK XTwo answers that sound contradictory and are both true. The wiring is PHB, the slowest of the on-board link types. Peer to peer is still OK across it. A budget two-GPU node has no NVLink and still gets direct card-to-card transfers, which is exactly why losing them costs you 4 GB/s instead of nothing.
One more group of lines, and then the log has no secrets left.
NCCL does not send one big stream. It splits the work into channels, parallel lanes that run at the same time, and it arranges the ranks into a ring so each one sends to the next.
localhost:24133:24153 [0] NCCL INFO Channel 00/04 : 0 1
localhost:24133:24153 [0] NCCL INFO Channel 01/04 : 0 1
localhost:24133:24153 [0] NCCL INFO Channel 02/04 : 0 1
localhost:24133:24153 [0] NCCL INFO Channel 03/04 : 0 1
localhost:24133:24153 [0] NCCL INFO Ring 00 : 1 -> 0 -> 1
localhost:24133:24153 [0] NCCL INFO Trees [0] 1/-1/-1->0->-1 [1] -1/-1/-1->0->1 [2] 1/-1/-1->0->-1 [3] -1/-1/-1->0->1Channel 00/04 reads as c
Practice Reading NCCL's Mind in a real Linux terminal at The Linux Camp. Progress is verified automatically as you type commands on the machine.