Every µs Matters: Achieving Near Speed-of-Light Latency in GPU Collectives

Siyuan Shen; Anton Korzh; John Bachan; Tiancheng Chen; Arnav Goel; Ludwig Schneider; Pouya Kousha; Zhenhao He; Sylvain Jeaugey; Kamil Iskra; Nishank Chandawala; Jeff R. Hammond; Torsten Hoefler

TL;DR

In long-context, small-batch LLM decode with tensor parallelism, every layer runs small AllReduce operations on the critical path of each token, so collective latency, not bandwidth, limits speed and cost. On 4 GB200 GPUs, cutting small-message AllReduce latency from 11.0 µs (NCCL ring) to 2.37 µs lowers Llama-3.1-70B inter-token latency by 8.7%.
The paper defines a speed-of-light floor for AllReduce, the least data movement the hardware allows, and measures it at 1.404 µs on GB200.
The largest avoidable cost is the memory barrier: each one takes 0.85 to 1.68 µs, and many kernels use two.
It removes barriers with LL flags, sentinel values, double buffering and a new two-shot LL128 atomic algorithm, packaged as low-latency APIs on NCCL’s device-side interface. The best kernel comes within 7% of the floor at 2 GPUs, and in vLLM inter-token latency drops by 7–13% on 4 GPUs and 9–11% on 8.

Tensor parallelism splits each layer’s weight matrices across GPUs, so every GPU ends a layer holding partial results that must be summed: an AllReduce after attention and another after the MLP. For Llama-3.1-70B, with 80 layers, that is at least 160 AllReduces per generated token.

Long contexts make each one smaller. The KV cache grows with sequence length, so the batch shrinks to fit in memory, and each AllReduce carries a small message: at a batch of 8, Llama-3.1-70B’s 8192-wide hidden state in 16-bit precision is 128 KiB. At that size the time depends mostly on how many times the GPUs must synchronize, not on how many bytes move. A ring AllReduce needs a number of rounds that grows with the GPU count and a tree needs a logarithmic number, while one-shot and two-shot designs inside one NVLink domain need a constant one or two. The paper works in that setting: GPUs in a single scale-up domain, where tensor-parallel groups usually run. See multi-GPU communication for the interconnects and NCCL for the library the work extends.

The speed-of-light floor

To know how far existing kernels are from optimal, the paper estimates the least data movement an AllReduce of one 128-byte cache line can need. The fastest shape is a one-shot push: each GPU loads its input from L2, stores it into every peer’s scratch buffer, and loads what the peers stored. L2 is the point of coherency between GPUs, so the floor is

L_{SoL} = 2L_{L2\,RTT} + L_{remote\,store}

It measures the L2 round trip as the latency of one __threadfence() (0.306 µs), and the remote store from a two-GPU ping-pong (0.792 µs), which gives 1.404 µs. The estimate assumes stores to all peers go out at the same time, so the floor does not depend on the number of GPUs, and it ignores computation and instruction scheduling. It is a lower bound, not a prediction. See GPU memory hierarchy for where L2 sits.

Barriers are the tax

The one-shot and two-shot kernels in vLLM, TensorRT-LLM, NVSHMEM and MSCCL++ use memory barriers to tell peers that data is ready. Each GPU propagates a flag to every peer and then waits until it has seen every other GPU’s flag. Using NCCL’s ncclLsaBarrierSession, a unicast barrier on GB200 takes 0.85 µs at 2 GPUs and 1.68 µs at 32. A barrier built on NVLink SHARP multicast starts higher, at 1.12 µs, but grows more slowly and costs only 1.25 µs at 32 GPUs. Many kernels need two barriers, so synchronization alone can take longer than the entire floor. The paper estimates that for a 5 µs AllReduce on 4 GPUs, two barriers account for about 40% of the time.

Signalling without barriers

The barrier exists so the reader knows its data has landed. The paper combines four techniques that let the data announce itself, all in push mode:

LL (low latency, from NCCL’s LL protocol) packs an 8-byte flag with 8 bytes of data and sends both in one atomic 16-byte store. The reader polls the flag. This halves the payload bandwidth and doubles the scratch space, so it suits very small messages.
Sentinel fills the receiving buffer with a value real data will never take, such as -NaN, and the reader polls until it changes. It keeps the full payload, but the buffer must be reset before reuse and valid data must never equal the sentinel.
Double buffering with bidirectional exchange handles messages too large for one pass. Each GPU alternates between two scratch buffers and only reuses one after it has received the peer’s data from the current round, so every receive acts as permission for the next send, like credit-based flow control. Ring AllReduce cannot use this, because each rank must receive before it forwards, so it still needs barriers.
Two-shot LL128 atomic is the paper’s new algorithm. In the ReduceScatter phase, groups of 8 threads each cover one 128-byte cache line. One thread moves the first element of the line to shared memory, puts a 1 in its place, and the group atomically adds the whole line into the target GPU’s buffer. NVLink applies 128-byte adds atomically, so when the first slot reads N, all N GPUs have contributed and the sum is complete. The AllGather phase restores the moved elements and writes the output.

LL128 atomic needs only D/N of scratch space per iteration and gives up 4 bytes in 128 for FP32 (2 in 128 for half precision). It has real limits: it needs NVLink for the cache-line atomics, supports only single- and half-precision types, only addition, and it is not deterministic because floating-point atomics can land in any order. Its error stays within the standard summation bound; for FP32 over 64 GPUs the worst-case error coefficient is about 3.8 × 10^-6.

