1 Introduction↩︎

The advent of agentic systems [1]–[3] and reasoning models [4] has made low-latency LLM token generation critical. In human-facing LLM applications such as chatbots, each token is read by users, so reducing latency beyond human reading speed provides diminishing returns. In contrast, these systems generate intermediate reasoning tokens that are not consumed by humans, so per-token latency compounds across steps and directly impacts end-to-end response time. Recent proposals for hardware accelerators tailored for reasoning [5] and industry platforms like NVIDIA Groq LPU [6] further emphasize the importance of optimizing for low-latency inference. In this work, we focus on reducing token generation latency in GPU-based inference.

The Low-Latency Challenge. The first step towards reducing token generation latency is to minimize batching, since larger batches delay individual token generation by (1) increasing the size of KV Cache and (2) activating more model parameters (for MoE [7]). With minimal batching, loading model weights and KV Cache from HBM becomes the primary bottleneck. Tensor Parallelism (TP) addresses this memory-bandwidth bottleneck by partitioning the model weights and KV Cache across multiple GPUs and loading them in parallel with aggregate bandwidth. However, scaling across GPUs to improve bandwidth introduces additional bottlenecks, including kernel launch overheads [8], on-chip delays [9], HBM under-utilization [10], and communication overheads.

All-Reduce: The Emerging Bottleneck. Recent advances in system software mitigate many non-communication bottlenecks. Megakernels [10], [11] eliminate kernel launch overheads and improve HBM utilization by fusing the forward pass, while Programmatic Dependent Launch [12] enables prefetching across kernels. As these optimizations reduce compute-side overheads, All-Reduce communication emerges as the dominant cost.

To quantify this trend, we measure time per output token (TPOT) for Llama 3.1 8B on H200 GPUs using a Megakernel implementation [10]. As we scale from TP=1 to TP=8, TPOT reduces from 2.73 ms to 1.43 ms, and the fraction of time in All-Reduce increases from 0% to 30% (corresponding to a 43% throughput gain). Thus, All-Reduce becomes a growing contributor to latency and is poised to dominate as other bottlenecks diminish. Despite this, optimizing All-Reduce for low-latency inference remains largely unexplored.

Limitations of Prior Work. Prior approaches to reducing All-Reduce overhead, such as kernel fusion [13]–[15] and microbatching [16]–[18], rely on overlapping communication with computation from large batches. However, these techniques are ineffective for low-batch inference, where there is insufficient compute to hide communication. The goal of this work is to directly optimize All-Reduce to minimize its latency.

All-Reduce Algorithms. The two most common All-Reduce algorithms in LLM inference are oneshot [19] and twoshot [20]. Oneshot completes reduction in a single round trip, where each GPU pulls the entire buffer from all peers and reduces it locally. Twoshot decomposes reduction into a reduce-scatter followed by a broadcast: GPUs first reduce disjoint \(K/N\) portions of the payload (\(K\): payload size, \(N\): number of GPUs), then broadcast the results so each GPU receives the full reduced output. Compared to oneshot, this reduces data transfer but requires two round trips.

Key Findings. Both oneshot and twoshot require a synchronization barrier before and after the data transfer phase. We find that for small payloads typical of low-batch inference, these barriers account for 32-50% of oneshot latency and 49-62% of twoshot latency (1 (a)). Our key insight is to minimize All-Reduce overheads by reducing the cost of these expensive barriers. To that end, we propose Synchronization-Free All-Reduce (SiFAR)¹.

Removing Bottom Barrier via Dual Buffering. The bottom barrier in oneshot ensures all peers have finished reading the payload buffer before the owning GPU overwrites it with new data, enforcing a write-after-write dependency. We eliminate this barrier using dual buffering. The inference engine alternates between two payload buffers across successive All-Reduce operations, ensuring that the payload is never overwritten before all peers finish reading. While dual buffering is well known, it is not widely used in current LLM serving systems because communication libraries like NCCL [21] are decoupled from model execution and cannot make assumptions about buffer reuse patterns. SiFAR overcomes this by co-designing communication with model execution, enabling dual buffering in systems like vLLM [22] and Megakernels [10] without changes to their memory management (1 (b)).

However, removing the bottom barrier is insufficient. Oneshot scales poorly with GPU count, making it slower than twoshot beyond TP=2. At TP=8 with an 8 KB payload, oneshot takes 4.01 \(\mu\)s for data transfer compared to 1.96 \(\mu\)s for twoshot. While twoshot has lower data transfer latency, it requires a bottom barrier that cannot be eliminated because it enforces a read-after-write dependency after the broadcast phase of twoshot. Ideally, we want oneshot latency to be comparable to twoshot so that dual buffering is effective.

Reducing Data Transfer via In-Switch Computation. To achieve this, we leverage in-switch computation available in modern NVIDIA switches via the multimem.ld_reduce instruction [23]. When a GPU issues ld_reduce, the switch pulls data from all peers, performs the reduction, and returns the result to the requesting GPU. In existing systems, this primitive is primarily used in the reduce-scatter phase of twoshot, where each GPU reduces a disjoint \(K/N\) portion of the payload at the switch. However, we find that ld_reduce can reduce oneshot latency beyond its conventional use in the reduce-scatter phase of twoshot. Reduce-scatter latency remains flat up to 64 KB and grows only slightly up to 256 KB. At TP=8, a 256 KB reduce-scatter means each GPU reduces 32 KB at the switch, which is comparable to the 4-32 KB payloads in low-batch inference. Our insight is that if the switch can reduce 32 KB per GPU with low latency, each GPU can issue ld_reduce over the entire 4–32 KB payload instead of a disjoint partition. We call this redundant pull: each GPU reduces the full \(K\)-byte payload at the switch and receives only \(K\) bytes instead of \(K \times (N-1)\) in oneshot, ensuring dual buffering remains effective (1 (c)).

