Appendix

Glossary

Every term, defined once

The vocabulary of this book, each entry quoted verbatim from the source that introduced it. Search covers both the terms and their definitions, so the concept works even when the name does not come to mind.

A

acceptance rate: The acceptance rate βx<t, given a prefix x<t, is the probability of accepting xt ~ q(xt|x<t) by speculative sampling. E(β) is then a natural measure of how well Mq approximates Mp.
algorithm cascading: Combine sequential and parallel reduction. Each thread loads and sums multiple elements into shared memory. Tree-based reduction in shared memory.
arithmetic intensity: Arithmetic intensity of a computation workload is commonly defined as the average number of computation operations performed per byte of data accessed from memory.
asynchronous copy: With cuda::memcpy_async, the thread block no longer stages data through registers, freeing the thread block from the task of moving data and freeing registers to be used by computations.
application replay: In Application Replay, all metrics requested for a specific kernel launch in NVIDIA Nsight Compute are grouped into one or more passes. In contrast to Kernel Replay, the complete application is run multiple times, so that in each run one of those passes can be collected per kernel.
atom: An "Atom" is the smallest collection of threads and data that must participate in the execution of a hardware-accelerated math or copy operation.
automatic prefix caching: Automatic Prefix Caching (APC in short) caches the KV cache of existing queries, so that a new query can directly reuse the KV cache if it shares the same prefix with one of the existing queries, allowing the new query to skip the computation of the shared part.
AccVGPR: The matrix core has its own VGPR file: the Accumulation ("Acc") GPRs. This is separate from the normal (Architectural, or "Arch") VGPRs in the original SIMD.

B

blocked program: We specifically revisit traditional "Single Program, Multiple Data" (SPMD) execution models for GPUs, and propose a variant in which programs, rather than threads, are blocked.
bank conflict: For a shared memory tile of 32 × 32 elements, all elements in a column of data map to the same shared memory bank, resulting in a worst-case scenario for memory bank conflicts: reading a column of data results in a 32-way bank conflict.
buffer rotation: By allocating duplicate buffers for each input tensor, with a total size at least twice the L2 cache capacity (and using at least two buffers per tensor), we ensure that after each kernel iteration, the next computation accesses data not resident in the cache.
block-sparse FlashAttention: we implement block-sparse FlashAttention, a sparse attention algorithm that is 2-4x faster than even FlashAttention, scaling up to sequence length of 64k.
backend scale-out network: The backend scale-out network often forms its own layer-3 subnet and is not usually connected directly to the frontend network.

C

collective: Collective communication routines are common patterns of data transfer among many processors.
cubin: A CUDA binary (also referred to as cubin) file is an ELF-formatted file which consists of CUDA executable code sections as well as other sections containing symbols, relocators, debug info, etc.
constrained decoding: To ensure that text generated by large language models (LLMs) is in an expected format, constrained decoding methods propose to enforce strict formal language constraints during generation.
chunked prefill: Sarathi-Serve introduces chunked-prefills which splits a prefill request into near equal sized chunks and creates stall-free schedules that adds new requests in a batch without pausing ongoing decodes.
cluster launch control: Blackwell introduces cluster launch control (CLC) for dynamic scheduling.
constant memory: Constant memory has grid scope and is accessible for the lifetime of the application. The constant memory resides on the device and is read-only to the kernel.
compute capability: All NVIDIA GPUs have a compute capability (CC) which is a two part identifier in the form major.minor. The major version identifies the GPU generation, while the minor number identifies the version within that generation.
collective builder: CollectiveBuilder accepts CUTLASS 2.x equivalent input template arguments, and attempts to build the best performing CollectiveMma from the given parameters.
critical batch size: Transformer matmuls are compute-bound iff the per-replica token batch size is greater than B_crit.

D

device: A GPU and the memory directly connected to it are referred to as the device and device memory, respectively. CUDA applications execute some part of their code on the GPU, but applications always start execution on the CPU.
draft model: Generating a short draft of length K. This can be attained with either a parallel model or by calling a faster, auto-regressive model K times. We shall refer to this model as the draft model, and focus on the case where it is auto-regressive.
dynamic sparse attention: We determine the optimal pattern for each attention head offline and dynamically build sparse indices based on the assigned pattern during inference.
distributed shared memory: Thread Block Clusters, introduced in compute capability 9.0 and facilitated by Cooperative Groups, provide the ability for threads in a thread block cluster to access shared memory of all the participating thread blocks in that cluster. This partitioned shared memory is called Distributed Shared Memory.
decoupled look-back: Our method embodies a decoupled look-back strategy that performs redundant work to dissociate local computation from the latencies of global prefix propagation.

