RCCL vs NCCL: same API, different transport
RCCL is AMD's NCCL fork — same collective primitives, same NCCL_* env vars, different intra-node fabric (xGMI vs NVLink), different multi-node transport (UCX-flavoured RDMA). What you actually have to change in your scripts and your containers.
help for the full list, or solutions for copy-paste fix recipes.If you've spent a year debugging NCCL — rail binding, NCCL_IB_HCA, peermem load order, NVL5/NVL18 ring construction — you have already done 90% of the learning needed for RCCL. RCCL is a fork of NCCL. AMD took the NCCL source, ran it through their compatibility shim, swapped CUDA calls for HIP calls, swapped NVLink transport for xGMI, and shipped it. The API is the same. The environment variables are mostly the same. The error messages are mostly the same. What's different is the underlying transport, which means the failure modes you'll hit are slightly different — and the diagnostic tools are slightly different.
This page is the translation table. It assumes you already know NCCL — see NCCL multi-node tuning if you don't.
What RCCL is
RCCL (pronounced "Rickle") is the ROCm Collective Communication Library. The ROCm-side equivalent of NCCL.
| Property | RCCL | NCCL |
|---|---|---|
| Source heritage | Forked from NCCL ~2018 | NVIDIA original |
| Language | C++ + HIP | C++ + CUDA |
| API | rccl.h — almost identical to nccl.h | nccl.h |
| Symbols | ncclXxx(...) — same names as NCCL | ncclXxx(...) |
| Library file | librccl.so | libnccl.so |
| Intra-node transport | xGMI / Infinity Fabric, PCIe | NVLink, NVSwitch, PCIe |
| Inter-node transport | InfiniBand Verbs, RoCE, TCP/IP sockets | InfiniBand Verbs, RoCE, TCP/IP sockets |
| Test repo | ROCm/rccl-tests (fork of nccl-tests) | NVIDIA/nccl-tests |
| GPU binding env | HIP_VISIBLE_DEVICES | CUDA_VISIBLE_DEVICES |
| RDMA peer-memory | Built into amdgpu kernel module | nvidia_peermem (separate module) |
The fact that RCCL keeps the ncclXxx symbol names is deliberate and operationally enormous. It means that PyTorch, DeepSpeed, Megatron-LM, vLLM, Horovod — anything that links against libnccl.so — works on AMD by swapping the library path. The application doesn't know which one it's calling. PyTorch's distributed backend, in particular, uses RCCL transparently when the build is ROCm and NCCL when the build is CUDA; you don't switch a flag.
What RCCL is NOT
A few NVIDIA features have no AMD counterpart and you should know up front:
- No SHARP equivalent. NVIDIA's CollNet topology uses Mellanox Quantum / Spectrum switches with in-network reduction (SHARPv1/v2/v3). AMD has no equivalent in-network reduction; RCCL on Quantum/Spectrum switches will not engage SHARP even if the switch supports it. For workloads where SHARP gave you 30%+ collective speedup at scale, RCCL won't match.
- No NVLink-SHARP (NVLS).
NCCL_NVLS_ENABLEis a no-op in RCCL. The "intra-node SHARP" path that Hopper added doesn't exist on AMD. - No NVSwitch fabric domain. With NVSwitch, all 8 GPUs (or all 72 in NVL72) are non-blocking all-to-all. AMD's xGMI is point-to-point mesh — you have full bandwidth between any pair, but uneven traffic can starve some links while others bottleneck. For balanced AllReduce, this is not visible. For uneven workloads (e.g. heavy AllToAll on MoE expert routing), the difference can show up.
Transports: where xGMI and Infinity Fabric replace NVLink
A typical 8-GPU MI300X OAM platform vs a typical 8-GPU H100 SXM5 platform:
| Aspect | 8× H100 SXM5 (NVLink + NVSwitch) | 8× MI300X OAM (xGMI mesh) |
|---|---|---|
| GPU-to-GPU all-pairs | 900 GB/s via NVSwitch | 896 GB/s aggregate, point-to-point |
| Topology | Star through 4× NVSwitch chips | Fully-connected mesh (each GPU has 7 xGMI links) |
| Bandwidth fairness | Non-blocking — all pairs simultaneously full BW | Per-link 128 GB/s; 8 simultaneous full-BW pairs use all links |
| Adverse traffic patterns | Switch can rebalance dynamically | Fixed link allocations — uneven traffic degrades |
| Failure modes | NVSwitch port flap, fabric manager failure | Single xGMI link degradation, IF link errors |
| Diagnostic | nvidia-smi nvlink -s, dcgmi | rocm-smi --showtopo, rocm-smi --showxgmierr |
To check xGMI link health on an MI300X box:
$ rocm-smi --showtopo
============================ ROCm System Management Interface ============================
======================================== Topology ========================================
GPU0 GPU1 GPU2 GPU3 GPU4 GPU5 GPU6 GPU7
GPU0 X XGMI XGMI XGMI XGMI XGMI XGMI XGMI
GPU1 XGMI X XGMI XGMI XGMI XGMI XGMI XGMI
GPU2 XGMI XGMI X XGMI XGMI XGMI XGMI XGMI
...
XGMI everywhere off-diagonal means the mesh is healthy. If you see PIX or PCIe between any pair of GPUs, that link has dropped — communication between those two GPUs will fall back to PCIe, costing you ~7× bandwidth. This is the AMD equivalent of seeing "NV1" instead of "NV18" on nvidia-smi topo -m.
For per-link error counters:
$ rocm-smi --showxgmierr
GPU[0]: XGMI Link Error Count: 0
GPU[1]: XGMI Link Error Count: 0
...
Anything non-zero, especially growing during a workload, is suspect. It's the analog of NVLink CRC errors.
Inter-node transport: same RDMA, different glue
RCCL uses the same multi-node transport options as NCCL:
| Transport | NCCL behavior | RCCL behavior |
|---|---|---|
InfiniBand Verbs (IB) | Uses libibverbs, requires nvidia_peermem | Uses libibverbs, peermem is in amdgpu itself |
| RoCE v2 | Same Verbs path, GID config | Same |
| TCP/IP sockets | Fallback if RDMA unavailable | Same |
| UCX | Optional plugin | Recommended on AMD — ships as librccl-net-ucx.so plugin |
The biggest practical difference: AMD recommends UCX as the network plugin for RCCL, especially on RoCE fabrics. UCX is the OpenUCX framework also used by Open MPI; it does the QP setup, GID selection, and connection management. NVIDIA NCCL uses its own native verbs path by default; AMD's recommendation is to load the UCX plugin, which gives more uniform behavior across HCAs.
To enable UCX in RCCL:
# Tell RCCL to use the UCX network plugin
export NCCL_NET=UCX
export NCCL_PLUGIN_P2P=ucx
# Or load the plugin explicitly
export LD_PRELOAD=/opt/rocm/lib/librccl-net-ucx.so:$LD_PRELOAD
Note that even with UCX, you still set NCCL_IB_HCA, NCCL_IB_GID_INDEX, NCCL_IB_TC etc — the env vars haven't changed.
Environment variables: what carries over, what changes
The vast majority of NCCL env vars work in RCCL. Here is the operator's cheat sheet:
| Variable | Works in RCCL? | Notes |
|---|---|---|
NCCL_DEBUG=INFO | yes | Same output format |
NCCL_DEBUG_SUBSYS=INIT,NET,GRAPH | yes | Same |
NCCL_DEBUG_FILE=/tmp/rccl.%h.%p.log | yes | Same %h hostname, %p pid templating |
NCCL_IB_DISABLE=0 | yes | Same default |
NCCL_IB_HCA=mlx5_0,mlx5_1,... | yes | Same HCA filtering |
NCCL_IB_GID_INDEX=N | yes | Same RoCE v2 GID selection |
NCCL_IB_TC=N | yes | Same DSCP/TOS — see RoCE |
NCCL_IB_TIMEOUT=N | yes | Same retry timeout semantics |
NCCL_IB_QPS_PER_CONNECTION=N | yes | Works; more impactful on AMD because of UCX plugin behavior |
NCCL_SOCKET_IFNAME=eth0 | yes | Same |
NCCL_NET=UCX | yes | AMD-specific recommendation |
NCCL_NVLS_ENABLE=1 | no-op | NVLink-SHARP doesn't exist on AMD |
NCCL_COLLNET_ENABLE=1 | no-op | No SHARP equivalent |
NCCL_TOPO_FILE=/path/topo.xml | yes | Same XML format; rccl-tests produces equivalent dumps |
NCCL_NET_GDR_LEVEL=PIX | yes | Same — AMD does GPUDirect RDMA the same way |
NCCL_P2P_DISABLE=0 | yes | Same |
NCCL_PROTO=Simple | yes | Same protocols (LL, LL128, Simple) |
NCCL_ALGO=Ring,Tree | yes | CollNet won't be selected |
NCCL_BUFFSIZE=N | yes | Same |
NCCL_MIN_NCHANNELS=N | yes | Same |
NCCL_MAX_NCHANNELS=N | yes | Same |
There are a handful of RCCL-specific env vars too. The important ones:
| Variable | Purpose |
|---|---|
RCCL_TREES | Define a custom tree topology — alternative to ring |
HSA_FORCE_FINE_GRAIN_PCIE=1 | Force fine-grained PCIe traffic — needed on some IOMMU configurations to make GPUDirect RDMA work |
NCCL_MSCCL_ENABLE=1 | Enable Microsoft Collective Communication Library (MSCCL) algorithms — Microsoft's contribution, can give 1.2–1.5× on certain AllReduce sizes |
RCCL_MSCCLPP_ENABLE=1 | Enable MSCCL++ (the C++/header-only successor) |
HSA_FORCE_FINE_GRAIN_PCIE=1 is the one that bites people. On boxes where the IOMMU is in passthrough mode, GPUDirect RDMA traffic between the NIC and GPU may need this flag set to avoid bouncing through host RAM. Setting it without need is harmless. Not setting it when needed costs you 50%+ of inter-node bandwidth without any error message — rccl-tests simply runs at half speed.
rccl-tests vs nccl-tests
ROCm/rccl-tests is a fork of NVIDIA/nccl-tests. Same flags, same output format, just compiled against HIP and RCCL.
# Build
git clone https://github.com/ROCm/rccl-tests
cd rccl-tests
make MPI=1 MPI_HOME=/opt/ompi HIP_HOME=/opt/rocm
# Run (single-node, 8 GPUs)
./build/all_reduce_perf -b 8 -e 8G -f 2 -g 8
# Run (multi-node via mpirun, 4 nodes × 8 GPUs)
mpirun -np 32 -hostfile hosts \
-x NCCL_DEBUG=INFO \
-x NCCL_IB_HCA=mlx5_0,mlx5_1,mlx5_2,mlx5_3,mlx5_4,mlx5_5,mlx5_6,mlx5_7 \
-x NCCL_NET=UCX \
-x HSA_FORCE_FINE_GRAIN_PCIE=1 \
./build/all_reduce_perf -b 8 -e 8G -f 2 -g 8
What you expect on a healthy 8× MI300X box, single-node AllReduce at 1 GB:
| Implementation | algbw | busbw | Notes |
|---|---|---|---|
nccl-tests on 8× H100 SXM5 | ~370 GB/s | ~640 GB/s | Reference NVIDIA baseline |
rccl-tests on 8× MI300X | ~340 GB/s | ~590 GB/s | Healthy MI300X. Within 10% of H100. |
rccl-tests on 8× MI300X (degraded) | ~50 GB/s | ~90 GB/s | One xGMI link down — fall back to PCIe |
Multi-node, 4 nodes × 8 GPUs, 8× 400 Gb/s NICs per node:
| Implementation | busbw | Notes |
|---|---|---|
nccl-tests on 4× HGX H100 | ~360 GB/s | Healthy IB fabric, peermem, rail-bound |
rccl-tests on 4× MI300X | ~330 GB/s | Healthy IB fabric, UCX, fine-grain-pcie set |
rccl-tests on 4× MI300X (no UCX) | ~280 GB/s | Sub-optimal but functional on default verbs path |
rccl-tests on 4× MI300X (no fine-grain-pcie + IOMMU on) | ~150 GB/s | Bouncing through host RAM |
If you measure significantly below these numbers, look at troubleshooting before tearing into the application.
HIP vs CUDA in your container
Your training/inference container has to be the right one for the GPU. You cannot run a CUDA-built PyTorch on AMD silicon — it links against libcudart.so, which doesn't exist there. There are three options:
1. Use AMD's pre-built PyTorch image
FROM rocm/pytorch:latest
# pytorch already built with ROCm/RCCL
Easiest. The rocm/pytorch image ships a PyTorch built against ROCm, with RCCL bundled. import torch; torch.distributed.init_process_group(backend="nccl") — yes, "nccl" is still the right backend name; PyTorch translates that to RCCL when running on ROCm.
2. Use AMD's vLLM image for inference
FROM rocm/vllm:latest
vLLM has been ROCm-supported since 0.3 and well-supported since 0.5. Performance on MI300X for Llama-3.1-70B / 405B is competitive with H100/H200.
3. Build your own
If you have to build your own (custom kernels, internal forks), the relevant flags:
FROM rocm/dev-ubuntu-22.04:7.0
# Tell PyTorch which gfx targets to compile for
ENV PYTORCH_ROCM_ARCH="gfx942" # MI300X
# Or for multi-target, semicolon-separated:
# ENV PYTORCH_ROCM_ARCH="gfx90a;gfx942;gfx950"
# RCCL is built in as part of ROCm
RUN pip install torch --index-url https://download.pytorch.org/whl/rocm6.2
The biggest mistake when self-building: PYTORCH_ROCM_ARCH set wrong. If you compile for gfx90a (MI250X) and run on gfx942 (MI300X), every kernel JIT-recompiles on first call. Your training run starts, sits at 0% utilization for 20 minutes, then either eventually starts running (slowly) or crashes with hipErrorNoBinaryForGpu.
Common gotchas porting a NCCL workload to RCCL
If you're moving an existing CUDA/NCCL workload onto AMD, here are the failure modes ranked by frequency:
1. CUDA_VISIBLE_DEVICES doesn't work — use HIP_VISIBLE_DEVICES
PyTorch on ROCm honours both, but other libraries (Horovod, DeepSpeed sometimes, custom kernels) only respect HIP_VISIBLE_DEVICES. Always set both:
export CUDA_VISIBLE_DEVICES=0,1,2,3
export HIP_VISIBLE_DEVICES=0,1,2,3
2. nvidia_peermem is irrelevant; check amdgpu GPUDirect status
On NVIDIA, you lsmod | grep peermem and panic if it's missing. On AMD, peermem is built into amdgpu. To verify GPUDirect RDMA is functional, check the InfiniBand side:
# /proc/driver/nvidia/peer_mem doesn't exist on AMD — instead:
ibv_devinfo | grep -i 'Atomic\|Peer'
# and a quick traffic test:
ib_write_bw -d mlx5_0 -F --use_cuda=0 ... # Note: still flags 'cuda' but works for AMD
3. Warp size mismatch in custom kernels
CUDA warp size = 32. AMD wavefront size = 64. If your custom CUDA kernel hard-codes 32 anywhere (__shfl_sync, shared memory tiling, voting ops), hipify-perl will translate the API call but won't fix the assumed-32 logic. You'll get correctness bugs that are hard to spot.
4. NCCL_DEBUG output looks the same — until it doesn't
NCCL_DEBUG=INFO prints almost identical lines on RCCL: ring construction, channel assignment, transport selection. The thing to watch for: messages mentioning "NVLS" (no-op on AMD), "CollNet" (no-op on AMD), or "peermem" (different on AMD). If your operator runbook greps for "peermem registered" as a health check, that string won't appear in RCCL logs.
5. The container build chain
You cannot reuse the same container image. CUDA→HIP requires recompilation of every device-side kernel. There is no "binary-compatible" path. Plan for parallel container build pipelines if you operate a mixed fleet.
Reading NCCL_DEBUG=INFO output on RCCL
The output looks almost identical to NCCL. The annotations show what each line means and what the AMD-specific differences are:
my-node-0:1234:1234 [0] NCCL INFO Bootstrap : Using eth0:10.20.30.4<0>
^^^^^^^^^^^^^^^^
# NCCL_SOCKET_IFNAME pick. Same as NCCL — pick the control-plane interface, not IPoIB.
my-node-0:1234:1234 [0] NCCL INFO NET/IB : Using [0]mlx5_0:1/RoCE [1]mlx5_1:1/RoCE [2]mlx5_2:1/RoCE [3]mlx5_3:1/RoCE
^^^^
# 4 HCAs picked. On a healthy 8-NIC node you want 8 here. If you see fewer,
# NCCL_IB_HCA filtering is missing some, or some HCAs are down.
my-node-0:1234:1234 [0] NCCL INFO Topology detection: 8 GPUs, 4 NUMA nodes, gfx942
^^^^^^
# Confirms RCCL identified the AMD silicon correctly.
my-node-0:1234:1234 [0] NCCL INFO Channel 00 : 0[0] -> 1[1] [send] via NET/IB/0/GDRDMA
^^^^^^^^
# GPUDirect RDMA confirmed for this channel. If you see "via NET/IB/0" without
# GDRDMA, traffic is going via host-bounce-buffer and you'll lose ~50% bandwidth.
# Investigate IOMMU and HSA_FORCE_FINE_GRAIN_PCIE.
my-node-0:1234:1234 [0] NCCL INFO Connected all rings
# Bootstrap done. If init hangs, it gets stuck before this line.
my-node-0:1234:1234 [0] NCCL INFO comm 0x... rank 0 nranks 32 cudaDev 0 nvmlDev 0 busId 18000 commId ...
^^^^^^^^ ^^^^^^^^
# On RCCL these still say "cudaDev" / "nvmlDev" — the field names didn't change
# when AMD forked NCCL. The values are HIP/ROCr device IDs.
Things you will NOT see in RCCL output that you would in NCCL:
NVLS Tree/NVLS Ring(NVLink-SHARP not present)CollNet Chain(no SHARP-equivalent)peer mapping donefornvidia_peermem(RDMA peermem is inamdgpu, no separate registration line)
Things you will see in RCCL that you wouldn't in NCCL:
ROCm version: 6.2.0— first init line for self-identificationxGMI bandwidth detected: 128 GB/s per link, 7 links per GPU— topology autodetectMSCCL kernel selected for AllReduce— when MSCCL is enabled (not in stock NCCL)
If you build operator runbooks that grep NCCL logs for specific strings, audit them when porting to AMD.
Bandwidth budget at the cluster level
A practical translation table for a 4-node × 8-GPU cluster, comparing achievable bandwidth on each platform:
| Workload | 4× HGX H100 (8× 400 Gb/s NICs) | 4× UBB MI300X (8× 400 Gb/s NICs) |
|---|---|---|
| Intra-node AllReduce, 1 GB | ~640 GB/s busbw | ~590 GB/s busbw (~92%) |
| Inter-node AllReduce, 1 GB | ~360 GB/s busbw | ~330 GB/s busbw (~92%) |
| Inter-node AllReduce, 4 KB | ~6 µs latency | ~7 µs latency |
| AllToAll, 256 MB (8 ranks/node) | ~280 GB/s busbw | ~260 GB/s busbw |
| Broadcast, 1 GB | ~50 GB/s alg / ~395 GB/s bus | ~46 GB/s alg / ~360 GB/s bus |
Rule of thumb: well-tuned RCCL achieves 90–95% of the equivalent NCCL number on AMD silicon. The gap is mostly in the absence of SHARP and NVLS for collective acceleration. For workloads that don't depend heavily on those features (most LLM training), the gap is operationally invisible.
What changes for your launch script
Most of your mpirun / torchrun / Slurm launch script doesn't change. Here's a side-by-side of a typical NCCL launch and the RCCL equivalent:
NVIDIA / NCCL launch
#!/bin/bash
export NCCL_DEBUG=INFO
export NCCL_IB_HCA=mlx5_0,mlx5_1,mlx5_2,mlx5_3,mlx5_4,mlx5_5,mlx5_6,mlx5_7
export NCCL_IB_GID_INDEX=3
export NCCL_IB_TC=106
export NCCL_SOCKET_IFNAME=bond0
export NCCL_TOPO_FILE=/etc/nccl/topo.xml
mpirun -np 32 -hostfile hosts \
-x NCCL_DEBUG -x NCCL_IB_HCA -x NCCL_IB_GID_INDEX -x NCCL_IB_TC \
-x NCCL_SOCKET_IFNAME -x NCCL_TOPO_FILE \
python train.py
AMD / RCCL launch
#!/bin/bash
# Identical NCCL_* env vars
export NCCL_DEBUG=INFO
export NCCL_IB_HCA=mlx5_0,mlx5_1,mlx5_2,mlx5_3,mlx5_4,mlx5_5,mlx5_6,mlx5_7
export NCCL_IB_GID_INDEX=3
export NCCL_IB_TC=106
export NCCL_SOCKET_IFNAME=bond0
export NCCL_TOPO_FILE=/etc/rccl/topo.xml # different path, but same XML format
# AMD-specific additions
export NCCL_NET=UCX # use UCX plugin for inter-node
export HSA_FORCE_FINE_GRAIN_PCIE=1 # avoid IOMMU bounce buffer
export HIP_VISIBLE_DEVICES=0,1,2,3,4,5,6,7 # belt-and-suspenders with CUDA_VISIBLE_DEVICES
export CUDA_VISIBLE_DEVICES=0,1,2,3,4,5,6,7
# PyTorch arch list (analog of TORCH_CUDA_ARCH_LIST)
export PYTORCH_ROCM_ARCH=gfx942
mpirun -np 32 -hostfile hosts \
-x NCCL_DEBUG -x NCCL_IB_HCA -x NCCL_IB_GID_INDEX -x NCCL_IB_TC \
-x NCCL_SOCKET_IFNAME -x NCCL_TOPO_FILE -x NCCL_NET \
-x HSA_FORCE_FINE_GRAIN_PCIE -x HIP_VISIBLE_DEVICES -x CUDA_VISIBLE_DEVICES \
-x PYTORCH_ROCM_ARCH \
python train.py
Three additions: NCCL_NET=UCX, HSA_FORCE_FINE_GRAIN_PCIE=1, the dual visible-devices export. Everything else carries over.
When MSCCL is worth enabling
MSCCL (Microsoft Collective Communication Library) is a contribution AMD ships in RCCL by default in 6.x+. It adds algorithmic optimizations for specific message size + topology combinations, picked from a library of pre-tuned algorithms.
| When MSCCL helps | Typical gain |
|---|---|
| AllReduce on 8-GPU UBB, 4–256 MB messages | 10–30% |
| AllToAll on MoE expert routing | 15–25% |
| Ring AllReduce at large scale (>64 nodes) | 5–15% |
| Small messages (<1 MB) | minor — sometimes negative |
| Large messages (>1 GB) | minor |
To enable:
export NCCL_MSCCL_ENABLE=1
# Optionally point to a custom algorithm config:
export RCCL_MSCCL_ALGO_DIR=/opt/rocm/share/rccl/mscclpp
Things to know:
- MSCCL is on by default in some RCCL builds and off by default in others. Check by setting
NCCL_DEBUG=INFOand grep forMSCCL kernel. - It interacts with
NCCL_ALGOandNCCL_PROTO. If you've manually pinned an algorithm, MSCCL won't override. - For training where you've already tuned
NCCL_PROTO=Simplefor stability, just leave MSCCL on default. For inference with very specific batch sizes, benchmark both.
Multi-node topology files for RCCL
RCCL accepts the same topology XML format as NCCL. AMD ships canned topology files for some platforms (/opt/rocm/share/rccl/topo/), but for OEM UBBs you may need to generate your own:
# Dump autodetected topology
NCCL_TOPO_DUMP_FILE=/tmp/rccl-topo.xml \
NCCL_DEBUG=INFO \
./build/all_reduce_perf -b 1G -e 1G -g 8
# Inspect
xmllint --format /tmp/rccl-topo.xml | head -100
# Edit if needed (e.g. fixing PCIe link speed reports), then deploy
sudo cp /tmp/rccl-topo.xml /etc/rccl/topo.xml
export NCCL_TOPO_FILE=/etc/rccl/topo.xml
The XML structure mirrors NCCL's, with one practical difference: GPU dev IDs in the AMD topology file map to gfx-target indices, not CUDA ordinals. If you copy a topology XML from an NVIDIA box and try to use it on AMD, the GPU references will be silently wrong. Always generate the topology from the AMD box itself.
Side-by-side: troubleshooting an underperforming AllReduce
The same problem (AllReduce running at half expected bandwidth on 4 nodes), debugged on each platform:
| Step | NVIDIA / NCCL | AMD / RCCL |
|---|---|---|
| 1. Verify all GPUs see each other | nvidia-smi topo -m | rocm-smi --showtopo |
| 2. Verify peer-to-peer DMA | nvidia-smi nvlink -s (NV18 active) | rocm-smi --showtopo (XGMI off-diagonal) |
| 3. Verify peermem | `lsmod | grep nvidia_peermem` |
| 4. Verify HCA selection | NCCL_DEBUG=INFO ⇒ "Using NIC mlx5_0" | same |
| 5. Verify GDR | log: "via GPU Direct RDMA" | log: "via GPU Direct RDMA" (same string) |
| 6. IOMMU passthrough | `dmesg | grep -i iommu` |
7. Run *_perf baseline | ./all_reduce_perf -b 1G -e 1G -g 8 | same with rccl-tests build |
| 8. Compare to reference | NCCL Whitepaper / DGX baselines | AMD reference + your own healthy-node baselines |
The 90% case in both worlds: rail binding wrong, or one link degraded. The other 10%: env-var typo, IOMMU misconfigured, or someone shipped the wrong container.
See also
- NCCL multi-node tuning — the NCCL details that almost all carry over
- ROCm stack — RCCL ships as part of ROCm
- AMD troubleshooting — diagnosing slow or failed RCCL collectives
- GPUDirect RDMA — AMD does this through
amdgpu, NVIDIA throughnvidia_peermem - InfiniBand vs RoCE — fabric choices apply identically to both vendors
- PCIe topology — rail binding and PCIe constraints are vendor-agnostic