Reducing Top Barrier Overhead via Speculative Reduction. The top barrier ensures all GPUs have produced valid partials before communication begins, enforcing a RAW dependency that cannot be simply eliminated. However, we observe that LLM decoding issues multiple All-Reduce operations per token. As a result, GPUs remain tightly synchronized after the first All-Reduce, as they execute identical operations across successive All-Reduces. Thus, the overhead of the top barrier is primarily due to flag exchange, not divergence between GPUs. This leads to our third optimization: speculate and verify. When a GPU enters an All-Reduce, it speculatively begins fetching data from peers, assuming their data is ready, thereby avoiding explicit flag exchange. To ensure correctness, each GPU writes a flag to a validation buffer that is reduced alongside the payload. If the result equals \(N \times \text{flag}\) after reduction, all GPUs had valid data and speculation succeeded. Otherwise, SiFAR reruns the All-Reduce. In practice, speculative reduction substantially reduces top-barrier overhead (1 (d)).

Results. We integrate SiFAR into Megakernels [10], a state-of-the-art for low-latency LLM inference, and evaluate it on H200s. SiFAR reduces All-Reduce latency by up to 52% for payloads up to 32 KB and improves throughput by 18.6% for Llama-3.1-8B and 13.1% for Qwen3.5-397B-17B at TP=8. Our contributions are as follows:

1. We identify that synchronization barriers dominate All-Reduce latency in low-latency tensor-parallel LLM inference, accounting for up to 62% of total All-Reduce time.

2. We propose Synchronization-Free All-Reduce (SiFAR), which eliminates the bottom barrier via dual buffering, reduces data transfer via redundant pull, and minimizes the top barrier via speculative reduction.

3. We demonstrate SiFAR reduces All-Reduce latency by up to 52% and improves end-to-end throughput by up to 18.6%.

2 Background and Motivation↩︎

2.1 LLM Inference: System Design and Metrics↩︎

LLM inference consists of two phases: prefill and decode. During the prefill phase, the model processes all input prompt tokens in parallel, making it compute-bound. In contrast, the decode phase loads all the model weights from memory to produce the next token, making it memory bandwidth-bound. To amortize the cost of repeatedly loading model weights, most LLM inference engines batch requests from multiple users together during inference, trading off latency for higher aggregate throughput.

Performance Metrics. The performance of LLM inference systems is measured using two latency metrics: Time-to-First-Token (TTFT) and Time-per-Output-Token (TPOT). TTFT measures the time taken to process the prompt and produce the first output token, while TPOT measures how quickly subsequent tokens are generated. Both metrics are critical for user experience: TTFT determines the initial responsiveness, and TPOT governs the perceived generation speed. In this work, we focus on optimizing TPOT, as it dominates end-to-end latency in many use cases of LLM-based applications.

2.2 Low Latency LLM Serving↩︎

Traditional LLM-based applications, such as chatbots, are designed to be human-facing, where each generated token is read by a human. In such settings, reducing token generation latency (TPOT) beyond a certain point yields diminishing returns, since humans can only read 4-7 tokens/second [24]. However, newer reasoning models [4] and agentic systems [1]–[3] operate under a different paradigm. Here, tokens are either consumed by other LLMs or used internally as intermediate reasoning steps rather than being read by humans. As a result, per-token latency compounds across steps, making TPOT the dominant contributor to end-to-end response time. Industry platforms such as the NVIDIA Groq LPU [6], which stores model weights in on-chip SRAM to improve memory bandwidth, as well as recent hardware proposals for reasoning [5], further highlight the growing importance of fast token generation. In this work, we focus on reducing TPOT in GPU-based inference.

2.3 Techniques for Low-Latency Serving↩︎

LLM inference systems face a fundamental tradeoff between throughput and interactivity. On one end, high-throughput systems maximize aggregate tokens generated per second across all users by batching hundreds of requests together. On the other end, low-latency systems minimize TPOT by reducing batch sizes, prioritizing responsiveness over throughput [25].

Batching requests from multiple users amortizes the cost of loading model weights across tokens. However, batching also increases token generation latency: larger batches require loading more KV cache during attention and activating more parameters in MoE models. Prior work [26] shows that TPOT increases substantially as batch size grows. We observe a similar trend when running Llama-3.1-8B and Qwen3.5-397B-17B at TP=8 on H200 GPUs using vLLM 0.18.1, where TPOT increases by more than 3.5\(\times\) when scaling from BS=1 to BS=256. Thus, minimizing batch size is essential for reducing TPOT in latency-sensitive settings.

Tensor Parallelism (TP) to Reduce TPOT. With minimal batching, autoregressive decoding becomes memory bandwidth-bound, since each token generation step repeatedly streams large model weights from GPU memory. TP [27] is traditionally used to fit models that exceed the memory capacity of a single GPU by sharding weight matrices across devices. However, TP also increases aggregate memory bandwidth, as each GPU loads only a fraction of the weights. Thus, after reducing batching, TP can further improve TPOT for bandwidth-bound LLM inference, even when the model fits on one GPU. 2 illustrates how TP operates in LLMs. TP shards weight matrices across GPUs, requiring each layer to perform two All-Reduce operations: one after attention and one after the feed-forward network. In each All-Reduce kernel all GPUs exchange partial results and reduce them to produce the final output.

Other parallelism strategies are less effective for reducing TPOT. Pipeline parallelism [28] does not increase memory bandwidth, as each stage still loads full model weights sequentially. Expert parallelism [29] introduces load imbalance across experts, leading to uneven bandwidth utilization. As a result, neither approach is optimal for reducing TPOT. Thus, in this work, we focus on TP.

2.4 Communication Bottlenecks in Tensor Parallelism (TP)↩︎

