# Glossary Every term of art in this book, defined once. Search covers both the terms and their definitions, so the concept works even when the name does not come to mind. ## A acceptance rate — We show that the expected acceptance rate of draft tokens is sufficient to offset the overhead of the drafting process for large language models, resulting in an effective and practical method for reducing sampling latency without the need for modifying the target model or biasing the sample distribution. 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. ## 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. ## C collective — NCCL is a stand-alone library of standard communication routines for GPUs, implementing all-reduce, all-gather, reduce, broadcast, reduce-scatter, as well as any send/receive based communication pattern. 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 prefills — 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. ## 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 — The latency of parallel scoring of short continuations, generated by a faster but less powerful draft model, is comparable to that of sampling a single token from the larger target model. 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 — A thread block can read from, write to, and perform atomics in shared memory of other thread blocks within its cluster. This is known as Distributed Shared Memory. ## 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. Nvidia uses the term math bandwidth for this delineation, which technically exists per kernel but can be abstracted to exist for groups of operations. 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. ## 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. ## 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. ## 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 sampling, the transformer performs self-attention, which requires the k and v 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, at the cost of some storage. 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. ## L layout algebra — CuTe provides an "algebra of Layouts" to support combining layouts in different ways. ## 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). ## N native sparse attention — Combining coarse-grained token compression with fine-grained token selection to preserve both global context awareness and local precision. NVTX — When you add NVTX markers and ranges to your application, the Timeline View shows when your CPU threads are executing within those regions. ## 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. ## 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 — Each inference request undergoes two phases: a compute-intensive prompt computation phase and a memory-intensive token generation phase, each with distinct latency, throughput, memory, and power characteristics. persistent kernel — Persistent clusters, or Workers, can stay on the GPU throughout kernel execution and process multiple tiles, hiding prologue and epilogue costs. ## 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 — Our system maintains an LRU cache of the KV cache for all requests within a radix tree. 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. ## 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. ## 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. All thread blocks in a cluster are executed in a single GPC, and threads in different blocks but within the same cluster can communicate and synchronize with each other. 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 — TF32 (TensorFloat-32) is a compute mode that processes float32 matrix multiplications using Tensor Cores with 19-bit mantissa precision, the same exponent range as IEEE float32 but with a truncated significand. 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. ## 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 1.0 Specification enables 200G per lane scale-up connection for up to 1,024 accelerators within an AI computing pod, delivering the open standard interconnect for next-generation AI cluster performance. Ultra Ethernet — Deliver an Ethernet based open, interoperable, high performance, full-communications stack architecture to meet the growing network demands of AI & HPC at scale. ## V virtual 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. ## 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.