E

execution configuration: The number of threads that will execute the kernel in parallel is specified as part of the kernel launch. This is called the execution configuration.

F

flop bound: Flop bound would then mean that there is time when nothing is being passed through memory, and memory bound would mean that no floperations are occuring.
FP8: An 8-bit floating point (FP8) binary interchange format consisting of two encodings - E4M3 (4-bit exponent and 3-bit mantissa) and E5M2 (5-bit exponent and 2-bit mantissa).
fair queueing: Fair queueing ensures that each client will get their "fair share". In the simplest case, if there are n clients sharing the same resource, the fair share is at least 1/n of the resource.
fast_p: A new evaluation metric fast_p, which measures the percentage of generated kernels that are functionally correct and offer a speedup greater than an adjustable threshold p over baseline.
feature-level drafting: Autoregression at the feature (second-to-top-layer) level is more straightforward than at the token level.
fatbinary: nvcc organizes its device code in fatbinaries, which are able to hold multiple translations of the same GPU source code. At runtime, the CUDA driver will select the most appropriate translation when the device function is launched.
frontend network: The frontend network is the operational network in datacenters that connects all compute nodes to the outside world (e.g., other datacenters or end customers on the Internet).

G

grid: Thread blocks are organized into a grid. All the thread blocks in a grid have the same size and dimensions. A grid may consist of millions of thread blocks, while the GPU executing the grid may have only tens or hundreds of SMs.
goodput: The maximum request rate that can be served adhering to the SLO attainment goal (say, 90%) for each GPU provisioned.
grouped-query attention: Grouped-query attention (GQA), a generalization of multi-query attention which uses an intermediate (more than one, less than number of query heads) number of key-value heads.
global memory: Global memory (also called device memory) is the primary memory space for storing data that is accessible by all threads in a kernel. It is similar to RAM in a CPU system.
GPU Speed Of Light Throughput: High-level overview of the throughput for compute and memory resources of the GPU. For each unit, the throughput reports the achieved percentage of utilization with respect to the theoretical maximum.
GpSimd Engine: GpSimd Engine (GpSimdE) is intended to be a general-purpose engine that can run any ML operators that cannot be lowered onto the other highly specialized compute engines discussed above efficiently, such as applying a triangular mask to a tensor.

H

host: The CPU and the memory directly connected to it are called the host and host memory, respectively. A GPU and the memory directly connected to it are referred to as the device and device memory, respectively.
HIP: HIP is a C++ runtime API and kernel language for AMD GPUs. It is part of AMD's ROCm platform and lets developers create applications that run on heterogeneous systems, using CPUs and AMD GPUs from a single source code base.
HIPIFY: HIPIFY is a collection of tools that automatically translate CUDA code to HIP code.

I

instruction set architecture: An instruction set architecture (ISA) is the specification for what instructions a processor can execute, their format, the behavior of those instructions, and their binary encodings.
inference gateway: A proxy/load-balancer that has been coupled with the EndPointer Picker extension. It provides optimized routing and load balancing for serving Kubernetes self-hosted generative Artificial Intelligence (AI) workloads.
interactivity: The horizontal axis is interactivity, measured in tokens per second per user. This is how fast the AI system response streams back to each query.

K

kernel: The code an application executes on the GPU is referred to as device code, and a function that is invoked for execution on the GPU is, for historical reasons, called a kernel. The act of starting a kernel running is called launching the kernel.
KV cache: In the sampling, the transformer performs self-attention, which requires the kv values for each item currently in the sequence. These vectors are provided a matrix known as the kv cache, aka past cache. The purpose of this is to avoid recalculations of those vectors every time we sample a token.
kernel replay: In Kernel Replay, all metrics requested for a specific kernel instance in NVIDIA Nsight Compute are grouped into one or more passes. For the first pass, all GPU memory that can be accessed by the kernel is saved.
KV cache quantization: The key cache should be quantized per-channel, i.e., group elements along the channel dimension and quantize them together. In contrast, the value cache should be quantized per-token.
KV cache reuse: Blocks containing KV state computed for previous requests are stored in a radix search tree as soon as they are filled. A search is performed when a new request is added, and matched blocks are reused instead of calculated.