Improving TPOT by scaling to more GPUs via TP introduces several overheads, including communication overheads from All-Reduce, kernel launch overheads, low HBM utilization, and GPU on-chip latency [10]. These overheads prevent ideal linear scaling of TPOT with increasing GPU count, which would otherwise be expected for bandwidth-bound, non-communication kernels. Recent system optimizations mitigate many of these non-communication bottlenecks. For example, Megakernels [10], [11] fuse the forward pass into a single kernel to eliminate launch overheads and improve HBM utilization, while Programmatic Dependent Launch (PDL) [12] enables weights prefetching across kernel boundaries. Thus, as non-communication overheads are mitigated, All-Reduce becomes the dominant bottleneck in low-latency inference.

To quantify the impact of All-Reduce on TPOT, we use a Megakernel [10] implementation, which fuses the forward pass into a single CUDA kernel, to evaluate Llama-3.1-8B [30], and Qwen3.5-397B-17B [31]. 3 shows throughput (tokens/s) across different TP degrees, comparing execution with and without All-Reduce (i.e., we feed the input of the All-Reduce directly to the next operation instead). For Llama-3.1-8B, removing All-Reduce improves performance by 18% at TP=2 and 43% at TP=8. For Qwen3.5-397B-17B, removing All-Reduce improves performance by 16% at TP=4 and 21% at TP=8. Thus, using TP to achieve lower TPOT makes All-Reduce a significant bottleneck and will become increasingly dominant as other overheads are reduced.

2.5 Limitations of Prior Work↩︎

Microbatching. Prior works such as TokenWeave [16], DeepEP [17], and NanoFlow [18] use microbatching to overlap computation and communication by dividing large batches into smaller microbatches and pipelining their execution. This approach is effective at overcoming communication latencies when the workload is compute-bound, which occurs during prefill or decode with large batch sizes. However, in low-latency inference, batch sizes are intentionally kept small to minimize latency. As a result, microbatching provides little benefit since there is insufficient computation to overlap with communication. Moreover, microbatching primarily improves throughput at the expense of single token latency, which is the key objective in our setting.

Fine-Grained Compute-Communication Overlap. Kernel fusion techniques [13]–[15] fuse the All-Reduce kernel to the preceding GEMM (matrix multiply) kernel and overlap compute and communication by sending partial results tile-by-tile. While effective in hiding All-Reduce overheads for compute-bound GEMMs with large batch sizes (see 4 (left)), fusion provides little benefit in low-latency inference (see 4 (right)). With small batches, each GEMM only issues a small number of threadblocks, which can execute concurrently on the GPU, making computation insufficient to hide any communication latency. We quantify this using a fused GEMM implementation in 6.4.

2.6 Goal of this Work↩︎

The goal of this work is to minimize All-Reduce communication overheads during the decode phase of low-latency LLM inference using Tensor Parallelism. Unlike prior approaches that rely on compute-boundedness from large batch sizes to overlap communication with computation, our solution directly reduces All-Reduce overheads in the low-latency inference regime, where batch sizes are minimal, and there is not enough computation to hide any communication overheads.

3 Towards Synchronization-Free All-Reduce↩︎

In this section, we first describe the two All-Reduce algorithms used in LLM inference engines. Next, we show that synchronization overheads dominate the overall runtime of All-Reduce operations for small payloads. Finally, we present our insights that enable us to reduce the impact of these overheads by (1) eliminating the bottom-barrier via dual buffering, (2) using in-switch computation to reduce data transfer during All-Reduce, and (3) reducing the top-barrier overhead via speculative reduction.

3.1 All-Reduce: Oneshot & Twoshot Algorithms↩︎

Most LLM inference engines, including vLLM [22], SGLang [32], and TRT-LLM [33], use the oneshot or the twoshot All-Reduce algorithm (or one of their variants) for low-latency serving. The choice depends on payload size and TP degree. In this work we use the best of the two algorithms as our baseline.

Oneshot. The oneshot All-Reduce algorithm completes the entire reduction in a single communication round. Each GPU pulls the entire buffer from all peers and performs reduction locally. It consists of three phases: (1) a top barrier to ensure all GPUs have produced the partials before communication begins. (2) each GPU pulls the partials from all other GPUs and performs element-wise reduction (summation) on received tensors to produce the final result. (3) a bottom barrier to ensure all peers have finished reading the payload buffer before the owning GPU overwrites it.

Twoshot. The twoshot All-Reduce algorithm decomposes reduction into two communication rounds: reduce-scatter followed by broadcast (all-gather). It consists of 4 phases: (1) a top barrier to ensure all GPUs have produced the partials before communication begins. (2) each GPU pulls a portion of the payload and reduces it. (3) each GPU broadcasts its reduced partial to all peers, allowing every GPU to reconstruct the full reduced result. (4) a bottom barrier to ensure that the broadcast writes are visible to all peers before moving on to the next operation.

Data Transfer Comparison. The key difference between oneshot and twoshot lies in data movement. In oneshot, each GPU sends and receives \(K \times (N - 1)\) bytes (where \(K\) is the payload size and \(N\) is the number of GPUs), scaling linearly with GPU count. In contrast, twoshot keeps data transfer constant by decomposing reduction into two steps: reduce-scatter and broadcast. The data transfer during twoshot can be further reduced using in-switch computation capabilities [20], [23] of modern switches. With in-switch computation, across both steps of twoshot, each GPU sends \(K + K/N\) and receives \(K\) bytes. In this work, we focus on this in-network implementation of twoshot. Oneshot is preferred when the payload size is small and the single round trip is more beneficial. Twoshot becomes more efficient at larger scales where the reduction in data transfer outweighs the cost of the additional round trip.

3.2 The Cost of Synchronization↩︎

While necessary for correctness, the barriers around All-Reduce introduce significant overhead. To quantify their impact on All-Reduce latency, we run the oneshot and twoshot All-Reduce implementations from vLLM [19], [20] with and without barriers across different payload sizes and TP configurations on an H200 node. 7 shows the results, breaking down total latency into data transfer time and synchronization time.

