Parallelism

One GPU is never enough for a frontier model. The whole art of scaling comes down to a single question: what do you split, and what does it cost to stitch back together?

Why split at all

A big model runs into two walls on one GPU. It may not fit: the weights, gradients, and optimizer state overflow the memory. And it may be too slow: one GPU cannot chew through a large batch or a long sequence fast enough.

Every strategy is a different answer to what to split across GPUs. Split the batch, and each GPU trains on different data. Split a layer, and each GPU holds a slice of its matmul. Split the layer stack, the experts, or the sequence. The catch is always the same: once you split, the GPUs must talk, and communication is the tax you pay.

Pick a strategy below. The colored artifact is the one being cut, one color per GPU. Anything hatched is copied on every GPU. Then read the price: it is always some flavor of communication.

Splits the batch.

GPU 0GPU 1GPU 2GPU 3 replicated

Batch

Model

Sequence

Communication all-reduce (gradients) once per step, after the backward pass

Saves memory on nothing: every GPU keeps a full copy of the model

The cost each GPU stores the whole model, gradients, and optimizer state

↳ Composes as: the outer loop, wrapped around everything else. DDP is the efficient PyTorch version.

Data parallel (DP / DDP) splits the batch. Communication: all-reduce (gradients), once per step, after the backward pass. It saves nothing: every GPU keeps a full copy of the model. The cost: each GPU stores the whole model, gradients, and optimizer state.

Show as data

Parallelism strategies: what each splits and what it costs
Strategy	Splits	Communication	Cost
DP / DDP	the batch	all-reduce (gradients)	each GPU stores the whole model, gradients, and optimizer state
FSDP / ZeRO	the batch and the model state	all-gather (params) + reduce-scatter (gradients)	more communication than DDP, gathering each layer just in time
TP	each layer's matmul	all-reduce (activations), twice per layer	chatty, latency-sensitive communication that must be fast
PP	the stack of layers	point-to-point (activations between stages)	the pipeline "bubble": stages sit idle while it fills and drains
EP	the experts	all-to-all (route tokens to their expert and back)	all-to-all traffic and load imbalance (see the MoE scene on the kernels page)
SP / CP	the sequence	ring / all-gather (passing K and V blocks around)	extra communication to pass keys and values around the ring

The five things you can split

Underneath the names, there are only a handful of axes. Everything else is a combination.

The batch → data parallel (DP / DDP). Every GPU keeps a full copy of the model and trains on a different slice of the batch. After the backward pass they average gradients with an all-reduce. Simple and fast, but the whole model sits on every GPU.
The model state → FSDP / ZeRO. The one people forget. It is still data parallel (the batch is split), but it also shards the params, gradients, and optimizer state across the same GPUs, gathering each layer's weights just in time. This is how a model that will not fit under plain DDP still trains. The price is more communication.
A layer's matmul → tensor parallel (TP). Cut each weight matrix across GPUs (Megatron splits attention heads and MLP columns). Every GPU does part of every layer and they all-reduce the activations. Very chatty, so it wants the fast links inside a single node.
The stack of layers → pipeline parallel (PP). Put the first layers on GPU 0, the next on GPU 1, and so on. Activations flow stage to stage. The cost is the bubble: while the pipeline fills and drains, some stages sit idle. Micro-batches shrink it.
The experts → expert parallel (EP). For mixture of experts: put different experts on different GPUs and all-to-all the tokens to wherever their expert lives. This is what DeepEP accelerates.
The sequence → sequence / context parallel (SP / CP). Split the sequence itself. SP splits the parts TP leaves whole (the norms and residuals) to save activation memory; CP (ring attention) splits attention over an ultra-long context, passing keys and values around a ring.

Composing them: 3D parallelism

Real training runs stack several at once, because the axes are independent. The classic recipe is 3D parallelism: tensor parallel inside a node (over fast NVLink), pipeline parallel across nodes, and data parallel (usually FSDP) wrapped around the whole thing. Add expert parallel for a mixture-of-experts model and context parallel when the sequences get very long, and you reach the 4D and 5D setups the largest runs actually use.