L

layout algebra: CuTe provides an "algebra of Layouts" to support combining layouts in different ways.
local memory: Local memory is so named because its scope is local to the thread, not because of its physical location. In fact, local memory is off-chip. Hence, access to local memory is as expensive as access to global memory.
launch bounds: Applications can optionally aid these heuristics by providing additional information to the compiler in the form of launch bounds that are specified using the __launch_bounds__() qualifier in the definition of a __global__ function.
Local Data Share: The AMD CDNA 4 LDS is also more efficient than the prior design and supports loading data directly from the L1 data cache, thereby reducing vector register usage and latency.
load generator: The LoadGen is a traffic generator for MLPerf Inference that loads the SUT and measures performance.

M

multi-head latent attention: MLA guarantees efficient inference through significantly compressing the Key-Value (KV) cache into a latent vector.
memory coalescing: The important thing to remember is that to ensure memory coalescing we want to map the quickest varying component to contiguous elements in memory.
microscaling: The core concept behind micro-scaling is enabling hardware support for a scale factor that is shared across a block of data elements (typically 32) within a tensor, rather than just a single scale factor for the entire tensor.
memory bandwidth utilization: Memory-bound typically means the achieved device memory bandwidth utilization (MBU) is close to 100% (60%+ is considered good in practice).
mbarrier phase bit: The parity qualifier, whose use entails providing a phase bit, indicates that the thread sleeps until that phase bit of the mbarrier flips.
micro-tensor scaling: The Blackwell Transformer Engine utilizes advanced dynamic range management algorithms and fine-grain scaling techniques, called micro-tensor scaling, to optimize inference performance, accuracy, and enable FP4 AI.
MFMA: The core operation implemented inside the matrix core is the 4 x 1 times 1 x 4 outer matrix product, yielding 16 output values.
M2N communication: the arbitrary parallelism configuration of the attention and FFN modules transforms the original All2All communication between them for token routing into M2N communication, where M and N represent the number of senders and receivers, respectively

N

native sparse attention: Combining coarse-grained token compression with fine-grained token selection to preserve both global context awareness and local precision.
NVTX: NVTX is a cross-platform API for annotating source code to provide contextual information to developer tools. The NVTX API is written in C, with wrappers provided for C++ and Python.
non-matmul FLOPs: the A100 GPU has a max theoretical throughput of 312 TFLOPs/s of FP16/BF16 matmul, but only 19.5 TFLOPs/s of non-matmul FP32. Another way to think about this is that each non-matmul FLOP is 16x more expensive than a matmul FLOP.

O

occupancy: Occupancy is the ratio of the number of active warps per multiprocessor to the maximum number of possible active warps.
online normalizer calculation: Calculates both the maximum value m and the normalization term d in a single pass over input vector with negligible additional cost of two operations per vector element. It reduces memory accesses from 4 down to 3 per vector element for the Softmax function evaluation.
overlap scheduler: The Overlap Scheduler maximizes GPU utilization by hiding CPU-bound latency behind GPU computation.
operational intensity: operations per byte of DRAM traffic

P

PTX: Parallel thread execution (PTX) is a virtual machine instruction set architecture that has been part of CUDA from its beginning. You can think of PTX as the assembly language of the NVIDIA CUDA GPU computing platform.
phase splitting: We propose Splitwise, a model deployment and scheduling technique that splits the two phases of LLM inference requests on to separate machines. Splitwise enables phase-specific resource management using hardware that is well suited for each phase.
persistent kernel: CUTLASS has adopted a software technique named persistent kernels. Persistent clusters, or Workers, can stay on the GPU throughout kernel execution and process multiple tiles, hiding prologue and epilogue costs.
prefix scan: Given a set of input elements and a binary reduction operator, a prefix scan produces a corresponding output list where each output is computed to be the reduction of the elements occurring earlier in the input.
preemption: Due to the autoregressive nature of transformer architecture, there are times when KV cache space is insufficient to handle all batched requests. In such cases, vLLM can preempt requests to free up KV cache space for other requests. Preempted requests are recomputed when sufficient KV cache space becomes available again.
performance ceiling: you cannot break through a ceiling without performing the associated optimization