We observe that synchronization overhead accounts for 32-50% of oneshot and 49-62% of twoshot latency for small payloads (8-32 KB) across all TP configurations. For example, at TP=4 with 8 KB payload, oneshot takes 4.66 \(\mu\)s with barriers but only 2.33 \(\mu\)s without, a 2\(\times\) slowdown purely from synchronization. For twoshot at TP=8 with 8 KB payload, barriers account for 3.15 \(\mu\)s out of 5.11 \(\mu\)s total, representing 62% overhead. Twoshot exhibits higher synchronization overhead because its bottom barrier must ensure that broadcast writes are visible before the next operation, which requires a more expensive acquire-release barrier that enforces write ordering. These sizes of 4-8 KB represent the All-Reduce payloads for most leading models [4], [7], [30], [34]–[36] at batch size 1. Thus, during autoregressive decoding, every All-Reduce operation encounters this synchronization bottleneck, spending more time waiting at barriers than transferring and reducing data.

3.3 Insight 1: Removing Bottom-Barrier via Dual Buffering↩︎

The bottom barrier in oneshot prevents Write-After-Write (WAW) hazards by ensuring all peers finish reading the payload buffer before the owning GPU overwrites it with new data. However, WAW is an avoidable dependency that can be eliminated through redundancy. Despite this, dual buffering is not widely used in current LLM serving systems because communication is handled by application-agnostic libraries like NCCL, which cannot make assumptions about buffer usage patterns and must include a bottom barrier to ensure correctness.

Eliminating the Bottom-Barrier. Our insight is to use dual buffering to eliminate the bottom-barrier from All-Reduce. By co-designing All-Reduce with the LLM inference engine, we can safely remove this barrier. 8 illustrates our approach. We modify the inference engine to allocate two input buffers and alternate between them across successive All-Reduce operations. This ensures that the payload buffer is never overwritten until the subsequent All-Reduce finishes, eliminating the need for a bottom-barrier.

Dual buffering requires only 4-8 KB per token of additional HBM per GPU, which at low batch sizes is negligible compared to multi-gigabyte model weights. Despite this, it is not widely used for two reasons. First, communication libraries like NCCL [21] are application-agnostic and cannot make assumptions about buffer reuse patterns, so they must include a bottom barrier for correctness. Second, LLM serving systems like vLLM [22] employ automatic memory management that reuses tensors aggressively to minimize HBM consumption. Since the input buffer of All-Reduce is the same size as the output of the following GEMM, the memory manager reuses it as the output buffer for the following GEMM, creating the WAW dependency shown in 8 (left). Co-designing communication with model execution allows us to allocate a second buffer and alternate between them, breaking this dependency.

Dual Buffering Alone is Insufficient. While dual buffering eliminates the bottom barrier, oneshot still scales poorly with GPU count since each GPU pulls the entire buffer from all peers. At higher TP degrees, twoshot outperforms oneshot despite requiring two round trips, because its data transfer remains constant (see 7). However, the bottom barrier of twoshot enforces a Read-After-Write (RAW) dependency, ensuring broadcast writes are visible before the next operation. Unlike the WAW dependency in oneshot, this RAW dependency cannot be eliminated through dual buffering. Thus, to benefit from dual buffering, we need a oneshot design with comparable data transfer latency to twoshot.

3.4 Insight 2: Using In-Switch Computation for Oneshot↩︎

Modern NVIDIA switches support in-network reduction via the ld_reduce instruction [23]. Each GPU maps the same virtual address range to its local buffer using symmetric memory. When a GPU issues ld_reduce on a symmetric address, the switch pulls data from all GPUs, reduces it, and returns the result. Currently, ld_reduce is only used for reduce-scatter, the first step of twoshot, where each GPU reduces a disjoint \(K/N\) portion of the payload (K: payload size, N: number of GPUs). However, we find that ld_reduce can be used beyond reduce-scatter to lower oneshot latency.

Observation: Constant Latency for Small Payloads. 9 shows that reduce-scatter latency remains nearly constant up to 64 KB across all TP configurations, because latency is dominated by round-trip latency between GPUs rather than data volume. At TP=8 with a 64 KB payload, reduce-scatter means each GPU reduces 8 KB at the switch, comparable to the 4-8 KB payload sizes of most leading models. Even up to 256 KB, reduce-scatter latency does not increase significantly. Our insight is that if the switch can reduce small payloads from each GPU with low latency during reduce-scatter, then each GPU can issue ld_reduce over the entire payload (\(K\) bytes) instead of a disjoint partition (\(K/N\) bytes) and achieve data transfer latency similar to twoshot.

Redundant Pull. We exploit this observation to propose redundant pull. Instead of pulling raw data from all peers and reducing locally as in oneshot, each GPU issues ld_reduce over the entire buffer. The switch reduces data from all peers and returns the full result. We call this redundant pull because every GPU requests reduction on the same addresses, making the switch perform redundant reductions. 10 illustrates this design. Compared to oneshot, redundant pull reduces data received per GPU from \(K \times (N-1)\) to \(K\). For small payloads, redundant pull achieves data transfer latency comparable to reduce-scatter, which is only one phase of twoshot, while retaining the ability to benefit from dual buffering.

Latency Comparison. 11 compares oneshot, twoshot, oneshot with dual buffering (DB), and redundant pull with DB. Oneshot with DB removes the bottom barrier, but its data transfer still grows with GPU count, making it slower than twoshot at TP=8. Redundant pull with DB achieves the lowest latency by combining reduced data transfer with the elimination of the bottom barrier.

3.5 Insight 3: Reducing Top-Barrier Overheads via Speculative Reduction↩︎

The top barrier enforces a Read-After-Write (RAW) dependency, ensuring all GPUs have produced valid partials before communication begins. Unlike the WAW dependency eliminated by dual buffering, this RAW dependency cannot be removed. However, we find that its overhead can be substantially reduced through speculate and verify. To explain this, we first break down the cost of the top barrier.