The ordering is not arbitrary. It follows the communication cost: put the chattiest strategy (TP) on the fastest link, and the cheapest (DP, one sync per step) on the slowest. Llama 3's 405B run is a clean example: 4D parallelism ordered TP, then CP, then PP, then DP from innermost to outermost, with TP capped at 8 so it never leaves a single NVLink node.

The device mesh

Frameworks arrange all of this as a device mesh: a grid of GPUs where each axis runs one strategy. The factors along the axes must multiply to your total GPU count. The mesh's order is what decides which groups land on fast NVLink inside a node and which cross the slower network between nodes.

Set the hardware and a factor per axis. Then pick an axis to see its communication groups: green cells that share a color talk to each other. Watch whether a group stays inside one node (fast) or spills across nodes (slow). The leftover GPUs become plain data-parallel replicas.

GPUs / node

Nodes

= 8 GPUs

2 auto

node 0

node 1

Show groups for

FSDP (Sharded data parallel) shards params, gradients, and optimizer state. It forms 4 groups of 2, talking by all-gather + reduce-scatter. Every group sits inside one node, so it runs on fast NVLink.

This is HSDP. FSDP shards the model inside each node, and DP replicates it across nodes, keeping the heavy shard traffic on NVLink and paying the slow link only for a cheap gradient sync.

Show as data

Device mesh: 8 GPUs, 2 x 4. Placement inner to outer.
Dim	Factor	Groups	Links	Comm
TP	2	4	intra-node	all-reduce
CP	1	—	—	—
FSDP	2	4	intra-node	all-gather + reduce-scatter
DP	2	4	inter-node	all-reduce (grads)
PP	1	—	—	—

Two compositions worth knowing

HSDP (hybrid sharded data parallel) is FSDP and DP stacked on one 2D mesh. FSDP shards the model within a group (say, one node), and the whole sharded model is replicated across groups (across nodes). Why bother? FSDP's all-gather and reduce-scatter get slower as the group grows, so at large scale pure FSDP collectives dominate. HSDP keeps those heavy collectives on fast intra-node NVLink and pays the slow inter-node link only for a cheap gradient all-reduce, once per step. Pure FSDP is just HSDP with the replicate factor set to 1. You saw it above: shard inside the node, replicate across.

FSDP + EP is how a mixture-of-experts model shards its experts. The order is experts first, then FSDP: expert parallel splits the set of experts across GPUs so each holds only a few, and then FSDP shards each of those local experts again across the remaining data-parallel group. The result is a 3D mesh (an expert-shard axis, an FSDP-shard axis, and a replicate axis). Sharding expert-first keeps each expert's all-to-all token routing local, then reclaims memory with FSDP on top. PyTorch's TorchTitan is the reference implementation for stacking all of these.

Cheat sheet

Parallelism strategies at a glance
Strategy	Splits	Communication	Reach for it when
DP / DDP	the batch	all-reduce (grads)	the model fits; you want speed
FSDP / ZeRO	batch + model state	all-gather + reduce-scatter	the model barely fits
TP	each layer's matmul	all-reduce (activations)	a layer is too big; you have NVLink
PP	the stack of layers	point-to-point	the model spans many nodes
EP	the experts	all-to-all	it is a mixture of experts
SP / CP	the sequence	ring / all-gather	the context is very long