R

ring attention: Ring Attention with Blockwise Transformers (Ring Attention), which leverages blockwise computation of self-attention and feedforward to distribute long sequences across multiple devices while fully overlapping the communication of key-value blocks with the computation of blockwise attention.
RadixAttention: This section introduces RadixAttention, a novel technique for automatic and systematic KV cache reuse during runtime. Unlike existing systems that discard the KV cache after a generation request finishes, our system retains the cache for prompts and generation results in a radix tree, enabling efficient prefix search, reuse, insertion, and eviction.
redundant experts: We adopt a redundant experts strategy that duplicates heavy-loaded experts. Then, we heuristically pack the duplicated experts to GPUs to ensure load balancing across different GPUs.
register spilling: Register spilling means that values currently stored on-chip in registers must be written out to global memory and then read back later to make space for other values.
register pressure: Register pressure occurs when there are not enough registers available for a given task. Even though each multiprocessor contains thousands of 32-bit registers, these are partitioned among concurrent threads.
ridge point: The ridge point is the point at which the memory bandwidth boundary meets the peak performance boundary.
recomputation (attention backward pass): We store the softmax normalization factor from the forward pass to quickly recompute attention on-chip in the backward pass, which is faster than the standard approach of reading the intermediate attention matrix from HBM.

S

streaming multiprocessor: The GPU can be considered to be a collection of Streaming Multiprocessors (SMs) which are organized into groups called Graphics Processing Clusters (GPCs). Each SM contains a local register file, a unified data cache, and a number of functional units that perform computations.
SIMT: In SIMT, all threads in the warp are executing the same kernel code, but each thread may follow different branches through the code. That is, though all threads of the program execute the same code, threads do not need to follow the same execution path.
shared memory: All threads of a thread block are executed in a single SM. This allows threads within a thread block to communicate and synchronize with each other efficiently. Threads within a thread block all have access to the on-chip shared memory, which can be used for exchanging information between threads of a thread block.
scaling factor: Higher precision values need to be multiplied with a scaling factor prior to their casting to FP8 in order to move them into a range that better overlaps with the representable range of a corresponding FP8 format.
selective batching: We suggest selective batching, which applies batching only to a selected set of operations.
salient weights: Protecting only 1% salient weights can greatly reduce quantization error. To identify salient weight channels, we should refer to the activation distribution, not weights.
SMEM (TPU): SMEM is a low-latency memory that supports random access, but lets you only read and write 32-bit values with a single instruction.
SparseCore: Each TPU v4 includes SparseCores, dataflow processors that accelerate models that rely on embeddings by 5x-7x yet use only 5% of die area and power.
software pipelining: To mitigate the effects of memory latency, CUTLASS uses software pipelining to overlap memory accesses with other computation within a thread. CUTLASS accomplishes this by double buffering.
structured outputs: vLLM supports the generation of structured outputs using xgrammar or guidance as backends.
SBUF and PSUM: Both SBUF and PSUM are considered two-dimensional memories with 128 partitions each, i.e., one SBUF partitions has 192KiB of memory while one PSUM partition has 16KiB.
Sync Engine: The Sync Engine is most commonly used to trigger DMA transfers without interfering with compute engine instruction scheduling and ordering.
scale-up network: Scale-up networks are typically very specialized short-range interconnects that often come with only a single tier of switches or possibly no switch at all.
server scenario: The server scenario represents online applications where query arrival is random and latency is important.
SOL-Score: a metric that grades custom kernel performance based on the theoretical roofline of a NVIDIA B200 GPU (obtained analytically with SOLAR)

T