Breakdown of Barrier Costs. The cost of the top barrier consists of two components: divergence and flag exchange. Divergence is the time the fastest GPU waits for the slowest GPU to reach the barrier. Flag exchange is the time spent exchanging synchronization flags across GPUs to confirm that all participants have arrived.

Divergence is Negligible in Practice. Modern LLM inference engines either launch all kernels for a forward pass within a single CUDA Graph [22], [32], or fuse them into a single Megakernel [10]. In both cases, the forward pass executes entirely on the GPU without host interaction after the initial launch. Once GPUs are synchronized by the first All-Reduce, they remain tightly synchronized throughout the remainder of the forward pass, since all GPUs execute identical kernels with identical input and output sizes. 12 illustrates this behavior using a representative CUDA Graph execution from Nsight profiler [37], showing that divergence is near-zero after the first All-Reduce. Most barrier time is therefore spent on flag exchange, where GPUs simply idle while exchanging synchronization flags to confirm all participants have arrived.

Eliminating Flag-Exchange Overhead. Since GPUs are implicitly synchronized during the forward pass, we can speculate that all GPUs are already synchronized and begin the data transfer phase of All-Reduce immediately, avoiding explicit flag exchange. When a GPU enters an All-Reduce, it speculatively begins fetching data from peers, assuming their data is ready. To ensure correctness, each GPU writes a flag to a validation buffer that is reduced alongside the payload. If the result equals \(N \times \text{flag}\) after reduction, all GPUs had valid data and speculation succeeded. Otherwise, SiFAR reruns the All-Reduce. We explain the detailed design of speculative reduction and the verification mechanism in 4.

4 SiFAR: Design and Implementation↩︎

4.1 Design Overview↩︎

In this section, we present the design and implementation of Synchronization Free All-Reduce (SiFAR) which uses the optimizations described in the previous section to minimize synchronization overheads in All-Reduce. 13 illustrates our approach.

Redundant Pull Design. SiFAR replaces the oneshot pull-and-reduce with redundant pull, where each GPU issues ld_reduce over the entire payload buffer. The switch reduces data from all peers and returns the full result, reducing the data transfer of oneshot.

Dual Buffering Design. SiFAR uses dual buffering to remove the bottom barrier by maintaining two buffers for consecutive All-Reduce operations. Thus, the input buffer to an All-Reduce is not reused until after the next All-Reduce, where the GPUs are again synchronized, safely eliminating the bottom barrier. We explain how SiFAR integrates dual buffering with the automatic memory management of vLLM in 4.2.

Speculative Reduction Design. SiFAR minimizes the overhead of the top barrier using speculative reduction. When a GPU enters the All-Reduce kernel, it speculatively begins issuing ld_reduce to fetch reduced data from remote peers, assuming all peer GPUs have already finished the previous operation. To facilitate validation, SiFAR allocates a small validation buffer per threadblock. After the preceding operation writes its result to the payload buffer, each GPU writes a flag to the validation buffer. The flag is reduced alongside the payload. If the reduced validation result equals ngpus\(\times\)flag, all GPUs had valid data and speculation succeeded. Otherwise, SiFAR reruns the All-Reduce until validation succeeds.

4.2 Implementation Details↩︎

Redundant Pull. We replace oneshot pull-and-reduce with redundant pull by changing how each threadblock performs reduction:

# Oneshot: pull from each peer and reduce locally
for peer in range(ngpus):
    data = load(peer, payload_buf)
    result += data
# Redundant Pull: switch reduces and returns result
result = mutlimem.ld_reduce(payload_buf)

Dual Buffering. Implementing dual buffering requires controlling the allocation and reuse of buffers. However, frameworks like vLLM and SGLang rely on automatic memory management in Pytorch, which eagerly reclaims buffers. Since the input buffer to All-Reduce matches the output size of the next GEMM, it is quickly reused:

buf = redundant_pull(input_buf0)
input_buf0 = GEMM(buf, weight)  # Buffer reused
...
buf = redundant_pull(input_buf1)

To implement dual buffering without modifying framework internals, we modify the All-Reduce function signature to accept an additional argument referencing the previous input buffer:

buf = redundant_pull(input_buf0)
...
buf = redundant_pull(input_buf1, input_buf0)

By passing input_buf0 as an additional argument to the second All-Reduce, the memory manager is forced to keep input_buf0 alive until after the second All-Reduce executes, since it cannot guarantee that input_buf0 will be unused. The additional argument is retained solely for its memory management side-effect and is not used in the computation itself. This approach requires only minimal changes to model execution code and works across frameworks without modifying their memory management systems.

Speculative Reduction. To implement speculative reduction, we build the SiFAR All-Reduce kernel using redundant pull for the payload and oneshot pull-and-reduce for validation. We show a pseudocode implementation below:

int* valid_buf;  # Dual buffered
valid_buf[blockIdx.x] = flag;
__syncthreads();
result = switch_reduce(payload_buf);  # warps 0..k
valid_out = local_reduce(valid_buf);  # warp k+1
while (valid_out != ngpus x flag) {
  result = switch_reduce(payload_buf);
  # Use .cg to bypass L1 cache for valid_buf
  valid_out = local_reduce_cg(valid_buf);
    __syncthreads();
}
flag = flag + 1

The payload buffer is reduced using switch_reduce, which performs reduction in the switch. For the validation buffer, we instead use pull-and-local-reduce to compute valid_out, as it empirically outperforms using switch_reduce for this purpose. Each threadblock runs with one extra warp to process the validation buffer while the remaining warps handle the payload. Before beginning the All-Reduce, each GPU writes the flag value into its validation buffer. After reduction, we check if the result matches ngpus\(\times\)flag. If validation fails, we re-execute until validation succeeds.

