DeepSpeed-Inference: Enabling Efficient Inference of Transformer Models at Unprecedented Scale

ArXiv: 2207.00032

🎯 Pitch

DeepSpeed Inference serves a trillion-parameter model in under 25ms by exploiting over 100 TB/s of aggregate memory bandwidth across hundreds of GPUs. It also runs models 25× larger than GPU memory on a single commodity workstation at over 50% of peak throughput.

---

1. Executive Summary

This paper introduces DeepSpeed Inference, a comprehensive system for efficient transformer model inference that spans single-GPU kernel optimization, multi-GPU dense and sparse model parallelism, and heterogeneous CPU+NVMe memory offloading. The system combines three named mechanisms—Deep-Fusion (a tile-granularity kernel fusion strategy that fuses reductions, data layout transformations, and GeMMs beyond element-wise operations), inference-optimized pipeline parallelism (a hybrid scheduling scheme that varies micro-batch counts between prompt processing and token generation phases, plus activation offloading to CPU), and ZeRO-Inference (a heterogeneous memory solution that streams model layers from CPU or NVMe into GPU memory to support large batch sizes)—evaluated on GPT-style decoder models from 1.5B to 530B parameters and MoE models up to 2 trillion parameters on up to 256 NVIDIA A100 GPUs. DeepSpeed Inference reduces latency by up to 1.9× over FasterTransformer for dense models and 7.3× for a 1-trillion-parameter sparse MoE model (serving it under 25 ms by leveraging 128 TB/sec aggregate memory bandwidth), while ZeRO-Inference enables inference of models 25× larger than GPU-only solutions and delivers 84 TFLOPS—over 50% of A6000 peak—on a single GPU workstation, establishing that massive model inference is achievable on resource-constrained hardware only when the system aggressively exploits heterogeneous memory to free GPU capacity for large batch sizes.

2. Context and Motivation

The Core Problem: Transformer Inference Is Fragmented Across Competing Constraints

The paper addresses a fundamental tension in deploying large transformer models: inference requirements vary dramatically across model architectures, deployment scales, and performance targets, yet no unified system existed that could handle this diversity efficiently. As the authors state in Section I, the landscape is fractured along multiple axes simultaneously — model size spans from hundreds of millions to trillions of parameters, architecture type ranges from dense to sparse Mixture-of-Experts (MoE), the performance goal may be latency-critical (interactive applications) or throughput-oriented (batch processing), and the available hardware could be anything from a single consumer GPU to hundreds of datacenter GPUs.

This fragmentation matters because each axis demands different, often conflicting, optimization strategies:

• Latency-critical inference at small batch sizes is bounded by how fast model weights can be read from GPU memory to registers, not by how fast the computations execute. The core challenge is memory bandwidth utilization, not compute throughput. As the paper notes, "inference latency of a model is lower bounded by the time it takes to load all the model parameters from memory to registers."
• Throughput-oriented inference at large batch sizes shifts the bottleneck to compute utilization, but with the added complication that autoregressive generative models introduce token-by-token dependencies absent in training — the key-value activations for every previously generated token must be cached, consuming precious GPU memory.
• Sparse MoE models are substantially larger than quality-equivalent dense models (up to 8×, citing Fedus et al., 2021; Kim et al., 2021), requiring vastly more aggregate memory bandwidth for comparable latency, while their conditional computation structure demands different parallelism strategies entirely.
• Resource-constrained environments — where most data scientists actually work — simply lack enough GPUs to fit models like MT-NLG 530B, which requires ~1 TB of GPU memory just for inference, necessitating over three DGX-2 nodes. The paper emphasizes this accessibility gap: "Most data scientists simply do not have access to such GPU resources needed for inference of these massive models."

Without a unified solution, practitioners were forced to cobble together different tools — or abandon large models entirely — depending on which constraint dominated their particular setting. The paper's core contribution is demonstrating that these seemingly disparate problems share common underlying bottlenecks (memory bandwidth, communication overhead, kernel launch overhead) that can be addressed through a layered system architecture.

Why This Problem Is Important

The stakes are both practical and foundational.

Practical importance: Transformer models were growing 10× per year in parameter count at the time of writing (Section II), from hundreds of millions (BERT, RoBERTa, GPT) to hundreds of billions (GPT-3 175B, Gopher 280B, MT-NLG 530B). This exponential growth showed no sign of stopping. Meanwhile, the nascent MoE technique (Lepikhin et al., 2020; Fedus et al., 2021) enabled models to scale to trillions of parameters with sublinear increase in compute cost by conditionally activating only a fraction of parameters per input. This meant that the memory capacity and memory bandwidth demands of state-of-the-art models were outpacing hardware improvements, making efficient inference an existential bottleneck for deploying these models in real applications. If inference remained prohibitively slow or expensive, the immense investment in training large models would be wasted — you cannot serve a trillion-parameter model in an interactive application if generating a single token takes seconds.

Theoretical importance: The paper identifies a structural mismatch between the optimization landscape for training and inference. Training kernels (cuBLAS, CUTLASS) are designed to maximize compute utilization at very large batch sizes where arithmetic intensity is high. But inference at small batch sizes operates in a fundamentally different regime — memory bandwidth-bound, with arithmetic intensity near zero. As Section I states, "Inference kernels must... achieve high memory bandwidth utilization and high compute utilization at small batch sizes, whereas training kernels simply need to achieve high compute utilization at much larger batch sizes." This means the entire toolkit developed for training — including parallelism strategies like tensor, pipeline, and expert parallelism — could not simply be reused for inference without significant re-engineering. The paper's thesis is that inference is not a solved problem just because training is well-optimized; it requires a ground-up rethinking of kernels, parallelism schedules, and memory management.

Democratization imperative: The paper explicitly frames access to large model inference as a fairness concern. Training MT-NLG 530B required massive GPU clusters accessible only to well-funded organizations, but even inference on that trained model required similar resources. This created a two-tier ecosystem: organizations with hundreds of GPUs could deploy large models, while individual researchers and smaller teams — who might benefit enormously from fine-tuning or evaluating these models — were locked out. ZeRO-Inference's stated goal is to "democratize access to large transformer inference by enabling them on systems with limited GPU resources." This is not just a performance optimization; it is an attempt to change who can participate in large model research.

Prior Approaches and Where They Fall Short

The paper organizes prior work into four categories and identifies specific deficiencies in each.

1. Parallelism Strategies Designed for Training

Tensor parallelism (Megatron-LM, Shoeybi et al., 2019), pipeline parallelism (GPipe, PipeDream, Huang et al., 2018; Harlap et al., 2018; Narayanan et al., 2021), expert parallelism (Fedus et al., 2021), and 3D parallelism combining all three (DeepSpeed, Rajbhandari et al., 2020) had been extensively developed for training. However, these strategies make assumptions that break during inference:

• Tensor parallelism incurs all-reduce communication in every transformer layer. At small inference batch sizes, this communication overhead dominates because there is insufficient computation to hide it. Moreover, tensor parallelism cannot scale efficiently across node boundaries where inter-node bandwidth (e.g., InfiniBand) is a fraction of intra-node bandwidth (e.g., NVLink/NVSwitch). The paper notes that "tensor parallelism is often restricted to groups of GPUs sharing the high-bandwidth interconnect within a node."