Thread: The smallest unit of work on a GPU. One thread runs one instance of the kernel on one lane.
Warp: A group of 32 threads that execute the same instruction together, in lockstep. The scheduling unit of the GPU.
SIMT: Single Instruction, Multiple Threads. All 32 threads of a warp run one shared instruction over their own data.
Lockstep: All threads in a warp advance together on the same instruction at the same time.
Coalescing: When a warp reads neighbouring addresses so the hardware serves them in as few memory transactions as possible.
Sector: The 32-byte unit the hardware fetches from global memory. A warp wants its data packed into as few sectors as possible.
HBM: High Bandwidth Memory, the large off-chip global memory. Biggest and slowest tier, hundreds of cycles away.
Shared memory: Fast on-chip scratchpad private to one thread block. Split into 32 banks.
Register: Per-thread on-chip storage. The fastest memory, about one cycle to access.
L2 cache: On-chip cache shared by all SMs, sitting between the per-SM L1 caches and global HBM.
SM: Streaming Multiprocessor. A core building block of the GPU that runs thread blocks. An A100 has 108 of them.
Sub-partition: One of the four processing blocks inside an SM. Each has its own warp scheduler and execution units.
Warp scheduler: The unit that picks one eligible warp each cycle and issues its next instruction. Four per SM.
Eligible: A warp that is ready to issue this cycle, not waiting on memory or a dependency.
Stalled: A warp that cannot issue yet because it is waiting, usually on a memory load.
Latency hiding: Keeping the machine busy during long waits by running other ready warps while one warp stalls.
Occupancy: The number of active warps on an SM divided by the maximum it can hold. More warps give the scheduler more to switch to.
Bank: One of the 32 slots shared memory is split into. Consecutive 4-byte words map to consecutive banks (word w lands in bank w mod 32), and each bank serves one word per cycle.
Bank conflict: When two or more threads in a warp want different words in the same bank. Their reads serialize.
Tiling: Loading a small block of a matrix into shared memory once so every thread in the block reuses it many times.
Data reuse: Using a value staged in fast memory many times before fetching new data, so slow global memory is touched as little as possible.
GEMM: General matrix multiply, C = A times B. The workhorse operation behind neural networks and the main thing GPUs are tuned for.
cp.async: An asynchronous copy from global memory straight into shared memory, without stalling the thread or passing through registers.
Software pipelining: Overlapping the load of the next tile with the compute on the current one, so memory latency hides behind useful work.
Double buffering: Using two shared-memory buffers that take turns: one is being computed on while the other is being filled by the next load.
Prefetch: Starting a load early, before the data is needed, so it has arrived by the time you use it.
Thread block: A group of threads that run on one SM and share its shared memory. Also called a CTA.
Register file: The pool of registers on an SM, shared out among all resident threads. About 256 KB on an A100.
Register pressure: How many registers a kernel needs per thread. High pressure means fewer warps fit, which lowers occupancy.
Register spilling: When a thread needs more registers than it has, the extra values spill to local memory, which actually lives in slow global memory.
Local memory: Per-thread memory that, despite the name, lives in slow off-chip global memory. Registers spill here when they run out.
TMA: Tensor Memory Accelerator. A Hopper copy engine that moves whole tensor tiles between global and shared memory from a single descriptor, so one thread issues the load.
Tensor map descriptor: A small host-built struct (128 bytes) that tells TMA the tensor base, shape, strides, tile size, element type, and swizzle. One thread passes it to issue a bulk copy.
mbarrier: An asynchronous barrier in shared memory. TMA signals it when a tile lands and waiting threads wake, handing each buffer stage from producer to consumer.
wgmma: Warpgroup matrix multiply. A Hopper instruction where 128 threads issue one asynchronous tensor-core matmul that reads its operands from shared memory.
Warpgroup: Four contiguous warps, 128 threads, the granularity Hopper wgmma operates on.
Warp specialization: Giving different warps different jobs: producer warps issue TMA loads while consumer warps run wgmma, overlapping load and compute.
Thread block cluster: A Hopper group of blocks co-scheduled on one GPC that can read each other’s shared memory (distributed shared memory).
Multicast: A TMA mode that broadcasts one global load into several blocks’ shared memory in a cluster, so a shared operand crosses the bus only once.
Tensor Memory: A dedicated on-SM memory on Blackwell (256 KB) that holds the MMA accumulator, so the register file no longer has to feed the tensor cores at FP4 rates.
tcgen05: Blackwell’s 5th-generation tensor-core MMA. A single thread issues the matmul for the whole block, reading operands from shared memory and TMEM.
Microscaling: Storing a block of low-precision values (say FP4) with one shared low-precision block scale (E8M0 for MXFP4, FP8 for NVFP4), so tiny formats stay accurate. NVFP4 uses a block of 16.
FP4: A 4-bit floating-point format (E2M1). Packs eight values per 32 bits and roughly quadruples tensor-core peak over FP16, at the cost of range.
Accumulator: The running sum D in D = A × B that a matmul builds up across its K steps. Where it lives (registers or TMEM) is a recurring bottleneck.
CuTe: The layout layer under CUTLASS. Expresses thread-to-data mappings as Shape ⊗ Stride and composes, tiles, and swizzles them at compile time.
CUTLASS: NVIDIA’s open template library for peak-performance GEMM and related kernels, built on CuTe layouts.
Layout: A CuTe object, Shape ⊗ Stride, that maps a logical coordinate to a linear memory offset. Change the stride to re-lay-out data without moving it.
Stride: How far apart, in memory, consecutive elements along an axis sit. A stride of 1 means contiguous, which is what makes a read coalesce.
Swizzle: A layout that permutes shared-memory addresses so a tile reads back bank-conflict-free and in the order the tensor cores want.
Roofline: A plot of attainable throughput against arithmetic intensity. An op is memory-bound under a bandwidth roof until enough reuse lifts it to the flat compute roof.
Arithmetic intensity: FLOPs performed per byte moved from memory. Low intensity is memory-bound; tiling raises it until the op becomes compute-bound.
Exponent: A float's exponent bits set how big or small it can get. More exponent means more dynamic range.
Mantissa: A float's mantissa bits set how finely it resolves values between powers of two. More mantissa means more precision.
BF16: Brain float 16: 8 exponent bits and 7 mantissa bits. Same range as FP32 with less precision, which is why it is the training default.
FP8: An 8-bit float in two flavors: E4M3 (more precision, forward pass) and E5M2 (more range, gradients). Introduced for tensor cores on Hopper.
Quantization: Storing weights or activations in fewer bits than they were trained in, usually with a shared scale factor to recover the real magnitudes.
INT8: An 8-bit integer format with evenly spaced steps and a shared scale. Cheap and tight when a tensor has no wild outliers.
Ternary (BitNet b1.58): Weights restricted to three values, -1, 0, and +1, about 1.58 bits each. The matmul becomes addition and subtraction with no multiplies; trained from scratch, not compressed after.
NVFP4: NVIDIA's 4-bit float (E2M1) with one FP8 scale per 16 values plus a per-tensor FP32 scale. Finer blocks than MXFP4 for better accuracy.
MXFP4: The open OCP microscaling 4-bit float: one power-of-two scale (E8M0) shared across every block of 32 values.
GGUF: The llama.cpp file format that packs a model plus metadata for local inference. Holds k-quant tensor types like Q4_K_M.
k-quant: A GGUF quantization scheme: weights in super-blocks of 256, split into sub-blocks of 32, with a two-level (super-block and sub-block) scale.
Attention: The operation that lets each token weigh every other token: score queries against keys, softmax the scores, then blend the values.
Softmax: Turns a row of numbers into positive weights that sum to 1, so they can act as attention weights or class probabilities.
Token: One chunk of a model’s input or output, roughly a word or word-piece. Sequences are measured in tokens.
Gradient: The correction signal used to update a model during training. Gradients span a huge range of magnitudes, which is why their number format needs range.
Tensor core: A dedicated unit inside the SM that multiplies a small matrix in one instruction, far faster than ordinary threads doing it multiply by multiply. First shipped on Volta.
MMA: Matrix multiply-accumulate: the tensor core’s core operation, D = A × B + C. wgmma (Hopper) and tcgen05.mma (Blackwell) are MMA instructions.
FLOP: One floating-point operation, a single multiply or add. Throughput is measured in FLOPs per second (FLOP/s).
Activation: The data flowing through a model (the inputs and intermediate results), as opposed to the fixed weights. Activations carry outliers, which makes them harder to quantize.
W4A16: A quantization recipe: 4-bit weights, 16-bit activations. Weight-only, so it saves memory but computes in 16-bit. GPTQ, AWQ, and GGUF are W4A16.
W4A4: A quantization recipe: 4-bit weights and 4-bit activations. The matmul runs on low-precision tensor cores, saving compute too, but activation outliers make it hard. NVFP4 is W4A4.
SDPA: PyTorch's scaled_dot_product_attention: a dispatcher that auto-picks a fused FlashAttention-style backend (flash, memory-efficient, or cuDNN) for your shapes and dtype.
FlexAttention: A PyTorch API for writing custom attention masks and score modifications as a small function that still compiles to one fused FlashAttention-style kernel.
SageAttention: Quantized attention: runs the attention matmuls in INT8 or FP4 on tensor cores, smoothing outlier channels first. The W4A4 idea applied to attention.
SonicMoE: A mixture-of-experts kernel library that packs routed tokens into contiguous groups so each expert’s grouped-GEMM tile is full and efficient.
Data parallel: Replicate the whole model on every GPU, split the batch, and average gradients with an all-reduce. DDP is the efficient PyTorch version.
FSDP / ZeRO: Sharded data parallel: split the batch AND shard params, gradients, and optimizer state across GPUs, gathering each layer just in time. Trades communication for memory.
Tensor parallel: Split each layer's weight matrices across GPUs (Megatron splits heads and MLP columns) and all-reduce the activations. Chatty, so it wants NVLink.
Pipeline parallel: Put different layers on different GPUs as stages; activations flow stage to stage. The idle time while the pipeline fills and drains is the "bubble."
Expert parallel: Spread a mixture-of-experts layer’s experts across GPUs and all-to-all the tokens to wherever their expert lives. DeepEP accelerates the shuffle.
Context / sequence parallel: Split the sequence across GPUs. Sequence parallel saves activation memory on the non-matmul parts; context parallel (ring attention) splits attention over ultra-long contexts.
All-reduce: A collective that sums a value across all GPUs and hands every GPU the total. Used to average gradients (data parallel) and combine activations (tensor parallel).
All-to-all: A collective where every GPU sends a different piece to every other GPU. Used to route tokens to their expert in expert parallelism.
KV cache: The stored Keys and Values of every past token, kept so they are not recomputed each step. It dominates long-context memory and grows with the number of KV heads, head size, layers, and sequence length.
RoPE: Rotary position embedding: encodes position by rotating the query and key vectors by an angle proportional to position, so the attention score depends on relative distance. Applied every layer, not added to embeddings.
NoPE: No positional encoding: a decoder-only causal model can infer position from the causal mask alone (a counting signal), with no explicit position input. Works only in the causal setting and has a finite usable range.
Sliding window attention: Each token attends only to the last W tokens (a local band), capping cost and local KV cache. Stacking layers still compounds the reach, so the model is not limited to W.
Attention sink: A learned bias that keeps the always-important first tokens in the softmax denominator, letting a very small sliding window (like gpt-oss’s 128) stay stable.
MHA: Multi-head attention: every query head has its own Key and Value head. Best quality, largest KV cache.
MQA: Multi-query attention: all query heads share a single Key/Value head. Smallest KV cache, but the hard sharing can cost quality.
GQA: Grouped-query attention: query heads are split into groups, each sharing one Key/Value head. Near-MHA quality at close to MQA memory, and the mainstream default.
MLA: Multi-head latent attention (DeepSeek): compress every head's K and V into one small shared latent, cache only that, and reconstruct per-head K/V on the fly. GQA-level cache at MHA quality, with a small decoupled RoPE key.
DSA: DeepSeek Sparse Attention (V3.2): a lightning indexer scores past tokens and each query attends only to its top-k (2048), turning attention from order N squared into order N times k. Built on MLA.
YaRN: A RoPE context-extension method: scale the rotation frequencies (NTK-style) and add an attention-temperature correction, so a model trained at one length works at a longer one.