Verification Logic Correctness. To maintain correctness across multiple All-Reduce invocations, we increment the flag value after each All-Reduce, ensuring each All-Reduce has a unique validation result. We apply dual buffering to the validation buffer to prevent race conditions where a fast GPU might overwrite its flag before all peers have finished reading it. Since remote loads are cached in L1 [38], upon mis-speculation we use the .cg modifier to load the validation buffer to bypass L1 and fetch fresh data. The payload switch_reduce does not require this, since data is fetched from the switch and not cached locally. The correctness of the verification logic relies on GPU monotonicity; see 8.1 for details.

Megakernel Integration. In Megakernels [10], the entire forward pass is fused into a single kernel, and successive operations within the kernel communicate completion via flags. SiFAR reuses these existing operator flags for validation by reading and reducing the flag from all remote GPUs, instead of explicitly writing to a separate validation buffer. This eliminates the overhead of writing to the validation buffer before each All-Reduce, further reducing latency.

5 Evaluation Methodology↩︎

Serving Framework. We evaluate SiFAR by integrating it into the Megakernel framework [10]. Megakernels fuse the entire forward pass into a single CUDA kernel, eliminating kernel launch overheads and improving HBM utilization. Within the Megakernel, weights of the next operation are prefetched into shared memory whenever possible, including during All-Reduce, further reducing latency. We evaluate on Megakernels rather than serving systems like vLLM [22] because Megakernels represent the state-of-the-art for low-latency inference, where non-communication overheads are already minimized, making the impact of All-Reduce optimization most visible. We extend the original Megakernel implementation to support Llama-3.1-8B and Qwen3.5-397B-17B with multi-GPU tensor parallelism. Our Megakernel baseline without SiFAR already achieves 1.3\(\times\) lower TPOT than vLLM 0.18.1 for Llama-3.1-8B and 2.2\(\times\) for Qwen3.5-397B-17B at TP=8, confirming that non-communication overheads are substantially reduced. We run both models on batch size 1 unless otherwise specified.

Models. We evaluate SiFAR on two models: Llama-3.1-8B [30], a dense transformer, and Qwen3.5-397B-17B [31], a Mixture-of-Experts (MoE) model that activates only 17B parameters per token but requires storing all 397B parameters in memory. We run both models in FP8 precision. Due to model size, Qwen3.5-397B-17B is evaluated only at TP=4 and TP=8. We evaluate SiFAR using input context lengths from 1K to 16K tokens.

Hardware. We use a single node with 8\(\times\) NVIDIA H200 GPUs, each with 141 GB of HBM3e memory. We use CUDA 12.9 with default clock frequencies. We evaluate TP across 2, 4, and 8 GPUs.

Metrics. We measure standalone All-Reduce latency to isolate the impact of SiFAR. For end-to-end impact, we measure throughput (inverse of TPOT) in tokens per second during the decode phase. We report speedup relative to the baseline implementation.

Baselines. Our baseline uses the best of oneshot and twoshot All-Reduce for each configuration. We also compare against two variants of oneshot: TRT-LLM oneshot [33] and Lamport oneshot [39]. We describe these in detail in 6.

6 Evaluation↩︎

In this section, we evaluate SiFAR. We first analyze its impact on All-Reduce latency, then measure end-to-end LLM inference improvements, and finally study the contribution of each optimization.

6.1 Impact on All-Reduce Latency↩︎

15 compares standalone All-Reduce latency across five implementations: oneshot, twoshot, TRT-LLM, Lamport, and SiFAR for TP configurations of 2, 4, and 8 GPUs with payloads from 8 to 32 KB. The baseline for each configuration is the best of oneshot and twoshot, which is the default in vLLM and SGLang. We execute All-Reduce back-to-back within a Megakernel to avoid measuring kernel launch overheads. The measured latency includes any mis-speculation overheads incurred by SiFAR.

SiFAR consistently outperforms all other implementations across every configuration. At TP=2, where oneshot is the best baseline, SiFAR reduces latency to 2.36 \(\mu\)s compared to 3.36 \(\mu\)s for oneshot at 8 KB, a 30% reduction. At TP=4, twoshot becomes the best baseline and SiFAR achieves a 45% reduction in latency from 4.38 \(\mu\)s to 2.39 \(\mu\)s. Finally, the benefits of SiFAR are largest at TP=8, where the baseline is twoshot with 5.11 \(\mu\)s latency at 8 KB, and SiFAR reduces it to 2.44 \(\mu\)s , a 2\(\times\) improvement. Next, we compare SiFAR against two variants of oneshot: TRT-LLM oneshot and Lamport oneshot.

TRT-LLM (Oneshot). We compare SiFAR against the TRT-LLM oneshot implementation, which operates in 3 phases: (1) each GPU copies its input data into a separate local buffer, making the payload buffer private to the All-Reduce kernel so the bottom barrier can be safely removed. (2) A top barrier ensures all GPUs have produced valid partials. (3) Each GPU pulls the full buffer from all peers and performs local reduction. Since TRT-LLM still pulls and reduces like oneshot, receiving \(K \times (N-1)\) bytes per GPU, it does not scale well to higher TP degrees. SiFAR outperforms TRT-LLM by 25-28% at TP=2, 43-44% at TP=4, and 48-57% at TP=8.

Lamport (with In-Network Multicast) [39]. Finally, we compare SiFAR against an optimized implementation (using In-Network Multicast) of a Lamport-style All-Reduce, which is a single-round, push-based algorithm. It operates in 3 phases: (1) each GPU pushes its partial results to all peers via multimem.st, which multicasts data through NVSwitch so each GPU sends only \(K\) bytes. (2) A bottom barrier, similar to the one in twoshot, ensures that all multicast writes are visible locally. (3) Each GPU then loads all received partials from its local buffer and reduces them. Lamport avoids the top barrier and transfers the same total data as redundant pull, but in the opposite direction: it sends \(K\) bytes and receives \(K \times N\) bytes into local memory. In contrast, redundant pull sends \(K \times N\) bytes and receives only \(K\) bytes of reduced data directly into registers via multimem.ld_reduce. SiFAR outperforms Lamport by 32-33% at TP=4 and 34-43% at TP=8 by avoiding the acquire-release bottom barrier and the on-chip latency of loading the full payload from local memory before reduction.