• Pipeline parallelism introduces pipeline bubbles — idle periods where some pipeline stages wait for others. Training hides these bubbles through large numbers of micro-batches, but inference with autoregressive generation has data dependencies between tokens (each token depends on the previous one's key-value activations), making it impossible to arbitrarily batch across sequence positions. Section IV-B explains: "These data dependencies induce frequent pipeline bubbles that degrade inference performance."

• Expert parallelism requires all-to-all communication between all expert-parallel GPUs to route tokens to their assigned experts. The latency of all-to-all scales linearly with the number of devices — an unacceptable cost when scaling to hundreds of GPUs for massive MoE models.

The paper does not reject these parallelism strategies but argues they are "mainly designed for maximizing training throughput and their effectiveness can be limited during inference because of insufficient parallelism with small batch sizes."

2. Existing Transformer Kernel Libraries

Several works had developed optimized transformer kernels, but with critical gaps:

• Stochastic transformer kernels (DeepSpeed's BERT training record) fused operators and reduced activation memory, but targeted training throughput at large batch sizes, not inference latency at small batch sizes.
• Ivanov et al. (2021) analyzed data movement patterns in transformer training and fused elementwise and reduction operators, but again focused on training.
• TurboTransformers (Fang et al., 2021) fuses elementwise and reduction operators specifically for inference, which is the closest prior work. However, the paper distinguishes Deep-Fusion by noting it "fuses a wider variety of operators, such as head-wise transformation that requires additional data layout transformation and layers beyond the self-attention sublayers, such as the intermediate layers and MoE specific layers." Crucially, TurboTransformers "do not consider support for KV-caching," which is essential for autoregressive generation.
• E.T. (Chen et al., 2021) combines fusion, custom GeMM, and pruning to accelerate inference. The paper later benchmarks against E.T. (Figure 12) and shows 1.7× lower latency for DistilBERT and 1.4× for BERT, attributing the gain to Deep-Fusion fusing more operators. But the larger gap is scope: E.T. targets small encoder models on single GPUs, while DeepSpeed Inference targets billion-to-trillion-parameter decoder and MoE models across hundreds of GPUs.

3. General DNN Inference Frameworks

Frameworks like TVM (Chen et al., 2018), ONNXRuntime, and TensorRT provide compilers and runtimes for deploying DNN models efficiently. The paper acknowledges these but identifies a scaling limitation: "These platforms have been mostly focused on optimizing DNN models that can fit in a single GPU, such as small transformers with a few hundreds millions of parameters." They do not address the multi-GPU memory and bandwidth challenges of models with tens or hundreds of billions of parameters.

4. Model Compression Techniques

Quantization, distillation, and sparsification can reduce computation time and memory consumption. The paper positions itself as complementary: "Our work is complimentary to these model compression techniques and can be combined together to boost performance further."

The Most Direct Baseline: FasterTransformer

The paper's primary comparison point for dense models is FasterTransformer (FT), an NVIDIA library that had become the de facto standard for efficient transformer inference. FT supports multi-GPU tensor and pipeline parallelism, KV-caching for autoregressive generation, and FP16 computation. This makes it the most feature-complete prior system and the strongest baseline.

However, FT has several shortcomings that the paper exploits:

• Suboptimal small-batch kernel performance: FT relies on cuBLAS for GeMM operations, which the paper argues is "not well tuned for extremely small batch sizes, and cannot achieve good memory-bandwidth utilization" (Section III-A). The paper's custom SBI-GeMM specifically targets this regime.
• No operator fusion beyond elementwise: FT does not fuse the wider range of operations — layernorm, attention transformations, residual additions — that Deep-Fusion targets, resulting in higher kernel launch and data movement overhead.
• No inference-optimized pipeline schedule: FT uses training-derived pipeline schedules that do not account for autoregressive dependencies between generated tokens or the differing computational characteristics of prompt processing versus token generation.
• No INT8 inference support for decoders: At the time of writing, FT's INT8 support was limited to encoder models like BERT, not GPT-style decoders. DeepSpeed Inference provides INT8 support for decoder models, achieving up to 1.95× further latency reduction over FP16.
• No sparse MoE support: FT targets dense transformers only. The PyTorch-based MoE baseline used for comparison lacks expert-slicing, the parallelism-coordinated all-to-all communication optimization, and the sparse kernel optimizations that DeepSpeed-MoE provides.

The paper's evaluation (Figure 6) shows DeepSpeed Inference's FP16 implementation achieves up to 1.55× lower latency than FT-FP16 at small batch sizes and up to 1.57× at larger batch sizes, with INT8 pushing these gains to 1.95× and 1.93× respectively. This is measured with identical tensor parallelism configurations, so the differences are purely attributable to kernel and scheduling optimizations.

A Gap Beyond Performance: Accessibility

Even with optimal multi-GPU inference, large models remained inaccessible to most practitioners. The paper quantifies this concretely: MT-NLG 530B requires ~1 TB of GPU memory for inference, demanding over three DGX-2 nodes with more than two dozen A100 40GB GPUs. Prior work on memory offloading — most notably ZeRO-Infinity (Rajbhandari et al., 2021) — had demonstrated that training could leverage heterogeneous CPU+NVMe memory to break the GPU memory wall. However, inference has different requirements than training: inference activations are smaller (no optimizer states, no gradients), but the autoregressive KV-cache grows with sequence length and batch size, creating a different memory pressure pattern.

The paper adapts ZeRO-Infinity's offloading techniques to inference with a crucial design decision: rather than pinning as many weights as possible in GPU memory (the intuitive strategy), ZeRO-Inference streams layers from CPU/NVMe on demand and uses the freed GPU memory to support much larger batch sizes. The counterintuitive insight is that for massive models, the small fraction of weights that would fit in GPU memory anyway provides negligible latency benefit, while the larger batch sizes enabled by offloading amortize the PCIe transfer cost through increased compute time per load. This design choice — prioritizing batch size over weight caching — is what enables 84 TFLOPS (over 50% of peak) on resource-constrained hardware, and it represents a genuinely different optimization philosophy from prior work.

How This Paper Positions Itself

The paper's intellectual contribution is not a single technique but a layered system architecture where each layer addresses a distinct constraint while remaining compatible with the others. This architecture directly reflects the problem decomposition:

1. Layer 1: Inference-optimized single-GPU kernels — address the memory bandwidth bottleneck at small batch sizes through Deep-Fusion (tile-granularity operator fusion) and SBI-GeMM (custom GeMM with cooperative-group reduction and cache-line-aware data layout). This layer does not assume or require multiple GPUs.

2. Layer 2: Many-GPU dense inference — leverage tensor parallelism for aggregate memory bandwidth within a node, combined with inference-optimized pipeline parallelism (hybrid micro-batch scheduling, activation offloading, communication-aware scheduling) for scaling across nodes. This layer builds on Layer 1's kernels.

3. Layer 3: Massive-GPU sparse inference — orchestrate tensor, data, and expert parallelism for MoE models, with the novel parallelism-coordinated communication (PCC) optimization that reduces all-to-all latency from O(p) to O(p/L) + O(L) by exploiting data replication across tensor-parallel ranks.

4. Layer 4: Heterogeneous inference (ZeRO-Inference) — when GPU resources are insufficient, leverage CPU and NVMe memory to stream model layers, using the freed GPU memory for large batch sizes. This layer does not require the parallelism of Layers 2-3 but achieves complementary goals (accessibility and throughput on limited hardware).

The paper explicitly presents this as a "three-layered system architecture" (Section I) but ZeRO-Inference adds the fourth layer addressing resource constraints. The key claim is that this decomposition is both necessary and sufficient for the diversity of transformer inference: necessary because each layer addresses a fundamentally different constraint that no single technique could solve, and sufficient because the layers compose cleanly (kernels are used by all parallelism strategies; pipeline parallelism extends tensor parallelism; ZeRO-Inference works independently or alongside GPU parallelism).

This contrasts with prior work, which the paper characterizes as solving fragments of the problem without a unified framework. FasterTransformer provides kernels and parallelism but not MoE optimization or heterogeneous memory; PyTorch-MoE provides expert parallelism but not optimized kernels or communication; ZeRO-Infinity provides memory offloading but was designed for training, not inference. DeepSpeed Inference's contribution is the integration — showing that these can be combined into a single system that achieves state-of-the-art results across all regimes simultaneously.

3. Technical Approach

3.1 Reader Orientation

DeepSpeed Inference is a layered inference serving system — think of it as a runtime engine that takes a trained transformer model and executes it efficiently across whatever hardware is available (one GPU, many GPUs, or even a single GPU augmented with CPU and NVMe storage) to generate text in response to user prompts. The fundamental problem it solves is that transformer model inference spans a fragmented landscape where the optimal strategy depends on model size (millions to trillions of parameters), model architecture (dense vs. sparse Mixture-of-Experts), performance target (low latency for interactive chat vs. high throughput for batch processing), and hardware budget (single consumer GPU vs. hundreds of datacenter GPUs). Rather than building separate systems for each regime, DeepSpeed Inference provides a single composable stack where each layer — optimized single-GPU kernels, multi-GPU dense parallelism, massive-GPU sparse parallelism, and heterogeneous memory offloading — addresses a distinct constraint while reusing the layers below it. The insight is that these constraints (memory bandwidth, communication overhead, GPU memory capacity) are orthogonal enough that solving each in isolation and then composing the solutions yields a system that is simultaneously state-of-the-art for single-GPU latency, for 256-GPU sparse model inference, and for running a 530B model on a single A6000 workstation.

3.2 Big-Picture Architecture (Diagram in Words)

The system is organized as four layers stacked vertically, where higher layers add capabilities for specific regimes while depending on lower layers for their core computation:

• Layer 1: Inference-Optimized Single-GPU Kernels — takes a single transformer layer's operations on a single GPU and produces the most efficient execution schedule for the given batch size. Two mechanisms work together: Deep-Fusion, which tiles the computation space to fuse operators that have cross-tile data dependencies (reductions, data layout transformations, GeMMs) into a single kernel launch, eliminating intermediate writes to GPU global memory; and SBI-GeMM, a custom matrix multiplication kernel tuned for the extremely small (batch size 1–16) and "skinny" (small hidden dimension) GeMMs characteristic of inference, using cooperative-group reduction and a transposed weight layout to saturate memory bandwidth. A CUDA Graph wrapper further eliminates CPU-side kernel launch overhead by recording the kernel trace and replaying it.

• Layer 2: Many-GPU Dense Inference System — takes a dense transformer model that is too large for a single GPU's memory or too slow on a single GPU's memory bandwidth, partitions it across multiple GPUs using tensor parallelism (within a node, to leverage aggregate memory bandwidth) and pipeline parallelism (across nodes, to increase memory capacity), and coordinates execution. The novel component is inference-optimized pipeline parallelism: a hybrid scheduling scheme that uses different numbers of micro-batches for prompt processing (many micro-batches to hide pipeline bubbles) versus token generation (few micro-batches to reduce total weight-loading time), activation offloading to CPU memory when the KV-cache exhausts GPU memory, and a communication-aware PCIe offload schedule where odd-numbered GPUs offload odd layers and even-numbered GPUs offload even layers to avoid bus contention.

• Layer 3: Massive-GPU Sparse Model Inference System — takes a Mixture-of-Experts transformer model (which contains both dense attention layers and sparse expert feed-forward layers) and partitions it across potentially hundreds of GPUs using a combination of tensor parallelism (for dense components and to split individual experts), expert parallelism (to distribute different experts across GPUs), and data parallelism (to replicate non-expert parameters at zero communication cost). The key innovation is Parallelism-Coordinated Communication (PCC): when tensor-parallel operators precede expert-parallel operators, the data replication inherent in tensor parallelism's all-reduce means the subsequent all-to-all communication does not need to involve all GPUs — only one representative from each tensor-parallel group needs to participate, reducing all-to-all latency from scaling with total GPU count $p$ to scaling with $p/L$ where $L$ is the tensor-slicing degree. Additionally, sparse kernel optimizations replace sparse einsum scatter/gather operations with dense table-based data layout transformations that avoid $E - 1$ out of $E$ multiplications with zero per token ($E =$ number of experts), reducing complexity from $S \times E \times M \times c_e$ to $S \times M \times c_e$, where $S$ is the number of tokens, $M$ is the hidden dimension, and $c_e$ is the expert capacity.

• Layer 4: ZeRO-Inference (Heterogeneous Memory Inference) — takes a massive model that will not fit in aggregate GPU memory (e.g., 530B parameters on a single A6000) and streams it through the available GPU(s) by keeping the full model in CPU DRAM or NVMe storage, loading one or a few layers into GPU memory at a time for computation, and using the freed GPU memory to support very large batch sizes. Prefetching overlaps layer transfer with computation of the current layer, and in multi-GPU settings each GPU fetches only a partition of each layer, using the fast GPU-GPU interconnect (NVLink) to aggregate partitions, effectively increasing the PCIe bandwidth by the number of GPUs.

Information flows as follows: an input prompt arrives → Layer 1's fused kernel executes the first transformer layer's operations on the current GPU(s) → if multiple GPUs are used, Layer 2's tensor or pipeline parallelism coordinates the layer's distribution and communication → if the model is sparse, Layer 3's MoE orchestration routes tokens to the correct expert GPUs using the PCC-optimized all-to-all → if GPU memory is insufficient, Layer 4's offloading streams the next layer's weights from CPU/NVMe while the current layer computes → the process repeats for all layers, generating output tokens autoregressively, with the KV-cache managed according to Layer 2's activation offloading policy.

3.3 Roadmap for the Deep Dive

• First, the Deep-Fusion kernel fusion mechanism, because every subsequent layer depends on the single-GPU kernel performance that Deep-Fusion enables. We need to understand how the system solves the memory bandwidth wall before we can scale it to multiple GPUs.
• Second, the SBI-GeMM custom matrix multiplication kernel, which works in concert with Deep-Fusion to maximize memory bandwidth utilization at small batch sizes. These two together constitute Layer 1.
• Third, inference-optimized pipeline parallelism — the hybrid scheduling, activation offloading, and communication optimization — because this is the primary mechanism that distinguishes DeepSpeed Inference's multi-GPU strategy from training-oriented parallelism (Layer 2).
• Fourth, the MoE inference system: the parallelism orchestration across tensor, expert, and data dimensions, the PCC communication optimization, and the sparse kernel optimizations (Layer 3). This builds naturally on the dense parallelism of Layer 2.
• Fifth, ZeRO-Inference's heterogeneous memory design (Layer 4) — the design decision to prioritize batch size over weight caching, the prefetching and multi-GPU PCIe aggregation mechanisms, and why this achieves >50% of peak TFLOPS on resource-constrained hardware.

3.4 Detailed, Sentence-Based Technical Breakdown

This is primarily a systems engineering paper whose core idea is a layered architecture that solves the transformer inference problem by decomposing it into orthogonal constraints — memory bandwidth, communication overhead, GPU memory capacity — and addressing each with a dedicated mechanism that composes cleanly with the others. The following subsections walk through each mechanism in the order they appear in the system stack, from the lowest-level kernel optimizations to the highest-level memory offloading strategy.

---

Deep-Fusion: Tile-Granularity Operator Fusion for Inference Kernels

The inference latency of a transformer layer — particularly at small batch sizes — is not dominated by the raw floating-point operations but by two overheads: (1) the CPU-side cost of launching many small GPU kernels (each layernorm, each bias addition, each GeMM is a separate kernel launch), and (2) the repeated writes to and reads from GPU global memory (HBM) between kernels, since each kernel must write its output to global memory and the next kernel must read it back. Standard operator fusion — merging element-wise operations into a single kernel — addresses these overheads for simple operators like ReLU or dropout because element-wise operations have no cross-element dependencies; each output element depends only on the corresponding input element, so the computation can be arbitrarily tiled without communication between tiles. However, transformer layers contain operators with cross-element data dependencies: reductions (layernorm computes mean and variance across the hidden dimension), data layout transformations (transposing the QKV outputs into per-head format), and GeMMs (where each output element is a dot product across the entire input, requiring a reduction that spans all input elements). These operators cannot be fused by standard element-wise fusion because a naive fusion would require different tiles to exchange partial results, which on a GPU requires a global memory synchronization — exactly the overhead fusion aims to avoid.

Deep-Fusion solves this by tiling the computation space along dimensions that have no cross-tile data dependencies and keeping data-dependent dimensions within a single tile. The key insight is that for most transformer operations, certain iteration dimensions are "embarrassingly parallel" while others require reduction. Concretely:

• In a layernorm, the computation can be tiled along the token dimension (the batch and sequence axes): each tile processes all hidden-dimension elements for a subset of tokens. Within a tile, the thread block computes the mean and variance across the hidden dimension (a reduction), then normalizes and applies the affine transformation. Tiles do not need to communicate with each other because layernorm is applied independently to each token. The paper states: "all micro-operations in a layer-norm can be tiled along the token dimension, while the reduction dimensions are processed within a tile. This allows all the micro-operations inside a layernorm to be fused into a single kernel despite consisting of multiple reduction operations."

• In a GeMM followed by a bias addition and activation function (e.g., the intermediate FF layer), the output dimension is tiled: each tile computes a subset of output elements by performing the full input-weight dot product (a reduction within the tile) and then immediately applying the bias and GELU activation. Again, no cross-tile communication is needed because each output element's computation is independent.

• Fusing across operation boundaries requires that each output tile of the first operator maps to exactly one input tile of the second operator with no cross-tile dependencies. When this condition holds — which the paper ensures by carefully choosing tile boundaries that align across adjacent operations — intermediate data can stay in registers or shared memory rather than being written to global memory, eliminating the data movement overhead.

The paper articulates the condition formally: "two operators can be fused using Deep-Fusion if each tile of the second operator depends on exactly one output tile of the first operator."

What Deep-Fusion is not: The paper explicitly contrasts this with the standard GPU execution model where "if data produced by a thread-block is consumed by a different one, a global memory synchronization is needed which invokes a new kernel." Standard approaches accept this kernel boundary; Deep-Fusion eliminates it by ensuring the same thread-block that produces data also consumes it.

Fusion schedule for small-batch inference (Figure 1c). The paper shows the specific fusion regions for a transformer layer:

1. QKV GeMM + input layernorm: The layernorm (which normalizes the input) and the QKV projection matrix multiplication are fused into a single kernel. The layernorm's tiling along the token dimension aligns with the GeMM's tiling along the output (QKV) dimension.

2. Transposition + attention computation: The Q, K, V outputs from the GeMM are reshaped from a flat QKV dimension into per-head format (adding a head dimension and transposing), and then attention scores (QK^T), softmax, and attention context (scores × V) are computed. These are fused into a single kernel region.

3. Post-attention layernorm + intermediate GeMM: The output of the attention block is normalized and then projected through the first feed-forward layer. To enable fusion here, the paper makes a deliberate tradeoff: "we broadcast the input batch across the SMs and perform the same operations that come before GeMM, so that there is no need of communicating data between SMs for adding the GeMM schedule." This means that every SM redundantly computes the layernorm on the entire batch, even though each SM will only process a subset of the GeMM output. The paper notes this is still a net performance win for very small batch sizes because the replicated computation time is less than the time saved by avoiding a separate layernorm kernel launch and global memory round-trip.

4. Bias and residual addition: The bias additions after each linear layer and the residual connections are fused together.

Large-batch kernel variant. For larger batch sizes where the GeMM becomes compute-bound rather than memory-bandwidth-bound, the same fusion regions are used but the GeMM operations are left unfused and executed using CUBLAS. The rationale is that at large batch sizes, CUBLAS achieves excellent compute utilization and the kernel launch overhead of the non-GeMM operations is a smaller fraction of total latency, making aggressive GeMM fusion less important. The paper states: "We follow the same fusion strategy as discussed above, with the difference that we use CUBLAS for GeMM operations, and keep them unfused."

CUDA Graph integration. Even with Deep-Fusion reducing kernel count, there remains CPU-side overhead from launching the fused kernels for each forward pass. For small to moderate model sizes with small batch sizes, the paper observes that "as we reduce the actual execution time of the kernels, the main latency bottleneck shifts from kernel execution to the kernel launch overhead on the CPU side." To eliminate this, the system records the entire sequence of kernel launches during the first forward pass using CUDA Graphs, constructing a computation graph that is replayed for all subsequent requests. This effectively removes the per-token kernel launch overhead from the critical path.

Why this approach over alternatives: The alternative of not fusing — executing each operation as a separate kernel — would incur a kernel launch overhead (typically microseconds) and a global memory write-read cycle for every operation. For a transformer layer with roughly 15–20 distinct operations (layernorms, GeMMs, bias additions, activation functions, attention operations, residual additions), this could mean 15–20 kernel launches per layer. For a 96-layer model like GPT-3 175B at batch size 1, the cumulative overhead would dominate the actual computation time. Deep-Fusion reduces this to 4 fused kernel regions per layer, a 4–5× reduction in launch count. The alternative of using a compiler-based approach (TVM, XLA) could achieve some fusion automatically, but the paper's manual tiling strategy exploits domain knowledge about transformer data dependencies that general-purpose compilers might miss — specifically, the observation that for small batch sizes, replicating work across SMs is cheaper than the communication needed to avoid replication.

---

SBI-GeMM: Custom Matrix Multiplication for Memory-Bandwidth-Bound Inference

Standard GPU matrix multiplication libraries (cuBLAS, CUTLASS) are optimized for throughput at moderate to large batch sizes where each matrix dimension is large enough to saturate GPU compute units and amortize data loading. However, inference at small batch sizes produces "skinny" GeMMs — matrix multiplications where one dimension (the batch size) is extremely small (1–16) while the other dimensions (hidden size, number of heads) are moderate. The paper identifies that "neither cuBLAS nor CUTLASS GeMM libraries are well tuned for extremely small batch sizes, and cannot achieve good memory-bandwidth utilization." The primary bottleneck shifts from compute to memory bandwidth: the time to load the weight matrix from GPU memory dominates, and the key metric is how efficiently those loads saturate the available memory bandwidth.

SBI-GeMM (the paper does not expand the acronym, but the context — "Small Batch Inference GeMM" — and the described features suggest this reading) is a custom GeMM implementation designed around three design principles:

1. Tiling Strategies (Figure 1a). The computation is tiled along the output dimension (the weight matrix's column dimension, typically the hidden size or FF intermediate size). This means each thread block computes a contiguous subset of output elements by performing the full dot product between the input vector and a partition of the weight matrix. The reduction across the input dimension is kept entirely within a thread block, eliminating cross-block communication. For very small models where the output dimension is too small to create enough parallel tiles to occupy all SMs (and thereby achieve high memory bandwidth), the paper additionally tiles the input dimension and implements the GeMM as a two-kernel sequence: the first kernel computes partial dot products, and the second kernel reduces across the input-dimension tiles. The paper notes: "For small models, where the output dimension is too small to create enough parallel tiles to achieve good memory bandwidth, we tile the input dimension as well and implement GeMM as two kernels to allow for reduction across tiles."

2. Cooperative-Group Reduction. After each thread block computes its partial outputs, a reduction across warps within the block is needed to sum the partial contributions to each output element. Standard practice is a binary tree reduction in shared memory requiring multiple __syncthreads() barriers — one per tree level — which serializes warp execution and creates a performance bottleneck. SBI-GeMM avoids this by first performing a data layout transpose in shared memory such that partial results for the same output element become contiguous in memory, then using a single warp to perform the final reduction directly in registers using cooperative-group collectives (CUDA's shfl_down_sync or equivalent). The paper describes the process: "At the end, the first thread of each warp holds the final result and writes it to shared memory. The results in shared memory are contiguous, allowing for a coalesced write to global memory." This reduces the reduction latency from O(log W) synchronization points to O(1).

3. Leveraging Full Cache-Line Utilization. GPU L1 cache lines are 128 bytes. A standard FP16 coalesced memory access with one element per thread per warp reads 32 × 2 = 64 bytes — only half a cache line. Reading two elements per thread along the output dimension would use the full cache line but would reduce the number of parallel tiles (since each tile now processes twice as many outputs), hurting parallelism. SBI-GeMM's solution is to transpose the weight matrix offline during initialization so that $M$ rows for each column are stored contiguously in memory, allowing each thread to read $M$ elements along the input dimension — not the output dimension — thereby consuming the full 128-byte cache line without reducing tile parallelism. The paper sets $M = 2$ for FP16 (2 × 16 bits × 32 threads = 128 bytes) and $M = 4$ for INT8 (1 × 8 bits × 32 threads × 4 elements = 128 bytes).

Why this design? The three techniques together address the specific pathology of small-batch GeMMs: the computation has low arithmetic intensity (few FLOPs per byte loaded), so performance is entirely determined by how fast the weight matrix can be streamed from HBM to registers. Standard libraries are optimized for high arithmetic intensity and leave memory bandwidth on the table at low intensity. The tiling strategy ensures enough parallel work to saturate memory channels; the cooperative-group reduction eliminates synchronization overhead in the critical path; and the cache-line-aware layout ensures that no memory bandwidth is wasted on partial cache-line reads. The composability with Deep-Fusion — the GeMM can be fused with preceding layernorm and following bias/activation — is also a design requirement that standard libraries do not satisfy.

Support for INT8. For INT8 inference, the paper uses CUTLASS INT8 GeMM implementations tuned for different batch sizes. Deep-Fusion adds a quantization operation before the GeMM (fused into the preceding kernel region) and the de-quantization after the GeMM is handled using CUTLASS's epilogue functionality (an operation applied to the GeMM output before writing to memory). This INT8 path provides an additional ~1.25–1.5× latency reduction over FP16 by halving the weight data volume and exploiting faster INT8 tensor core throughput.

---

Inference-Optimized Pipeline Parallelism for Autoregressive Generation

When a dense transformer model exceeds the memory capacity of a single GPU node, pipeline parallelism (PP) is used to split the model's layers across multiple nodes. During training, PP is relatively straightforward: large numbers of micro-batches (independent training examples) can be pipelined through the stages to hide the "bubble" where some stages are idle waiting for others. But autoregressive inference introduces three complications that make training-derived PP schedules inefficient:

Complication 1: Token-by-token data dependencies (Figure 2). In training, all tokens in a sequence are processed in parallel with causal masking. In autoregressive inference, each generated token becomes the input for generating the next token. This creates a data dependency across pipeline stages: the last stage produces the token, which feeds back as input to the first stage for the next generation step. If the pipeline depth is $P$, and we have $B$ ongoing sequences, the naive schedule would complete all $B$ sequences' token $t$ before starting token $t+1$, creating a pipeline bubble of size $P-1$ between tokens. The paper observes: "These data dependencies induce frequent pipeline bubbles that degrade inference performance."

The solution: dynamic micro-batch queuing. The paper sets the number of micro-batches equal to the pipeline depth $P$, which is the minimum needed to keep all stages busy, but avoids increasing it further due to "latency and memory costs of the larger batch size." Crucially, the system does not wait for all sequences to complete token $t$ before starting token $t+1$. Instead, when a stage finishes processing a micro-batch and produces its output, that micro-batch immediately advances to the next pipeline stage, and a new micro-batch (the next token for a different sequence) can enter the first stage. This effectively amortizes the pipeline bubble over all generated tokens across all sequences: the system maintains a steady flow of micro-batches at different stages of generation, and the pipeline bubble only occurs at the very beginning (startup) and end (drain) of processing a batch of sequences, rather than between every token. The paper describes: "We avoid intermediate pipeline bubbles by dynamically queuing micro-batches of generated tokens until the sequences terminate. The resulting schedule amortizes the pipeline bubble over all generated tokens without allocating extra activations from a larger batch size."

Complication 2: Prompt processing vs. token generation have different performance characteristics. When processing the initial prompt (e.g., 128 tokens), the model sees many tokens simultaneously, making the computation compute-bound and capable of saturating GPU cores. The number of micro-batches affects the pipeline bubble but not the per-micro-batch execution time. However, during token generation, each micro-batch contains only a single token per sequence, making the computation memory-bandwidth-bound: the execution time is dominated by loading model weights, not by computing. The paper observes: "the execution time for a micro-batch does not change much with change in the size of micro-batch as most of the time is spent in fetching model parameters." This means that during token generation, total latency is proportional to the number of micro-batches, because each micro-batch requires reloading the model weights from memory — there is essentially no reuse of weights across micro-batches. Efficient token generation therefore demands minimizing micro-batch count, but efficient prompt processing demands enough micro-batches to hide the pipeline bubble.

The solution: hybrid scheduling (Figure 3). DeepSpeed Inference uses different numbers of micro-batches for the two phases. During prompt processing, a larger number of micro-batches is used to minimize the pipeline bubble (since the per-micro-batch time is small relative to the bubble, and compute saturation matters). During token generation, the number of micro-batches is reduced — the paper suggests to the absolute minimum that still avoids idle stages, which is $P$ — to reduce the total weight-loading time. The paper uses an illustration in Figure 3 to show this: "We use larger number of micro-batches during the prompt processing stage to minimize the pipeline bubble, while during the token generation phase, we reduce the number of micro-batches to reduces the overall execution time."

Complication 3: KV-cache memory pressure. Autoregressive inference caches the key and value activations of every transformer layer for every previously generated token, so that future tokens can attend to the full sequence without recomputing. This KV-cache memory scales with (number of layers × number of sequences × sequence length × hidden dimension × 2 for keys and values × bytes per element). For large models with many sequences and long generation lengths, the KV-cache can exhaust GPU memory even when the model weights fit.

The solution: activation offloading to CPU with communication-aware scheduling. When allocated GPU activation memory exceeds a threshold, the system offloads KV-cache activations to CPU memory. The paper exploits the predictable reuse pattern: "The activations of sequence $s_i$ will not be used again until generating the next token of $s_i$." Unlike training activations which are needed for backpropagation, inference activations are needed only once more — for the next token's attention computation — making them ideal candidates for offloading.

However, naively offloading over PCIe introduces a new bottleneck: PCIe bandwidth is much lower than GPU HBM bandwidth, and multiple GPUs in a node typically share PCIe links. The paper's architecture-aware optimization has a specific rule: "odd-numbered GPUs offload activations for odd-numbered layers, while even-numbered GPUs offload activation for even-numbered layers." The rationale is that most system architectures share a single PCIe link across two GPUs. If both GPUs attempt to offload simultaneously, they contend for the bus, halving effective bandwidth. By partitioning layers such that GPUs sharing a link never offload simultaneously (odd GPUs handle odd layers at odd time steps, even GPUs handle even layers at even time steps), each GPU gets the full PCIe bandwidth when it needs to offload. The paper states: "Scheduling odd and even layer offloading across GPUs prevents contention on the PCIe link, allowing each GPU to fully leverage the PCIe bandwidth when it needs to offload." Additionally, the offloading is overlapped with computation so that the GPU is not stalled waiting for PCIe transfers.

Why these optimizations matter for inference but not training: In training, the backward pass is the dominant cost, and it intrinsically recomputes or retrieves activations through the computation graph. The forward pass activations are typically kept in memory for the backward pass, but training batch sizes are limited by the need to store optimizer states and gradients anyway. In inference, there is no backward pass, so the forward pass is the entirety of the cost, and memory pressure comes entirely from the KV-cache and model weights. This fundamentally changes the tradeoff: inference can afford to offload activations because the cost of retrieving them for the next token is small compared to reloading all model weights, whereas in training, offloading activations would stall the critical path of the much more expensive backward pass.

---

Parallelism Orchestration and Communication Optimization for Sparse MoE Models

Mixture-of-Experts models introduce sparsity into the transformer architecture by replacing the dense feed-forward (FF) block in each transformer layer with a Position-wise MoE layer containing $E$ experts, where a learned gating function selects the top-$k$ experts to process each token (typically $k = 1$ or $k = 2$). This allows scaling the total parameter count with sublinear increase in computation — because each token only activates $k$ out of $E$ experts — but introduces two new system challenges for inference:

1. Sparse models are much larger than quality-equivalent dense models (up to 8× larger by the paper's citation of prior work), requiring much higher aggregate memory bandwidth to achieve comparable latency. Every token incurs the latency of loading the chosen expert's parameters from memory, but the total number of expert parameters is $E \times (\text{FF size})$, which can be hundreds of billions for large $E$. This demands scaling across many more GPUs than dense models.

2. Sparse models have a fundamentally different computational structure: transformer layers contain both dense components (self-attention, layernorm) and sparse components (MoE FF). The dense components benefit from tensor parallelism (splitting individual operators across GPUs for aggregate memory bandwidth), while the sparse components require expert parallelism (placing different experts on different GPUs and routing tokens to the appropriate GPU).

Parallelism orchestration (Figure 4). The paper uses a combination of three parallelism strategies operating simultaneously:

• Tensor parallelism (tensor-slicing) splits individual operators — both the dense attention/FF parameters and individual expert parameters — across multiple GPUs within a tensor-parallel group. This leverages the aggregate memory bandwidth of those GPUs. For the non-expert (dense) parameters, the paper calls this "tensor-slicing" and it uses standard all-reduce communication to aggregate partial activations. For the expert parameters, the paper calls this "expert-slicing" — splitting each individual expert across multiple GPUs — which requires communication within the expert-slicing group but increases the memory bandwidth available to load any given expert's parameters.

• Expert parallelism places different experts on different GPUs. Tokens are routed via an all-to-all communication: each GPU sends the tokens assigned to experts that reside on other GPUs and receives tokens assigned to experts it hosts. Unlike tensor parallelism, expert parallelism does not reduce the computation granularity of individual operators (each expert processes its assigned tokens as a single batch), allowing it to scale to hundreds of GPUs without the fine-grained communication overhead that would plague tensor-slicing at that scale.

• Data parallelism replicates the non-expert (dense) parameters across data-parallel groups. Because these parameters are identical across the data-parallel replicas and process different subsets of tokens, no communication is needed between data-parallel GPUs — they each compute independently on their token subset. This allows scaling the dense components to match the number of GPUs used by expert parallelism at zero communication cost.

The paper provides a concrete example in Figure 4: with 16 total GPUs, 8 total experts, expert-slicing degree 2 (each expert split across 2 GPUs), expert-parallel degree 8 (8 distinct expert placements), tensor-slicing degree 4 (dense parameters split across 4 GPUs), and data-parallel degree 4 (4 replicas of dense parameters). The tensor-slicing groups are GPUs {0,1,2,3} and {12,13,14,15}; the expert-slicing pairs are {GPU 0, GPU 8}, {GPU 1, GPU 9}, ..., {GPU 7, GPU 15}; and the expert-parallel set spans all 16 GPUs.

Parallelism-Coordinated Communication (PCC) for All-to-All (Figure 5). The standard all-to-all communication required by expert parallelism scales poorly: with $p$ GPUs, the latency is $O(p)$ because each GPU must communicate with every other GPU. For hundreds of GPUs, this latency alone can exceed the target inference budget.

The paper makes a crucial observation: when tensor-slicing (or expert-slicing) is used alongside expert parallelism, the all-reduce inherent in tensor-slicing replicates data across the tensor-parallel ranks. This means that after a tensor-parallel operation, every GPU in a tensor-parallel group holds identical data. Therefore, the subsequent all-to-all for expert parallelism does not need to involve all GPUs — only one representative from each tensor-parallel group needs to participate, because the data is already replicated.

The paper formalizes this as a four-step PCC protocol (illustrated in Figure 5):

Step 1: Local transformation and splitting. The data on each GPU is transformed into a format that separates the contributions for different experts and splits the data along the tensor-parallel dimension. This is a local operation requiring no communication.

Step 2: Intra-model-parallel (MP) and inter-MP all-to-all. Within each tensor-parallel group (intra-MP), a small all-to-all redistributes data such that each rank holds data for a disjoint subset of experts. Then, across tensor-parallel groups (inter-MP), an all-to-all exchanges data between the representative ranks of different groups. Because each tensor-parallel group has $L$ GPUs (the tensor-slicing degree), the inter-MP all-to-all involves only $p/L$ GPUs rather than all $p$ GPUs. The paper states: "the all-to-all can happen within just the subset of devices that share the same tensor-slicing rank, since the data across tensor-parallel ranks are replicated."

Step 3: Intra-MP allgather. After the inter-MP all-to-all, the data for each expert now resides on one GPU per tensor-parallel group. An allgather within each tensor-parallel group replicates this data to all $L$ GPUs in the group, satisfying the need for expert-slicing to have the full expert parameters distributed.

Step 4: Local transformation. A final local transformation rearranges the data into the format expected by the subsequent computation.

Latency reduction analysis. Without PCC, the all-to-all latency is:

$T_{\text{baseline}} = p \cdot C_1 + C_2$

where $p$ is the total number of GPUs, $C_1$ is the per-peer latency (dominated by point-to-point communication setup), and $C_2$ is the bandwidth-dependent term (proportional to message size divided by bandwidth). With PCC and tensor-slicing degree $L$, the latency becomes bounded by:

$T_{\text{PCC}} \approx \frac{p}{L} \cdot C_1 + C_2$

for the inter-MP all-to-all step, plus an additional $O(L)$ term for the intra-MP allgather:

$T_{\text{PCC}} = \frac{p}{L} \cdot C_1 + C_2 + O(L)$

Concrete example. The paper gives the case of 128 GPUs with 8-way tensor-slicing ($L = 8$) and 128-way expert parallelism ($p = 128$). Without PCC: $T_{\text{baseline}} = 128C_1 + C_2$. With PCC: $T_{\text{PCC}} \approx (128/8)C_1 + C_2 + O(8) = 16C_1 + C_2 + O(8)$. This is a factor of ~8× reduction in the dominant linear latency term.

Why this works for inference specifically. The PCC optimization exploits a property that is always true when tensor and expert parallelism are combined — that tensor-parallel all-reduce replicates data — but it requires that the tensor-parallel and expert-parallel operations appear in a specific order (tensor-parallel before expert-parallel for the forward pass, or expert-parallel before tensor-parallel for the backward/final forward). In training, the backward pass reverses this order, making the optimization symmetric in the opposite direction. In inference, only the forward pass exists, simplifying the analysis but also making the optimization only applicable in one direction per transformer layer. The paper accounts for this by noting that after expert-parallel operators, a similar PCC scheme applies but with an allgather following the all-to-all to re-establish replicated data for subsequent tensor-sliced operators.

---

Sparse Kernel Optimizations: Replacing Sparse Einsums with Dense Table-Based Transformations

The MoE gating and token routing involves several operations that, if implemented naively with sparse tensor representations, introduce substantial overhead:

1. Gating function: For each token, compute a softmax over all $E$ experts, select the top-$k$, and produce a mapping from tokens to experts. The naive sparse representation is a one-hot vector of size $E$ per token, which is $E-1$ out of $E$ zeros.

2. Cumulative sum (cumsum): Convert the token-to-expert mapping into an expert-to-tokens mapping — for each expert, which token IDs are assigned to it. This requires scanning the token-to-expert assignments.

3. Scatter: Distribute tokens to their assigned experts. Naively implemented as a sparse einsum between the expert-to-token mapping matrix and the input token matrix, with complexity $S \times E \times M \times c_e$, where $S$ is the number of tokens, $E$ is the number of experts, $M$ is the hidden dimension, and $c_e$ is the expert capacity (typically very small). Because each token is assigned to only $k$ out of $E$ experts (usually $k = 1$), this sparse einsum performs $(E - 1)$ multiplications with zero per token — an enormous waste of computation.

4. Gather: After each expert processes its assigned tokens, redistribute the processed tokens back to their original sequence ordering. Again, a sparse einsum with the same wasteful complexity.

The paper identifies that "the sparse tensor representation in the gating function and sparse einsum operators introduce a significant latency overhead" and that the einsums "have a complexity of $S \times E \times M \times c_e$ ... ($E-1$) out of $E$ operators for each token are multiplications and additions with zeros."

Optimization 1: Replace one-hot vectors with dense tables. Instead of representing the token-to-expert mapping as a one-hot vector of size $E$ per token (mostly zeros), the system uses a table data structure: for each token, store only the integer ID of the assigned expert. This reduces the memory for the mapping from $O(S \times E)$ to $O(S)$ and eliminates all operations that processed the zeros. The paper describes: "we replace the one-hot representation of the token to expert mapping using a table data-structure, greatly reducing the memory overhead from eliminating all the zeros in the one-hot vectors."

Optimization 2: Parallel scan for inverse mapping. To create the expert-to-tokens mapping from the token-to-expert table, the system simply scans through the token-to-expert table in parallel: each GPU thread examines a subset of tokens, checks which expert each token is assigned to, and appends the token ID to that expert's list. This is a $O(S)$ operation rather than the $O(S \times E)$ cumsum over a sparse matrix. The paper states: "we create the inverse mapping (expert-to-tokens mapping table) from the tokens-to-expert mapping table by simply scanning though the token-to-expert table in parallel."

Optimization 3: Data-layout transformation replaces sparse einsum scatter. Rather than performing a sparse matrix multiplication to scatter tokens to experts, the system uses the expert-to-tokens mapping table to directly copy tokens: for each expert, look up which token IDs are assigned to it from the table, then copy those tokens' hidden states from the input buffer to the appropriate expert's input buffer. The paper calls this a "data-layout transformation" because it is essentially a gather operation driven by the mapping table rather than a matrix multiplication. The complexity drops from $S \times E \times M \times c_e$ to $S \times M \times c_e$: each token's hidden state is read exactly once and written exactly once, rather than being multiplied by a masking matrix with $E-1$ zero entries.

Optimization 4: Data-layout transformation replaces sparse einsum gather. Similarly, after expert computation, the processed tokens are copied back to their original sequence positions using the inverse mapping table, again at $O(S \times M \times c_e)$ cost.

Fusion. All steps except the final data-layout transformation (the gather after expert computation) are fused into a single kernel "using basic fusion principles." The paper notes these optimizations together "result in over 6× reduction in MoE kernel-related latency."

Why the sparse approach exists at all. The sparse einsum approach is the "obvious" generalization of dense matrix operations to sparse layout: it represents conditional computation as a sparse matrix multiply with a learned gating matrix. This is elegant from a mathematical perspective and easy to implement in frameworks like PyTorch that support sparse tensors. However, the elegance comes at a steep performance cost because the sparsity pattern — exactly 1-out-of-E non-zero per token — is so extreme that the sparse representation is almost entirely overhead. The paper's key insight is that the sparsity pattern is known and structured (it is produced by the gating function), so the routing can be implemented as an explicit table lookup and copy rather than a general sparse matrix operation.

---

ZeRO-Inference: Heterogeneous Memory Inference for Resource-Constrained Systems

When a model does not fit in aggregate GPU memory and the user lacks access to enough GPUs, the fallback is ZeRO-Inference, which leverages CPU DRAM and NVMe storage as an extension of the memory hierarchy. This is not simply a fallback — the paper argues it can be more efficient than GPU-only inference in some regimes because the freed GPU memory supports larger batch sizes.

The core design decision: pin weights vs. pin activations. The intuitive strategy for heterogeneous inference is to keep as many model weights as possible pinned in GPU memory, fetching only the remainder from CPU/NVMe when needed. This minimizes data transfer volume. The paper rejects this strategy for two reasons: "(i) it allows only small batch sizes which hurts efficiency, and (ii) the latency savings for hundred-billion parameter models are negligible since only a small fraction of the weights can fit in GPU memory anyway."

Instead, ZeRO-Inference adopts the opposite strategy: pin model weights in CPU DRAM or NVMe, and stream each layer into GPU memory only when it is needed for computation. The GPU memory is used primarily for activations, KV-cache, and the current layer's weights and computation. The paper explains the reasoning:

"Despite the latency of fetching model weights over PCIe, ZeRO-Inference is able to achieve high efficiency for two reasons. First, by limiting GPU memory usage of the model to one or a few layers of weights, ZeRO-Inference is able to use large batch sizes for inference. Second, a large model layer requires significant amount of compute, especially given their long input sequence length (e.g., 2048). For example, one GPT3-175B layer requires about 7 TFlops to process an input of batch size 1. Therefore, large batch sizes cause compute time to dominate the latency of fetching model weights, which ultimately improves efficiency."

Mathematical intuition. Let the time to compute a layer for a micro-batch be $T_{\text{compute}}$ and the time to transfer the layer's weights from CPU/NVMe to GPU be $T_{\text{transfer}}$. At batch size 1, $T_{\text{compute}}$ is small (the layer is memory-bandwidth-bound with low arithmetic intensity), and $T_{\text{transfer}}$ dominates. At large batch sizes, $T_{\text{compute}}$ grows (more computation per weight byte loaded), eventually exceeding $T_{\text{transfer}}$. The system's efficiency is $\frac{T_{\text{compute}}}{T_{\text{compute}} + T_{\text{transfer}}}$ — the fraction of time the GPU is computing rather than waiting for weights. ZeRO-Inference's design increases $T_{\text{compute}}$ by increasing batch size, driving efficiency toward 1.

Prefetching. To further overlap transfer and computation, ZeRO-Inference prefetches a configurable number of layers ahead. While layer $i$ is computing on the GPU, layer $i+1$ (and possibly $i+2$, etc.) is being transferred from CPU/NVMe to a GPU buffer. This trades increased GPU memory consumption (to hold the prefetched layers) for reduced idle time waiting for transfers. At small batch sizes where $T_{\text{compute}}$ is small, prefetching provides significant benefit by hiding the transfer latency. At large batch sizes where $T_{\text{compute}}$ dominates, prefetching provides diminishing returns because the transfer is already well-hidden.

Multi-GPU PCIe bandwidth aggregation. When multiple GPUs are available, ZeRO-Inference leverages them differently than DeepSpeed Transformer. Rather than partitioning the model across GPUs (tensor or pipeline parallelism), each GPU independently streams the full model but only fetches a partition of each layer's weights over PCIe. The GPUs then share their partitions over the fast GPU-GPU interconnect (NVLink, typically 600 GB/s vs. PCIe Gen4's ~32 GB/s for a single GPU), reassembling the full layer for computation. This effectively multiplies the aggregate PCIe bandwidth by the number of GPUs: with 2 GPUs sharing a layer, each transfers half the weights over its dedicated PCIe link in parallel, then they exchange halves over NVLink. The paper notes: "In multi-GPU scenarios, the aggregate PCI-e bandwidth is used to reduce the layer transfer time by having each GPU only fetch a partition of the layer and then aggregating partitions over the much faster GPU-GPU interconnect."

NVMe I/O optimizations. For models stored on NVMe, the paper implements bulk read/write requests with asynchronous completion, aggressive parallelization of I/O requests, work scheduling to avoid idle time, memory pinning to avoid copy overhead, and zero-copy data paths where possible. The paper does not claim novelty on these specific optimizations, citing them as inherited from ZeRO-Infinity (Rajbhandari et al., 2021), but their integration into the inference pipeline is new.

Why this achieves >50% of peak TFLOPS. The key to understanding ZeRO-Inference's efficiency (84 TFLOPS = 54% of A6000 peak 156 TFLOPS for FP16 tensor core operations, though the paper quotes 158.4 TFLOPS theoretical peak) is the ratio of compute to data transfer for large transformer layers at large batch sizes. A single GPT-3 175B transformer layer with hidden size 12,288 processes inputs of shape $[B, S, 12288]$ where $B$ is batch size and $S$ is sequence length. The layer's weight matrices total roughly $12 \times 12288 \times 12288 \approx 1.8 \times 10^9$ bytes (counting QKV, attention output, FF intermediate, and FF output). The computation for this layer is approximately:

• Attention: $O(B \times S^2 \times 12288 + B \times S \times 12288^2)$ FLOPs
• Feed-forward: $O(B \times S \times 12288 \times (4 \times 12288))$ FLOPs

At $B=1, S=128$ (small batch, short prompt), the FLOP count is ~3.5 TFLOPs per layer, and the compute time on A6000 (~38 TFLOPS effective for these GEMM shapes) is small compared to transferring 1.8 GB of weights over PCIe. At $B=32, S=2048$ (large batch, long prompt), the FLOP count scales by ~$32 \times (2048/128)^2 = 32 \times 256 = 8192$ for the attention term (which is $O(S^2)$) and by $32 \times 16 = 512$ for the feed-forward term (which is $O(S)$). The total compute time grows dramatically while the transfer time stays roughly constant (the same 1.8 GB per layer, minus the fraction that can be cached). Once $T_{\text{compute}} \gg T_{\text{transfer}}$, the GPU spends most of its time computing, and efficiency approaches the GPU's peak.

The "democratization" claim concretely. The paper states that ZeRO-Inference can inference a 530B model on a single A6000 GPU, which is "25× larger than the largest model that can be inferenced with a GPU-only solution." The A6000 has 48 GB of GPU memory. An FP16 530B model requires ~1 TB of parameters. Even with INT8, it requires ~530 GB. The GPU-only limit on 48 GB would be a model around 20B parameters (INT8, ~20 GB for weights, leaving ~28 GB for activations and KV-cache). ZeRO-Inference offloads all but one or a few layers to CPU/NVMe, needing only enough GPU memory for the active layer (~2–4 GB for weights of a large layer) plus activations. The remaining 40+ GB of GPU memory can be used for a large batch KV-cache, enabling the high throughput that makes the system practical despite the PCIe transfer overhead. This is not just a memory capacity trick — it is a deliberate tradeoff that converts unused GPU memory into batch size, which in turn converts PCIe latency into hidden overhead, yielding net efficiency that approaches GPU-only performance.

4. Key Insights and Innovations

Innovation 1: Inference Is Not Degraded Training — It Requires a Fundamentally Different Optimization Target

The paper's most important conceptual move is not any single technique but the reframing of inference as an optimization problem with a different objective function than training. Prior to this work, the dominant assumption — implicit in the reuse of training kernels (cuBLAS, CUTLASS), training parallelism strategies (Megatron-LM tensor parallelism, GPipe pipeline schedules), and training memory management (ZeRO-Infinity applied naively to inference) — was that inference is essentially "training's forward pass without the backward pass," and therefore that optimizations developed for training should transfer with minor adaptation. The paper systematically dismantles this assumption by identifying three structural mismatches that make training-optimized systems suboptimal for inference.

Mismatch 1: The arithmetic intensity regime is inverted. Training operates at large batch sizes where GeMM operations are compute-bound, with high arithmetic intensity (many FLOPs per byte loaded). cuBLAS and CUTLASS are engineered for this regime. Inference at small batch sizes — the regime for interactive applications with latency constraints — operates at near-zero arithmetic intensity, where the bottleneck is not compute but memory bandwidth: how fast model weights can be streamed from HBM to registers. The paper's diagnosis (Section III-A) is that "neither cuBLAS nor CUTLASS GeMM libraries are well tuned for extremely small batch sizes, and cannot achieve good memory-bandwidth utilization." This is not an incremental tuning gap — it is a categorical mismatch. A library optimized for compute throughput will systematically underutilize memory bandwidth when the work per weight byte is small, because its tiling strategies, wave scheduling, and reduction algorithms are designed for the opposite regime.

Mismatch 2: Autoregressive dependencies break pipeline schedule assumptions. Training pipeline parallelism schedules (GPipe, PipeDream) assume that micro-batches are independent, so pipeline bubbles can be amortized over many micro-batches without penalty. In autoregressive inference, this independence does not hold: "the inputs to the model inference are previously-generated outputs" (Section IV-B), creating a data dependency where the last pipeline stage produces a token that feeds back as input to the first stage. The paper identifies that this dependency "induces frequent pipeline bubbles that degrade inference performance" — a problem that simply does not exist in training because all tokens are processed in parallel with causal masking. The earlier prior section describes how DeepSpeed Inference solves this with dynamic micro-batch queuing, but the insight is that the pipeline scheduling problem is structurally different for inference, not just differently parameterized.

Mismatch 3: KV-cache creates a unique memory pressure pattern. Training stores activations for the backward pass; inference stores KV-cache activations for future token generation. While both consume GPU memory, the reuse patterns differ fundamentally. Training activations are consumed once in the backward pass with a fixed order; KV-cache activations have a predictable but different reuse pattern — "the activations of sequence $s_i$ will not be used again until generating the next token of $s_i$" (Section IV-C2) — which makes them amenable to a specific offloading strategy (CPU offload with PCIe scheduling) that would be disastrous for training because it would stall the backward pass's critical path. The paper's contribution is not just the offloading mechanism but the identification that inference memory pressure is qualitatively different from training memory pressure, enabling optimizations that are invalid for training.

Significance beyond performance. This reframing has implications beyond DeepSpeed Inference. It establishes a design principle: when building inference systems, start from the inference workload's characteristics (batch size distribution, memory access patterns, dependency structure), not from training infrastructure. The paper's layered architecture — where kernels are redesigned for memory bandwidth, pipeline schedules are redesigned for autoregressive dependencies, and memory management is redesigned for KV-cache access patterns — is a concrete instantiation of this principle. The speedups over FasterTransformer (1.55× to 1.93× at small batch sizes, 1.51× throughput for GPT-3 175B) are evidence that the principle yields practical gains, but the principle itself is the deeper contribution: it tells future system builders not to treat inference as an afterthought to training infrastructure.

This is a fundamental conceptual shift, not an incremental refinement. The paper is essentially arguing that the optimization target for inference is memory bandwidth utilization (for latency) and the ratio of compute to data transfer (for throughput), whereas the optimization target for training is compute utilization. Systems that confuse these two targets will leave performance on the table regardless of how well they execute.

---

Innovation 2: Tile-Granularity Fusion with Deliberate Work Replication Breaks the Element-Wise Fusion Barrier

Operator fusion — merging multiple operations into a single GPU kernel to eliminate kernel launch overhead and intermediate global memory writes — was a well-established technique before this paper, used in XLA, TVM, ONNXRuntime, and TensorRT. However, prior fusion was almost entirely restricted to element-wise operators (ReLU, dropout, bias addition, scaling) because these operators have no cross-element data dependencies: each output element depends only on the corresponding input element, so the computation can be arbitrarily tiled without communication between tiles. This restriction meant that transformer inference — where reductions (layernorm), data layout transformations (attention head transpositions), and GeMMs dominate the operation count — benefited only marginally from fusion, because the expensive operations remained unfused.

Deep-Fusion's conceptual contribution is the recognition that tiling along dimensions with no cross-tile dependencies while keeping reduction dimensions within a single tile enables fusing operators far beyond element-wise boundaries. The idea is simple in retrospect: for a layernorm, tile along the token dimension so that each tile independently computes its own mean, variance, and normalization across the hidden dimension. For a GeMM followed by bias and activation, tile along the output dimension so that each tile computes its own dot products and immediately applies bias and activation. The condition for fusing two operators is that "each tile of the second operator depends on exactly one output tile of the first operator" (Section III-B) — a generalization of the element-wise condition that does not require global synchronization.

The deliberate work replication tradeoff. The paper makes an explicit, non-obvious design choice that distinguishes Deep-Fusion from compiler-based fusion approaches: when tiling constraints require communication between tiles to avoid redundant computation, Deep-Fusion sometimes chooses to replicate work across SMs rather than communicate. The specific case is the post-attention layernorm + intermediate GeMM fusion (Figure 1c, region 3): "we broadcast the input batch across the SMs and perform the same operations that come before GeMM, so that there is no need of communicating data between SMs." This means every SM redundantly computes the layernorm on the entire batch, even though each SM will only process a subset of the GeMM output. For very small batch sizes, the paper observes that "in spite of replicating the work across SMs, we still gain performance benefit compared to the non-replicated, non-fused kernel implementation."

Why this is intellectually distinctive. Compiler-based fusion (XLA, TVM) typically optimizes for minimal total FLOP count — it would never deliberately replicate work because that increases total computation. Deep-Fusion optimizes for a different objective: minimizing global memory round-trips, even at the expense of redundant computation. This makes sense only in the specific regime where memory bandwidth is the bottleneck and compute is cheap — exactly the small-batch inference regime the paper targets. A compiler optimizing for FLOP minimization would reject this fusion as inefficient; Deep-Fusion recognizes that FLOPs are abundant and memory bandwidth is scarce.

This insight is fundamental rather than incremental because it changes the optimization criterion for kernel fusion. It implies that the right fusion strategy depends on the workload's arithmetic intensity: at high intensity (training, large-batch inference), avoid redundant computation; at low intensity (small-batch inference), avoid redundant memory traffic even if it means redundant computation. Future systems targeting both regimes may need to dynamically select fusion strategies based on batch size — exactly the hybrid approach the paper takes (Deep-Fusion schedule for small batches, CUBLAS GeMMs with unfused GeMMs for large batches).

---

Innovation 3: Expert-Parallel Communication Can Be Decoupled from GPU Count by Exploiting Tensor-Parallel Data Replication

Scaling Mixture-of-Experts inference to hundreds of GPUs faces a fundamental communication bottleneck: expert parallelism requires an all-to-all communication where every GPU must exchange tokens with every other GPU to route tokens to their assigned experts. The latency of this all-to-all scales as O(p) where p is the number of expert-parallel GPUs. For a 128-expert model on 128 GPUs, this means 128 pairwise communication rounds per transformer layer, each incurring setup latency proportional to point-to-point network latency. At the inference latency targets the paper targets (25 ms for a trillion-parameter model), this communication alone could consume the entire budget.

Prior work on MoE training (GShard, Switch Transformer, DeepSpeed-MoE) accepted this O(p) scaling as inherent, because the all-to-all is necessary to route tokens to the correct experts. The typical mitigation was to keep expert parallelism within a single node where NVLink provides higher bandwidth and lower latency, but this limits model scale to the number of GPUs per node. To scale to hundreds of GPUs across nodes, the O(p) latency becomes prohibitive.

The paper's Parallelism-Coordinated Communication (PCC) introduces a genuinely novel insight: when tensor parallelism (tensor-slicing or expert-slicing) is used alongside expert parallelism, the all-reduce operation inherent in tensor parallelism replicates data across all GPUs within each tensor-parallel group. This replication means that after a tensor-parallel operation, every GPU in a tensor-parallel group of size L holds identical data. Therefore, the subsequent all-to-all for expert routing does not need to involve all p GPUs — only one representative from each tensor-parallel group needs to participate, because the data is already replicated across the other L-1 GPUs in the group.

The reduction in communication latency is dramatic: from O(p) to O(p/L) + O(L), where L is the tensor-slicing degree. For 128 GPUs with 8-way tensor-slicing, this reduces the dominant linear term from 128 to 16 — an 8× reduction. The paper provides this concrete analysis in Section V-B: "when scaling to 128 GPUs with 8-way tensor-slicing and 128-way expert parallelism, this approach reduces the latency overhead of the all-to-all from (128C1 + C2) to (16C1 + C2) due to 8-way tensor-slicing."

Why this is conceptually distinctive, not just an engineering trick. The insight is not about implementing all-to-all more efficiently — it is about recognizing that the combination of parallelism strategies creates information redundancy that can be exploited to reduce communication. Tensor parallelism and expert parallelism are typically treated as independent design choices: tensor parallelism for dense components, expert parallelism for sparse components, with communication handled separately for each. PCC recognizes that they interact through data replication: the all-reduce that tensor parallelism requires (a cost) also produces replicated data (a benefit) that can be leveraged to reduce the cost of expert parallelism's all-to-all. This is a complementarity between parallelism strategies that prior work had not identified because it analyzed each parallelism dimension in isolation.

The significance extends beyond MoE inference. Any system that combines multiple parallelism strategies where one strategy replicates data and another requires all-to-all communication can potentially apply the same principle. The paper does not explore this generalization, but the pattern — identifying redundant communication created by one optimization and using it to accelerate another — is a transferable design principle for distributed systems, not just MoE transformers.

This is a fundamental contribution to the design space of distributed model parallelism. It changes the analysis from "tensor parallelism costs all-reduce, expert parallelism costs all-to-all" to "tensor parallelism costs all-reduce but also provides replicated data that can partially substitute for all-to-all." The cost-benefit calculus for choosing parallelism strategies shifts when this interaction is accounted for.

---

Innovation 4: For Memory-Constrained Inference, Freeing GPU Memory for Batch Size Beats Caching Model Weights

When a model does not fit in GPU memory, the intuitive strategy — and the default in many systems — is to cache as many model weights as possible in GPU memory and fetch only the overflow from CPU or NVMe. This minimizes total data transfer volume and, one would assume, maximizes throughput. The paper explicitly rejects this approach for inference and articulates a counterintuitive alternative that is arguably the deepest systems insight in the paper.

The rejected approach: "One approach is to pin as much of the model weights as possible into GPU memory, and fetch the remainder (from DRAM or NVMe) when needed for computation." The paper identifies two flaws: "(i) it allows only small batch sizes which hurts efficiency, and (ii) the latency savings for hundred-billion parameter models are negligible since only a small fraction of the weights can fit in GPU memory anyway" (Section VI-A).

The adopted approach (ZeRO-Inference): pin model weights entirely in CPU DRAM or NVMe, stream each layer into GPU memory on demand, and use the freed GPU memory to support very large batch sizes. The conceptual leap is recognizing that for large transformer layers, the compute time at large batch sizes can dominate the weight transfer time, making the transfer latency effectively free from a throughput perspective. The paper gives a concrete example: "one GPT3-175B layer requires about 7 TFlops to process an input of batch size 1." At batch size 32 with sequence length 2048, the compute time per layer grows dramatically (the feed-forward computation scales linearly with batch size and sequence length), while the transfer time for the layer's weights (~1.8 GB) remains constant. Once compute time exceeds transfer time, the GPU is compute-bound and the weight streaming happens in the background, hidden by computation.

Why this is a non-obvious insight. The default systems intuition — cache what you can to minimize data movement — is correct when compute time is proportional to data volume (e.g., element-wise operations, small models). But transformer inference at scale has a nonlinear relationship between batch size and compute time per weight byte. The feed-forward layers have compute that scales as O(B × S × d_model × d_ff), while the weight data volume is constant per layer. At sufficiently large batch sizes, the ratio of compute to weight data becomes arbitrarily large. In this regime, the bottleneck shifts from "how fast can I load weights" to "how much computation can I do per weight load," and the way to maximize computation per load is to maximize batch size — which requires maximizing available GPU memory for activations, not for weight caching.

This insight is fundamental rather than incremental because it inverts the optimization criterion for heterogeneous memory systems. Prior work (ZeRO-Infinity, and indeed most systems that use CPU offloading) optimizes for minimal data transfer. ZeRO-Inference optimizes for maximal compute-to-transfer ratio, which leads to the opposite design decision (offload everything, not cache anything) when the model is large enough that the GPU-resident fraction of weights is negligible anyway. The paper's evidence — 84 TFLOPS on a single A6000 (54% of peak) and the ability to serve models 25× larger than GPU-only solutions — demonstrates that this inversion is not just theoretically interesting but practically transformative.

Connection to the broader scaling landscape. This insight has implications beyond the specific ZeRO-Inference implementation. It suggests that as models continue to grow, the economic sweet spot for inference may shift from "enough GPUs to hold the model" toward "a few GPUs with enough batch size to hide PCIe latency." The paper's finding that ZeRO-Inference "offers over 50% better throughput than the GPU-only solution" even for models that fit in GPU memory (because it can use larger batch sizes) is a provocative result: it implies that streaming weights from CPU can be faster than keeping them in GPU memory if the batch size gain is large enough. This challenges the assumption that GPU memory capacity is always the binding constraint for inference performance and suggests a regime where PCIe bandwidth and GPU compute, not GPU memory size, are the relevant hardware limits.

---

Innovation 5: Inference Pipeline Scheduling Must Be Phase-Aware Because Prompt Processing and Token Generation Occupy Different Performance Regimes

Autoregressive transformer inference has two distinct computational phases: prompt processing (where the entire input prompt is processed in parallel to generate the first token) and token generation (where subsequent tokens are generated one at a time, each depending on the previous token's KV-cache). Prior work on pipeline parallelism for training treated all forward-pass micro-batches as homogeneous — each has the same computational characteristics, so a single pipeline schedule with fixed micro-batch count works uniformly. The paper identifies that this assumption breaks for inference because the two phases have categorically different performance characteristics:

• Prompt processing: many tokens per sample (e.g., 128), making the computation compute-bound. The per-micro-batch execution time is large enough that pipeline bubbles are the primary efficiency concern. More micro-batches → smaller bubble fraction → better efficiency.

• Token generation: a single token per sample, making the computation memory-bandwidth-bound. "The execution time for a micro-batch does not change much with change in the size of micro-batch as most of the time is spent in fetching model parameters. However, the overall execution time is proportional to the number of micro-batches, as the forward pass on each micro-batch requires fetching the weights all over again" (Section IV-C1). Fewer micro-batches → less total weight-loading time → lower latency.

The two phases therefore demand opposite micro-batch count strategies: prompt processing wants many micro-batches (to hide bubbles), token generation wants few micro-batches (to minimize weight reloads). A single schedule optimized for one phase will be suboptimal for the other.

The hybrid scheduling solution — using different micro-batch counts for prompt processing and token generation — is conceptually simple once the diagnosis is made, but the diagnosis itself is the contribution. The paper is identifying a structural heterogeneity in inference workloads that prior systems ignored because they imported training's homogeneous-micro-batch assumption. FasterTransformer, the primary baseline, does not implement this phase distinction, which is part of why DeepSpeed Inference achieves 1.51× higher throughput for GPT-3 175B (Figure 8).

Why this is a diagnostic contribution, not just a scheduling trick. The paper is essentially arguing that inference pipeline scheduling has an additional degree of freedom — the micro-batch count per phase — that training scheduling does not have (or does not need, because all training micro-batches are homogeneous). Recognizing this degree of freedom requires understanding that the performance regime (compute-bound vs. memory-bandwidth-bound) is not static across the lifetime of an inference request — it changes when the system transitions from prompt processing to token generation. This is a genuinely new characterization of the inference scheduling problem, and it implies that future inference systems should be designed with phase awareness as a first-class concept, not as an afterthought.

The prompt processing latency results (Figure 13: 1.18× and 3.06× speedup over FasterTransformer for different parallelism configurations) provide evidence that this diagnosis translates to real gains, but the conceptual value is in the framework: it tells system designers that when optimizing inference pipelines, they should analyze the two phases separately, characterize each phase's bottleneck (compute vs. memory bandwidth), and tune micro-batch counts accordingly, rather than searching for a single compromise setting.

This is an incremental refinement in mechanism (it is a scheduling policy, not a new computational primitive) but a fundamental contribution in diagnosis: it identifies a previously unrecognized degree of freedom in inference scheduling that arises from the autoregressive nature of decoder models. Prior systems missed it because they treated inference as a minor variant of training; the paper's reframing of inference as its own optimization target (Innovation 1) is what makes this diagnosis possible.

5. Experimental Analysis

Evaluation Methodology

• Dataset. The paper evaluates on GPT-style autoregressive transformer decoder models, parameterized by hidden dimension, number of layers, and attention heads following the GPT-3 family and its publicly available variants (Table I). For dense models, evaluated configurations span from 1.5B parameters (GPT-2) to 530B parameters (MT-NLG). For sparse Mixture-of-Experts models, configurations range from a 52B-parameter MoE (1.3B base + 128 experts) to a 2-trillion-parameter MoE (47B base + 128 experts, Table II). Latency-sensitive experiments use an input prompt of 128 tokens generating 8 tokens (dense) or 100 tokens (sparse) with varying batch sizes. Throughput-oriented experiments use an input prompt of 512 tokens generating 50 tokens. Resource-constrained experiments measure the maximum batch size possible for generating a single token.

• Base model(s). The system is model-agnostic in principle; evaluation covers GPT-2, GPT-Neo (1.5B–2.7B), GPT-J (6B), GPT-13B, GPT-NeoX-20B, GPT-50B, GPT-87B, LM-175B, and LM-530B for dense models, and a series of MoE models with base sizes of 1.3B, 2.4B, 8B, 24B, and 47B parameters each combined with 128 experts for sparse evaluation. The diversity of hidden dimensions (1,600 to 20,000), layer counts (16 to 96), and attention heads (20 to 128) is intentional: the paper states it aims to "cover a wide range of model configurations and different number of parameters." This is not a standard benchmark but a system performance evaluation across model scales.

• Metrics. Three metrics are used: (1) latency — end-to-end output generation time for a batch of input prompts, measured in milliseconds; (2) token throughput — tokens-per-second processed; and (3) compute throughput — TFLOPS per GPU. For resource-constrained evaluation, compute throughput at maximum batch size is the primary metric. The latency metric for sparse models is reported as per-token latency (100 tokens generated per measurement). Memory bandwidth utilization is reported as a percentage of theoretical peak HBM bandwidth aggregated across GPUs (e.g., "33% of peak memory bandwidth" for the 2T MoE model on 256 GPUs).

• Baselines. The primary dense model baseline is FasterTransformer (FT) [31], described as "an efficient implementation of transformer models provided by NVIDIA." Both FT and DeepSpeed Inference use identical tensor parallelism (TP) configurations where applicable, so all latency differences come from kernel implementations, not parallelism strategy. FT supports FP16 but, at the time of writing, did not support INT8 for decoder models ("FasterTransformer only supports INT8 computation for Transformer models with just the encoders, e.g., BERT, but not decoders"). For sparse MoE models, the baseline is a full-featured distributed PyTorch implementation that supports both tensor and expert parallelism [37], using the same data/tensor/expert parallelism configuration as DeepSpeed-MoE but without expert-slicing, PCC, or MoE-specific kernel optimizations. For smaller encoder models (DistilBERT, BERT), the paper also compares against E.T. [27], a state-of-the-art transformer kernel with fusion, custom GeMM, and pruning.

• Generation budget / compute accounting. The paper does not define a universal "generation budget" the way the example paper does with "number of generated solutions." Instead, compute is measured in several ways appropriate to the system: (1) latency at fixed workload (fixed prompt length, fixed number of output tokens, fixed batch size), which is the standard for latency-critical evaluation; (2) throughput at maximum feasible batch size, where each system is allowed to use the largest batch that fits in memory — this gives throughput-oriented systems full advantage of their memory optimization; (3) TFLOPS per GPU as a hardware efficiency metric normalized against theoretical peak. For multi-GPU experiments, parallelism configurations (TP degree, PP stages, EP degree) are specified per model in Tables I and II. The paper explicitly states that for the FT comparison, "Both the baseline and DeepSpeed Inference uses identical TP strategy so all the latency differences in these results come from the differences in kernel implementations." For throughput comparisons, each system uses the batch size that gives its best performance.

• Cross-validation / statistical protocol. The paper does not describe any cross-validation, statistical significance testing, or confidence intervals. The experiments are performance measurements of a deterministic system on fixed hardware; there is no model training or hyperparameter tuning being evaluated. The primary source of measurement variance would be hardware-level noise (GPU clock variation, network contention), which the paper does not discuss. All results appear to be single-run measurements. The lack of error bars or variance reporting is a limitation in assessing whether the reported speedups (e.g., "up to 1.55×") are statistically reliable or subject to run-to-run variation.

Main Quantitative Results

Latency-Critical Inference for Dense Models (Figure 6)

Headline claim: DeepSpeed-FP16 achieves up to 1.55× lower latency than FasterTransformer-FP16 at small batch sizes, and DeepSpeed-INT8 achieves up to 1.95× over the FP16 baseline. At larger batch sizes, the gains are up to 1.57× (FP16) and 1.93× (INT8).

The paper evaluates eight dense model configurations on 1, 8, and 16 batch sizes (Figure 6). In the single-batch (batch size 1) configuration:

• For the smallest models (GPT-2, GPT-Neo-2.7B), DeepSpeed-FP16 latency is visibly lower than FT-FP16, with the gap narrowing as batch size increases. The paper attributes this to the largest kernel-launch overhead and worst cuBLAS memory bandwidth utilization in the smallest GeMMs — precisely the regime where Deep-Fusion and SBI-GeMM provide the most benefit.
• For GPT-J-6B at batch size 1, FT-FP16 shows approximately 95 ms latency while DeepSpeed-FP16 shows approximately 60 ms — roughly 1.58× speedup (read from the bar chart).
• For GPT-13B at batch size 1, DeepSpeed-FP16 achieves roughly 40 ms vs. FT's 62 ms — roughly 1.55× speedup.
• For GPT-NeoX-20B, GPT-50B, and GPT-87B, the pattern continues with DeepSpeed-FP16 consistently lower than FT-FP16 across all batch sizes.
• For LM-175B (requiring 16-way tensor parallelism per Table I), DeepSpeed-FP16 shows approximately 500 ms at batch size 1 vs. FT-FP16's 800 ms — roughly 1.6× speedup. At batch size 8 and 16, the absolute latency gap narrows but the relative advantage persists.

The DeepSpeed-INT8 results (shown only where applicable for models where INT8 is supported) demonstrate a further reduction: for GPT-2 at batch size 1, INT8 latency is roughly half that of FT-FP16, consistent with the claim of up to 1.95× over the FP16 baseline. The paper notes that FT does not support INT8 for decoder models, so the INT8 comparison is against FT-FP16, not FT-INT8.

Throughput results on the same figure: The paper overlaid token throughput (tokens-per-second, right y-axis) on the same bar charts. DeepSpeed-FP16 achieves higher throughput than FT-FP16 across all configurations, with the absolute gap widening at larger batch sizes where the non-GeMM data-movement reduction from Deep-Fusion becomes a larger fraction of overall latency. DeepSpeed-INT8 roughly doubles the throughput of FT-FP16 for the smallest models.

Key interpretation point: The paper states that the INT8 improvements "further improve upon the DeepSpeed-FP16 performance by utilizing the higher peak of the INT8 tensor-cores compared to FP16." This means the INT8 gains come from both reduced weight data volume (half the bytes) and higher tensor-core throughput (INT8 tensor cores have higher TFLOPS than FP16). The paper does not separately ablate these two factors.

Latency-Critical Inference for Sparse MoE Models (Figure 7)

Headline claim: DeepSpeed-MoE achieves up to 7.3× reduction in latency compared to the PyTorch-MoE baseline, serving a 1-trillion-parameter MoE model under 25 ms on 256 GPUs, while achieving 33% of peak memory bandwidth utilization (128 TB/sec aggregate).

Figure 7 shows single-token-generation latency and per-GPU throughput for five sparse model sizes (52B to 2T parameters) on up to 256 GPUs:

• For the 52B MoE model (1.3B base + 128 experts) on 128 GPUs, DeepSpeed-MoE achieves roughly 5-6 ms latency vs. the PyTorch baseline's ~40 ms — approximately 7.3× reduction. The throughput improvement is correspondingly large (the bar chart shows DeepSpeed-MoE throughput approximately 6-7× higher).
• For the 107.7B model (2.4B base + 128 experts) on 128 GPUs, latency drops from roughly 15 ms (PyTorch) to roughly 3 ms (DeepSpeed-MoE) — approximately 5× reduction.
• For the 349B model (8B base + 128 experts) on 128 GPUs, DeepSpeed-MoE achieves approximately 2 ms vs. ~10 ms baseline — approximately 5× reduction.
• For the 1064.9B model (24B base + 128 experts) on 256 GPUs, latency is roughly 6-7 ms vs. roughly 50 ms baseline — approximately 7× reduction.
• For the 2024B (2T) model (47B base + 128 experts) on 256 GPUs, latency is roughly 15-18 ms vs. ~35 ms baseline — approximately 2× reduction.

The paper highlights that achieving this latency for a trillion-parameter model required leveraging an aggregate GPU memory bandwidth of 128 TB/sec, which they note is "33% of peak memory bandwidth." The paper characterizes this as "an unprecedented result due to the intensive communication required in such scenarios," explicitly distinguishing this from compute-bound applications where 33% utilization would be "fairly low."

Critical detail: The paper states that "the configuration for data/tensor/expert parallelism is the same for both the baseline and DeepSpeed Inference-MoE." The 7.3× speedup therefore comes exclusively from (1) expert-slicing (splitting individual experts across GPUs for aggregate memory bandwidth), (2) the PCC all-to-all optimization, and (3) the sparse kernel optimizations replacing sparse einsums with dense table-based transformations. The paper does not ablate the individual contributions of these three optimizations to the 7.3× figure, reporting only the combined effect.

Throughput-Oriented Massive Model Inference (Figure 8)

Headline claim: DeepSpeed Inference achieves 1.51× higher throughput than the best FasterTransformer configuration for GPT-3 175B on 16 GPUs, and 1.53× for MT-NLG 530B on 40 GPUs.

Figure 8 shows throughput (tokens-per-second per GPU, left y-axis) and compute efficiency (TFLOPS per GPU, right y-axis) for two configurations:

• LM-175B on 16 GPUs (2 nodes of 8× A100): The paper compares four configurations: FT with TP-only, DeepSpeed Inference with TP-only, FT with TP+PP, and DeepSpeed Inference with TP+PP. DeepSpeed Inference with TP+PP achieves roughly 17 tokens/sec/GPU and 72 TFLOPS/GPU vs. FT's TP+PP at roughly 11.5 tokens/sec/GPU and 69 TFLOPS/GPU. The paper reports "over 1.51× throughput improvement." The TP-only configurations show a smaller gap. The paper states that the improvement "comes from our improved pipeline parallelism schedule, and ability to run much larger batch sizes using memory optimization and communication minimization strategies."

• LM-530B on 40 GPUs (5 nodes of 8× A100): The paper could not run FT with TP+PP ("without crashing"), so the comparison is DeepSpeed Inference TP+PP vs. FT TP-only. DeepSpeed achieves roughly 6 tokens/sec/GPU and 58 TFLOPS/GPU vs. FT TP-only at roughly 48 TFLOPS/GPU and lower throughput. The paper reports "over 1.53× throughput improvement running on 5 nodes." This is a weaker comparison because FT TP-only is not the best FT configuration for this model — the crash prevents the full TP+PP comparison. The paper does not explain the cause of the FT crash.

Important caveat: The throughput numbers are measured at the best batch size for each configuration, meaning DeepSpeed Inference's memory optimizations (activation offloading, communication-aware PCIe scheduling) that enable larger batch sizes are factored into the comparison. The paper is not claiming 1.51× higher throughput at identical batch size — it is claiming that DeepSpeed Inference's ability to use larger batch sizes (due to memory savings) combined with its improved pipeline schedule yields 1.51× higher maximum throughput. This is a valid system-level comparison but means the gain is not purely from scheduling.

Democratizing Large Model Inference with ZeRO-Inference (Figure 9)

Headline claim: ZeRO-Inference can inference a 530B model on a single A6000 GPU (25× larger than GPU-only), achieving 84 TFLOPS (54% of peak) and offering >50% better throughput than GPU-only solutions even for models that fit in GPU memory.

Figure 9 evaluates three aspects:

Model scale (Figure 9b): The bar chart shows maximum inferable model size for three configurations: GPU-only, CPU-only, and ZeRO-Inference. GPU-only reaches approximately 20B parameters (limited by A6000's 48 GB); CPU-only reaches approximately 50B; ZeRO-Inference reaches 530B. The 25× claim is relative to GPU-only (530B vs. ~20B); the 10× claim is relative to CPU-only (530B vs. ~50B).

Inference throughput across models (Figure 9b): For models that fit in GPU memory (GPT-Neo-2.7B, GPT-J-6B, GPT-NeoX-20B), ZeRO-Inference achieves higher throughput than the GPU-only solution — "over 50% better throughput" for models that fit in single GPU memory. The paper explains this counterintuitive result: ZeRO-Inference can support much larger batch sizes by offloading parameters and using GPU memory for activations, and "the benefit of larger batch size is shown in Fig. 9(a)." For models that fit in CPU memory, ZeRO-Inference offers "over 25× higher throughput than the CPU-only solution."

The peak throughput number — 84 TFLOPS, 54% of theoretical peak (158.4 TFLOPS for A6000 as quoted, though the standard A6000 FP16 tensor-core peak is 156 TFLOPS with sparsity; the paper's number may include sparsity or a different SKU) — is measured at very large batch sizes where compute dominates weight transfer time.

Scalability (Figure 9c): For the GPT-50B model on DGX2 V100 (16 GPUs), ZeRO-Inference achieves near-linear throughput scaling: single-GPU throughput of approximately 67 TFLOPS (53% of V100 peak), scaling to roughly 16× on 16 GPUs. The paper attributes this to leveraging aggregate PCIe bandwidth across GPUs (each GPU fetches only a partition of each layer over its dedicated PCIe link, then aggregates over NVLink). The near-perfect linear scaling suggests that PCIe bandwidth is not the bottleneck when distributed across 16 GPUs, and that the NVLink aggregation effectively hides the per-GPU transfer time.

Performance Breakdown and Analysis

The paper provides several finer-grained analyses to decompose the contributions of individual optimizations.

Dense GPU kernel performance breakdown (Figure 10a): The ablation compares PyTorch baseline, PyTorch + Deep-Fusion, and PyTorch + Deep-Fusion + custom GeMM for the GPT-2 model. The baseline latency (unitless y-axis, relative scale) is highest. Adding Deep-Fusion reduces latency substantially, and adding the custom SBI-GeMM on top provides further reduction. The paper states that "deep-fusion offers a significant reduction in latency by reducing kernel launch and data movement overheads, while our custom GeMM implementation offers further reduction for small batch sizes by increasing memory bandwidth utilization of GeMM." This is the only direct ablation of the kernel optimizations in the paper.

Throughput breakdown for massive model GPU inference (Figure 10b): This shows the incremental throughput gain from adding optimizations for the 530B model: starting from the optimized dense kernel baseline, adding inference-optimized scheduling, then memory optimizations (enabling larger batch sizes), then communication optimizations. Each step increases the throughput bar. The paper does not provide exact numbers for each increment in the text, relying on the bar chart.

Prompt latency improvement with hybrid scheduling (Figure 13): For GPT-3 175B, the hybrid scheduling achieves 1.18× speedup over FasterTransformer with PP+TP configuration, and 3.06× speedup with TP-only configuration (batch size 24). The larger speedup in TP-only mode suggests that FT's TP-only scheduling is particularly suboptimal for prompt processing, while FT's PP+TP scheduling already achieves some of the benefit. The paper notes an unexplained issue: "we set the batch size to 24, because the latency dramatically increases when the batch size is larger than 24. We suspect this is related to an issue in the AllReduce kernel in Pytorch."

Memory bandwidth scalability for sparse MoE (Figure 11): For the 52B MoE model scaling from 8 to 128 GPUs, DeepSpeed-MoE achieves significantly higher per-GPU memory bandwidth than the PyTorch baseline at every scale, and the gap widens at larger GPU counts. On 8 GPUs, DeepSpeed-MoE achieves roughly 65% of peak memory bandwidth vs. baseline at roughly 40%. On 128 GPUs, DeepSpeed-MoE achieves roughly 50% vs. baseline's ~15%. The paper attributes this to "the combined effect of MoE kernels and all-to-all optimizations."

Impact of prefetching on ZeRO-Inference throughput (Figure 10c): Prefetching improves throughput at small batch sizes where compute time is small and weight transfer time dominates. At larger batch sizes, the benefit diminishes because compute time grows large enough to hide the transfer latency anyway. This is an expected result confirming the paper's model of the compute-to-transfer ratio.

Comparison with E.T. for small encoder models (Figure 12): On DistilBERT with batch size 1 and sequence length 128 on A100, DeepSpeed Inference achieves approximately 0.4 ms latency vs. E.T.'s 0.7 ms — a 1.7× speedup. On BERT-base, the numbers are approximately 0.9 ms vs. 1.25 ms — a 1.4× speedup. The paper attributes this to Deep-Fusion fusing more operators, leading to lower kernel invocation overhead and higher memory bandwidth utilization. The paper notes that while DeepSpeed Inference is faster for these small models, its scope is "much broader than E.T., where DeepSpeed Inference supports encoder, decoder, and sparsely gated MoE models at much larger scale."

Ablation Studies and Robustness Checks

Deep-Fusion vs. no fusion (Figure 10a): The GPT-2 kernel performance breakdown shows that Deep-Fusion alone (without custom GeMM) provides a substantial latency reduction over the PyTorch baseline. The custom GeMM adds a further, smaller reduction on top. This confirms that both mechanisms contribute, with Deep-Fusion providing the larger absolute gain.

Custom GeMM vs. cuBLAS (implicit in Figure 6): The small-batch latency advantage of DeepSpeed-FP16 over FT-FP16 (up to 1.55×) is largely attributable to SBI-GeMM vs. cuBLAS for the GeMM operations, since both systems use Deep-Fusion-level optimization for non-GeMM operators. The paper does not isolate this component explicitly but the gap is most pronounced at the smallest batch sizes and smallest models — exactly where SBI-GeMM's memory-bandwidth-optimized design (cooperative-group reduction, cache-line-aware layout) provides the most benefit over cuBLAS's compute-optimized design.

INT8 vs. FP16 (Figure 6): DeepSpeed-INT8 provides approximately 1.25-1.3× further latency reduction over DeepSpeed-FP16 for the models where it is evaluated (GPT-2, GPT-Neo-2.7B, GPT-J-6B, GPT-13B). Since FT does not support INT8 for decoder models at the time of comparison, this is a capability advantage, not just an optimization advantage. The paper does not evaluate INT8 accuracy or compare against any INT8 baseline (since none exists for decoders in FT).

Hybrid scheduling vs. fixed micro-batch count (Figure 13): The 1.18× (PP+TP) and 3.06× (TP-only) prompt processing speedups provide evidence that phase-aware micro-batch scheduling matters. The larger gain in TP-only mode is notable but partially explained by the AllReduce issue that forced batch size 24 — the paper acknowledges this is a limitation and leaves fixing it as future work.

Activation offloading enabling larger batch sizes (Figures 8 and 10b): The throughput improvements for GPT-3 175B and MT-NLG 530B depend on DeepSpeed Inference's ability to run larger batch sizes than FT, which comes from activation offloading and the communication-aware PCIe scheduling. The ablation in Figure 10b shows that "memory optimizations" and "communication optimizations" each contribute incremental throughput gains, but the paper does not report the exact batch sizes achieved by each configuration.

PCC all-to-all optimization (implicit in Figures 7 and 11): The MoE scaling results (Figure 11) show DeepSpeed-MoE maintaining much higher memory bandwidth utilization at large GPU counts than the baseline. Since the baseline uses the same expert and tensor parallelism degrees, the difference in scaling behavior is attributable to PCC (reducing all-to-all latency) and sparse kernel optimizations. The paper does not ablate these two components separately — we cannot attribute how much of the 7.3× speedup is from PCC vs. from sparse kernels.

Expert-slicing vs. no expert-slicing (not directly ablated): The PyTorch-MoE baseline does not use expert-slicing, so some fraction of the 7.3× MoE speedup comes from the aggregate memory bandwidth benefit of splitting individual experts across GPUs. The paper does not isolate this fraction from the PCC and kernel optimization contributions.

Prefetching depth (Figure 10c): The prefetching ablation shows the expected diminishing-returns pattern: large benefit at small batch sizes, minimal benefit at large batch sizes. The paper does not experiment with different prefetching depths (number of layers prefetched ahead) or analyze the optimal depth as a function of batch size and model layer size.

ZeRO-Inference design choice: streaming vs. caching (Figure 9b): The counterintuitive result that ZeRO-Inference outperforms GPU-only solutions even for models that fit in GPU memory is validated by the >50% throughput improvement for GPT-NeoX-20B on A6000. This is the key evidence that the design choice to offload everything and maximize batch size dominates the alternative of caching weights in GPU memory. The paper does not provide a direct ablation comparing the two strategies head-to-head for the same model at the same total memory budget, which would make the comparison more rigorous.

Multi-GPU ZeRO-Inference scaling (Figure 9c): Near-linear scaling from 1 to 16 GPUs for GPT-50B on DGX2 V100 confirms that PCIe bandwidth aggregation (each GPU fetching only a partition of each layer) effectively removes the PCIe bottleneck when enough GPUs contribute. It would be informative to see at what GPU count the scaling deviates from linear — i.e., where PCIe or NVLink bandwidth becomes a bottleneck — but the paper does not explore this saturation point.

E.T. comparison (Figure 12): The 1.7× (DistilBERT) and 1.4× (BERT) speedups over E.T. on small encoder models demonstrate that the kernel optimizations apply beyond the large-model regime the paper primarily targets. However, this comparison is for batch size 1 and sequence length 128 only — the paper does not evaluate a range of batch sizes or sequence lengths against E.T.

Critical Assessment

Claim 1: DeepSpeed Inference reduces latency by up to 1.9× over state-of-the-art for dense models.

This claim is supported by the data in Figure 6, but with several qualifications that narrow its scope:

• The 1.9× figure refers specifically to DeepSpeed-INT8 vs. FasterTransformer-FP16, not DeepSpeed-FP16 vs. FT-FP16. The FP16-to-FP16 comparison shows up to 1.55× (small batch) and 1.57× (larger batch). The INT8 comparison is against an FP16 baseline because FT did not support INT8 for decoder models — this means the 1.9× includes both software optimization and a datatype advantage. An INT8-to-INT8 comparison cannot be made, and the paper does not compare against any other INT8-capable inference system for decoders.

• The comparison uses identical tensor parallelism configurations. This isolates kernel performance but also means the results only apply to the TP degrees tested. For models that would benefit from different TP strategies, the relative advantage might differ.

• The evaluation covers eight dense model configurations but all are GPT-style decoders. There is no evidence that the same speedups would apply to encoder-only models (beyond the E.T. comparison for small BERT/DistilBERT), encoder-decoder models, or non-language transformers (vision, speech).

• The 1.55× speedup at small batch sizes degrades at the largest models (LM-175B shows a smaller relative gap in Figure 6 than GPT-2). The paper does not discuss whether this trend continues to even larger models, which is relevant since DeepSpeed Inference specifically targets massive-scale inference.

Claim 2: DeepSpeed Inference reduces latency by up to 7.3× for a trillion-parameter sparse MoE model.

The 7.3× figure appears to apply to the 52B MoE model (1.3B base + 128 experts) on 128 GPUs, not the trillion-parameter model. For the 2T model on 256 GPUs, Figure 7 suggests roughly a 2× reduction. The paper's abstract states "7.2× for sparse models (a 1T model under 25 ms)" — this conflates two separate results: the 7.2× speedup (likely from the 52B model at a specific configuration) and the <25 ms latency for the 1T model. The per-token latency for the 1T model under 25 ms is a legitimate achievement, but attributing a 7.3× speedup to that model size is misleading given Figure 7's data.

More fundamentally, the baseline is a PyTorch implementation that "does not" include expert-slicing, PCC, or optimized MoE kernels. This is a reasonable baseline for demonstrating the impact of these optimizations, but it is not the strongest possible baseline — a version of PyTorch-MoE with expert-slicing (but not PCC or kernel optimizations) would allow attributing the gain to specific components. The 7.3× number bundles three distinct optimizations (expert-slicing for higher memory bandwidth, PCC for reduced communication, and sparse kernel optimizations for reduced MoE computation overhead) into a single speedup figure without an ablation that isolates their contributions.

Additionally, the memory bandwidth scalability results (Figure 11) show DeepSpeed-MoE maintaining ~50% peak bandwidth at 128 GPUs vs. baseline ~15%. This difference accounts for roughly a 3.3× improvement in memory bandwidth utilization. The remaining factor of ~2.2× to reach 7.3× must come from other sources (PCC reducing communication time, sparse kernels reducing MoE computation time). Without separate ablations, we cannot know the relative importance of each optimization, which limits the paper's guidance for practitioners prioritizing which optimizations to implement first.

Claim 3: ZeRO-Inference enables inference of models 25× larger than GPU-only solutions while achieving >50% of peak hardware performance.

The "25× larger" claim is measured against the largest model that fits in A6000 GPU memory (~20B for INT8, given 48 GB capacity). This is a hardware limit, not a comparison against another system, so the claim is true by construction: if you can offload to CPU/NVMe, you can run larger models than if you cannot. The more meaningful claims are the throughput claims.

The headline 84 TFLOPS (54% of peak) is measured at very large batch sizes where compute dominates. The paper does not report the batch size at which this throughput was achieved, the sequence length, or the model used. This makes it impossible to assess whether this throughput is achievable for practical inference workloads (e.g., generating 50 tokens with a 128-token prompt) or only for the specific "maximum batch size for generating a single token" benchmark described in the evaluation methodology. The throughput may degrade substantially for workloads with longer generation lengths where the KV-cache consumes more memory, reducing the feasible batch size.

The >50% throughput advantage over GPU-only solutions for models that fit in GPU memory is the most provocative result. However, the paper does not provide a detailed comparison: what batch sizes does each configuration achieve? What is the GPU memory breakdown (weights vs. activations vs. KV-cache) for each? Without this detail, the result is suggestive rather than conclusive. A skeptical reader could wonder whether the GPU-only configuration was properly tuned (e.g., whether it used the maximum possible batch size for its memory budget) or whether the 50% gain partially reflects suboptimal GPU-only configuration rather than a fundamental advantage of offloading.

Claim 4: DeepSpeed Inference increases throughput by over 1.5× for throughput-oriented scenarios.

For GPT-3 175B on 16 GPUs, the 1.51× throughput improvement over FT is measured at each system's best batch size. This conflates two sources of gain: (1) DeepSpeed Inference's improved pipeline schedule (hybrid scheduling, dynamic micro-batch queuing) and (2) DeepSpeed Inference's ability to use larger batch sizes (due to activation offloading and memory optimizations). The paper's ablation (Figure 10b) shows that "memory optimizations" and "communication optimizations" each contribute nonzero throughput gains, but does not quantify what fraction of the 1.51× comes from each. A comparison at identical batch sizes would isolate the scheduling improvement, but this is not provided.

For MT-NLG 530B, the comparison is weakened by FT's inability to run with TP+PP ("without crashing"). The 1.53× figure is against FT TP-only, which is not the strongest FT configuration. The paper does not explain the crash, leaving open whether it reflects a fundamental limitation of FT or a configuration issue that might be resolved. A fairer comparison would require running FT at its best working configuration, even if that configuration differs from DeepSpeed Inference's.

Claim 5: DeepSpeed Inference provides state-of-the-art latency and throughput for all variations of transformer model inference.

This is the paper's broadest claim and the one least supported by the evidence. The evaluation covers:

• Dense models: GPT-style decoders from 1.5B to 530B parameters. No encoder-only models evaluated at scale (the E.T. comparison is for small BERT models only). No encoder-decoder models. No non-language transformers.
• Sparse models: MoE models with 128 experts and base sizes from 1.3B to 47B. No evaluation with different numbers of experts (e.g., 8, 64, 256) or different top-k gating values. The specific MoE architecture is fixed; results may not transfer to other MoE designs (e.g., Switch Transformer with k=1 vs. GShard with k=2).
• Batch sizes: Primary latency evaluation uses batch sizes 1, 8, 16 for dense models and batch size 8 for sparse models. No evaluation at batch sizes 2, 4, 32, or 64, which would show the transition from memory-bandwidth-bound to compute-bound regimes.
• Sequence lengths: Fixed at 128-token prompt, 8-token generation for latency; 512-token prompt, 50-token generation for throughput. No sensitivity analysis to sequence length, which matters because the KV-cache memory pressure scales with sequence length and could significantly affect the benefits of activation offloading.

Missing Experiments That Would Strengthen the Paper

1. Ablation of MoE optimizations separately: The PCC communication optimization and the sparse kernel optimizations are never ablated independently. Reporting how much of the 7.3× MoE speedup comes from each would guide practitioners on where to invest implementation effort.

2. Comparison against model compression baselines: The paper acknowledges that quantization, distillation, and sparsification are complementary, but does not combine DeepSpeed Inference with any of these techniques to show the combined effect. A comparison showing DeepSpeed-INT8 against a distilled model of equivalent accuracy, or against a sparsified model, would contextualize the pure-system-optimization approach.

3. Latency-throughput tradeoff curves: The paper reports latency at specific batch sizes and throughput at maximum batch size, but does not show the latency-vs-throughput Pareto frontier. For practitioners deciding between DeepSpeed Inference and alternatives, this curve is more informative than point comparisons.

4. Accuracy evaluation for INT8: The INT8 speedup numbers are impressive, but no accuracy comparison is provided. For decoder models, quantization can introduce token-level errors that compound autoregressively, potentially degrading output quality even if per-layer quantization error is small. Without accuracy evaluation, the INT8 gains are a performance result, not an end-to-end deployment result.

5. Breakdown of speedup sources at identical batch sizes: For the multi-GPU throughput comparisons (Figure 8), showing the throughput improvement at identical batch size would isolate the scheduling improvement. The memory-optimization benefit (larger batch sizes) should be reported separately as a capability gain, not bundled into the throughput speedup.

6. ZeRO-Inference sensitivity to sequence length: The 84 TFLOPS result is at "maximum batch size for generating a single token." How does throughput degrade when generating 50 tokens (as in the throughput benchmarks) or 100 tokens (as in the MoE latency benchmarks)? The KV-cache for long generations may consume enough memory to reduce batch size, which would reduce the compute-to-transfer ratio and potentially eliminate the throughput advantage over GPU-only solutions.

7. Comparison against CPU-only inference at equivalent model scales: The paper mentions that ZeRO-Inference offers "over 25× higher throughput than the CPU-only solution" for models that fit in CPU memory, but this result is not shown quantitatively in Figure 9 or elsewhere. CPU inference with optimized libraries (e.g., ONNX Runtime with Intel oneDNN) can achieve non-trivial throughput for large-batch inference — a head-to-head comparison would better contextualize the GPU+CPU heterogeneous approach.

8. Startup cost of ZeRO-Inference: Streaming layers from NVMe on-demand means the first token latency is much higher than subsequent tokens because no weights are pre-loaded. The paper does not report first-token latency vs. steady-state token latency for ZeRO-Inference, which is critical for interactive applications even if they are "less latency sensitive."

Conditional Boundaries of the Claims

The paper's results establish that DeepSpeed Inference is effective on GPT-style decoder models on NVIDIA A100/V100/A6000 GPUs for the specific workloads tested (128-token prompts, 8-100 token generation, batch sizes 1-16 for latency). Extrapolation to other model architectures, hardware platforms, or workload profiles is not supported by the presented data. Specifically:

• The benefits of Deep-Fusion and SBI-GeMM depend on the specific operator mix in the transformer layer. Models with different attention mechanisms (e.g., sparse attention, linear attention, multi-query attention) or different normalization schemes may not benefit equally from the described fusion schedules.
• The inference-optimized pipeline schedule exploits the compute-bound vs. memory-bandwidth-bound phase distinction in autoregressive generation. For encoder-only models (which process all tokens simultaneously with no generation phase) or for models where prompt processing and token generation are not clearly separated, the hybrid scheduling provides no benefit.
• The PCC all-to-all optimization depends on the specific interaction between tensor parallelism and expert parallelism. Models using different parallelism combinations (e.g., only expert parallelism without tensor slicing) cannot use PCC.
• ZeRO-Inference's throughput advantage depends on achieving sufficiently large batch sizes that compute time dominates PCIe transfer time. For models with smaller layers (less compute per weight byte), for deployments with slower NVMe drives, or for workloads requiring very long generation (where KV-cache limits batch size), the advantage may shrink or disappear.
• All results are on NVIDIA hardware with specific interconnect topologies (NVLink within node, InfiniBand across nodes, PCIe Gen4 for A6000). Performance on other GPU architectures (AMD, Intel) or with different interconnect bandwidth ratios is not evaluated.

6. Limitations and Trade-offs

Limitation 1: No Accuracy Evaluation for INT8 Inference on Decoder Models

The assumption or constraint. The paper presents INT8 inference as a key capability for latency reduction, achieving up to 1.95× speedup over the FP16 baseline for decoder models (Figure 6). However, the paper provides no evaluation of whether INT8 quantization degrades the model's output quality. This is a significant omission because autoregressive decoder models compound quantization errors across generated tokens — a small per-layer quantization error can accumulate over dozens of generation steps, potentially producing incoherent or factually incorrect outputs. The paper does not compare INT8 outputs against FP16 reference outputs for any model, any task, or any generation length, nor does it report any perplexity, accuracy, or human evaluation metric.

The consequence. A practitioner cannot determine whether the 1.95× latency reduction comes at an acceptable accuracy cost. The speedup figures are a pure performance result, not an end-to-end deployment characterization. If INT8 inference produces lower-quality outputs, the latency gain may be irrelevant for applications where output quality matters. The paper's comparison of DeepSpeed-INT8 against FasterTransformer-FP16 (which lacks INT8 support for decoders) further obscures this tradeoff: a portion of the 1.95× gain comes from the INT8 datatype advantage rather than software optimization, but we cannot assess whether that advantage is "real" if INT8 outputs are degraded.

What evidence exists in the paper. None. The INT8 evaluation in Figure 6 reports only latency and throughput. There is no mention of accuracy, perplexity, BLEU score, human evaluation, or any quality metric anywhere in the paper for INT8 decoder models. The paper acknowledges in Section I that model compression techniques like quantization involve "a small accuracy trade-off," but does not quantify this trade-off for its own INT8 implementation. The absence is complete — not even a single example of INT8 vs. FP16 output text is provided.

Mitigation status. The paper does not attempt to measure or mitigate the accuracy impact of INT8 quantization. It does not discuss calibration methodology (how the quantization scales were determined), per-tensor vs. per-channel quantization granularity, or whether the INT8 path uses post-training quantization or quantization-aware training. Future work would need to establish the accuracy-efficiency Pareto frontier for INT8 decoder inference across model scales, generation lengths, and tasks. Without this, the INT8 results are a lower bound on latency with an unknown upper bound on quality degradation.

---

Limitation 2: Difficulty Estimation Cost for Hetergeneous Memory Offloading (ZeRO-Inference Latency Characterization)

The assumption or constraint. The paper's ZeRO-Inference throughput claims (84 TFLOPS, 54% of peak) are measured at "maximum batch size for generating a single token" (Section VII-A3), and the throughput numbers in Figure 9 are presented as steady-state throughput. However, autoregressive inference with model streaming has a fundamentally different cost structure for the first token versus subsequent tokens. When weights are streamed from CPU/NVMe on demand, the first token must incur the full latency of loading every layer from cold storage. Once layers are streamed through, subsequent tokens may benefit from prefetching and overlap, but the initial cold-start penalty is substantial and is nowhere quantified in the paper.

The consequence. For any interactive application — even "throughput-oriented" ones that the paper positions as ZeRO-Inference's target — the latency of generating the first token of a request matters. Users experience the time-to-first-token, not the steady-state tokens-per-second. If ZeRO-Inference requires loading 96 layers from NVMe to produce the first token (at, say, 7 GB/s NVMe read bandwidth, ~1.8 GB per layer for GPT-3 175B, roughly 96 × 1.8 / 7 ≈ 25 seconds just for weight transfer), the first-token latency may be orders of magnitude worse than a GPU-only solution even if steady-state throughput is higher. The paper's claim that ZeRO-Inference is suitable for "applications that are less latency sensitive but resource constrained" (Section I) acknowledges this tension but provides no data to characterize the actual latency profile.

What evidence exists in the paper. The evaluation methodology (Section VII-A3) states that throughput is measured "using maximum batch size possible for generating a single token." The paper does not report first-token latency, cold-start overhead, or the latency-throughput tradeoff as a function of generation length. The prefetching ablation (Figure 10c) shows throughput at different batch sizes but does not measure latency. The throughput results in Figure 9 are aggregate numbers that do not distinguish between cold-start and steady-state performance.

Mitigation status. The paper partially acknowledges this by positioning ZeRO-Inference for "non-latency sensitive applications" (Section VI introduction), but does not define what latency sensitivity threshold makes it appropriate. The prefetching optimization (Section VI-B) attempts to hide transfer latency by overlapping with computation, but this only helps after the first few layers are loaded — the initial layers must still be loaded from cold storage before any computation can begin. The paper suggests future work on "pretraining or finetuning models to directly predict difficulty" for the search-based inference problem in the example paper; analogously, a latency model for ZeRO-Inference that characterizes first-token and steady-state latency as a function of model size, batch size, sequence length, and storage bandwidth would be needed to make informed deployment decisions. The paper provides none of this.

---

Limitation 3: Single Hardware Architecture and Interconnect Topology

The assumption or constraint. All experiments in the paper are conducted on NVIDIA GPUs with specific interconnect topologies: A100 40GB with NVLink and NVSwitch within DGX nodes, V100 32GB with NVLink within DGX-2, and A6000 48GB with PCIe Gen4. The parallelism strategies, communication optimizations, and memory bandwidth utilization claims depend intimately on the ratios of these interconnects' bandwidths to GPU HBM bandwidth, to PCIe bandwidth (for CPU/NVMe offloading), and to each GPU's compute throughput. The paper never evaluates on alternative hardware (AMD GPUs, Intel GPUs, non-NVLink interconnects, PCIe Gen3, different NVMe drive classes) and does not provide a performance model that would allow practitioners to predict performance on their own hardware.

The consequence. A practitioner with non-NVIDIA hardware, or with older NVIDIA hardware (e.g., V100 with PCIe-only interconnects, A100 with NVLink but not NVSwitch), or with different PCIe topology (e.g., single PCIe switch vs. dual-root), cannot determine whether DeepSpeed Inference's optimizations will provide similar benefits. Several optimizations are specifically designed for the evaluated topology: the communication-aware PCIe offload schedule ("odd-numbered GPUs offload activations for odd-numbered layers" — Section IV-C3) assumes a specific PCIe sharing pattern (two GPUs sharing a link) that may not hold on all systems. The PCC all-to-all optimization's reduction from O(p) to O(p/L) depends on the ratio of intra-node (NVLink) to inter-node (InfiniBand) bandwidth; on systems with different ratios, the optimal tensor-slicing degree L may differ, and the paper provides no guidance for choosing it.

What evidence exists in the paper. The hardware testbeds (Section VII-A4) describe three specific configurations: 256 A100 GPUs in DGX boxes, an A6000 workstation, and a DGX2 V100 server. No sensitivity analysis across hardware parameters is performed. The paper does not report effective bandwidth achieved for NVLink, InfiniBand, PCIe, or NVMe transfers — only aggregate throughput and latency. The paper notes that "modern clusters have heterogeneous network topology (eg. intra-node NVLink/NVSwitch and inter-node InfiniBand) which requires further consideration when developing parallelism strategies" (Section I), acknowledging the hardware-dependence, but evaluates only one topology.

Mitigation status. The paper does not attempt to provide a hardware-agnostic performance model. The optimizations are implemented for NVIDIA GPUs using CUDA, NCCL, and CUDA-specific features (CUTLASS, CUDA Graphs, cooperative groups). Porting to other hardware would require re-implementing the custom GeMM, Deep-Fusion tiling strategies, and communication collectives for the new platform's programming model and interconnect topology. The paper is transparent that this is a CUDA/NVIDIA system, but the claims of generality ("state-of-art latency and throughput at unprecedented scales for both dense and sparse transformer models") implicitly assume the hardware characteristics of the evaluated platforms. A deployment guide mapping model size, batch size, and hardware configuration to expected performance would make the results more actionable for practitioners with different hardware.

---

Limitation 4: Untested Generalization to Non-GPT, Non-Language Transformer Architectures

The assumption or constraint. The paper evaluates exclusively on GPT-style autoregressive decoder models for language (GPT-2, GPT-Neo, GPT-J, GPT-NeoX, LM-175B, LM-530B variants) and their MoE extensions. The kernel fusion schedules (Figure 1c) are specific to the standard transformer decoder architecture: self-attention followed by feed-forward, with layernorm, residual connections, and the specific QKV/attention/FNN operator sequence. The inference-optimized pipeline parallelism depends on the autoregressive token generation pattern with KV-caching that is characteristic of decoder-only models. The sparse MoE evaluation uses a specific architecture (128 experts, top-k gating, standard transformer blocks with MoE FF layers).

The consequence. The paper's claimed scope — "all variations of transformer model inference: dense or sparse, small or large batches, billions to trillions of parameters" (Section I) — is not supported for transformer architectures that deviate from the evaluated pattern. Encoder-decoder models (T5, BART) have a different computational structure (encoder processes all input tokens in parallel, decoder is autoregressive but with cross-attention to encoder outputs), which changes both the operator mix (cross-attention is not fused in the described schedule) and the memory access pattern (encoder KV-cache does not grow with generation length). Vision transformers (ViT, DeiT) process fixed-size image patches and lack autoregressive generation entirely, making the pipeline parallelism and KV-cache management optimizations irrelevant. Models with sparse attention patterns (BigBird, Longformer) or alternative positional encodings (relative position bias, rotary embeddings) may not fuse cleanly with the described tiling strategy. Even within language models, encoder-only architectures (BERT, RoBERTa) lack the prompt-processing/token-generation phase distinction that motivates hybrid scheduling.

What evidence exists in the paper. The only non-decoder evaluation is the small-scale comparison against E.T. for DistilBERT and BERT (Figure 12), which shows 1.4×–1.7× lower latency for these encoder models at batch size 1 and sequence length 128. This is a kernel-level comparison only (single GPU, small model) and does not evaluate any of the multi-GPU parallelism, pipeline scheduling, or MoE optimizations that constitute the bulk of the paper's contribution. The paper explicitly notes that "the scope of our work is also much broader than E.T." — but the broader scope is only demonstrated for decoder models.

Mitigation status. The paper does not claim evaluation on non-decoder architectures, but it does claim generality in its conclusions ("versatile solution capable of achieving state-of-art latency and throughput for all variations of transformer model inference"). The Deep-Fusion tiling strategy is described in general terms (tiling along dimensions with no cross-tile dependencies), suggesting it could be extended to other architectures, but the specific fusion schedule (Figure 1c) would need to be redesigned for each new architecture's operator graph. The inference-optimized pipeline parallelism is so tightly coupled to autoregressive generation that it provides no benefit for non-autoregressive models. The paper would benefit from explicitly scoping its claims to autoregressive decoder architectures and acknowledging that encoder and encoder-decoder models require separate validation.

---

Limitation 5: Unquantified First-Token Latency Overhead from CUDA Graph Capture

The assumption or constraint. The paper introduces CUDA Graph support to eliminate CPU-side kernel launch overhead (Section III-D). The mechanism records the kernel trace during the first forward pass and replays it for subsequent requests. This means the first request (or first request after any model or batch size change) incurs the full cost of kernel launches plus the graph capture overhead, while subsequent requests benefit from replay. The paper claims this "largely eliminates the kernel launching overhead and substantially improves the performance" but never measures the graph capture overhead or the first-request latency penalty.

The consequence. For serving systems with continuous traffic, the amortized overhead may be negligible. But in several common deployment scenarios, the first-request penalty matters: (1) cold-start in serverless deployments where the model is loaded on-demand per request; (2) dynamic batching systems that aggregate different numbers of requests into variable batch sizes, requiring graph recapture for each batch size; (3) A/B testing or model update scenarios where the model changes frequently. In these cases, the CUDA Graph capture overhead may either negate the launch overhead savings or force the system to use a fixed batch size to avoid recapture, sacrificing the throughput benefits of dynamic batching. The paper provides no guidance on whether DeepSpeed Inference supports dynamic batch sizes without graph recapture, or what the capture overhead is in absolute terms.

What evidence exists in the paper. None. Section III-D mentions CUDA Graph support without any latency measurement, ablation, or discussion of the capture cost. The evaluation in Figure 6 does not distinguish between first-request and steady-state latency. The paper does not describe the batching interface — whether requests are processed individually with fixed-size CUDA Graphs, or whether a graph is captured per batch size and reused.

Mitigation status. The paper does not address this limitation. CUDA Graphs are presented as a pure optimization with no discussion of their tradeoffs. Since the paper positions DeepSpeed Inference for both latency-critical interactive applications (where cold-start matters) and throughput-oriented applications (where dynamic batching matters), the uncharacterized CUDA Graph overhead is a practical gap. Future work could measure capture overhead as a function of model size and batch size, explore whether CUDA Graph "update" (rather than recapture) can handle variable batch sizes, or implement a hybrid approach where small-batch requests use graph replay and large-batch requests use direct kernel launch if capture overhead dominates.

---

Limitation 6: The Paper Does Not Ablate the Relative Contribution of MoE Optimizations, Making Investment Prioritization Impossible

The assumption or constraint. The sparse MoE evaluation reports a single combined speedup of up to 7.3× over the PyTorch baseline (Figure 7), which bundles three independent optimizations: expert-slicing (splitting individual experts across GPUs for aggregate memory bandwidth), PCC (parallelism-coordinated communication using tensor-parallel data replication), and sparse kernel optimizations (replacing sparse einsums with dense table-based transformations). These optimizations address different bottlenecks — memory bandwidth, communication latency, and computation overhead, respectively — and require different implementation investments. The paper never measures each optimization's individual contribution.

The consequence. A practitioner deciding whether to adopt DeepSpeed-MoE cannot determine which optimizations are most impactful for their specific bottleneck. If their model is communication-bound (many GPUs, low inter-node bandwidth), PCC may provide most of the gain and the sparse kernel optimizations may be marginal. If their model is compute-bound (small number of experts, large batch sizes), the sparse kernel optimizations may dominate and PCC may be unnecessary. The paper's single speedup figure forces an all-or-nothing adoption decision and provides no guidance for incremental implementation — precisely the kind of guidance that systems papers typically aim to provide through careful ablation.

What evidence exists in the paper. The memory bandwidth scalability graph (Figure 11) shows DeepSpeed-MoE maintaining higher per-GPU bandwidth than baseline at scale, which is attributable to expert-slicing (higher aggregate bandwidth) plus PCC (less communication overhead eating into bandwidth utilization). But this graph does not separate the two effects. The paper claims "over 6× reduction in MoE kernel-related latency" from the sparse kernel optimizations (Section V-C), but this 6× figure is for the MoE kernel component in isolation, not for end-to-end inference latency. Without an end-to-end ablation, we cannot determine how much of the 7.3× total speedup comes from each source — the 6× kernel speedup may translate to a much smaller end-to-end gain if kernel execution is not the dominant latency component.

Mitigation status. The paper does not provide any ablation of the three MoE optimizations. The evaluation methodology states that "the configuration for data/tensor/expert parallelism is the same for both the baseline and DeepSpeed Inference-MoE" and that "the main differences are optimizations that DeepSpeed Inference has, such as expert-slicing, parallelism coordinated all-to-all and MoE-specific kernels, but the PyTorch-MoE baseline does not." This bundles all differences into a single comparison. A standard systems evaluation would include incremental results: (1) baseline, (2) baseline + expert-slicing, (3) baseline + expert-slicing + sparse kernels, (4) baseline + expert-slicing + sparse kernels + PCC (the full system). The paper's failure to provide this decomposition is the most significant methodological weakness in the evaluation, because it prevents the very kind of bottleneck analysis that the paper's own diagnostic framework (separating memory bandwidth, communication, and computation bottlenecks) would recommend.

This limitation is particularly consequential given the paper's claim to provide a "comprehensive system solution" where each layer "addresses a unique aspect of the latency challenge." If the layers are truly independent, their contributions should be independently quantifiable. Without this quantification, the paper's architectural argument — that a layered approach is necessary and each layer is valuable — is asserted rather than demonstrated for the MoE system.

7. Implications and Future Directions

How This Work Changes the Landscape

This paper does not introduce a new architectural innovation for transformer models — it does not propose a novel attention mechanism, a new model family, or a fundamentally different training objective. Instead, it makes a systems-level argument that reshapes how the field should think about the relationship between model training and model deployment. The conceptual shift is: inference is not a simpler version of training's forward pass — it is a distinct computational regime with its own optimization target, memory access patterns, and scheduling constraints, and treating it as an afterthought to training infrastructure leaves enormous performance on the table.

This reframing matters because it challenges the prevailing workflow in large-scale NLP research, where training systems (Megatron-LM, DeepSpeed training, PyTorch distributed) are developed and optimized first, and inference is often handled by those same systems running in "eval mode" or by repurposing training-optimized libraries (cuBLAS, CUTLASS) without restructuring for the inference workload. The paper's layered architecture demonstrates that a ground-up inference design — where kernels are rebuilt for memory bandwidth rather than compute throughput, pipeline schedules are redesigned for autoregressive dependencies rather than independent micro-batches, and memory management is reconfigured for KV-cache access patterns rather than activation checkpointing — yields 1.5–7.3× improvements over systems that reuse training infrastructure.

The magnitude of this shift is a reframing, not a paradigm shift. The paper does not introduce a new conceptual framework for understanding computation (the way, say, the Chinchilla scaling laws reframed the pretraining compute allocation problem). Rather, it systematically demonstrates that a principle many researchers would have agreed with in the abstract — "inference has different requirements than training" — has been underappreciated in practice, and that acting on this principle yields engineering gains large enough to change deployment feasibility. The paper's influence is likely to be on what systems researchers optimize for rather than on what model architects build.

Reconciling prior contradictions. Several optimizations in the paper resolve tensions that were implicit in the literature but never explicitly articulated:

• Why do training-optimized parallelism strategies underperform for inference? Prior work on tensor, pipeline, and expert parallelism (Megatron-LM, GPipe, GShard) reported excellent training throughput but was sometimes disappointing when applied to inference at small batch sizes. The paper provides a clear diagnosis: training parallelism is designed to maximize compute utilization at large batch sizes where arithmetic intensity is high, but inference at small batch sizes operates in a memory-bandwidth-bound regime where these strategies introduce communication overhead that cannot be hidden behind computation. The inference-optimized pipeline schedule with hybrid micro-batching, the PCC all-to-all optimization that exploits tensor-parallel data replication, and the expert-slicing for aggregate memory bandwidth are all responses to this diagnosis — they restructure the parallelism to target memory bandwidth and communication latency rather than compute throughput.

• Why do existing inference libraries (FasterTransformer, TurboTransformers) still undershoot hardware limits? The 1.55–1.95× improvement over FasterTransformer at small batch sizes (Figure 6) suggests that even dedicated inference libraries — not just training repurposed for inference — were not fully exploiting the regime where memory bandwidth is the bottleneck. The paper attributes this to two gaps: (a) operator fusion was limited to element-wise operations, missing the reductions and GeMMs that dominate transformer layers, and (b) GeMM libraries (cuBLAS) optimized for compute throughput systematically underutilize memory bandwidth at the extremely small and skinny matrix dimensions of single-batch inference. Deep-Fusion and SBI-GeMM directly address these two gaps.

• Can massive models be served on resource-constrained hardware without sacrificing throughput? The conventional wisdom was that if a model does not fit in GPU memory, inference becomes impractically slow due to PCIe transfer overhead. ZeRO-Inference's counterintuitive finding — that streaming weights from CPU/NVMe and using freed GPU memory for larger batch sizes can actually outperform GPU-only inference (Figure 9b, >50% throughput improvement for models that fit in GPU memory) — rejects this conventional wisdom with compelling empirical evidence. The key insight is that for large transformer layers at sufficient batch sizes, compute time dominates transfer time, making the transfer effectively free from a throughput perspective.

What research directions become more attractive. The paper makes several previously speculative or niche research directions more concrete and actionable:

• Inference-specific compiler optimization becomes more attractive because Deep-Fusion demonstrates a manual version of what an inference-aware compiler could automate: tiling strategies that minimize global memory traffic at the expense of redundant computation, fusion schedules that adapt to arithmetic intensity (aggressive fusion at low intensity, selective fusion at high intensity), and data layout transformations (weight transposition for cache-line utilization) that a compiler could apply automatically if given the right cost model.

• Heterogeneous memory systems for inference are validated as a practical deployment strategy, not just a fallback. ZeRO-Inference's 84 TFLOPS on a single A6000 (54% of peak) demonstrates that streaming weights from NVMe is not just a "make it work" hack but a potentially optimal strategy for throughput-oriented serving of large models. This opens up research into more sophisticated memory hierarchy management for inference — e.g., learned prefetching policies, dynamic layer caching based on access patterns, or tiered storage where hot experts or layers are promoted to GPU memory.

• Communication-computation co-design for distributed inference is shown to have substantial headroom. The PCC all-to-all optimization demonstrates that the interaction between different parallelism strategies creates redundancy (tensor-parallel replication) that can be exploited to reduce communication in other strategies (expert-parallel all-to-all). This is a specific instance of a broader principle: in systems that combine multiple parallelism dimensions, the communication requirements are not independent, and joint optimization can reduce total communication below the sum of the individual costs.

What research directions become less attractive. Conversely, the paper casts doubt on several approaches:

• Simply scaling training infrastructure to handle inference becomes harder to justify. The 1.5× throughput gap between DeepSpeed Inference and FasterTransformer at the same parallelism configuration (Figure 8) quantifies the cost of treating inference as a variant of training — and this cost grows with model size and the complexity of the parallelism strategy. Organizations that have invested heavily in training infrastructure may be tempted to reuse it for inference; this paper's results suggest that dedicated inference optimization pays for itself quickly in reduced hardware requirements.

• Developing ever-more-sophisticated sparse matrix libraries for MoE routing becomes less urgent. The paper's sparse kernel optimizations achieve a 6× reduction in MoE kernel latency by replacing sparse einsums — the "obvious" mathematical representation of expert routing — with dense table-based data layout transformations. This suggests that for structured sparsity patterns like top-k expert selection (where exactly 1-out-of-E experts is nonzero per token), general sparse matrix operations are a poor fit and dedicated routing implementations are both simpler and faster. Research effort is better spent on optimizing the routing logic for specific sparsity patterns than on improving general sparse matrix libraries for this use case.

• Pinning model weights in GPU memory as the default for heterogeneous inference is shown to be potentially suboptimal. ZeRO-Inference deliberately offloads everything and uses GPU memory for activations — the opposite of the intuitive "cache what you can" strategy — and achieves better throughput. This suggests that future work on CPU-GPU heterogeneous inference should start from the compute-to-transfer ratio optimization rather than from data-transfer minimization, which is a different analytical framework.

---

Follow-Up Research This Work Enables

1. Automated inference-specific kernel fusion through learned cost models. Deep-Fusion's fusion schedule (Figure 1c) is manually crafted for the standard GPT-style transformer layer. The decision to replicate layernorm computation across SMs to enable GeMM fusion (Section III-B) is an explicit tradeoff — more FLOPs for less memory traffic — that a compiler optimizing for total FLOP count would reject. A natural follow-up is to build an inference-aware compiler that searches over fusion schedules using a cost model that weights memory bandwidth utilization above FLOP count when arithmetic intensity is below a threshold. The concrete experiment: take a transformer architecture not evaluated in the paper (e.g., T5 encoder-decoder, ViT, or a multimodal transformer like CLIP), use the compiler to automatically derive a fusion schedule, and compare against both a manual Deep-Fusion-style schedule and a cuBLAS/CUTLASS baseline. The key metric is whether the compiler can match or exceed the manual schedule's memory bandwidth utilization at batch size 1 across architectures, without requiring per-architecture hand-tuning. The paper's explicit fusion condition — "two operators can be fused using Deep-Fusion if each tile of the second operator depends on exactly one output tile of the first operator" — provides a formal constraint that an automated system could use as a search primitive.

2. Quantifying the accuracy-efficiency Pareto frontier for INT8 autoregressive decoder inference at scale. The paper reports INT8 latency improvements of up to 1.95× over FP16 (Figure 6) with no accuracy evaluation. Autoregressive decoders compound quantization errors across generated tokens — a small per-layer error can produce divergent outputs after dozens of steps. The follow-up experiment: for a range of model sizes (GPT-2 1.5B, GPT-J 6B, GPT-NeoX 20B, and ideally a model in the 100B+ range evaluable via ZeRO-Inference on limited hardware), measure both FP16 and INT8 inference on standard generation benchmarks (e.g., WikiText perplexity, LAMBADA accuracy, or a few-shot task suite like SuperGLUE) across generation lengths from 8 to 512 tokens. Plot the per-token quality degradation as a function of generation length to determine whether there is a threshold beyond which INT8 outputs become unreliable. The paper's claim that INT8 provides a latency reduction "by reducing the overall size of the parameters in half compared to FP16" (Section VII-B1) is a performance claim only; establishing the accuracy boundary would convert this into a deployable result. Strong negative findings (e.g., INT8 decoders diverge after ~50 tokens for models above a certain size) would be equally valuable, as they would guide practitioners toward FP16 for long-generation tasks and INT8 only for short-generation or classification tasks.

3. Joint optimization of model architecture and inference scheduling for deployment efficiency. The paper optimizes inference for fixed model architectures (GPT-style transformers). But the choice of architecture — layer count, hidden dimension, FF intermediate size, number of attention heads, number of experts — determines the arithmetic intensity, memory bandwidth demand, and communication patterns that the inference system must handle. A co-design experiment: given a fixed inference budget (latency SLA, hardware configuration, batch size distribution), search over both model architecture hyperparameters and DeepSpeed Inference configuration (tensor/pipeline/expert parallelism degrees, batch size, INT8 vs. FP16) to find the jointly optimal point. The concrete output would be a set of "inference-efficient" model architectures that achieve target accuracy with lower deployment cost than standard architectures optimized only for training efficiency. The paper's performance data across model scales (Table I) provides the raw material for building a performance model of DeepSpeed Inference as a function of architecture parameters, which could then be combined with scaling law estimates of model quality to perform the co-optimization. This would extend the paper's framework from "given a model, how to inference it efficiently" to "given a deployment target, what model should you train."

4. Characterizing the cold-start and batching dynamics of ZeRO-Inference for practical serving. The paper reports ZeRO-Inference throughput at steady state with maximum batch sizes (Figure 9), but real serving systems experience request arrival patterns that include cold starts (model not yet loaded), variable batch sizes (requests arrive over time), and generation length variability (some requests produce short outputs, others long). A latency characterization workload: implement a request arrival simulator that generates Poisson-distributed requests with varying prompt lengths (drawn from a realistic distribution), and measure the time-to-first-token and tokens-per-second distributions for ZeRO-Inference under different batching policies (dynamic batching with a maximum latency constraint, continuous batching with in-flight request addition/removal). Compare against a GPU-only baseline at the batch sizes it can support. The key question is whether ZeRO-Inference's throughput advantage (84 TFLOPS, 54% of peak) survives under realistic request patterns, or whether the cold-start latency from streaming layers from NVMe and the KV-cache memory pressure from long-generation requests erode the benefit. The paper's prefetching optimization (Figure 10c) provides a knob for trading memory for reduced transfer latency, but its interaction with dynamic batching is unexplored.

5. Extending PCC-style communication optimization to other parallelism combinations. The Parallelism-Coordinated Communication optimization (Section V-B) exploits a specific interaction: tensor-parallel all-reduce creates replicated data that reduces the scope of expert-parallel all-to-all. This principle — that data replication from one parallelism strategy can reduce the communication requirements of another — may apply to other combinations. For example: pipeline-parallel micro-batches generate intermediate activations that are only needed on adjacent stages; can this locality be exploited to reduce the scope of tensor-parallel all-reduce within each stage? When ZeRO-style data parallelism partitions optimizer states, can the partition boundaries be aligned with expert placement to reduce all-to-all for MoE layers? The concrete experiment: for each pair of parallelism strategies (tensor, pipeline, expert, data), identify whether a PCC-like optimization exists where the communication of one strategy satisfies or partially satisfies the communication requirement of another. Measure the communication volume reduction for a representative large MoE model (e.g., the 52B MoE configuration from Table II) on 128 GPUs. The paper's insight that communication requirements are not independent — that there is redundancy to exploit — suggests a systematic search space that the paper only samples at one point.

6. Benchmarking DeepSpeed Inference against model compression at equivalent quality. The paper positions DeepSpeed Inference as complementary to model compression techniques (distillation, quantization beyond INT8, structured pruning), but never evaluates the combined system or compares against compression alone. A critical comparison: for a target latency budget (e.g., 50 ms per token for GPT-3 175B on a single A100), compare two approaches: (a) the full-precision model with DeepSpeed Inference (FP16, optimized kernels, TP=8 as in Table I), and (b) a compressed model (distilled to a smaller student, or pruned to a sparser architecture) running on general-purpose inference infrastructure (FasterTransformer or ONNX Runtime) that meets the same latency budget. The metric is generation quality (perplexity or task accuracy) — not latency, since both are constrained to meet the target. This comparison would answer whether DeepSpeed Inference's systems optimizations are a substitute for or a complement to model compression: if DeepSpeed Inference + full model achieves better quality than compression + standard inference at the same latency, the case for systems investment is strengthened. The paper's E.T. comparison (Figure 12) and INT8 results partially address this, but the E.T. comparison is at a small scale (DistilBERT, BERT) where compression is standard practice, and the INT8 results lack accuracy evaluation. A head-to-head at the 175B+ scale where inference cost is prohibitive would be informative.

---

Practical Applications and Downstream Use Cases

1. On-device or single-workstation fine-tuning and evaluation of large models. The most immediate practical impact of ZeRO-Inference is enabling individual data scientists to run inference on models that previously required GPU clusters. The paper demonstrates running a 530B-parameter model on a single A6000 GPU (48 GB) — hardware that costs a few thousand dollars and fits in a workstation, not a datacenter rack. This means a researcher can download a trained 175B or 530B checkpoint, run evaluation on their custom dataset, perform prompt engineering, or generate training data for distillation — all without requesting cluster access. The concrete benefit: the paper's 84 TFLOPS on A6000 corresponds to generating roughly 20-30 tokens per second for a 175B model at batch size 8-16 (extrapolating from Figure 6's throughput numbers at larger batch sizes), which is fast enough for interactive iteration. Before ZeRO-Inference, this workflow required at minimum a multi-GPU node that most practitioners cannot access on demand. The "democratization" framing (Section VI) is not rhetorical — it changes who can participate in large model research by shifting the hardware requirement from "multiple datacenter GPUs" to "a single consumer-grade GPU with sufficient CPU RAM or NVMe storage."

2. Cost-efficient batch inference for content generation and data processing pipelines. Many production applications — offline query rewriting, document summarization, training data generation, bulk evaluation — prioritize throughput over latency. For these workloads, DeepSpeed Inference's throughput-oriented configurations (Figure 8) offer a direct cost reduction: the 1.51× throughput improvement for GPT-3 175B on 16 GPUs means the same batch job completes in 66% of the time, or equivalently requires 34% fewer GPU-hours. At cloud GPU pricing (~$3-4 per A100-hour at the time of writing), this translates to hundreds of dollars saved per large batch job. The activation offloading that enables larger batch sizes (Section IV-C2) is particularly valuable here: by freeing GPU memory for larger batches, DeepSpeed Inference amortizes the model weight loading cost over more tokens, increasing hardware efficiency. The paper's reporting of TFLOPS per GPU (72 TFLOPS for GPT-3 175B, 58 TFLOPS for MT-NLG 530B, Figure 8) gives practitioners a direct measure of hardware utilization they can use to estimate cost for their specific workloads. The key deployment question is whether the batch size that maximizes throughput also meets any application-specific latency constraints — the paper's hybrid scheduling (Figure 3) provides a mechanism to tune this tradeoff by varying micro-batch counts between prompt processing and token generation.

3. Serving sparse MoE models in latency-critical interactive applications. The paper's most striking latency result is serving a trillion-parameter MoE model under 25 ms per token on 256 GPUs (Figure 7), which the paper characterizes as "making it possible to serve such a massive model even in extremely interactive online applications." Before this result, MoE models at this scale were considered infeasible for interactive serving because the all-to-all communication for expert routing would introduce unacceptable latency. The PCC optimization and sparse kernel optimizations that enable this result (Sections V-B and V-C) are not model-specific — they apply to any MoE architecture with top-k gating and expert parallelism. This opens up a practical deployment path for the largest language models in applications like chatbots, code completion, and real-time translation, where users expect sub-100ms token generation latency. The concrete benchmark: for the 349B MoE model on 128 GPUs, DeepSpeed-MoE achieves ~2 ms per token (Figure 7), which means a 50-token response generates in ~100 ms — within the threshold for perceived interactivity. Without these optimizations, the same model at ~10 ms per token (baseline) would require ~500 ms for a 50-token response, which users experience as noticeably slow. The 5× latency gap between DeepSpeed-MoE and baseline at this scale is the difference between deployable and non-deployable for interactive use.

4. Scaling inference capacity linearly with available hardware for throughput-bound models. ZeRO-Inference's near-perfect linear scaling from 1 to 16 GPUs (Figure 9c, GPT-50B on DGX2 V100) provides a simple capacity planning model: if you need 2× the throughput, add 2× the GPUs, and the per-GPU efficiency (TFLOPS per GPU) remains roughly constant. This is not obvious for a system that streams model weights over PCIe — one might expect PCIe bandwidth to become a bottleneck that limits scaling. The paper's multi-GPU PCIe bandwidth aggregation (each GPU fetches only a partition of each layer over its dedicated PCIe link, then aggregates over NVLink in Section VI-B) explains why scaling is near-linear: the aggregate PCIe bandwidth grows with GPU count, so the per-GPU transfer time does not increase. This has practical implications for organizations deploying large models at scale: they can provision inference capacity incrementally without worrying about nonlinear scaling cliffs. The paper does not show at what GPU count the scaling deviates from linear (the DGX2 V100 has 16 GPUs with NVLink full-mesh connectivity — scaling to 32 or 64 GPUs with partial connectivity might show the limit), but the 16-GPU result is sufficient for node-level provisioning. For larger deployments, the 256-GPU MoE results (Figure 7) demonstrate that the system scales to datacenter scale for sparse models, though without the clean linear scaling characterization.

---

When to Prefer This Method

The paper explicitly positions DeepSpeed Inference against two alternatives: FasterTransformer for dense models and PyTorch-based distributed MoE for sparse models. The latency, throughput, and memory capacity results in Section VII provide the basis for a decision rule grounded in the paper's measured performance characteristics.

• Prefer DeepSpeed Transformer over FasterTransformer when: (1) You are running GPT-style autoregressive decoder models at batch sizes ≤ 16, where DeepSpeed-FP16 provides 1.3–1.55× lower latency than FT-FP16 (Figure 6) due to memory-bandwidth-optimized kernels (Deep-Fusion and SBI-GeMM); OR (2) You need INT8 inference for decoder models, which FT did not support at the time of writing (Section VII-B1), and the INT8 accuracy tradeoff (not quantified in the paper) is acceptable for your application; OR (3) You are running throughput-oriented inference with pipeline parallelism on models >175B parameters, where DeepSpeed Inference's hybrid scheduling and activation offloading enable 1.51–1.53× higher throughput than FT at each system's maximum batch size (Figure 8); OR (4) Your model exceeds the memory capacity of your GPU configuration and you need inference-optimized pipeline parallelism with activation offloading — the paper demonstrates 530B model inference on 40 GPUs with DeepSpeed Inference, while FT crashed when attempting TP+PP for the same model (Section VII-C).

• Prefer DeepSpeed-MoE over PyTorch-distributed MoE when: (1) You are serving MoE models at scale (>52B parameters, >32 GPUs), where the combined effect of expert-slicing, PCC all-to-all, and sparse kernel optimizations provides up to 7.3× lower latency than the PyTorch baseline (Figure 7); OR (2) Your deployment requires sub-25ms per-token latency for trillion-parameter MoE models, which the paper demonstrates is achievable only with PCC reducing all-to-all latency from O(p) to O(p/L) (Section V-B); OR (3) You are scaling MoE inference to hundreds of GPUs, where the PyTorch baseline's memory bandwidth utilization degrades to <20% of peak while DeepSpeed-MoE maintains ~50% (Figure 11).

• Prefer ZeRO-Inference over GPU-only deployment when: (1) Your model exceeds your available GPU memory — ZeRO-Inference supports models up to 25× larger than GPU-only on equivalent hardware (Figure 9b); OR (2) Your application is throughput-oriented (batch processing, data generation) rather than latency-critical, and you can achieve batch sizes large enough to hide PCIe/NVMe transfer latency — the paper demonstrates >50% throughput improvement over GPU-only for models that fit in GPU memory specifically because the larger batch sizes enabled by offloading increase the compute-to-transfer ratio (Section VI-A); OR (3) You have limited GPU budget (single workstation) but abundant CPU RAM or NVMe storage — ZeRO-Inference achieves 84 TFLOPS on a single A6000 (54% of peak), making large model inference economically accessible without GPU clusters.

• Prefer GPU-only (FasterTransformer or equivalent) over ZeRO-Inference when: (1) Your model fits comfortably in aggregate GPU memory AND your workload is latency-sensitive with small batch sizes — the paper acknowledges ZeRO-Inference is designed for "non-latency sensitive applications" (Section VI), and the cold-start latency of streaming weights from NVMe is not characterized but likely much higher than GPU-resident inference; OR (2) Your application requires very long generation sequences (hundreds or thousands of tokens), where the KV-cache memory pressure may limit ZeRO-Inference's batch size advantage — the paper only evaluates generation of 1-100 tokens and does not characterize throughput degradation at long generation lengths.

Important caveat: These preferences assume NVIDIA GPU hardware with NVLink intra-node interconnect and either InfiniBand (for multi-node) or PCIe Gen4 (for heterogeneous memory). The paper does not evaluate on AMD GPUs, non-NVLink interconnects, or alternative CPU/NVMe configurations, so the decision rules do not generalize to those platforms without independent validation. Additionally, the INT8 recommendation is conditional on acceptable accuracy, which the paper does not quantify — practitioners should benchmark INT8 output quality on their specific task and generation length before committing to INT8 deployment.