thread block: When an application launches a kernel, it does so with many threads, often millions of threads. These threads are organized into blocks. A block of threads is referred to as a thread block. All threads of a thread block are executed in a single SM.
thread block cluster: Clusters are a group of thread blocks which, like thread blocks and grids, can be laid out in 1, 2, or 3 dimensions. Because the thread blocks are scheduled simultaneously and within a single GPC, threads in different blocks but within the same cluster can communicate and synchronize with each other using software interfaces provided by Cooperative Groups.
tile programming: In tile programming, the programmer writes code at the level of an entire thread block, describing operations on multidimensional collections of data called tiles. The compiler maps these operations to the individual threads of the block.
tile block: Tile IR models the GPU as a tile-based processor. In Tile IR, each logical thread (tile block) computes over partial fragments (tiles) of multi-dimensional arrays (tensors).
tiling: We restructure the attention computation to split the input into blocks and make several passes over input blocks, thus incrementally performing the softmax reduction (also known as tiling).
TF32: Non-tensor operations continue to use the FP32 datapath, while TF32 tensor cores read FP32 data and use the same range as FP32 with reduced internal precision, before producing a standard IEEE FP32 output. TF32 includes an 8-bit exponent (same as FP32), 10-bit mantissa (same precision as FP16) and 1 sign-bit.
time to first token: It is defined as the time taken between arrival and first output token generated by system for each request. TTFT includes both scheduling delay and prompt processing time.
time per output token: It is defined as total time taken to generate all output tokens divided by the number of output tokens generated.
time between tokens: It is defined as the time taken between two consecutive output tokens generated by system for each request.
tree attention: Using a tree-based attention mechanism, Medusa constructs multiple candidate continuations and verifies them simultaneously in each decoding step.
tile grid: Tile IR allows tile blocks to be grouped into a tile grid, similar to CUDA C++, enabling users to launch sets of tile blocks that execute in parallel.
Tensor Memory Accelerator: TMA (Tensor Memory Accelerator) is a new feature introduced in the NVIDIA Hopper architecture for doing asynchronous memory copy between a GPU's global memory (GMEM) and the shared memory (SMEM) of its threadblocks (i.e., CTAs).
tcgen05.mma: Blackwell SM100 introduces tcgen05.mma instructions. tcgen05.mma instructions support all legacy types (tfloat32_t, half_t, bfloat16_t, int8_t, uint8_t) and the new 4, 6, and 8-bits floating point datatypes with and without scale factors.
Tensor Memory (TMEM): A new on-chip memory dedicated to the accumulator (and, optionally, operand A). tcgen05 MMA reads and writes the accumulator in TMEM directly, freeing the register file for other work.

U

unified memory: Unified memory is a feature of the CUDA runtime which lets the NVIDIA Driver manage movement of data between host and device(s).
UALink: The UALink 200G 1.0 Specification defines a low-latency, high-bandwidth interconnect for communication between accelerators and switches in AI computing pods.
Ultra Ethernet: The Ultra Ethernet (UE) specification covers a broad range of software and hardware relevant to AI and HPC workloads: from the API supported by UE-compliant devices to the services offered by the transport, link, and physical layers, as well as management, interoperability, benchmarks, and compliance requirements. UE does not require or mandate changes to the network layer or the Ethernet PHY and link layers.
unified data cache: The unified data cache provides the physical resources for shared memory and L1 cache. The allocation of the unified data cache to L1 and shared memory can be configured at runtime.

V

virtual architecture: GPU compilation is performed via an intermediate representation, PTX, which can be considered as assembly for a virtual GPU architecture. Contrary to an actual graphics processor, such a virtual GPU is defined entirely by the set of capabilities, or features, that it provides to the application.
VMEM: VMEM is fairly large for such a low-level memory hierarchy (16MB+), making it possible to use large window sizes.
vectorized load: Using vectorized loads reduces the total number of instructions, reduces latency, and improves bandwidth utilization.

W

warp: Within a thread block, threads are organized into groups of 32 threads called warps. A warp executes the kernel code in a Single-Instruction Multiple-Threads (SIMT) paradigm.
warp divergence: If some threads within a warp follow a control flow branch in execution while others do not, the threads which do not follow the branch will be masked off while the threads which follow the branch are executed. When different threads in a warp follow different code paths, this is sometimes called warp divergence.
warp group: Note that a warp = 32 threads, so 128 threads will comprise 4 warps. A group of 4 warps is called a warp-group in Hopper architecture.
warp shuffle: Warp shuffle functions exchange a value between non-exited threads within a warp without the use of shared memory.
warp scheduler: A warp scheduler can issue one warp instruction per cycle.