Comparison with NCCL 2.30. We compare SiFAR against NCCL 2.30, which provides custom low-latency kernels that use the LL protocol [40]. We run NCCL with its default auto-selection policy, which uses a heuristic to choose a communication strategy based on the network topology and payload size. 16 shows that NCCL auto is only better than the best of oneshot and twoshot for the 8 and 16 KB payloads at TP=8. Finally, despite the custom low-latency kernels, SiFAR achieves 46–62% lower latency than NCCL across evaluated TP degrees and payload sizes.

6.2 Impact on End-to-End LLM Inference↩︎

17 shows end-to-end throughput across TP degrees for Llama-3.1-8B and Qwen3.5-397B-17B, with input and output sequence length of 1000. SiFAR improves throughput consistently across all configurations with best of oneshot and twoshot as the baseline. For Llama-3.1-8B, the improvement grows with TP degree: 9.1% at TP=2, 14.9% at TP=4, and 18.6% at TP=8. For Qwen3.5, we observe a similar trend with 9.2% improvement at TP=4 and 13.1% at TP=8. This scaling behavior is expected because the impact of All-Reduce increases with TP degree, making the optimizations of SiFAR more impactful.

Impact of Input Sequence Length (ISL). 18 shows the impact of SiFAR when varying ISL for Llama-3.1-8B and Qwen3.5-397B-17B at TP=8. For Llama-3.1, SiFAR provides a 18.6% improvement at ISL=1000 and 14.6% at ISL=16K. For Qwen3.5, the improvement ranges from 13.1% at ISL=1000 to 10.2% at ISL=16K. As ISL increases, the fraction of time spent in attention grows, which does not benefit from SiFAR. However, even at ISL=16K, SiFAR provides meaningful improvements because All-Reduce remains a significant fraction of total execution time.

Impact of Batch Size. 19 evaluates SiFAR across BS=1, 2, and 4 on Llama-3.1-8B at TP=8. SiFAR improves throughput by 18.6% at BS=1, 15.0% at BS=2, and 10.2% at BS=4. The gain is largest at BS=1 and shrinks as the batch grows. Larger batches increase the GEMM and attention compute, so All-Reduce is a smaller fraction of TPOT. SiFAR targets the BS=1 regime because it is latency-critical. Although larger batches raise throughput, they worsen per-request TPOT by 45% from BS=1 to BS=4.

Comparison with Other All-Reduce Algorithms. 20 compares the throughput of SiFAR against the TRT-LLM oneshot variant and the Lamport-style push-based. At TP=8 for Llama-3.1-8B, TRT-LLM and Lamport reach only 85.6% and 90.3% of SiFAR throughput, so SiFAR outperforms even these optimized oneshot-style implementations end-to-end.

Impact of SiFAR on Tail Latency. 21 reports TPOT percentiles at TP=8, normalized to the baseline p50. SiFAR lowers both median and tail latency: for Llama-3.1-8B it reduces p50 to 81.4% and p99 to 84.0%, and for Qwen3.5-397B-17B it reduces p50 to 86.3% and p99 to 87.8%. The speedup from SiFAR holds through p99.9, so mis-speculation retries do not inflate tail latency.

6.3 Impact of Individual Optimizations↩︎

To understand the individual contributions of our optimizations, we conduct an ablation study where we enable each optimization incrementally: redundant pull alone, redundant pull with dual buffering, and SiFAR with all three optimizations including speculative reduction. 22 shows the contributions of each of our optimizations.

For Llama-3.1-8B, redundant pull contributes 6.7% at TP=8 by reducing data transfer through in-switch reduction. Dual buffering builds on this by eliminating the bottom barrier, contributing an additional 3.4% at TP=8. Speculative reduction reduces latency by minimizing the top-barrier overhead, contributing 3.8% at TP=2 and up to 8.5% at TP=8. Together, the three optimizations are complementary and address distinct sources of All-Reduce overhead.

Mis-Speculation and Verification Overhead. We compare SiFAR against two idealized configurations in 23: perfect speculation (no retries) and redundant pull without barriers. The gap between SiFAR and perfect speculation isolates the cost of mis-speculation retries, while the gap between perfect speculation and redundant pull without barriers isolates the verification overhead.

For Llama-3.1-8B at TP=8, SiFAR achieves a 18.6% improvement. With perfect speculation, the improvement increases to 20.2%, meaning mis-speculation retries cost only 1.6% despite a 32.2% average mis-speculation rate (24). Removing barriers entirely yields 21.6%, so verification overhead accounts for an additional 1.4%. Together, mis-speculation and verification account for only 3.0% of lost improvement, showing that speculative reduction captures most of the available benefit.

Retry Overhead. 25 shows the SiFAR latency for the no-retry case and the added latency from each retry. We measure this by forcing SiFAR to retry a fixed number of times. Even with retries, the average SiFAR latency stays well below the baseline (oneshot or twoshot) across all evaluated TP degrees and payload sizes. Each retry costs less than the first attempt, which pays the one-time startup cost of beginning the All-Reduce.

6.4 Comparison with Compute-Comm. Fusion↩︎

26 compares fine-grained compute-communication fusion from ParallelKittens [41] against an unfused implementation for the down-projection GEMM of Llama-3.1-8B at TP=8. At small batch sizes, the fused GEMM has similar or worse latency, because there is not enough computation to overlap with communication. Fusion helps only at very large batch sizes, such as BS=2048, where the GEMM has enough work to hide communication. In the regime that SiFAR targets, overlap provides little benefit, so SiFAR reduces All-Reduce latency directly instead of trying to hide it.

7 Related Work↩︎

We discuss related work in four categories: (1) speculation and relaxation techniques to reduce synchronization overheads, (2) low-latency LLM inference optimizations, (3) collective communication optimizations, and (4) in-network push-based reduction.

