-
Notifications
You must be signed in to change notification settings - Fork 0
AI Performance
collective aggregation operations like all-reduce, all-to-all, and all-gather, which are extensively during model training and inference
-
Measure goodput: Look beyond raw FLOPS or utilization. Instead, measure the ratio of time the GPU spends performing useful work (e.g., forward/backprop computations) versus waiting on data or other overhead. Use NVIDIA Nsight Systems/Compute or PyTorch profiler to measure this ratio. Strive to improve this ratio, as goodput focuses on effective, useful GPU utilization.
-
Prefer skillful engineering optimizations instead of brute-force spending: More hardware isn’t a silver bullet. Clever software and system optimizations can bridge the gap when hardware is limited, enabling results that would otherwise require far more expensive infrastructure. We saw this with DeepSeek’s achievement. By optimizing communication and hardware usage, their engineers trained a frontier model on restricted H800 GPUs (with limited interconnect bandwidth) at a fraction of the cost. The DeepSeek model matched the performance of frontier models trained on far more powerful hardware. In other words, skillful engineering outperformed brute-force spending.
-
Look for order-of-magnitude impact with incremental optimizations: At scale, even a small-percentage efficiency gain can save millions of dollars. Said differently, small inefficiencies such as redundant computations and slow data pipelines can silently increase costs as the system scales.
-
Approach performance tuning with a profile-driven mindset: Use data and profiling tools to guide optimizations. Use profilers to identify the true bottlenecks—whether it’s compute utilization, memory bandwidth, memory latency, cache misses, or communication/network delays. Then apply targeted optimizations for that bottleneck.
-
Maintain a holistic view: Improving AI systems performance spans hardware, including the GPU, CPU, memory, and network—as well as software such as algorithms and libraries. A weakness in any layer can bottleneck the whole. The best performance engineers consider hardware-software codesign: sometimes algorithm changes can alleviate hardware limits, and sometimes new hardware features enable new algorithms.
-
Stay informed on the latest hardware, software, and algorithms: Modern AI hardware and software are evolving rapidly. New capabilities such as unified CPU-GPU memory, faster interconnects, and novel numerical-precision formats can change the optimal strategies. A good performance engineer keeps an eye on these and updates their mental models accordingly to eliminate bottlenecks quickly. Additionally, the MLPerf benchmark suites are a great resource to understand AI hardware performance for various models.
- cuPyNumeric: dropin replacement for numpy.
When a researcher submits a training job, the scheduler reserves nodes, the OS provides the GPU devices and memory allocations using the NVIDIA driver, and the container provides the correct software environment (including the optimized, hardware-aware CUDA libraries). The user code (e.g., PyTorch, TensorFlow, JAX) uses these CUDA libraries, which ultimately communicate with the driver and hardware.
A NUMA node is a logical grouping of CPUs, GPUs, network interface con‐ trollers (NICs), and memory that are physically close to one another.
By binding a process to a CPU on the same NUMA node as its GPU, we can avoid this extra overhead. For instance, you can use numactl --cpunodebind= --membind= to bind both CPU threads and memory allocations to the GPU’s local NUMA node.
The key idea is to keep CPU execution and memory access local to the GPU that it’s serving
By default, processes may be migrated across NUMA nodes. This will lead to additional latency caused by remote memory accesses. As such, it’s important to explicitly bind processes and memory to the same NUMA node as the local GPU. You can do this using numactl, taskset, or cgroups
To explicitly specify NUMA-affinity, you need to “pin” processes or threads to spe‐ cific CPUs that are connected to the same NUMA node as the GPU. This type of CPU affinity is called CPU pinning.
If you launch eight training processes, one per GPU, you should bind each training process to a CPU core—or set of CPU cores—connected to the same NUMA node as the GPUs
Linux provides tools to do this, including numactl --cpunodebind= --membind= , which launches a process pinned to the given NUMA node. You can also use taskset to pin processes to specific core IDs. Here is an example using numactl to bind the train.py script to a CPU running in the same NUMA
node 1 as GPU 4:
numactl --cpunodebind=1 --membind=1 \
python train.py --gpu 4This assumes we know the NUMA node ID and that we are binding the script to only one GPU. Binding the train.py to multiple GPUs to an unknown NUMA node is a bit more complicated. The following script dynamically queries the topology using nvidia-smi topo and binds the script to GPUs using the local NUMA node:
#!/bin/bash
for GPU in 0 1 2 3; do
# Query NUMA node for this GPU
NODE=$(nvidia-smi topo -m -i $GPU \
| awk '/NUMA Affinity/ {print $NF}')
# Launch the training process pinned to that NUMA node
numactl --cpunodebind=$NODE --membind=$NODE \
bash -c "CUDA_VISIBLE_DEVICES=$GPU python train.py --gpu $GPU"
doneMany deep learning frameworks also let you set thread affinities programmatically. For instance, PyTorch’s DataLoader exposes worker_init_fn so you can set CPU affinity for each worker process during initialization
import os
import re
import glob
import subprocess
import psutil
import ctypes
import torch
import torch.distributed as dist
from torch.nn.parallel import DistributedDataParallel as DDP
from torch.utils.data import DataLoader, Dataset
from functools import partial
# Optional: NVML is preferred for GPU↔NUMA mapping
try:
import pynvml as nvml # pip install nvidia-ml-py3
_HAS_NVML = True
except Exception:
_HAS_NVML = False
# --- libnuma for memory binding
_libnuma = ctypes.CDLL("libnuma.so")
if _libnuma.numa_available() < 0:
raise RuntimeError("NUMA not available on this system")
_libnuma.numa_run_on_node.argtypes = [ctypes.c_int]
_libnuma.numa_set_preferred.argtypes = [ctypes.c_int]
def parse_physical_cpu_list(phys_str: str):
"""Parse '0-3,8-11' -> [0,1,2,3,8,9,10,11]."""
cpus = []
if not phys_str:
return cpus
for part in phys_str.split(','):
part = part.strip()
if not part:
continue
if '-' in part:
start, end = map(int, part.split('-'))
cpus.extend(range(start, end + 1))
else:
cpus.append(int(part))
return cpus
def get_numa_cpus_for_node(node: int):
"""Read /sys/devices/system/node/node{node}/cpulist."""
path = f"/sys/devices/system/node/node{node}/cpulist"
with open(path, "r") as f:
return parse_physical_cpu_list(f.read().strip())
def get_numa_cpus_and_memory():
"""Return (current_cpu_mask, preferred_node) from numactl --show."""
out = subprocess.run(
["numactl", "--show"],
capture_output=True,
text=True
).stdout
phys = re.search(r"physcpubind:\s*([\d,\-\s]+)", out).group(1)
cpus = parse_physical_cpu_list(phys)
node = int(
re.search(r"preferred node:\s*(-?\d+)", out).group(1)
)
return cpus, node
def get_gpu_numa_node(device: int) -> int:
"""
Determine NUMA node for a GPU (prefer NVML; fall back to sysfs;
final fallback to current preferred node).
"""
# NVML path (preferred)
if _HAS_NVML:
try:
nvml.nvmlInit()
props = torch.cuda.get_device_properties(device)
pci = props.pci_bus_id
# Normalize to 8-hex-digit domain if needed for NVML
try:
domain, bus, devfn = pci.split(':')
if len(domain) < 8:
domain = domain.rjust(8, '0')
pci8 = f"{domain}:{bus}:{devfn}"
except ValueError:
pci8 = pci
try:
handle = nvml.nvmlDeviceGetHandleByPciBusId_v2(pci8)
except AttributeError:
handle = nvml.nvmlDeviceGetHandleByPciBusId(pci8)
# Direct NUMA ID if driver exposes it
try:
numa_id = nvml.nvmlDeviceGetNUMANodeId(handle)
if isinstance(numa_id, int) and numa_id >= 0:
return numa_id
except Exception:
pass
# Derive from NVML CPU affinity
cpu_count = psutil.cpu_count(logical=True)
elems = (cpu_count + 63) // 64
mask = nvml.nvmlDeviceGetCpuAffinity(handle, elems)
cpus = []
for i, m in enumerate(mask):
m = int(m)
for b in range(64):
if m & (1 << b):
cpu_id = i * 64 + b
if cpu_id < cpu_count:
cpus.append(cpu_id)
# Build CPU→NUMA map from sysfs and choose majority node
cpu2node = {}
for node_path in sorted(
glob.glob("/sys/devices/system/node/node*")
):
node_id = int(
os.path.basename(node_path).replace("node", "")
)
with open(os.path.join(node_path, "cpulist"), "r") as f:
for c in parse_physical_cpu_list(f.read().strip()):
cpu2node[c] = node_id
counts = {}
for c in cpus:
n = cpu2node.get(c)
if n is not None:
counts[n] = counts.get(n, 0) + 1
if counts:
return max(
counts.items(),
key=lambda kv: kv[1]
)[0]
except Exception:
pass
# sysfs fallback
try:
props = torch.cuda.get_device_properties(device)
pci = props.pci_bus_id
sysfs_path = f"/sys/bus/pci/devices/{pci}/numa_node"
with open(sysfs_path, "r") as f:
val = int(f.read().strip())
return val if val >= 0 else 0
except Exception:
pass
# Last resort: current preferred node
_, node = get_numa_cpus_and_memory()
return node if node >= 0 else 0
def set_numa_affinity(node: int):
"""Bind current process to CPUs and memory of the given NUMA node."""
cpus = get_numa_cpus_for_node(node)
# IMPORTANT: CPUs of target node
psutil.Process(os.getpid()).cpu_affinity(cpus)
_libnuma.numa_run_on_node(node)
_libnuma.numa_set_preferred(node)
print(
f"PID={os.getpid()} bound to NUMA node {node} "
f"(CPUs={cpus})"
)
return cpus
def _worker_init_fn(worker_id: int, node: int, cpus: list):
"""Reapply binding in each DataLoader worker (no CUDA calls here)."""
psutil.Process(os.getpid()).cpu_affinity(cpus)
_libnuma.numa_run_on_node(node)
_libnuma.numa_set_preferred(node)
print(
f"Worker {worker_id} "
f"(PID={os.getpid()}) "
f"bound to NUMA node {node}"
)
# ----- Example usage below -----
class MyDataset(Dataset):
def __len__(self):
return 1024
def __getitem__(self, idx):
return torch.randn(224 * 224 * 3, device="cpu")
def main():
# DDP setup
dist.init_process_group(
backend="nccl",
init_method="env://"
)
device = torch.cuda.current_device()
# Determine GPU's NUMA node and bind this process
gpu_node = get_gpu_numa_node(device)
cpus = set_numa_affinity(gpu_node)
# Build dataloader with closure-based worker_init_fn
dataset = MyDataset()
init_fn = partial(
_worker_init_fn,
node=gpu_node,
cpus=cpus
)
dataloader = DataLoader(
dataset,
batch_size=32,
num_workers=4,
pin_memory=True,
persistent_workers=True,
worker_init_fn=init_fn,
prefetch_factor=2,
)
# Model and DDP
model = torch.nn.Linear(
224 * 224 * 3,
10,
bias=True
).to("cuda")
ddp_model = DDP(
model,
device_ids=[device],
static_graph=True
)
for batch in dataloader:
batch = batch.to(
"cuda",
non_blocking=True
)
out = ddp_model(batch)
# ... loss, backward, optimizer ...
if __name__ == "__main__":
main()
This script binds the main training process and each DataLoader worker process to the GPU’s local NUMA node to prevent cross-NUMA memory access. In the DataLoader, we pass a closure-based worker_init_fn that reapplies the precomputed NUMA binding inside each worker. And we do this without touching any CUDA APIs in the worker.
Because some launchers, container runtimes, or kernels do not reliably propagate NUMA policy to children, we explicitly reapply and verify the binding in every forked worker. It’s not safe to rely on inheritance alone.
Remember to set pin_memory=True and use non_blocking=True on H2D copies so that page-locked host buffers stay on the correct NUMA node. Prefer persis tent_workers=True to avoid re-forking workers and losing their affinity between epochs. And do not call torch.cuda.* in worker_init_fn. Instead, pass the GPU index using a closure or environment variable. The result is that data preparation and batch loading happen entirely in local mem‐ ory. This way, your GPUs stay busy and never need to pause for a remote‐NUMA hop. With this code, you get robust, topology‐aware affinity on any Linux server with libnuma and numactl installe
By default, numactl applies its CPU and memory policy to a process and is docu‐ mented to inherit that policy to all forked children. In practice, however, threads spawned by Python frameworks or exec’d subprocesses don’t always pick up the same settings on every kernel or Linux distribution. When using framework-managed worker processes, you should explicitly reassert the CPU and memory policy inside of each worker.
With a superchip architecture like Grace Blackwell (and Vera Rubin), the CPU and GPU are coherent using NVLink-C2C. How‐ ever, Linux still models CPU DRAM and GPU HBM as separate pools. Binding CPU threads to the local CPU NUMA node still remains beneficial for locality
In practice, pinning can eliminate unpredictable CPU scheduling behavior. It ensures that a critical thread such as a data-loading thread for your GPU doesn’t suddenly get migrated by the OS to a core on a different NUMA node in the middle of training or inferencing. In practice, it’s possible to see 5%–10% training throughput improve‐ ments just by eliminating cross-NUMA traffic and CPU core migrations. This also tends to reduce performance jitter and variance.
Many high-performance AI systems evaluate CPU simultaneous multithreading (SMT), or hyperthreading as it’s often called—and sometimes disable it for more pre‐ dictable per-core performance, but the benefit is workload-dependent. These systems may also reserve a handful of cores exclusively for OS background tasks by setting the isolcpus kernel parameter to isolate them from the general scheduler. You can also use Kubernetes CPU isolation for system daemons. This ensures that the remaining cores are dedicated entirely to training and inference threads and doing useful work.
By default, a process will allocate memory from the NUMA node of the CPU it’s cur‐ rently running on. So if you pin a process to NUMA node 0, its memory will natu‐ rally come from NUMA node 0’s local RAM, which is ideal. However, if the OS scheduler migrates threads—or if some memory was allocated before you did the pinning—you could end up with the nonideal scenario in which a process running in NUMA node 0 is using memory from NUMA node 1. To avoid this, the numactl --membind option forces memory allocation from a spe‐ cific NUMA node The general rule is to keep memory close to the CPU, which is close to the GPU.
As such, any worker subprocesses forked by your training script will automatically use the same NUMA memory binding. How‐ ever, they must be created using a fork-based model. If you switch to a spawn start method, or otherwise exec a new program, those child processes do not inherit the parent’s memory policy.
In addition, pinned memory, also called page-locked memory, is essential for efficient and direct GPU access. When memory is pinned, the OS won’t swap or move it. This leads to faster direct memory access (DMA) transfers.
You can test the data-transfer bandwidth between CPU memory and GPU memory using bandwidthTest --memory= from the installed CUDA utilities.
Deep learning frameworks provide options to use pinned memory for data loaders. For example, PyTorch’s DataLoader has a flag pin_memory=True
n short, you should make sure that your data loader uses pinned memory (e.g., pin_memory=True in PyTorch DataLoader) and that GPUDirect RDMA and GDS are enabled for supported hardware. This will reduce data transfer latency.
If you plan to use large, pinned buffers, ensure the ulimit value is high—or set it to unlimited. Otherwise the allocation might fail. Typically, one sets it to unlimited for large AI workloads and highperformance computing (HPC) applications.
Linux memory management typically uses 4 KB pages, but managing millions of tiny pages is inefficient when you have processes using tens or hundreds of gigabytes of memory, as in the case of deep learning data‐ sets, prefetched batches, model parameters, etc. Hugepages—2 MB or even 1 GB pages—can reduce the overhead of virtual memory management by making memory chunks bigger. The main benefits are fewer page faults and less pressure on the translation lookaside Buffer (TLB).
In scenarios with very large memory pools (e.g., preallocated pinned buffers for I/O), you may also consider explicit hugepages using vm.nr_hugepages or hugetlbfs for more deterministic performance.
Remember that, when using large, pinned memory regions, you should raise the ulimit -l setting (max locked memory) to a high value or unlimited. If this limit is too low, your attempt to pin memory can fail, leading to fallback on swappable memory—or out-of-memory (OOM) errors.
It’s important to note that THP’s background compaction can introduce unpredicta‐ ble pauses that are disastrous for latency-sensitive LLM inference workloads. Linux is configured by default to use THP to automatically allocate 2 MB pages whenever pos‐ sible. This is often sufficient, but it’s worth testing for your workload.
You can disable THP, but you will need to manually allocate and control hugepages. This will incur extra complexity, but it might be needed for low-latency workloads like inference. With THP disabled, your system will avoid stalls caused by kerneldriven defragmentations.
The modern consensus is to enable THP for most GPU-based training workloads in which throughput is important and to disable THP completely (transparent_hugepage=never)—or use madvise—for workloads like inference in which latency is important. This is also true for distributed training workloads in which many ranks (GPUs) allocate memory simultaneously.
On a busy system, you want to make sure that important threads such as datapipeline threads aren’t interrupted frequently. Linux by default uses the Completely Fair Scheduler (CFS) that works well for most cases.
But if you have a very latency-sensitive thread that feeds the GPU with data, for example, you could consider using real-time first in, first out (FIFO) or round-robin (RR) priority scheduling for that thread. This would make sure that the high-priority thread runs without being preempted by normal-priority threads. In practice, however, if you’ve pinned your threads to dedicated cores, you often don’t need to mess with real-time thread priorities.
Another option is to isolate cores or create separate CPU partitions to further reduce interruptions on these dedicated compute resources. To do this, you can use cset, kernel parameters like isolcpus and nohz_full, or cgroup cpuset isolation. With isolation, the OS scheduler leaves those CPU cores for you to use as you wish.
cgroup CPU and memory affinity is strongly recommended in production environments. Using these, each AI workload is isolated on its own physical cores and memory regions. This will prevent cross-workload contention and NUMA penalties. Tools like cpuset cgroups or container runtimes (docker --cpuset-cpus) should be used to enforce this.
You can assign each device’s hardware interrupts to cores on the same NUMA node. This will prevent cross-node interrupt handling that would otherwise incur extra latency and evict useful cache lines on a remote node.
In practice, performance-sensitive systems often disable the default irqbalance daemon or run it with bespoke rules. The other option is to manually set each interrupt’s affinity mask using /proc/irq/*/smp_affinity. By pinning every GPU and NIC interrupt to the nearest cores, you guarantee that those device interrupts are always serviced on the optimal NUMA node.
In short, the combination of dedicated cores, appropriate scheduling priorities, and NUMA-aware hardware interrupt bindings can help minimize jitter for data loading threads that are feeding the GPUs.
you should always try to avoid memory swapping to avoid seeing a catastrophic, multiple-orders-of-magnitude slowdown.
We recommend setting vm.swappiness=0, which tells Linux to avoid swapping except under extreme memory pressure. It effectively isolates your training job’s memory with cgroup limits to prevent any swapping.
You should use cgroups v2 through Docker or Kubernetes to pin memory and CPUs to the AI process. This will enforce NUMA affinity and no-swap policies in containerized environments.
You can also use sudo swapoff -a to temporarily disable all swap devices and files until the next reboot. Just make sure you have enough RAM for your workload—or put limits to prevent overcommit. Otherwise, the OOM killer may reap the process. Monitor swap usage using vmstat or free -m to make sure swap stays at zero.
A best practice for large training jobs is to write frequent checkpoints to disk in case you need to restart a failed job from a known good checkpoint. During checkpointing, however, huge bursts of data might fill up the OS page cache and cause stalls.
For storage, you can adjust vm.dirty_ratio and vm.dirty_background_ratio to tune the page-cache size for buffering writes. For example, with multi-GB check‐ points, using a higher dirty ratio lets the OS batch more data in RAM before flushing to disk. This will smooth out large checkpoint writes and reduce stalls in your train‐ ing loop.
Another option is to perform checkpointing in a separate thread. A more recent option in PyTorch is to write distributed checkpoint partitions from nodes across the cluster. In this case, the checkpoint partitions will be combined when the checkpoint is loaded after a failed-job restart.
In latency-sensitive training workflows, it’s best to bypass the page cache entirely. For example, open checkpoint files with O_DIRECT or use Linux’s io_uring for asynchro‐ nous I/O to avoid page-cache stalls. After writing each checkpoint, call posix_fadvise (fd, 0, 0, POSIX_FADV_DONTNEED) to immediately drop those pages from cache and prevent memory pressure on subsequent iterations.
By default, many compute nodes will run CPUs in a power-saving mode, which either downclocks a CPU or puts it to sleep when it’s idle this can ause extra latency when the system wakes the CPUs up again when new work arrives.
Configure the governor to performance. This can be done using cpupower frequency-set -g performance or in the Basic Input/Output System (BIOS).
Likewise, disabling deep C-states can keep cores from going into a low-power sleep state. C0 is active; everything above C0 represents a deeper state of sleep
In practice, many server BIOS/UEFI (Unified Extensible Firmware Interface) offer a high-performance profile that automatically sets the CPU governor to “Performance” and disables deep C-states
Bubbles are periods of time when the GPU is wait‐ ing for the CPU to resume data processing. By keeping the CPU ready, we reduce such hiccups. Many BIOSes for servers have a setting to disable C-states—or at least limit them
You should always turn off anything in your system that might introduce unpredictable latency, such as excess context switching, CPU frequency scaling, and memory-to-disk swapping. The result should be that your CPUs deliver data to the GPUs as fast as the GPUs can consume it, without the OS scheduling things on the wrong core or taking CPU cycles away at the wrong tim
On a well-tuned GPU server, CPU usage may not be very high since GPUs handle most of the computation. However, CPU usage should remain steady and in lockstep with GPU activity. The CPUs must stay busy preparing each incoming batch while the current batch is being processed by the GPU.
By tuning your host’s memory allocator (jemalloc or tcmalloc), you can eliminate unpre‐ dictable pauses in data preparation.
After tuning, you should see each GPU’s utilization hover near 100% and drop only at required synchronization barriers. The GPUs should never stall for data due to CPUside delays. With jemalloc, you can shard allocations into per-CPU arenas (narenas), enable background_thread for off-path purging, and lengthen dirty_decay_ms/muzzy_decay_ms so that freed pages aren’t immediately returned to the OS. This will minimize lock contention and fragmentation.
You can tune jemalloc with the MALLOC_CONF environment variable as follows:
export MALLOC_CONF="narenas:8,dirty_decay_ms:10000,muzzy_decay_ms:10000
,background_thread:true"
Similarly, tcmalloc benefits from tuning the TCMALLOC_MAX_TOTAL_THREAD_ CACHE_BYTES and TCMALLOC_RELEASE_RATE environment variables. These will provide larger per-thread caches so that small allocations avoid global locks and syscalls keeping CPU threads ready to feed the GPU with low, predictable latency. You can do this as follows:
export TCMALLOC_MAX_TOTAL_THREAD_CACHE_BYTES=$((512*1024*1024))
export TCMALLOC_RELEASE_RATE=16
In short, optimizing the allocator can reduce allocator overhead and fragmentation. This will keep CPU threads consistently fast and avoid unexpected stalls feeding the GPU. Experiment with these environment variables and tune them for your specific workload and environment
GPU persistence mode, the partitions of MPS, MIG, and a few other considerations like clock settings, ECC memory, and out-of-memory behavior.
Persistence mode is enabled by running the nvidia-persistenced daemon. This keeps the GPU driver loaded and the hardware in a ready state even when no application is active.
systemctl enable nvidia-persistenced
In Kubernetes environments, the NVIDIA GPU Operator can be configured to enable persistence mode on all GPUs automatically.
Normally, when multiple processes share a single GPU, the GPU’s scheduler timeslices between them. NVIDIA’s MPS is a feature that creates a sort of umbrella under which multiple processes can run on the GPU concurrently and without strict time-slicing. With MPS, the GPU can execute kernels from different processes at the same time as long as the GPU resources (streaming multiprocessors [SMs], Tensor Cores, etc.) are available. MPS essentially merges the contexts of the processes into one scheduler context. This way, you don’t pay the full cost of switching and idling between independent processes.
If you enable MPS for these inference jobs, the GPUs can interleave their work so that while one job is waiting on memory, another job’s kernel might fill the GPU, etc. The result is higher overall GPU utilization. In practice, if two processes each use 40% of a GPU, with MPS you might see the GPU at 80%–90% utilization serving both.
Setting up MPS involves running an MPS control daemon (nvidia-cuda-mpscontrol), which then launches an MPS server process that brokers GPU access. On modern GPUs, MPS is more streamlined as clients (the processes) can talk directly to the hardware with minimal interference from the compute node itself.
Another feature of MPS is the ability to set an active thread percentage per client. This limits how many SMs (GPU cores, essentially) a client can use. This can be useful if you want to guarantee quality of service (QoS) where two jobs, for example, each get at most 50% of the GPU’s execution resources. In this case, you can set CUDA_MPS_ACTIVE_THREAD_PERCENTAGE=50 to cap a client to about 50% of SM execution capacity. If not explicitly set, the jobs will just compete and use whatever GPU resources they can.
Note that MPS does not partition GPU memory, so all processes will share the full GPU memory space. MPS is mainly about compute sharing and scheduling. The issue is that one process could request a massive amount of GPU RAM, cause an OOM error on the GPU, and result in terminating all of the other processes running on the GPU. This is very disruptive.
By default, all MPS clients must run as the same Unix user since they share a context. In multiuser clusters, this means MPS is usually set up at the scheduler level such that only one user’s jobs share a GPU at a time. Otherwise, you can configure a system-wide MPS that’s shared by all users, but understand that the jobs are not isolated from a security standpoint. Prefer MIG when strong isolation is required.
One specific alternative to MPS is a feature for time-slicing GPUs in Kubernetes. Time-slicing on Kubernetes allows the device plugin to schedule different pods on the same GPU by time. For instance, if you configure a single GPU with a time-slicing replication factor of four, four pods on that GPU can each receive a time share.
Kubernetes time-slicing is sort of an automated time-sharing algorithm that doesn’t require MPS. However, this doesn’t overlap execution. Instead, it just switches more rapidly than the default driver would. Time-slicing may be useful for interactive workloads where you prefer isolation at the cost of some idle time. For hight hroughput jobs, overlapping with MPS or splitting the GPU with a MIG is usually better than fine-grained time-slicing.
Modern GPUs can be partitioned at the hardware level into multiple instances using MIG. MIG is a form of virtualization but done in hardware. This way, the overhead is very low—maybe a few percentdue to the loss of some flexibility. If one instance is idle, it can’t lend its resources to another, as they are hard partitioned. MIG allows a GPU to be sliced into as many as seven smaller logical GPUs--each with its own dedicated portion of memory and compute units, or SMs.
Administrators can enable or disable only the supported MIG profiles (e.g., 1g.23gb, 2g.45gb, 4g.90gb, etc.) on each GPU using tools like nvidia-smi -mig or using the NVIDIA Kubernetes GPU Operator’s nvidia.com/mig.config config map. Reconfiguring MIG requires draining workloads and invoking MIG’s dynamic reconfiguration capability to apply the changes.
Once a GPU is in MIG mode, modern GPUs can create and destroy MIG partitions dynamically without rebooting the entire system. You can adjust MIG instances on the fly after draining existing workloads, but to enable or disable MIG mode itself on a GPU, a reset of that GPU is needed.
Each MIG instance acts like a separate GPU from the perspective of software since it has its own memory, its own SMs, and even separate engine contexts.
For large-scale model training jobs and inference servers that span many GPUs, MIG is typically not useful since we want access to the full set of GPUs. On the other hand, for multitenant, small-model inference servers that can run smaller GPU partitions, MIG and its isolation features could be useful.
As of this writing, when a GPU is in MIG mode, GPU-to-GPU peer-to-peer communication (including VLink) is disabled. This applies to both NVLink and PCIe P2P across GPUs. MIG instances cannot engage in P2P with other GPUs. CUDA IPC across MIG instances is also limited. This can reduce distributed training throughput.
In short, enable MIG only when you need to run multiple independent jobs on the same GPU with strong isolation. Do not use MIG for large-scale distributed training or inferencing that spans GPUs, as you want access to the full power of the GPUs and their fast interconnects.
NVIDIA GPUs have something called GPU Boost, which automatically adjusts the core clock within power and thermal limits. Some users like to lock the clocks for consistency so that the GPU always runs at a fixed maximum frequency. Fixing the clock is extremely important when performing benchmarks since later runs may be throttled due to excessive heat.
This is mostly relevant during benchmarking to achieve deterministic and reproducible results.
ECC memory on GPUs is another consideration. ECC ensures that if there’s a singlebit memory error caused by cosmic rays, for example, the memory can be corrected on the fly. And if there’s a double-bit error, the error is detected and will throw an error to the calling code. ECC is usually enabled by default on NVIDIA data center GPUs. Disabling ECC can free up a small amount of memory since ECC requires extra bits for error checking. This might yield a marginal performance gain by reducing the overhead associated with on-the-fly error checking, but typically just a few percent. However, turning off ECC also removes critical memory-error protection, which can lead to system instability or undetected data corruption
For NVIDIA’s data center GPUs, including Hopper and Blackwell, ECC comes enabled by default and is intended to remain enabled to ensure reliable, errorcorrected computation and data integrity.
The only time you’d possibly consider turning it off is in a research setting where you are fine with taking the risk because you need that extra sliver of memory for your model to fit into your limited-memory GPU cluster.
Unlike CPU RAM, by default there is no such thing as GPU “swap” memory. If you try to allocate more GPU memory than available, you will get an unfriendly OOM error along with an even-unfriendlier process crash
There are a couple of mecha‐ nisms to mitigate this issue: allow memory to grow dynamically, embrace unified memory across CPU and GPU, and utilize memory pools and caching allocators.
By default, some frameworks (e.g., TensorFlow) grab all of the available GPU mem‐ ory at startup to avoid fragmentation and improve performance. If you don’t know this, it can be very bad in scenarios where you are sharing the GPU. PyTorch, by default, allocates GPU memory only as needed.
CUDA’s Unified Memory system lets you allocate memory without predefining whether it resides on the CPU or GPU. The CUDA Runtime handles moving pages as needed. Modern NVIDIA GPUs like Hopper and Blackwell include hardware support for on-demand paging using the Page Migration Engine (PME)
However, while PME provides flexibility, relying on it can introduce performance penalties compared to having enough GPU memory for your workload. This mechanism is mostly a convenience for practitioners trying to run models that don’t fit into GPU RAM.
If you run into the GPU OOM error, which you surely will at some point, it’s likely caused by memory fragmentation or excessive memory caching. You can try to clear the cache using PyTorch’s Torch.cuda.empty_cache(), but it almost always means your workload legitimately needs that much memory.
PyTorch also provides tools like torch.cuda.memory_stats() and torch.cuda.mem ory_summary() to help diagnose fragmentation by showing allocated versus reserved memory. NVIDIA’s Nsight Systems also shows GPU memory usage patterns to help identify memory leaks, long-lived allocations that correlate with leaks, CPU-GPU interconnect activity, and GPUDirect Storage timeline tracing. Additionally, the Nsight Compute profiler provides low-level kernel analysis, including occupancy, throughput, and NVLink usage. We’ll cover all of these in the upcoming chapters.
Docker provides the --gpus flag to select and expose GPUs to a container, but it does not support setting a GPU memory limit. If you need hard isolation for GPU memory or compute, use MIG to partition the device or use Multi-Process Service (MPS) with active thread percentage for fair sharing. Configure limits in Kubernetes using MIG resources like nvidia.com/mig-2g.45gb when you require strict partitioning.
In general, running out of GPU memory is something you can manage at the application level. For instance, you can reduce the data batch size, model weight precision, or even the model parameter count, if that’s an option.
A best practice is to monitor GPU memory usage with nvidia-smi or NVML APIs during model training and inferencing. If you’re close to the memory limit, consider workarounds like reducing batch size, using activation checkpointing for training, or other techniques to lower memory usage.
Also, you should ensure that your CPU memory isn’t being swapped, as this would indirectly hurt your GPU utilization and goodput because each time your GPU tries to fetch something from the CPU host, but the host memory page has been swapped to disk, your performance will be bottlenecked by the much slower disk I/O.
The only time you want to unload the GPU driver is for troubleshooting or upgrading the driver.
For modern GPUs running with the latest NVIDIA Container Toolkit, GPU performance within a properly configured environment is virtually identical (< 2% difference) to running the code directly on the bare-metal host outside of the container.
One challenge when using containers with GPUs is making sure that the CUDA libraries inside the container match the driver on the host.
The general rule is that the host’s NVIDIA driver version must be at least as recent as the minimum driver version required by the CUDA version inside the container.
Alternatively, NVIDIA’s container runtime can actually inject the host driver libraries into the container at runtime, so you don’t even need to ship the NVIDIA driver inside the image.
The main difference when running in a Docker container versus running directly on the host might be in I/O. Containers often use a union filesystem that transparently overlays multiple underlying filesystems, like the host filesystem and the container filesystem, into a single, unified view.
In a union filesystem such as OverlayFS, files and directories from multiple sources will appear as if they belong to one filesystem. This mechanism is especially useful for containers, where the read-only filesystem from the base image layer is combined with a writable container layer. There is some overhead when using an overlay filesystem, however.
Model training often involves heavy I/O operations when reading datasets, loading a model, and writing model checkpoints. To work around this, you can mount a host directory—or network filesystem into the container using bind mounts.
Bind mounts bypass the overlay and therefore perform similarly to disk I/O directly on the host. If the host filesystem is something like an NVMe SSD or an NFS mount, you get the full performance of that underlying storage device. We purposely do not package a multi-terabyte dataset inside the image. Instead, we bring the data in through the mounts.
In fact, it’s a best practice to avoid heavy data reads/writes against the container’s writable layer. Instead, mount your data directory and output directory from the host into the container. You want to ensure that I/O is not bottlenecked by the overhead of the container’s CoW mechanism.
The NVIDIA device plugin for Kubernetes is a lightweight component that advertises GPU hardware to the schedu‐ ler. It mounts those device nodes into your pods when you request nvidia.com/gpu under resources.limits and optionally under resources.requests, if you want to set both explicitly. This way, when you deploy a container on Kubernetes with this device plugin, Kubernetes takes care of making the GPUs available to the container. The device plugin is topology aware, as well. This means it can prefer to allocate mul‐ tiple GPUs from the same NVLink Switch or the same NUMA node for a given pod.
When using Kubernetes to orchestrate GPU-based containers, you want it to allocate resources to containers in a manner that is aware of the hardware topology, including the NUMA node and network bandwidth configurations. However, by default, Kubernetes is not topology-aware. It treats each GPU as a resource but doesn’t know if GPU 0 and GPU 1 are on the same NUMA node or if they use the same NVLink interconnect. This could make a big difference.
To avoid resource contention, you should try to either reserve the resources that you need or request the entire node for your job. For the container/pod placements, you should align pods with CPU affinities and NUMA nodes using the Kubernetes Topol‐ ogy Manager component to bind the container’s CPUs to the same NUMA node as the GPUs that the container was allocated.
Kubernetes Topology Manager can provide detailed topology information. For example, it can detect that GPU 0 is connected to NUMA node 0, NVLink domain A, and PCIe bus Z. The Kubernetes scheduler can then use this information to allocate containers to GPUs in an optimal way for efficient processing and communication.
Topology-aware GPU scheduling is still maturing. In many clusters, administrators explicitly label nodes using Kubernetes labels to capture the GPU and system topology. These labels ensure that multi-GPU pods land on servers whose GPUs share the same NVLink interconnect or reside within the same NUMA domain.
For our purposes, if you’re running multi-GPU jobs in Kubernetes, make sure to enable topology-aware scheduling. This typically involves configuring --topologymanager-policy to best-effort, restricted, or, in some cases, single-numanode. This policy configuration helps multi-GPU and CPU + GPU workloads achieve lower latency by avoiding remote memory access. This complements the OS-level NUMA tuning.
When scheduling collective-heavy training, prefer placements that keep traffic inside the fast NVLink domain before crossing the slower network fabric.
Commonly, the Simple Linux Utility for Resource Management (SLURM) is used for training clusters, while Kubernetes is typically favored for inference clusters. However, hybrid solutions have emerged that integrate SLURM with Kubernetes. The open source Slinky project is an example solution to simplify cluster management across training and inference workloads
If not properly set, a scheduler might treat all GPUs as identical and provide nonideal allocations for your multi-GPU container requests. Proper configuration can avoid unnecessary cross-NUMA-node and cross-NVLink GPU communication overhead.
This single-node constraint can cause pods to never run—even if the combined MIG resources can be found across different nodes of the cluster. The request can be satisfied only if the requested MIG resources are available on a single node
Using host networking allows a container to access the InfiniBand interconnect exactly as the host does—without any additional translation or firewall layers. This is particularly useful for MPI jobs because it eliminates the need to configure port mappings for every MPI rank.
However, if host networking is not an option due to security policies, you must ensure that your Kubernetes container network interface (CNI) and any overlay network can handle the required traffic. In such cases, you may need to open specific ports to support the handshake of NCCL and data exchange, using environment variables like NCCL_PORT_RANGE and NCCL_SOCKET_IFNAME to help establish connections.
When using a Kubernetes environment and you want to enable RDMA, consider installing the Kubernetes RDMA device plugin from Mellanox.
To safeguard against resource contention, Kubernetes lets you define resource requests and limits for pods.
These limits are enforced using Linux cgroups, so if your container exceeds its allocation, it can be throttled or even terminated by the OOM killer. It’s common practice to use resource requests—and optionally the CPU Manager feature to pin cores—to ensure that performance-critical jobs get exclusive access to the necessary CPU resources so that other processes cannot steal CPU time from your reserved cores.
Ideally, a GPU node is fully dedicated to your job. However, if it’s not, you should ensure that the node is carefully partitioned using Linux cgroup controllers for I/O and CPU so that other workloads don’t interfere.
Fortunately, Kubernetes supports CPU isolation, which ensures that pods get the dedicated CPU cores and memory they request—and prevents other pods from being scheduled on the same CPU core as yours. This avoids extra overhead from context switching and resource contention.
In practice, performance-sensitive Kubernetes jobs should request all of the CPUs and GPUs of a given node so that nothing else interferes or contends with the jobs’ resources. Easier said than done, but this is the ideal job configuration from a performance and consistency standpoint.
If an unbounded container uses too much memory on the host, the infamous Linux “OOM killer” will start killing processes—and potentially your Kubernetes job—even if your job wasn’t the one using too much memory.
In Kubernetes, a Pod with no requests/limits is treated as BestEf fort and is the most likely to be evicted. To obtain Guaranteed QoS, every container must set requests == limits for both CPU and memory. Setting a high limit alone will result in a Burstable QoS, not Guaranteed
As of this writing, Kubernetes does not offer native, first-class I/O isolation out of the box, unfortunately. While Linux does support I/O controls using cgroup controllers, Kubernetes itself does not automatically enforce I/O limits in the same way it does for CPU and memory.
If you need to ensure that heavy I/O workloads on a GPU node don’t interfere with one another, you might need to manually configure I/O controls at the node level. This can involve adjusting the cgroup v2 I/O controller or using other OS-level con‐ figurations to partition I/O resources.
It’s a good idea to always ensure that the host machine is tuned since containers can’t change kernel parameters like hugepage settings or CPU governor limits.
By tuning bucket sizes and scheduling these transfers appropriately, one can achieve a higher degree of overlap and prevent communication delays from stalling the com‐ pute pipeline. Tools such as the PyTorch profiler and NVIDIA Nsight Systems offer insight into whether your computation and communication are overlapping, allowing engineers to adjust these parameters for maximal efficiency
By combining larger batch sizes, gradient accumulation, asynchronous transfers, compression, and bucketing into one cohesive strategy, large, distributed AI models can overcome network limitations and reduce idle time. This design minimizes syn‐ chronization events while achieving high throughput and optimal GPU utilization
AI frameworks hide most of this complexity. PyTorch’s DistributedData Parallel automatically installs hooks on the backward pass so that each gradient bucket triggers an asynchronous NCCL all-reduce on a dedicated communication CUDA stream, while the default CUDA stream continues computing gradients for subsequent layers
To maintain proper overlapping, avoid unnecessary synchronization points with torch.cuda.synchronize() or inadvertently triggering a full device sync by moving tensors to the CPU with torch.Tensor.item().
Bucketing, as implemented in PyTorch’s Distributed Data Parallel (DDP) communication mechanism, also reduces per-call overhead by grouping many small tensors into larger messages. However, bucket sizing is a trade-off. Very large buckets maxi‐ mize bandwidth utilization but delay the start of communication since you wait for more gradients to accumulate before kicking off the all-reduce. Very small buckets start transfers earlier but incur more overhead due to many small NCCL calls. As of this writing, the default bucket size in PyTorch DDP is 25 MB. This is a balance that overlaps well in most cases. However, if you have a model with very large layers, you might increase this to reduce overhead. If you have a model with many small layers, you might actually benefit from smaller buckets to start communication sooner. Ultimately, achieving maximal overlap may require profiling different bucket sizes to see which yields the best iteration time.
The takeaway is that a well-tuned DDP should overlap most of the gradient commnication with computation.
PyTorch’s DDP’s default overlap strategy is often described as wait-free back-propagation (WFBP), which bucketizes gradients and launches reductions as soon as each bucket is ready
Avoid operations that inadvertently move tensors from the GPU to the CPU (e.g., calling .item() on a tensor) until you’re sure that all asynchronous GPU work is finished. Otherwise, you will force a synchronization, stall the computation, and slow down your training or inference workload. This typically happens when adding print() or log() statements for debugging. These can be disastrous for performance.
Also, manual calls to torch.cuda.synchronize() should be minimized and used only for accurate benchmarking—or when required for correctness. Otherwise, they will serialize GPU work and negatively impact performance
DDP’s design and PyTorch’s operations are already asynchronous and handle dependencies correctly. Explicit synchronization is rarely needed in user code.
Prefer RDMA paths where available and verified. And always confirm (and continuously reconfirm) with logs and micro-benchmarks that the RDMA data path is active.
n container environments like Docker and Kubernetes, ensure the container has direct access to the host’s InfiniBand devices (e.g., /dev/infiniband). Otherwise, NCCL may silently fall back to TCP sockets instead of GPUDirect RDMA—and without any obvious errors to highlight the degradation. This results in throughput dropping from tens of GB/s to only a few Gb/s, with no obvious error messages.
A related container pitfall arises when the container’s GID assignments don’t match the host, as in some “rdma-shared” Docker images. This prevents GPUDirect registration and uses CPU-driven RDMA copies instead of using true GPU-based RDMA.
Always verify that it is true GPUDirect RDMA. Confirm that the kernel module is loaded with lsmod | grep nvidia_peermem, and check dmesg for initialization. For an end-to-end check, run NCCL with NCCL_DEBUG=INFO to confirm NET/IB paths and use RDMA perftests with --use_cuda to validate GPU-to-GPU transfers. Verifying will help prevent stealthy performance degradations
- Understand the topology: Use nvidia-smi topo -m to get a basic GPU interconnect view. For NVSwitch- and NVLink-based systems also use nvidiasmi nvlink or Nsight Systems to understand multihop switch fabric connectivity.
- Leverage NVLink Switch: Make sure jobs are placed within the same NVLink domain to fully utilize this ultrafast interconnect.
- Make sure to use RDMA
- Aggregate bandwidth with multiple NICs: NCCL can stripe traffic across multiple NICs (called multirail) to increase bandwidth.
- Utilize optimized “direct NIC”: Favor high-bandwidth, multirail NIC configurations that give each GPU or small groups of GPUs sufficient dedicated network bandwidth. With modern GPU systems, NCCL supports GPU-initiated networking with InfiniBand GPUDirect Async (IBGDA) and the direct NIC path without CPU intervention.
- Check for misconfiguration: A common pitfall is a mismatch in network configuration that causes a fallback to a slower path. Tools like NCCL’s debugging output and network interface counters (ibstat, ifstat) can help verify which interface is being used more heavily. For modern systems with large 200–400 Gbps paths, dropping to 10 Gbps would cause a severe bottleneck.
Gloo pythorch backend uses CPUs and TCP sockets. NCCL is the preferred backend for NVIDIA GPUs. If we were trying to allreduce 400MB, we would observe it taking 200ms (2GB/s)-- much lower than infiniband limit of 100GB/s, and also we would see near 100% CPU utilization. You can verify the backend in PyTorch by calling torch.distributed.get_backend(). In a production cluster environment with multiple NICs, you should explicitly set NCCL_SOCKET_IFNAME=ib0 so that NCCL’s initial TCP handshake runs over the InfiniBand host channel adapter (HCA). This ensures it bootstraps correctly and then hands off to GPUDirect RDMA on the fastest path. Ensure that all nodes can reach one another over the selected interconnect.
If you run PyTorch’s bundled NCCL (e.g., torch.cuda.nccl.version() == ()) against a different version of the system-installed libnccl, you will hang the system or fall back to a slower implementation. Make sure you have alignment by matching nvidia-nccl-cu* packages or rebuilding PyTorch against the system NCCL.
NCCL uses ephemeral TCP ports for its out‐of‐band setup, and if your OS’s net.ipv4.ip_local_port_range is too narrow, you can exhaust available ports, caus‐ ing failed or stalled handshakes. It’s recommended that you widen your port range in /proc/sys/net/ipv4/ip_local_port_range (e.g., 50000 51000) to avoid hidden bootstrap failures.
When profiling your workload under these unfortunate conditions, you will observe that scaling to multiple nodes significantly slows down training. In other words, the “per-GPU throughput” will drop. In this case, check the network links. Monitor the network throughput using nvidia-smi dmon, for instance, to collect NVLink/PCIe/Network statistics. You can also use built-in tools like ethtool -S or ip -s link show for byte/packet counters, or launch interactive monitors such as iftop or nload to watch live NIC throughput.
You can also try to utilize multiple interfaces, if available. If you’re saturating an 800 Gbps (100 GB/s) InfiniBand link, for instance, and your job needs more network throughput, consider enabling NCCL’s multi-NIC support—assuming that you have multiple NICs. Make sure that NCCL_NSOCKS_PERTHREAD and NCCL_SOCKET_NTHREADS are tuned, as these control how many parallel connections and threads NCCL uses for network transfers. In cases with multiple NICs, increasing these environment variable values from their platform-dependent defaults can help utilize both NICs
Remember that the product of threads and sockets should not exceed 64 per NVIDIA guidance since more threads mean more CPU usage. Increase these thread-related settings stepwise (e.g., 2 → 4 → 8), and continuously measure the throughput. Too many threads will contend for resources and potentially diminish returns.
In multinode training, the slowest node, or GPU, will determine the overall pace because synchronization needs to wait for every node and GPU to respond. Using monitoring tools like NVIDIA’s DCGM or InfiniBand counters on each node can help spot if one node has degraded performance due to NIC link flapping or GPU thermal throttling. It’s also useful to use collective profiling tools such as PyTorch’s torch.distributed.monitored_barrier to identify if a particular rank is consistently lagging
# barrier_straggler
import torch
import torch.distributed as dist
import os
import datetime
def run(rank, world_size):
dist.init_process_group(backend="nccl", init_method="env://")
local_rank = int(os.environ["LOCAL_RANK"])
torch.cuda.set_device(local_rank)
# ... your forward/backward work here ...
# Before syncing at end of iteration, use a monitored barrier:
try:
# Wait up to 30 seconds for all ranks
# if one lags, you’ll get a timeout on that rank
dist.monitored_barrier(timeout=datetime.timedelta(seconds=30))
except RuntimeError as e:
print(f"Rank {rank} timed out at barrier: {e}")
# Now proceed knowing all ranks are roughly in sync
dist.destroy_process_group()
Here, dist.monitored_barrier(timeout=datetime.timedelta(seconds=30)) will raise an error on any GPU that doesn’t arrive within 30s. This will help you pinpoint stragglers. Combine this with NCCL_DEBUG=INFO and NCCL_ASYNC_ERROR_HANDLING=1 to get both PyTorch and NCCL logs around which rank or link is slow.
PyTorch’s caching allocator holds onto GPU memory across iterations. In distributed settings using UCX/RDMA, these long-lived allocations can exhaust registration pools or fragment memory, causing sporadic allocation failures or performance cliffs. Monitoring torch.cuda.memory_reserved() versus memory_allocated() helps surface these edge cases:
NVIDIA NCCL is a many-to-many communication library for operations, called collectives, used by groups of GPUs to share data. NCCL underpins most multi-GPU training workloads in NVIDIA’s ecosystem.
NCCL provides optimized implementations of collective communication operations like all-reduce, all-gather, broadcast, and reduce-scatter that scale from a few GPUs to many thousands and, someday, millions.
While NCCL can use a simple pattern communication like ring all-reduce to commu‐ nicate with each link equally, it will automatically use a topology-aware hierarchical communication pattern to maximize communication performance. For systems with multiple NUMA node domains, for instance, NCCL might first do an intranode reduce, then a cross-node reduce, then an intranode broadcast, which is effectively a hierarchical all-reduce.
It is possible to override NCCL’s algorithm selection with the environment vari‐ able NCCL_ALGO (e.g., NCCL_ALGO=NVLS,NVLSTree,Tree,Ring,PAT, etc.), but generally NCCL does a good job of automatically choosing the best path based on the topology. Manual override is usually only for specific situations like troubleshooting, research experiments, and more.
Tools like NVIDIA Nsight Systems—or NCCL’s own traces with NCCL_DEBUG=INFO and NCCL_TOPO_DUMP_FILE=—will show if NVLink paths are being utilized fully.
Internally, NCCL can employ different communication algorithms depending on the size of data, number of GPUs, and topology. The primary algorithms NCCL uses for collectives are Ring, Tree, CollTree, CollNet, and Parallel Aggregated Tree (PAT).
- Ring: bandwidth-dominated workload
- Tree and NVLSTree: latency-dominated workload
- CollTree (hierarchical tree collectives): preferred when cross node latency dominates
- CollNet (hierarchical collectives across nodes): very large, multinode GPU clusters
- Parallel aggregated tree (PAT): The result is near–ring-level throughput for large data transfers plus tree‐level latency advantages for smaller segments.
As when choosing any communication algorithm, the choice of NCCL algorithm typ‐ ically comes down to message size and topology. Small messages (on the order of 10s of megabytes) favor tree algorithms since there are fewer steps. Large messages favor ring algorithms because they provide better bandwidth utilization.
if profiling your workload reveals suboptimal communication, such as unexpectedly high cross-node latency, you can override the communication algorithm on a case-by-case basis by setting the NCCL_ALGO environment variable. This will force NCCL to use a particular algorithm on that communicator. If setting this variable in code, make sure to do it before calling ncclCommInitRank()
The key is to overlap communication with computation at every level. This includes using NCCL for all-reduce and NIXL for one-to-one transfers. Using these mecha‐ nisms, you can scale to thousands and millions of GPUs with high efficiency
Other techniques like gradient accumulation and activation check‐ pointing are also critical at ultrascale to manage the memory foot‐ print without sacrificing throughput.
When scaling to multiple GPUs on a single node, PyTorch offers both data-parallel (split the data) and model-parallel (split the model) approaches at the framework level.
let’s compare two of the most basic data-parallel strategies from a systems performance standpoint: nn.DataParallel (DP) and torch.distributed.DistributedDataParallel (DDP). It’s important to understand their differences as choosing the wrong one can severely impact performance:
- Data parallelism (DP): DP is an easy-to-use API that involves a single process, or single Python thread, controlling multiple GPUs.
- Fully sharded data parallelism (FSDP): FSDP avoids full model replicas by sharding activations, gradients, and parameters across GPUs, greatly reducing memory overhead.
- Distributed Data Parallel (DDP): DDP uses one process per GPU device and relies on NCCL to communicate gradients. Like most simple data parallel strategies (FSDP being the exception), each process has its own copy of the model.
- Creating NCCL communicators too often: init_process_group call is designed to be called once at startup, you should avoid any design that reinitializes it on every iteration.
- Do not create and destroy NCCL communicators on every iteration: create the subcommunicators once at the beginning using PyTorch’s torch.distributed.new_group() and reuse these communicators. If you need to create multiple communicators because, for instance, you have a dynamic runtime membership scenario or a staged initialization, NCCL provides a C++ API to initialize multiple communicators together using ncclGroupStart(), ncclCommInitRank(...), and ncclGroupEnd(). PyTorch does not support fully dynamic membership changes at runtime without a full communicator teardown. All ranks must invoke creation and destruction calls in lockstep to prevent hangs.
- Avoid overtuning or disabling NCCL features with environment variables
- Verify CPU-GPU NUMA-node affinity for NCCL threads: The recommended approach is to bind each GPU process to the CPU cores for its NUMA domain and then set NCCL_IGNORE_CPU_AFFINITY=1 so that NCCL can fine-tune thread placement within those cores. PyTorch’s launch utilities handle much of this automatically, but it’s good to verify.
- Resist the temptation to ignore NCCL warnings and errors
- NCCL communicator hangs, errors, or shuts down completely: NCCL supports asynchronous error handling and failover for cases like network errors
NIXL was designed specifically to accelerate large-scale LLM distributed and disaggregated inference. NIXL is a core component of NVIDIA’s open source Dynamo inference engine. NCCL remains the standard for many-to-many collective operations common in large-scale training such as all-reduce. NIXL, however, targets one-to-one or one-to-few data transfers that are common in large-scale inference such as moving KV cache data.
The inference path of a transformer-based model is actually split into two different stages: prefill and decode.
- Prefill, is often compute bound as it uses many matrix multiplications to build the KV cache from the incoming request data (aka prompt).
- Decode, is often memory-throughput bound, as it needs to gather the model weights from GPU HBM memory to calculate the next set of tokens (aka completion or response).
This prefill/decode split is implemented in common inference engines vLLM, SGLang, and NVIDIA’s Dynamo and TensorRT-LLM. The prefill (prompt ingestion) creates the KV cache, and the decode (generation) uses this cache. NIXL specifically accelerates the transfer of the KV cache between nodes in this workflow.
The traditional setup has each GPU node handle both the prefill (compute-bound) and decode (memory-bound, I/O-bound) phases. The disaggregated serving configuration places the prefill workers in the GPU cluster and the decode workers in another GPU cluster. A GPU in the prefill cluster generates the KV cache for the input sequence and uses NIXL to transfer it to a GPU in the decode cluster. This specialization produces higher overall throughput and advanced scaling configurations.
In such cases, the KV cache, which can run into tens of gigabytes in a long prompt, must move seamlessly from one processing unit to another in near-real time. This way, the text generation happens at speeds that are unnoticeable to end users.
NIXL provides a direct channel for transferring data from one GPU to another or a small group of GPUs across compute nodes and even across racks. The system looks at the available pathways and always selects the one that gets the data there the quickest.
NIXL offers a straightforward API. You post a transfer request with a pointer to the data and a destination—either GPUs, CPUs, or storage targets like Amazon S3. NIXL will transfer that data as fast as possible. You register memory with regis terMem, obtain transfer descriptors with trim, prepare a nonblocking request with prepXfer, and submit it with postXfer. NIXL chooses whether to perform a direct PCIe or NVLink copy, an RDMA transfer, or a storage path such as GPUDirect Storage. The NIXL library is nonblocking and returns a request handle that you poll with checkXfer to detect completion.
A nixlAgent is NIXL’s core transfer object. It encapsulates the endpoint configuration, memory registrations, and backend selection. It also manages metadata, connection information, and asynchronous transfer requests to and from other agents. You need two agents for a transfer because each nixlAgent instance represents one endpoint in the transfer. The source agent (agentSrc) encapsulates the context, memory registrations, and backends for the origin of the data. The destination agent (agentDst) does the same for the receiver side.
NIXL is not a replacement for NCCL but a complement. NCCL still handles synchronized collectives for GPUs working on a single task/stage in parallel, such as an all-reduce split across multiple GPUs. NIXL, on the other hand, performs asynchronous data transfers between tasks/stages—or between distinct components (e.g., GPUs, CPUs, storage) in a distributed system.
- NCCL (collective communication) Primary use case: Many-to-many collectives (e.g., all-reduce, all-gather) for tightly coupled GPU groups in training
- NIXL (point-to-point communication) Primary use case: One-to-one or one-to-few transfers (e.g., sending large tensors or caches) for distributed inference or pipelining.