One-shot or two-shot

For N GPUs reducing M bytes, with D bytes reduced per iteration, the paper compares its designs as follows:

Algorithm	Data sent per GPU	Synchronizations	Scratch per iteration	Deterministic
One-shot (LL)	2(N−1)M	1	2ND	Yes
One-shot (sentinel)	(N−1)M	1	ND	Yes
Two-shot (LL)	4(N−1)M/N	2	2D	Yes
Two-shot (sentinel)	2(N−1)M/N	2	D	Yes
Two-shot LL128 atomic	≈2(N−1)M/N	2	≈D/N	No

One-shot needs one synchronization but sends every GPU’s whole input to every peer, so it wins for small messages. Two-shot adds a second synchronization but cuts the data each GPU sends by a factor of N/2, from O(N·M) to O(M), so it wins as messages grow, provided the scratch buffer holds the whole message in one pass.

The paper settles on 4 MiB of scratch for one-shot kernels and 64 MiB for two-shot and LL128 atomic. It picks kernels from measurements: at 4 GPUs, messages under 1 MiB use one-shot and messages from 1 to 2 MiB use two-shot.

The API

The techniques are packaged as experimental low-latency APIs on NCCL’s device-side communication interface. A device-side ncclLLBuffer wraps symmetric memory; one template parameter chooses LL or sentinel mode and another turns on NVLS multicast. It exposes thread-level send, recv, recvUnrolled, recvReduce, bcast and reset, plus advanceEpoch() to switch buffers between iterations. On the host, ncclCalcScratchBufferSize() sizes the buffer and ncclLLBufferInitSentinel() fills it for sentinel mode. A one-shot AllReduce then fits in a loop of a few lines:

code

for (int i = tid; i < nElts; i += nthreads) {
  float data = inputBuf[i];
  int slot = threadIdx.x + rank * blockDim.x;
  llBuf.template bcast<4, float>(team, slot, data);   // push to every peer
  float result = llBuf.template recvReduce<4, float, false>(
      threadIdx.x, nRanks, blockDim.x,                 // poll each peer's slot
      [](float v) { return v; },
      [](float a, float b) { return a + b; });
  outputBuf[i] = result;
  llBuf.advanceEpoch();                               // swap buffers, no barrier
}

LL128 atomic is not part of the API, because it needs 8 threads to act as a unit. In the paper’s modified NCCL, released with the code, the new kernels are selected with NCCL_SYM_KERNEL (AllReduce_LLBuffer, AllReduce_LLBuffer_Twoshot, AllReduce_LL128_Atomic), and NCCL_SYM_LLBUFFER_SYNC switches between LL and sentinel.

Results

The microbenchmarks run on a GB200 NVL72 system from 2 to 64 GPUs, against NCCL 2.29.1 (ring, tree and symmetric-memory kernels), NVSHMEM 3.5.21, MSCCL++ 0.8.0 and vLLM’s custom AllReduce.

The new one-shot kernel has the lowest latency for small messages at every GPU count. At 2 GPUs it is about 7% above the floor for a 128-byte message; at 64 GPUs the multicast one-shot variants stay within about 70%.
LL is slightly faster for the smallest messages, and sentinel takes over as message size and GPU count grow.
LL128 atomic gains over the standard two-shot design as the GPU count rises.
For large messages NCCL’s existing two-shot symmetric kernel stays best, so the new kernels complement it rather than replace it.

In vLLM, with 100k to 200k tokens of input context, 16k output tokens and a batch of 8, the best configuration cuts inter-token latency by 7–13% on 4 GPUs and 9–11% on 8, with similar throughput gains, across dense (Llama), mixture-of-experts (DeepSeek) and hybrid-attention (Qwen3-Next) models. Converted to cost at CoreWeave’s GB200 price, the savings exceed $11 per million output tokens for DeepSeek-V3 on 8 GPUs.

The same kernels speed up cuSOLVERMp’s generalized eigensolver on the Alps supercomputer (GH200 nodes) by 7.0% (m = 32768 on 2 GPUs) and 1.5% (m = 65536 on 4 GPUs). The gain is larger where communication is a bigger share of the runtime.

Why it matters

Collective libraries have mostly been tuned for bandwidth. This paper argues that for decode-heavy serving the right target is the latency floor, and that the largest avoidable cost is synchronization rather than data. Measuring the floor turns “fast enough” into a gap that can be quantified, 7% at best and larger as GPU counts grow. LL, sentinel and double buffering need only GPU-initiated writes to peer memory, polling and fences, so they are not tied to NVIDIA hardware. It complements NanoFlow: NanoFlow hides collective time behind computation, and this work shrinks the collectives themselves. It is a recent preprint and the APIs are marked experimental.

Megatron-LM: the paper that put two all-reduces in every tensor-parallel transformer layer
Collective Communication for 100k+ GPUs: Meta's host-driven, zero-copy collectives for jobs that span data-center buildings, the large-scale end of the same problem
NanoFlow: overlaps network-bound collectives with compute inside one GPU, where this paper makes the collectives shorter
PagedAttention (vLLM): the serving engine used in the paper’s inference experiments
Making Deep Learning Go Brrrr From First Principles: the overhead-bound regime, where fixed per-operation costs rather than bandwidth set the time
A Survey of Techniques for Optimizing Transformer Inference: the wider set of inference optimizations this work fits into