Using Speculation to Reduce Synchronization Overheads. A long line of work explores reducing synchronization latency by executing critical sections speculatively rather than waiting at locks or barriers. Early hardware proposals such as Transactional Memory (TM) [42], Speculative Lock Elision [43], and Speculative Synchronization [44] execute code regions in parallel while deferring conflict checks to commit time. TM provides architectural support for atomic sections whose updates are speculatively buffered and committed only when no conflicts are detected. Transactional Lock-Free Execution [45] extends this model using timestamp-based ordering to allow multiple threads to speculatively enter the same critical section. Subsequent work, such as TMbarrier [46] and Elastic Barriers [47], applies similar ideas to barrier synchronization by speculating across phase boundaries. Speculative Lock Forwarding [48] and Redundant Barrier Elimination [49] further reduce waiting by dynamically validating safety conditions at runtime.

While these techniques also use speculation to reduce synchronization costs, they primarily target removing locks and barriers placed conservatively for cases where threads may not actually conflict. In contrast, SiFAR addresses scenarios with true read-after-write dependencies where multiple GPUs must be synchronized for correctness. SiFAR exploits the observation that GPUs executing identical work remain largely synchronized, allowing it to speculate past the barrier and eliminate explicit synchronization latency.

Low-Latency LLM Inference Optimizations. Programmatic Dependent Launch (PDL) [12] reduces inter-kernel idle time by allowing a kernel to begin while the preceding one is still running. Megakernels [10] fuse the entire forward pass into a single kernel, eliminating launch overheads and enabling weight prefetching. Mirage [50] provides a compiler that automatically fuses operators into Megakernels, making them easier to implement. Speculative Decoding [51] reduces the number of sequential decode steps by generating multiple tokens in parallel. SiFAR is complementary to all of these: it reduces the latency of All-Reduce communication within each decode step, which these techniques do not address.

Collective Communication Optimizations. During training, All-Reduce communication accounts for a major fraction of training time due to the large volume of gradient data exchanged. Multiple works exploit tolerance to gradient errors by communicating only a subset of gradients [52]–[54], applying lossy compression [55], [56], or using quantization [57]–[59]. Our work is orthogonal as we focus on low-latency inference, where the payload size is very small, and the latency of the synchronization barrier is significant. During inference, PRESERVE [60] proposes prefetching model weights from HBM to L2 cache during All-Reduce operations to overlap memory access with communication. However, our evaluation setup using Megakernels already prefetches weights into shared memory, and for MoE models, prefetching during All-Reduce is not possible because expert selection depends on the All-Reduce output. SiFAR directly reduces All-Reduce latency rather than hiding it, providing benefits even when prefetching opportunities are exhausted.

In-Network Push-Based Reduction. Several works propose push-based in-network reduction where GPUs send data to a switch, which reduces and broadcasts without a top barrier. SwitchML [61] uses programmable switches for streaming aggregation during training. Klenk et al. [62] propose push-based reduction in NVSwitch fabrics, achieving up to 18\(\times\) speedup for small messages. CAIS [63] extends NVSwitch with compute-aware request merging that unifies push and pull. These approaches require hardware modifications to the switch. In contrast, SiFAR achieves barrier-free reduction on current systems by combining dual buffering and speculative reduction, without requiring push-based hardware support.

8 Discussion↩︎

8.1 Monotonicity Assumption and Correctness↩︎

SiFAR’s speculative reduction relies on a validation mechanism to ensure correctness, which depends on the monotonicity assumption. Consider the case where GPU1 is ahead of GPU0, so GPU1 reads GPU0’s payload before it finishes writing. In this case, the verification flag correctly indicates an error. However, if GPU1 slows down after reading the payload and only reads the verification flag after GPU0 has caught up, GPU1 would read an incorrect payload but a correct flag, causing incorrect verification. Thus, correctness assumes that if GPU1 is ahead of GPU0 at time \(T_1\), it remains ahead at time \(T_2 > T_1\), i.e., relative GPU progress is monotonic.

In LLM inference, this assumption holds for two reasons. First, CUDA Graphs or Megakernels eliminate host-side variability, the primary source of extreme GPU divergence. Second, GPU kernels typically execute with a run-to-completion model. Unlike CPU processes that can be easily preempted, GPU kernels run atomically to completion unless explicitly preempted [64], [65]. Thus, once a kernel begins executing, its relative progress with respect to other kernels is unlikely to change significantly, i.e., if GPU1 is ahead of GPU0 when reading the payload, it will likely remain ahead when reading the verification flag. To verify this independently of the validation mechanism itself, we compare the output of SiFAR against the twoshot All-Reduce. Across 100,000 decode iterations spanning multiple models and TP configurations, SiFAR matches this baseline, confirming that no silent error occurs.

9 Conclusion↩︎

With the rise of reasoning models and agentic systems, token generation latency has become critical for LLMs. To minimize latency, inference engines reduce batching and use Tensor Parallelism to leverage aggregate memory bandwidth across GPUs. However, All-Reduce emerges as the dominant cost as other bottlenecks diminish. Removing All-Reduce improves throughput by 43% at TP=8 for Llama 3.1 8B on H200 GPUs. We present Synchronization-Free All-Reduce (SiFAR), which minimizes All-Reduce overhead through three techniques: redundant pull to reduce data transfer via in-switch computation, dual buffering to eliminate the bottom barrier, and speculative reduction to minimize the top barrier. We integrate SiFAR into Megakernels, a state-of-the-art approach for low-latency inference, and evaluate it on Llama-3.1-8B and Qwen3.5-397B-17B. SiFAR reduces All-Reduce latency by up to 52% and improves end-to-end throughput by 18.6% for Llama-3.1-8B and 13.1% for Qwen3.5-397B-17B at TP=8.

References↩︎

---

1. Sifar means zero in Urdu, reflecting our goal of making All-Reduce overheads “zero”.↩︎