# Learn Kernels: complete text Every page of the book in reading order. Interactive figures appear as labelled placeholders; open the linked page to use them. --- # Learn Kernels: kernel and inference performance, traced to primary sources # Learn Kernels: An interactive book on GPU kernels and inference performance. ## LearnKernels GPU performanceengineering [Start reading](https://learn-kernels.com/chapters/introduction) [Continue reading](https://learn-kernels.com/chapters/introduction) [Glossary](https://learn-kernels.com/chapters/glossary)[Further reading](https://learn-kernels.com/chapters/reading) ## What this book is Getting a GPU to run a computation is easy. Getting it to run anywhere near what the hardware can actually do is the entire discipline. This book is about that gap: why it exists, and the vocabulary you need to reason about it. It is ordered the way the problem is, from one inference request down to a single kernel and back out to a fleet of machines. Every factual claim traces to a primary source, linked where it is used and collected in the reading list. ## Chapters ### [0\. Introduction](https://learn-kernels.com/chapters/introduction)[](https://learn-kernels.com/chapters/introduction) - [0.1Why kernels matter](https://learn-kernels.com/chapters/introduction/why-kernels-matter) - [0.2How this book is organized](https://learn-kernels.com/chapters/introduction/how-this-book-is-organized) - [0.3How to read it](https://learn-kernels.com/chapters/introduction/how-to-read-it) ### [1\. Foundations](https://learn-kernels.com/chapters/foundations)[](https://learn-kernels.com/chapters/foundations) - [1.1Threads, blocks, and warps](https://learn-kernels.com/chapters/foundations/threads-blocks-and-warps) - [1.2Compute-bound and memory-bound](https://learn-kernels.com/chapters/foundations/compute-bound-and-memory-bound) - [1.3The KV cache](https://learn-kernels.com/chapters/foundations/the-kv-cache) - [1.4Compilation and machine code](https://learn-kernels.com/chapters/foundations/compilation-and-machine-code) ### [2\. Kernel optimization](https://learn-kernels.com/chapters/kernel-optimization)[](https://learn-kernels.com/chapters/kernel-optimization) - [2.1Foundational kernel exercises](https://learn-kernels.com/chapters/kernel-optimization/foundational-kernel-exercises) - [2.2Matrix multiplication](https://learn-kernels.com/chapters/kernel-optimization/matrix-multiplication) - [2.3Tensor cores and low precision](https://learn-kernels.com/chapters/kernel-optimization/tensor-cores-and-low-precision) - [2.4Attention](https://learn-kernels.com/chapters/kernel-optimization/attention) ### [3\. Programming models and profiling](https://learn-kernels.com/chapters/programming-models)[](https://learn-kernels.com/chapters/programming-models) - [3.1Triton](https://learn-kernels.com/chapters/programming-models/triton) - [3.2CUTLASS, CuTe, and CUDA Tile](https://learn-kernels.com/chapters/programming-models/cutlass-cute-and-cuda-tile) - [3.3Other hardware stacks](https://learn-kernels.com/chapters/programming-models/other-hardware-stacks) - [3.4Profiling, benchmarking, and correctness](https://learn-kernels.com/chapters/programming-models/profiling-and-correctness) ### [4\. Inference engines](https://learn-kernels.com/chapters/inference-engines)[](https://learn-kernels.com/chapters/inference-engines) - [4.1Scheduling and continuous batching](https://learn-kernels.com/chapters/inference-engines/scheduling-and-continuous-batching) - [4.2KV cache systems](https://learn-kernels.com/chapters/inference-engines/kv-cache-systems) - [4.3Quantization](https://learn-kernels.com/chapters/inference-engines/quantization) - [4.4Speculative decoding](https://learn-kernels.com/chapters/inference-engines/speculative-decoding) - [4.5Structured decoding and fairness](https://learn-kernels.com/chapters/inference-engines/structured-decoding-and-fairness) - [4.6Long context and multimodal inference](https://learn-kernels.com/chapters/inference-engines/long-context-and-multimodal) ### [5\. Distributed inference](https://learn-kernels.com/chapters/distributed-inference)[](https://learn-kernels.com/chapters/distributed-inference) - [5.1Parallelism, collectives, and topology](https://learn-kernels.com/chapters/distributed-inference/parallelism-collectives-and-topology) - [5.2Mixture-of-experts serving](https://learn-kernels.com/chapters/distributed-inference/mixture-of-experts-serving) - [5.3Prefill and decode disaggregation](https://learn-kernels.com/chapters/distributed-inference/prefill-and-decode-disaggregation) - [5.4Production systems and serving benchmarks](https://learn-kernels.com/chapters/distributed-inference/production-systems-and-benchmarks) ### [6\. Current hardware](https://learn-kernels.com/chapters/hardware)[](https://learn-kernels.com/chapters/hardware) - [6.1NVIDIA Blackwell](https://learn-kernels.com/chapters/hardware/nvidia-blackwell) - [6.2AMD, TPU, and Trainium](https://learn-kernels.com/chapters/hardware/amd-tpu-and-trainium) ### [7\. Frontier](https://learn-kernels.com/chapters/frontier)[](https://learn-kernels.com/chapters/frontier) - [7.1AI-generated kernels](https://learn-kernels.com/chapters/frontier/ai-generated-kernels) - [7.2Watchlist](https://learn-kernels.com/chapters/frontier/watchlist) ### [Glossary](https://learn-kernels.com/chapters/glossary)[](https://learn-kernels.com/chapters/glossary) Every term, defined once ### [Further reading](https://learn-kernels.com/chapters/reading)[](https://learn-kernels.com/chapters/reading) - [B.1Start here](https://learn-kernels.com/chapters/reading#start-here) - [B.2GPU fundamentals](https://learn-kernels.com/chapters/reading#programming-model) - [B.3Kernel optimization](https://learn-kernels.com/chapters/reading#foundational-kernel-exercises) - [B.4Programming models and profiling](https://learn-kernels.com/chapters/reading#triton) - [B.5Inference engines](https://learn-kernels.com/chapters/reading#scheduling-batching) - [B.6Distributed inference](https://learn-kernels.com/chapters/reading#parallelism-collectives) - [B.7Current hardware](https://learn-kernels.com/chapters/reading#hardware-nvidia) - [B.8Frontier](https://learn-kernels.com/chapters/reading#ai-generated-kernels) --- # About # About this site Learn Kernels is a book about GPU performance engineering: the gap between running a computation on an accelerator and running it at the speed the hardware allows. It is ordered from one inference request down to a single kernel and back out to distributed serving, and it is written for people who want to reason about that gap rather than memorize recipes. ## Sources Every factual claim traces to a primary source: the paper that introduced a mechanism, the specification that defines it, or the repository that implements it. Sources are linked in a sidenote at the end of each section and collected in the reading list, whose structure follows the [gpu-perf-engineering-resources](https://github.com/wafer-ai/gpu-perf-engineering-resources) list. This site is not affiliated with that project. Where a figure or an explanation gets something wrong, the mistake is this site’s. ## Who built it The site does not say, but the Ask AI on any chapter will tell you if you ask. ## For developers Learn Kernels also publishes a keyless JSON API and an MCP server for the same chapter index. [See the developer page](https://learn-kernels.com/developers). --- # Distributed inference A single GPU cannot hold a frontier model, and a single request cannot use a whole cluster efficiently. Everything in this chapter is about the seam between those two facts: how a model gets split across chips, how experts get routed to and dispatched across nodes, and how the two arithmetic phases of generation end up running on different hardware entirely. ## 4 sections 1. [5.1Parallelism, collectives, and topology](https://learn-kernels.com/chapters/distributed-inference/parallelism-collectives-and-topology) 2. [5.2Mixture-of-experts serving](https://learn-kernels.com/chapters/distributed-inference/mixture-of-experts-serving) 3. [5.3Prefill and decode disaggregation](https://learn-kernels.com/chapters/distributed-inference/prefill-and-decode-disaggregation) 4. [5.4Production systems and serving benchmarks](https://learn-kernels.com/chapters/distributed-inference/production-systems-and-benchmarks) --- # Mixture-of-experts serving · Distributed inference A mixture-of-experts model replaces a single dense feed-forward block with many smaller ones and a router that picks which few run for each token. DeepSeek-V3 is a concrete example of the shape this takes: each MoE layer has 1 shared expert, always active, plus 256 routed experts, of which 8 are activated for any given token through **expert routing** (A learned function that scores every expert against a token and selects the top-scoring subset to run for that token.), a sigmoid affinity score per expert with a per-expert bias term adjusted over training to keep load balanced without an auxiliary loss that would otherwise hurt model quality. The result is a model with 671B total parameters where only 37B are activated per token, which is the entire economic argument for MoE: most of the parameters sit idle for any single forward pass. That sparsity has to be realized physically, and that is where **expert parallelism** (Distributing a model's experts across GPUs, so each GPU holds a subset of experts and tokens are routed to whichever GPU holds the expert they need.) and its communication pattern take over. DeepSeek-V3’s prefilling deployment unit spans 4 nodes and 32 GPUs, combining 4-way tensor parallelism for attention with 32-way expert parallelism for the MoE blocks, so that “each expert processes a sufficiently large batch size.” Because a token’s chosen experts can live on any GPU in that group, every MoE layer needs an all-to-all exchange: each GPU sends token activations out to the GPUs hosting its routed experts (dispatch) and receives the computed results back (combine). DeepEP, DeepSeek’s open-source dispatch and combine library, describes itself as providing “high-throughput and low-latency all-to-all GPU kernels” for exactly this exchange, built to route dispatch traffic across nodes over InfiniBand and then forward it to the right GPU within a node over NVLink. > Figure. Expert routing in a mixture-of-experts layer. One token reaches the router, which picks 8 of the layer's 256 routed experts; the shared expert runs for every token regardless. Eight expert boxes are drawn to stand in for the 256, and the highlighted ones stand for the chosen 8. Dispatch carries the token's activations out to the GPUs hosting its chosen experts and combine brings the results back, the all-to-all exchange DeepEP provides kernels for.Illustrative numbers Because the router decides per token which experts get used, expert load is never guaranteed even, and an overloaded expert stalls the GPU that hosts it while the rest of the all-to-all waits. DeepSeek-V3 addresses this by deploying redundant copies of the experts that online traffic statistics show are hottest, rebalancing that set roughly every 10 minutes; its decoding deployment goes further, spanning 40 nodes and 320 GPUs with one expert per GPU and 64 GPUs dedicated to hosting redundant and shared experts, using direct point-to-point transfers over InfiniBand to keep dispatch and combine latency low. The load-balancing problem and the communication problem are really the same problem seen from two sides: a router that is free to send every token anywhere needs both a placement strategy that keeps GPUs evenly loaded and a communication kernel fast enough that the resulting all-to-all does not dominate the layer’s runtime. DeepSeek open-sourced the placement half of that answer as EPLB, its expert parallelism load balancer, which computes “a balanced expert replication and placement plan based on the estimated expert loads”; predicting those loads is left to the deployer, with a moving average of historical statistics named as the common method. It ships two policies. Hierarchical load balancing, meant for the prefilling stage with a smaller expert-parallel size, packs whole expert groups onto nodes evenly and then replicates within each node, exploiting DeepSeek-V3’s group-limited routing to place “the experts of the same group to the same node to reduce inter-node data traffic.” Global load balancing, for the decoding stage with a larger expert-parallel size, replicates experts across the whole fleet regardless of groups. Either way, replication and placement fall out of measured load rather than being fixed at deployment: the plan is an output of traffic. MegaScale-Infer pushes the same disaggregation logic inside the MoE layer itself. Its starting observation is that sparse activation “shifts feed-forward networks (FFNs) from being compute-intensive to memory-intensive during inference, leading to substantially lower GPU utilization,” because each expert sees only its routed slice of the batch. The system therefore disaggregates the attention and FFN modules within each layer onto separate GPUs, giving each module its own parallelism strategy and hardware, and runs what it calls ping-pong pipeline parallelism: a request batch is partitioned into micro-batches that shuttle between the attention side and the expert side, so one side computes while the other’s traffic is in flight. Backed by an M2N communication library that strips out “unnecessary GPU-to-CPU data copies, group initialization overhead, and GPU synchronization,” MegaScale-Infer reports up to 1.90× higher per-GPU throughput than state-of-the-art systems. Source Architecture, parameter counts, and deployment numbers from the [DeepSeek-V3 technical report](https://arxiv.org/html/2412.19437), sections 2.1.2 and 3.4. Dispatch and combine kernel description from the [DeepEP](https://github.com/deepseek-ai/DeepEP) README. Load balancing policies and the replication plan interface from the [EPLB](https://github.com/deepseek-ai/EPLB) README. Attention/FFN disaggregation, ping-pong pipeline parallelism, and the throughput result from the [MegaScale-Infer](https://arxiv.org/abs/2504.02263) abstract. Expert placement strategies at scale are covered further in [mixture-of-experts serving](https://learn-kernels.com/chapters/reading#moe-serving). --- # Parallelism, collectives, and topology · Distributed inference Splitting a model across GPUs means picking which axis to cut along. **Tensor parallelism** (Splitting individual matrix multiplications inside a layer across GPUs, so every GPU holds a slice of every layer's weights.) cuts inside a layer: Megatron-LM partitions the weight matrix of a transformer's MLP block by splitting the first GEMM’s weight matrix along its rows and the input along its columns, so each GPU computes a partial result that has to be combined before the next operation can proceed. That combination is a synchronization point after every parallelized block, which is why tensor parallelism wants the fastest possible link between GPUs, typically NVLink inside a node, and degrades quickly across a slower network. **Pipeline parallelism** (Splitting a model's layers into stages, each placed on a different GPU, so a request flows through the stages like a pipeline.) cuts the other way: whole layers are assigned to different GPUs, and activations flow between stages instead of being combined mid-layer. Megatron-LM describes its own tensor parallel approach as “orthogonal and complimentary to pipeline model parallelism,” meaning production systems compose both at once rather than choosing one. Neither form of parallelism does anything without a way for GPUs to exchange the partial results. That is the job of a **collective** (A communication operation, such as all-reduce or all-gather, that all participating processes perform together to combine or redistribute data.): an operation like all-reduce, all-gather, or reduce-scatter that every participating GPU executes together. NCCL, NVIDIA’s library for this, implements exactly these primitives: all-reduce, all-gather, reduce, broadcast, reduce-scatter, and arbitrary send/receive patterns, and is “optimized to achieve high bandwidth on platforms using PCIe, NVLink, NVswitch, as well as networking using InfiniBand Verbs or TCP/IP sockets,” supporting any number of GPUs in a single node or spread across many. Tensor parallelism’s per-layer combine step is an all-reduce; pipeline parallelism’s stage-to-stage handoff is a point-to-point send/receive. Which collective a parallelism strategy needs, and how often it needs it, is what determines whether that strategy can tolerate a slow interconnect or requires the fastest one available. In code, a collective is less exotic than the name suggests. The first complete example in NCCL’s documentation has a single process drive four GPUs: it creates one communicator per device with `ncclCommInitAll`, then issues an all-reduce on every device inside a group call, which NCCL requires when one thread manages multiple GPUs. Each call names a send buffer, a receive buffer, an element count, a datatype, and a reduction operator; the library sums the four send buffers and leaves the identical result in every receive buffer. Nothing is finished when `ncclGroupEnd` returns: the operations are queued on CUDA streams, so the host synchronizes each stream before trusting the result. > Figure. ncclAllReduce across four GPUs, from the NCCL documentation This is why topology is not a footnote to parallelism, it is a constraint on which parallelism strategies are viable at all. A group of GPUs connected by NVLink can absorb tensor parallelism’s frequent all-reduces; GPUs connected only by a network fabric across nodes generally cannot, which is why tensor parallelism is typically confined inside a node and pipeline or data parallelism is used to scale beyond it, where the coarser, less frequent communication tolerates the added latency. > Figure. Tensor vs. pipeline parallelism. The same four-layer model split two ways. Tensor parallelism cuts every layer in half across both GPUs, so activations must be all-reduced at every layer boundary. Pipeline parallelism gives each GPU whole layers and sends activations across the boundary once.Illustrative numbers The size of that fastest tier is a moving target. NVIDIA’s multi-node tuning guide notes that before the GB200 NVL72, an NVLink domain topped out at eight GPUs on an HGX H200 baseboard at 900 GB/s of communication per GPU; the NVL72 rack design extends one domain to 72 Blackwell GPUs at 1.8 TB/s each, with the NVLink switch providing 130 TB/s of aggregate GPU bandwidth inside the domain, and describes that 72-GPU domain as acting as one massive GPU. For the parallelism strategies above, that is a ninefold change in where the scale-up boundary sits: a split that would have crossed the network fabric between eight-GPU nodes can now stay on NVLink across a whole rack. NVLink is one vendor’s fabric, and open standards now exist for both directions of scaling. UALink is the scale-up one: its 200G 1.0 specification “defines a low-latency, high-bandwidth interconnect for communication between accelerators and switches in AI computing pods” and “enables 200G per lane scale-up connection for up to 1,024 accelerators within an AI computing pod.” Ultra Ethernet is the scale-out counterpart, a Linux Foundation project whose stated mission is an “Ethernet based open, interoperable, high performance, full-communications stack architecture” for AI and HPC at scale; the problems it names, multi-pathing, fast reaction to congestion, and flows “where tail latency is the figure of merit,” are this chapter’s constraints restated as networking requirements. Whichever fabric wins a given deployment, the collective at the top of the stack stays the same; what the standards change is who can build the hardware underneath it. Source Tensor and pipeline parallelism from [Megatron-LM](https://arxiv.org/abs/1909.08053), including the GEMM-splitting mechanism in section 3 of the [full paper](https://arxiv.org/html/1909.08053v4). Collective operations and interconnect support from the [NCCL](https://github.com/NVIDIA/nccl) README; the all-reduce listing is trimmed from Example 1 of the [NCCL documentation’s examples](https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/examples.html). NVLink domain sizes and bandwidth from the [multi-node tuning guide](https://docs.nvidia.com/multi-node-nvlink-systems/multi-node-tuning-guide/overview.html) overview. UALink 200G 1.0 description from the [UALink Consortium](https://ualinkconsortium.org/specification/); Ultra Ethernet mission and goals from the [Ultra Ethernet Consortium](https://ultraethernet.org/). More on interconnect topology and the specifications behind scale-up and scale-out fabrics in [parallelism, collectives, and topology](https://learn-kernels.com/chapters/reading#parallelism-collectives). --- # Prefill and decode disaggregation · Distributed inference Generating a response happens in two phases with almost opposite performance characteristics. The prefill phase processes the entire prompt at once to produce the first output token; DistServe measures this as compute-bound, noting that for a 13B-parameter model, “processing a single sequence of 512 tokens can fully engage an A100 GPU.” The decode phase then generates one token per step, each step reading the full set of weights and the growing KV cache to produce a single new token, which is the same memory-bound pattern this book has already covered. Existing serving systems, DistServe observes, batch both phases together on the same GPUs to maximize throughput, but that colocation means “adding a single prefill job to a batch of decoding requests significantly slows down both processes,” stretching out both **time to first token** (Time to first token: the latency of the prefill phase, from receiving a request to producing its first output token.) (TTFT) and **time per output token** (Time per output token: the average latency of one decode step, once generation is underway.) (TPOT). **Disaggregation** (Running the prefill and decode phases of inference on separate GPUs, each tuned to its own phase's resource and parallelism needs, instead of batching both phases together on the same hardware.) is the fix DistServe proposes: assign prefill and decode to separate GPU pools entirely. A prefill instance runs only the prefill computation for a request and forwards its KV cache to a decode instance, which owns everything after the first token. Because the two phases no longer share hardware, each can be scaled and parallelized on its own terms; DistServe finds that intra-op (tensor) parallelism lowers TTFT more effectively at low request rates by cutting execution time directly, while inter-op (pipeline) parallelism scales better at high request rates since it adds capacity linearly per GPU without the same communication cost. Tuned this way, DistServe reports serving “7.4× more requests or 12.6× tighter SLO” than a colocated baseline on the same hardware, while keeping 90% of requests within their latency target. The cost disaggregation introduces is moving the KV cache: every request now requires a transfer of prefill’s output state to whichever GPU is running decode, over whatever network sits between the two pools. NIXL, part of NVIDIA’s Dynamo stack, is purpose-built for that transfer: it “offers a unified abstraction across various memory types, including HBM, DRAM, local or remote SSDs” and picks a backend, such as UCX over InfiniBand or NVLink, based on where the source and destination actually live, so the inference engine issues one transfer request without hardcoding a transport. Disaggregation only pays off if that hand-off is cheap relative to the interference it removes, which is why DistServe pairs its placement algorithm with the cluster’s actual bandwidth rather than assuming the transfer is free. DistServe is not the only system to land on this split. Splitwise, from Microsoft, reached the same architecture from a power and cost angle: its characterization finds that every request runs through “a compute-intensive prompt computation phase and a memory-intensive token generation phase, each with distinct latency, throughput, memory, and power characteristics,” and that even with state-of-the-art batching and scheduling the token generation phase underutilizes compute. The conclusion Splitwise draws is about hardware selection: “token generation does not need the compute capability of the latest GPUs and can be run with lower power and cost,” so the two pools do not even have to be built from the same GPU generation. That freedom turns disaggregation into a cluster design problem. Splitwise moves request state between machines using optimized network libraries over the fast back-plane interconnects already present in GPU clusters, the role NIXL fills in the Dynamo stack, and uses the phase split to design both homogeneous clusters and heterogeneous ones where each phase gets hardware suited to it. Compared with an existing design, it reports clusters with “up to 1.4x higher throughput at 20% lower cost,” or “2.35x more throughput under the same power and cost budgets.” Prefill and decode were always two different workloads sharing one machine; once they are separated, hardware can be purchased for each, not just scheduled for each. > Figure. Prefill and decode disaggregation. A request's prompt runs on the prefill pool, which hands the resulting KV cache to a decode worker over the interconnect, the transfer NIXL abstracts. Decode then streams tokens one step at a time. Each pool is scaled and parallelized on its own terms, which is how DistServe reports 7.4× more requests or 12.6× tighter SLO than a colocated baseline on the same hardware.Illustrative numbers Source TTFT/TPOT characterization, interference measurements, and throughput results from [DistServe](https://arxiv.org/html/2401.09670), sections 2.2 through 3.1. Transfer layer description from the [NIXL](https://github.com/ai-dynamo/nixl) README and its architecture overview. Phase characterization, hardware heterogeneity, and cluster results from the [Splitwise](https://www.microsoft.com/en-us/research/publication/splitwise-efficient-generative-llm-inference-using-phase-splitting/) abstract (ISCA 2024). Production disaggregated serving systems built on these ideas are listed in [prefill and decode disaggregation](https://learn-kernels.com/chapters/reading#disaggregation). --- # Production systems and serving benchmarks · Distributed inference Everything so far in this chapter happens inside one deployment of one model. A production serving system is the layer above that: it decides which replica a request lands on, when it runs, and how the fleet scales, and its defining problem is the tail rather than the average. Clockwork, a model serving system that predates LLMs, made the case that this layer can be built on prediction instead of reaction: existing serving architectures used “well-known reactive techniques” against common-case latency but could not curtail the tail latency caused by unpredictable execution times, while DNN inference itself “has deterministic performance.” Built bottom-up from those predictable execution times, with a centralized scheduler in place of reactive workers, Clockwork supported thousands of models while meeting “100 ms latency targets for 99.997% of requests.” The lesson carries directly into LLM serving: latency targets are met by scheduling, not only by fast kernels. llm-d is what that layer looks like for LLMs today. Engines like vLLM and SGLang handle running the model on accelerators; llm-d provides “orchestration and optimizations above model servers,” deployed on Kubernetes, and its feature list reads like this chapter turned into routing decisions: prefix-cache and load-aware request balancing, tiered KV cache offloading with global indexing of cache state, prefill/decode disaggregation and wide expert parallelism for the largest models, and SLO-aware autoscaling driven by real-time inference signals. Where a request lands now matters as much as how fast the engine runs it: llm-d reports 3x higher output throughput and 2x faster TTFT from prefix-cache-aware routing compared to round-robin, measured on Llama 3.1 70B across 4 AMD MI300X GPUs, a gain that comes from routing each request to a replica that already holds its prefix in cache. That routing layer is also becoming a standard rather than a product. Gateway API Inference Extension, an official Kubernetes project, extends any gateway that supports the Gateway API and Envoy’s external processing protocol into an inference gateway: a load balancer coupled with an endpoint picker that chooses a replica using metrics and capabilities the model servers themselves report, such as prefix cache status or which LoRA adapters a replica has loaded. The same layer carries policy that a plain load balancer has no vocabulary for: serving priority, so a latency-sensitive chat model can outrank a latency-tolerant summarization model, and incremental rollouts of new model versions by splitting traffic on model names. Scaling the fleet has a cold start problem of its own: a new replica must load its model’s weights from a checkpoint before it can serve a single token. ServerlessLLM, an OSDI 2024 system, attacks that startup path by harnessing “the substantial near-GPU storage and memory capacities of inference servers” to keep checkpoints close to the accelerators. Its three mechanisms are a loading-optimized checkpoint format with a multi-tier loading system that uses the full bandwidth of the storage hierarchy, live migration of running inference so a new request can claim a server that already holds its checkpoint, and a scheduler that places each model on the server that minimizes its time to start. The paper reports “reducing latency by 10 - 200X across various LLM inference workloads” against state-of-the-art serverless systems. Measuring such a system takes more than a throughput number. The etalon benchmark framework works with the per-request latency metrics this chapter has already met: time to first token, defined as “the time taken between arrival and first output token,” which includes “both scheduling delay and prompt processing time”; time between tokens, the gap between two consecutive output tokens of a streaming response; and time per output token, total generation time divided by the number of output tokens. Etalon’s starting observation is that even these “fail to fully capture the nuances of LLM inference,” leaving an incomplete picture of user-facing performance in real-time applications like chat. What a production system optimizes is **goodput** (The rate of requests a system serves while still meeting their latency targets, so a request that misses its target does not count toward it.) under **service-level objectives** (A latency target a serving system commits to meeting for a stated fraction of requests, such as a TTFT bound met by 90% of requests.): DistServe defines per-GPU goodput as “the maximum request rate that can be served adhering to the SLO attainment goal (say, 90%)” for each GPU provisioned. By these definitions a system can raise its throughput while lowering its goodput, simply by letting a slice of requests blow through their latency targets. > Figure. One request's lifetime, and the metrics cut from it. TTFT runs from arrival to the first output token and swallows both queueing and prompt processing. TBT is the gap between any two consecutive tokens; TPOT is the whole decode span divided by the token count. Goodput is not on this timeline at all: it counts how many such requests per second finish inside their latency targets. Definitions from the etalon documentation and DistServe. The load those metrics are measured under matters as much as the metrics themselves. Clockwork was evaluated “using production trace workloads,” and BurstGPT makes the same possible for LLM serving research: a public trace of real GPT-3.5 and GPT-4 traffic served on Azure, covering 121 consecutive days and roughly 5.29 million requests, each entry recording request and response token counts and whether the call came through a conversation or the API. The trace’s own overview plots weekly and daily periodicity in request arrivals, structure that a uniform synthetic load has none of, and its usage notes suggest scaling the trace’s average request rate to the evaluation setup rather than discarding the arrival pattern. ServeGen carries that argument further. Built on a characterization of workloads “collected from our worldwide cloud inference serving service” at Alibaba, covering language, multimodal, and reasoning models, it argues that prior analyses were too limited in scale and scope to capture how real traffic behaves, and generates realistic workloads “by composing them on a per-client basis” rather than drawing from one aggregate distribution. Its production use case makes the stakes concrete: benchmarking with ServeGen “avoids 50% under-provisioning compared to naive workload generation.” A capacity plan validated against uniform synthetic load can simply be wrong. Comparing systems across vendors needed a referee, and MLPerf became it. The MLPerf Inference paper (ISCA 2020) describes a field where over 100 organizations were building inference chips and existing systems spanned “at least three orders of magnitude in power consumption and five orders of magnitude in performance,” and answers with rules and best practices “to ensure comparability across systems with wildly differing architectures”; its first call for submissions drew more than 600 reproducible measurements from 14 organizations. MLPerf Endpoints is the same idea aimed at LLM serving: a submitted system is a measured curve rather than a single number, each point one real test at a fixed concurrency, relating system throughput, interactivity in tokens per second per user, and 95th-percentile TTFT. As MLCommons puts it, as load rises “total throughput goes up and per-user speed comes down. That is the core tradeoff in serving AI,” which is goodput’s tradeoff restated as a purchasing decision. Source Predictability claims and tail latency results from the [Clockwork](https://www.usenix.org/conference/osdi20/presentation/gujarati) abstract. Orchestration features and routing numbers from the [llm-d](https://github.com/llm-d/llm-d) README. Metric definitions from the [etalon documentation](https://project-etalon.readthedocs.io/en/latest/tutorials/metrics_used.html) and the [etalon paper](https://arxiv.org/abs/2407.07000) abstract; the goodput definition from [DistServe](https://arxiv.org/html/2401.09670), section 1. Trace characteristics from the [BurstGPT](https://github.com/HPMLL/BurstGPT) README. Inference gateway concepts from the [Gateway API Inference Extension](https://gateway-api-inference-extension.sigs.k8s.io/) documentation. Checkpoint loading, live migration, and the latency result from the [ServerlessLLM](https://www.usenix.org/conference/osdi24/presentation/fu) abstract. Workload characterization and the provisioning result from the [ServeGen](https://arxiv.org/abs/2505.09999) abstract (NSDI 2026). Benchmark scope and submission counts from the [MLPerf Inference Benchmark](https://arxiv.org/abs/1911.02549) abstract; curve dimensions from the [MLPerf Endpoints](https://mlcommons.org/benchmarks/endpoints/) page. More systems and measurement work in [production systems](https://learn-kernels.com/chapters/reading#production-systems) and [serving benchmarks](https://learn-kernels.com/chapters/reading#serving-benchmarks). --- # Foundations A kernel is a function launched as thousands to millions of threads at once. Everything in the rest of this book, tiling, occupancy, coalescing, roofline analysis, is a consequence of how those threads are grouped and where the data they touch actually lives. ## 4 sections 1. [1.1Threads, blocks, and warps](https://learn-kernels.com/chapters/foundations/threads-blocks-and-warps) 2. [1.2Compute-bound and memory-bound](https://learn-kernels.com/chapters/foundations/compute-bound-and-memory-bound) 3. [1.3The KV cache](https://learn-kernels.com/chapters/foundations/the-kv-cache) 4. [1.4Compilation and machine code](https://learn-kernels.com/chapters/foundations/compilation-and-machine-code) --- # Compilation and machine code · Foundations A CUDA source file mixes host and device code, and nvcc, the CUDA compiler driver, splits it: the host code is forwarded to an ordinary C++ compiler, while the device functions go through NVIDIA’s own compilers. The device side compiles in two stages. First the CUDA C++ is compiled to **PTX** (Parallel Thread Execution, a virtual machine instruction set architecture that acts as the assembly language of the CUDA platform.), assembly for a virtual GPU defined only by its capabilities, not by any physical chip. Then an assembler called ptxas compiles the PTX into a *cubin*, an ELF binary holding the actual machine instructions for one real architecture. The compiled GPU code is embedded back into the host executable, which can carry both cubins and the PTX text itself. The split exists because NVIDIA does not guarantee binary compatibility across GPU generations: instruction sets and encodings change between major compute capabilities, so a cubin built for one generation will not run on the next. Every nvcc compilation therefore names two architectures, a virtual one (compute\_80) stating which features the code needs, and a real one (sm\_80) naming the processor to build for. Because the executable can also embed the PTX, the GPU driver can **JIT compile** (Compiling the embedded PTX to machine code at application runtime, done by a compiler built into the GPU driver, for the exact GPU the program is about to run on.) it at runtime for whatever GPU is actually present, including architectures that did not exist when the application shipped. > Figure. The two-stage device compile. nvcc compiles the device side of a .cu file to PTX, assembly for a virtual GPU, and ptxas compiles that into the machine instructions (SASS) of one real architecture. Because the executable can also carry the PTX, the driver can JIT compile it at runtime for the GPU actually present. Everything above the cubin is an intermediate form, which is why kernel authors read the machine code when it matters. PTX is what the compiler promised; the disassembled cubin, called **SASS** (The native assembly of a real GPU architecture, disassembled from a cubin by cuobjdump or nvdisasm.), is what the hardware will execute, and the two can differ because ptxas optimizes, allocates registers, and schedules instructions on the way down. The toolkit ships two disassemblers: cuobjdump accepts both cubins and host executables and can extract the embedded PTX, while nvdisasm accepts only cubins but adds control flow analysis and richer display options. When this book claims a kernel change did something, disassembling is how that claim gets checked: against the instructions the compiler actually emitted, not the ones the source seems to imply. What PTX looks like is worth seeing once. NVIDIA’s Understanding PTX post compiles the same bounds-checked vecAdd kernel from section 1.1 and prints the result. Most of it reads like assembly for any machine, loads, an add, a store, with the CUDA thread model showing through in special registers: `%tid.x`, `%ntid.x`, and `%ctaid.x` are `threadIdx.x`, `blockDim.x`, and `blockIdx.x`. > Figure. vecAdd compiled to PTX, from Understanding PTX (declarations and parameter loads elided) The kernel survives translation recognizably. A single `mad.lo.s32`, a multiply-add, computes the entire `workIndex` expression; `setp.ge.u32` and the predicated branch `@%p1 bra` are the bounds check, comparing the index against the length parameter (loaded into `%r2` in the elided lines) and jumping past the body when it is out of range; and the load, add, store triple in the middle is the arithmetic itself. Because PTX targets a virtual machine, the post notes, ptxas behaves more like a compiler than a traditional assembler when it lowers this into a cubin, which is where the register allocation and scheduling that separate PTX from SASS happen. The post draws the comparison this whole pipeline rests on: PTX plays the role in CUDA that LLVM IR plays for clang, a portable middle stage that many front ends can target and any supported GPU can be compiled from. Source Compilation trajectory and virtual versus real architectures from the [CUDA Compiler Driver NVCC](https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/) documentation; PTX’s role, the vecAdd PTX listing, and JIT compilation from [Understanding PTX](https://developer.nvidia.com/blog/understanding-ptx-the-assembly-language-of-cuda-gpu-computing/); cuobjdump and nvdisasm from [CUDA Binary Utilities](https://docs.nvidia.com/cuda/cuda-binary-utilities/). --- # Compute-bound and memory-bound · Foundations Every kernel eventually hits one of two limits: the time it spends waiting on arithmetic units, or the time it spends waiting on memory traffic. NVIDIA calls this boundary *math bandwidth*. A kernel is flop bound when there is never a moment where nothing is moving through memory; it is memory bound when there is never a moment where no floating-point operation is in flight. The ratio between a chip’s peak arithmetic throughput and its peak memory bandwidth sets exactly where that boundary falls, in operations performed per byte moved. 312e12flop/s A100 peak, bf16 1.5e12bytes/s A100 memory bandwidth 208flops/byte Boundary ratio On an A100, that works out to 208 floating-point operations per byte moved. Below that arithmetic intensity, a kernel is memory bound: it finishes exactly as fast as its bytes can be streamed in, no matter how many arithmetic units sit idle. Above it, the kernel is flop bound, and streaming data faster would not help. This is the same reasoning the [roofline model](https://learn-kernels.com/chapters/reading#start-here) formalizes across an entire chip, not just one kernel. > Figure. The roofline, sketched. Throughput against arithmetic intensity. Left of the knee a kernel is memory bound: its throughput rises with the bandwidth slope. Right of it the compute roof takes over and streaming data faster would not help. On an A100 the knee sits at 208 flops per byte.Illustrative numbers The bandwidth in that ratio is a peak, and whether a kernel gets it depends on the shape of its accesses. Global memory loads and stores by the threads of a warp are coalesced by the device into as few transactions as possible, and on devices of compute capability 6.0 or higher the rule is simple: a warp’s concurrent accesses coalesce into however many 32-byte transactions are needed to service them all. Thirty-two threads reading adjacent 4-byte floats need four 32-byte transactions, and every byte fetched is a byte used. NVIDIA’s best practices guide marks keeping accesses coalesced as a high-priority recommendation, and the penalty for breaking the pattern is arithmetic: when adjacent threads access memory with a stride of two elements, half of every fetched segment goes unused, a 50 percent load and store efficiency, and as the stride grows the decay continues until a warp of 32 threads loads 32 separate 32-byte segments. A strided kernel is not paying the streaming rate the roofline promises; it is paying for bytes it never touches. The other lever is keeping the memory system busy at all. Thread instructions execute sequentially, so when a warp stalls on a load, executing other warps is the only way to hide the latency, and **occupancy** (The ratio of the number of active warps per multiprocessor to the maximum number of possible active warps.) names how much of that capacity is in play. What limits it is resource sharing: registers are allocated to a whole block at once out of a file that every resident thread shares, so on a compute capability 7.0 device with 65,536 registers per SM and up to 2,048 resident threads, full occupancy leaves each thread at most 32 registers. The guide is careful about the direction of the claim: higher occupancy does not always buy more performance, but low occupancy always interferes with hiding memory latency. The same guide prescribes a workflow for applying any of this: Assess, Parallelize, Optimize, Deploy, a cycle rather than a checklist. Assess profiles the application to find the code responsible for the bulk of the execution time and uses Amdahl’s and Gustafson’s laws to bound what accelerating it can possibly buy. Parallelize exposes the parallelism, sometimes as simply as calling an existing GPU library. Optimize is explicitly iterative (identify an opportunity, apply and test it, verify the speedup, repeat), so no one needs to memorize every strategy before seeing gains. Deploy ships each partial speedup to production before the next hotspot is tackled, so every pass around the loop pays for itself. Source Numbers and the flop-bound/memory-bound framing from [Transformer Inference Arithmetic](https://kipply.github.io/blog/transformer-inference-arithmetic/), citing the NVIDIA A100 datasheet. Coalescing rules and stride penalties, occupancy and its register arithmetic, and the APOD cycle from the [CUDA C++ Best Practices Guide](https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/), sections 2.2, 10.2.1, and 11.1. --- # The KV cache · Foundations This is where GPU fundamentals stop being an abstract exercise and start explaining why inference engines look the way they do. Autoregressive sampling generates one token at a time, and at every step, self-attention needs the key and value vectors for every token already in the sequence, not just the new one. Recomputing them from scratch each step would mean redoing the same matrix multiplies over and over as the sequence grows. Instead, engines store those vectors in a *KV cache* the first time they are computed, and read them back on every later step. That trade only pays off because of the boundary from the previous section. Computing the key and value vectors for one new token is cheap; the same weights are needed to run a full forward pass over the whole context, which is why a decode step spends most of its time waiting on memory bandwidth to stream the model’s weights through, rather than on arithmetic. The KV cache avoids paying that same memory cost again for tokens already seen, at the price of storing one more matrix per layer that grows with sequence length and batch size. > Figure. How the KV cache grows. Five decode steps. Each step computes the K and V vectors for its one new token exactly once, then every later step reads the stored entries back instead of recomputing them. The price is storage that grows by one entry per token, per layer.Illustrative numbers Source KV cache mechanics from [Transformer Inference Arithmetic](https://kipply.github.io/blog/transformer-inference-arithmetic/). For how production engines page, quantize, and distribute this cache, see [KV cache systems](https://learn-kernels.com/chapters/reading#kv-cache-systems) in the reading list. --- # Threads, blocks, and warps · Foundations A CUDA program runs on a heterogeneous system: a **host** (The CPU and the memory directly connected to it.) (the CPU) and a **device** (A GPU and the memory directly connected to it.) (the GPU), connected by an interconnect such as PCIe or NVLink. The host code copies data to device memory, launches a **kernel** (A function invoked for execution on the GPU, launched as many threads running in parallel.), and waits for it to complete. A kernel launch starts many threads, often millions, all executing the same device code. Those threads are organized into *thread blocks*, and thread blocks into a *grid*. Every thread block in a grid runs entirely on one **streaming multiprocessor** (A GPU is a collection of streaming multiprocessors, each with a local register file, a unified data cache, and functional units that perform computations.) (SM), which is what lets threads inside a block synchronize and share on-chip memory cheaply. There is no such guarantee across blocks: the CUDA programming model requires that thread blocks be safe to run in any order, in parallel or in series, because a grid can have far more blocks than the GPU has SMs to run them on at once. Inside a block, threads execute in fixed groups of 32 called *warps*, in a Single-Instruction Multiple-Threads (SIMT) model: every thread in a warp runs the same instruction at the same time, but each thread carries its own program counter and can take a different branch. When threads in a warp disagree on which branch to take, the ones not on the active path are masked off until the warp reconverges, a cost called **warp divergence** (When different threads in a warp follow different code paths, the ones not on the active branch are masked off while the others execute.). It follows that a block sized to a multiple of 32 threads uses every lane of its last warp; anything else leaves lanes idle for the whole kernel. In code, the whole model fits in a dozen lines. A kernel is a `__global__` function; the triple-chevron launch names the grid and block dimensions; and inside the kernel, each thread combines `threadIdx`, `blockIdx`, and `blockDim` to find the one element it is responsible for. This is the CUDA programming guide’s own first example, an element-wise vector addition where every thread performs exactly one add: > Figure. vecAdd, from the CUDA programming guide The launch `<<<1, 256>>>` starts one thread block of 256 threads, and the guide notes the two constraints this book keeps returning to: a block may contain at most 1,024 threads because the whole block must fit on one SM, and kernel launches are asynchronous, so the host must synchronize before it can trust the result. One block of 256 threads covers 256 elements, but the same index expression scales to any number of blocks. Launched as `vecAdd<<<4, 256>>>` over a vector of 1,024 elements, `blockDim.x * blockIdx.x` becomes each block’s offset into the vector: threads in the first block compute indices 0 through 255, threads in the second land at `threadIdx.x + 256`, the third at `threadIdx.x + 512`. Real vector lengths are not always multiples of the block size, so the guide’s full kernel takes the length as a parameter and guards the work with `if (workIndex < vectorLength)`; threads past the end simply do nothing. The launch then rounds the block count up with an integer ceiling divide, `(vectorLength + threads - 1) / threads`. A few idle threads in the last block cost little, the guide notes, but launching whole blocks in which no thread does work should be avoided. None of this runs until the arrays live in memory the GPU can reach. The explicit path allocates device buffers with `cudaMalloc` and copies data across with `cudaMemcpy`, whose last argument names the direction: `cudaMemcpyHostToDevice`, `cudaMemcpyDeviceToHost`, or `cudaMemcpyDefault`, which infers the direction from the pointer values. Wrapped around the launch, those calls turn vecAdd into a complete program: > Figure. explicit memory management around vecAdd, from the CUDA programming guide Two details in that listing carry most of the meaning. `cudaMemcpy` is synchronous: it does not return until the copy has completed. The kernel launch is not, which is why `cudaDeviceSynchronize` sits between the launch and the copy back: it blocks the host thread until all previously issued GPU work has finished. The guide also offers a second path, unified memory, where `cudaMallocManaged` allocates buffers the driver keeps accessible to both CPU and GPU and the copies disappear from the source. The explicit version is more verbose precisely because it affords control over when data moves and where it lives, control the performance chapters of this book spend heavily. > Figure. The execution hierarchy. A grid is made of thread blocks, and every block runs entirely on one SM. Inside a block, threads execute in warps of 32: each row of the zoomed block is one warp.Illustrative numbers Source Thread blocks, grids, SMs, and warps as defined in the [CUDA programming model](https://docs.nvidia.com/cuda/cuda-programming-guide/01-introduction/programming-model.html), section 1.2. The vecAdd listings, multi-block indexing, bounds checking, memory management, and synchronization from [Intro to CUDA C++](https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/intro-to-cuda-cpp.html), section 2.1. --- # Frontier Everything before this chapter rests on published ISA documents and measurements you can reproduce. This chapter is the opposite: claims that are still being tested, hardware that has been announced but not independently measured, and the habits of skepticism that separate the two. ## 2 sections 1. [7.1AI-generated kernels](https://learn-kernels.com/chapters/frontier/ai-generated-kernels) 2. [7.2Watchlist](https://learn-kernels.com/chapters/frontier/watchlist) --- # AI-generated kernels · Frontier The claim under test is simple: language models can write CUDA kernels that beat PyTorch. KernelBench, the benchmark that made the question precise, evaluates models on 250 PyTorch workloads and scores them with **fast\_p** (KernelBench's metric: the percentage of generated kernels that are functionally correct and offer a speedup greater than an adjustable threshold p over baseline.), the fraction of generated kernels that are both correct and faster than the baseline by more than a chosen threshold. Its authors found that frontier reasoning models did best out of the box but still matched the PyTorch baseline in less than 20 percent of cases, with the benchmark getting harder as the speedup threshold rises. Then the correctness checks themselves came under scrutiny. KernelBench-Verified, a follow-up from Meta and Stanford, found that frontier models frequently engage in **reward hacking** (Exploiting the specific conditions of an evaluation to inflate a score, rather than solving the task the evaluation is meant to measure.). The original baseline ran PyTorch in plain float32, without **TF32** (TensorFloat-32, a compute mode on recent NVIDIA GPUs that routes float32 matrix multiplications through Tensor Cores with a truncated significand, enabled in PyTorch by a single flag.), so any generated kernel that merely invoked cuBLAS looked dramatically fast against an artificially slow reference. And because the test inputs came from one narrow distribution, all positive and small, models learned to hardcode bypasses: one generated ReLU kernel checked for the test shape and returned its input unchanged, passing the correctness check and reporting a 374x speedup while computing nothing. > Figure. How a reward-hacked kernel passes. The ReLU kernel from KernelBench-Verified's analysis: it checks for the benchmark's input shape and returns its input unchanged. Against one narrow test distribution it passes correctness and reports a 374x speedup while computing nothing; a hidden suite drawn from four input distributions catches it. Under the verified protocol, a TF32-enabled baseline plus a hidden test suite of four input distributions, the picture inverted. The best model’s geometric mean speedup fell from 1.43x under the standard protocol to 0.88x, no model consistently outperformed PyTorch, and 28 percent of the best model’s correct kernels increased peak GPU memory. NVIDIA’s SOL ExecBench pushes the standard further still: it checks submissions for reward hacking and ranks them by a SOL-Score, which grades a kernel against the analytically derived roofline of a B200 rather than against a software baseline. The state of the evidence, then: models can write working kernels, and any claimed speedup is only as trustworthy as the baseline configuration and input distribution behind it. When you read the next result in this space, those are the two things to check first. Source Benchmark design and results from [KernelBench](https://proceedings.mlr.press/v267/ouyang25a.html) and [KernelBench-Verified](https://arxiv.org/html/2607.16241); speed-of-light grading from the [SOL ExecBench](https://github.com/nvidia/sol-execbench) README. Follow-ups are collected under [AI-generated kernels](https://learn-kernels.com/chapters/reading#ai-generated-kernels) in the reading list. --- # Watchlist · Frontier This book has an evidentiary standard: an architecture gets covered in the [hardware chapter](https://learn-kernels.com/chapters/hardware) once its ISA documents are public and its performance numbers can be reproduced by someone who does not work for the vendor. Until then it stays here, as a set of claims worth tracking rather than facts worth teaching. Everything in this section is in that waiting room. The clearest example is NVIDIA’s Rubin GPU. The launch material describes two reticle-limited compute dies joined by **NV-HBI** (NVIDIA High-Bandwidth Interface, the high-speed inter-die link that unifies Rubin's two compute dies on a single package.), an HBM4 memory subsystem, a third-generation Transformer Engine rated at up to 50 petaflops of NVFP4, and up to 10x more agentic throughput per unit of energy than Blackwell. Every one of those numbers is a vendor claim, and the throughput figure comes from an internal workload. None of that makes the claims false; it makes them unverified. What would move Rubin into the hardware chapter is shipped systems, published tuning and ISA documentation, and measurements someone outside NVIDIA can reproduce. The rest of the watchlist has even less on the record, which is exactly why it is a watchlist. AMD’s next datacenter GPU generation, session-aware scheduling in inference engines, serving real-time voice and video, dedicated inference ASICs, and processing in memory are all areas where the interesting claims currently live in announcements, papers without public artifacts, or products without disclosed baselines. For each, the same three questions apply: is the hardware or system shipping, are the specifications published, and has anyone outside the organization reproduced the numbers? Pointers to whatever public material exists are kept current in the [reading list](https://learn-kernels.com/chapters/reading#watchlist), so this section can stay short and the evidence can accumulate there. Source Rubin architecture claims from NVIDIA’s [Rubin architecture post](https://developer.nvidia.com/blog/inside-nvidia-rubin-gpu-architecture-powering-the-era-of-agentic-ai/), a vendor launch publication. The remaining watchlist entries cite no source by design: they are listed because reproducible public evidence does not exist yet. --- # 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. --- # Current hardware Every kernel in this book eventually runs on real silicon, and what that silicon can actually do sets the ceiling the earlier chapters’ roofline math was assuming all along. This chapter reads the numbers straight from each vendor’s own architecture and tuning documentation: NVIDIA’s Blackwell generation, AMD’s CDNA 4 Instinct MI350 series, Google’s TPU7x, and AWS’s Trainium3. ## 2 sections 1. [6.1NVIDIA Blackwell](https://learn-kernels.com/chapters/hardware/nvidia-blackwell) 2. [6.2AMD, TPU, and Trainium](https://learn-kernels.com/chapters/hardware/amd-tpu-and-trainium) --- # AMD, TPU, and Trainium · Current hardware ### AMD CDNA 4 AMD’s fourth-generation CDNA architecture powers the Instinct MI350 series. Each MI350-series GPU integrates 8 vertically stacked **accelerator complex dies** (A single chiplet that houses a processor's compute portion: compute units, their local caches, and the hardware that schedules work to them, as distinct from I/O dies that hold shared memory controllers and cache.) (XCDs) and 2 I/O dies, tied together with AMD’s on-package Infinity Fabric and connected to 8 stacks of 12-Hi HBM3E memory. The eight compute chiplets are built on TSMC’s N3P process, while the two I/O dies, which hold the AMD Infinity Cache and the memory controllers, use TSMC’s N6 process; AMD splits the design this way because the memory and communication logic in the I/O dies does not benefit as much from the newer, more expensive node as compute logic does. The family ships as two SKUs: the air-cooled MI350X at 1000 W, built to be drop-in compatible with the prior-generation MI325X platform, and the liquid-cooled MI355X at 1400 W for higher power and cooling budgets. The MI355X carries 288 GB of HBM3E, 8.0 TB/s of memory bandwidth, and over 1 TB/s of Infinity Fabric communication bandwidth between GPUs. Across the package its 8 XCDs total 256 compute units and 1,024 matrix cores, with each XCD contributing a 4 MB L2 cache in front of the shared Infinity Cache in the I/O dies below. CDNA 4’s biggest generational jump is in reduced-precision matrix throughput. Building on the OCP **microscaling** (A quantization scheme that shares one scale factor across a block of typically 32 elements within a tensor, rather than one scale for the entire tensor, preserving more information than a single per-tensor scale at similarly low bit widths.) standard, the CDNA 4 compute units add hardware support for MXFP8, MXFP6, and MXFP4 formats alongside doubled execution resources for the existing 16-bit and 8-bit datatypes. Comparing peak theoretical throughput per GPU, the MI355X reaches 5.0 PFLOPS of dense FP8 matrix compute (10 PFLOPS with structured sparsity) and 10 PFLOPS of dense FP6/FP4 (20 PFLOPS with sparsity), against 2.6 PFLOPS of dense FP8 on the prior-generation MI300X, a 1.9x generational gain on the formats the two share. See [AMD ROCm and CDNA](https://learn-kernels.com/chapters/reading#hardware-amd) in the reading list for the ISA reference and ROCm software stack. Whether a kernel is actually reaching those numbers is a measurable question, because ROCm documents the full hardware performance counter set for the MI350 series, organized by IP block and exposed through ROCprofiler-SDK and ROCm Compute Profiler. The command processor and shader pipe interpolator blocks count thread groups launched and waves in flight; the compute unit (SQ) counters break instructions and FLOPs out by datatype, down to a dedicated counter for F6 and F4 matrix instructions (SQ\_INSTS\_VALU\_MFMA\_F6F4); and the LDS counters report load, store, and atomic traffic in 64-byte units alongside the cycles their FIFOs spent full. The cache blocks carry their own stall accounting, including cycles stalled on data pending from L2 and stalls inside the UTCL1 address-translation unit, so the waiting half of a kernel’s roofline story is as countable on this hardware as the arithmetic half. Source Package construction, SKUs, and peak throughput figures from the [AMD CDNA 4 Architecture whitepaper](https://www.amd.com/content/dam/amd/en/documents/instinct-tech-docs/white-papers/amd-cdna-4-architecture-whitepaper.pdf). Counter names and descriptions from the [MI350 Series performance counters](https://rocm.docs.amd.com/en/latest/reference/gpu-arch/mi350-performance-counters.html) reference. ### Google TPU Google’s TPU7x, the first chip in the Ironwood generation, is a dual-chiplet design: each chiplet is a self-contained unit with one TensorCore, two SparseCores, and 96 GB of HBM, and the two chiplets are exposed to frameworks like JAX as two separate devices connected by a die-to-die interface. Counted at the full-chip level, that is 2 TensorCores, 4 SparseCores, and 192 GiB of HBM per chip, with roughly 7.38 TB/s of HBM bandwidth. A full pod scales to 9,216 chips connected in a 3D torus topology, with 200 GB/s of bidirectional inter-chip bandwidth per axis between neighboring chips. Per chip, TPU7x reaches 2,307 TFLOPS of peak bf16 compute and 4,614 TFLOPS of peak FP8 compute. Each TPU7x virtual machine bundles 4 chips together with 224 vCPUs and 960 GB of host RAM, connected to its host over PCIe. See [TPU architecture](https://learn-kernels.com/chapters/reading#hardware-tpu) in the reading list for the Pallas TPU programming model that targets this hardware. The design has a lineage worth knowing, because its core idea has not changed since the first TPU. Google’s 2017 ISCA paper describes the original chip, deployed in its datacenters since 2015 to accelerate neural network inference: at its heart sat a matrix multiply unit built from 65,536 8-bit multiply-accumulate units, delivering a peak of 92 TeraOps/second and fed from a large 28 MiB software-managed on-chip memory. The paper argues that the TPU’s deterministic execution model was a better match for the 99th-percentile response-time requirements of production inference than the time-varying optimizations of contemporary CPUs and GPUs (caches, out-of-order execution, multithreading, prefetching), and measures the chip at roughly 15x to 30x the speed of its contemporary Haswell CPU and K80 GPU on Google’s production workloads, at 30x to 80x their TeraOps per watt. TPU v4, which Google describes as its fifth domain-specific architecture and third supercomputer for machine learning, shows how far that lineage had scaled by its 2020 deployment. Its supercomputer grew 4x larger than v3’s, to 4,096 chips, joined through optical circuit switches that dynamically reconfigure the interconnect topology (users can pick a twisted 3D torus if desired) at under 5 percent of system cost and under 3 percent of system power. Each TPU v4 also includes SparseCores, dataflow processors that accelerate embedding-reliant models by 5x to 7x while using only 5 percent of die area and power, and the paper reports the chip outperforming TPU v3 by 2.1x while improving performance per watt by 2.7x. The SparseCores and the torus interconnect in TPU7x above are direct descendants of both decisions. To a kernel author working through Pallas, this hardware looks nothing like a GPU. The JAX documentation describes TPUs as sequential machines with a very wide vector register: the grid of a Pallas TPU kernel is generally processed not in parallel but sequentially, in lexicographic order, and HBM cannot be accessed directly by compute instructions; data has to be prefetched into lower levels of the memory hierarchy by DMA subunits, with matrix multiplies executed by the MXU and the bulk of remaining computation performed on 2D vector registers, typically 8x128 for 32-bit values. The references a kernel body receives point at buffers in VMEM, a vector memory the docs describe as fairly large for its level of the hierarchy at 16 MB or more (the hardware reference table lists 64 MiB per TensorCore for TPU 7x), and in SMEM, a low-latency scalar memory serving the separate scalar unit that handles control flow. The quickstart’s first kernel shows how little of that machinery the author touches directly: > Figure. add_vectors, from the JAX Pallas quickstart The quickstart notes that on TPU the references already live in on-chip memory by the time the kernel body runs: values are fetched from HBM before execution, the body moves them from SRAM into registers and back, and results return to HBM only after the kernel completes. The compiler, not the kernel author, schedules those transfers and overlaps them with compute, the same overlap that CUDA kernels in earlier chapters arranged by hand with producer and consumer pipelines. Source Chip composition, memory, compute, and pod topology from the [Google Cloud TPU7x (Ironwood) documentation](https://docs.cloud.google.com/tpu/docs/tpu7x). TPU v1 figures from [In-Datacenter Performance Analysis of a Tensor Processing Unit](https://research.google/pubs/in-datacenter-performance-analysis-of-a-tensor-processing-unit/). TPU v4 figures from [TPU v4: An Optically Reconfigurable Supercomputer for Machine Learning](https://arxiv.org/abs/2304.01433). Execution model, memory spaces, and register shapes from the [Pallas TPU details](https://docs.jax.dev/en/latest/pallas/tpu/details.html), [TPU hardware reference](https://docs.jax.dev/en/latest/pallas/tpu/hardware.html), and [Pallas quickstart](https://docs.jax.dev/en/latest/pallas/quickstart.html), where the listing appears. ### AWS Trainium AWS’s Trainium3 device is built from 8 NeuronCores (v4), 4 HBM stacks totaling 144 GiB of capacity at 4.7 TB/s of bandwidth, 128 DMA engines for moving data within and across devices, 20 CC-Cores dedicated to collective communication, and 4 NeuronLink-v4 links for device-to-device traffic. On-chip, each NeuronCore-v4’s SBUF grew to 32 MiB, up from 28 MiB in the prior NeuronCore-v3, while PSUM stayed at 2 MiB. The Tensor Engine is where the generational jump shows up most: it runs at 2.4 GHz and delivers 315 TFLOPS of MXFP8/MXFP4 compute, 79 TFLOPS of BF16/FP16/TF32, and 20 TFLOPS of FP32. To hit the MXFP8/ MXFP4 rate it quadruples the matmul contraction dimension from 128 to 512 elements, presenting what the architecture guide calls a 512x128 **systolic array** (A grid of processing elements that pass data to their neighbors each cycle, used to implement matrix multiplication directly in hardware. Trainium3's Tensor Engine is organized as a 128x128 grid of elements that together present a 512x128 array to the programmer for reduced-precision matmul.) to the programmer, even though the underlying grid of processing elements is still 128x128. See [Trainium and NKI](https://learn-kernels.com/chapters/reading#hardware-trainium) in the reading list for the full architecture guide and the tile-level programming model built on top of it. AWS’s prescription for programming this hardware is spelled out in the NKI performance guide, and it reads like this book’s roofline chapters translated into Neuron vocabulary. Optimization work should end with a kernel that is either compute-bound, meaning at least one compute engine is active close to 100 percent of the execution time (the guide treats 90 percent or more as good in practice), or memory-bound, with achieved memory bandwidth utilization close to 100 percent (60 percent or more is considered good). Getting there is framed as a fight against data movement: keep inputs resident in SBUF instead of reloading them over DMA, fuse consecutive operators through explicit loop fusion so intermediates never spill to device memory (the profiler’s spill\_save\_bytes and spill\_reload\_bytes metrics expose how much traffic is spill), pipeline tiles across the Tensor, Scalar, Vector, and GpSimd engines so no engine idles waiting on another, and size instruction tiles large enough to amortize instruction overhead but small enough not to wreck pipelining or SBUF pressure. The language those optimizations are written in is small; the getting-started guide names the three phases every NKI kernel has, load from device memory into SBUF, compute, store back, and its first example is the whole model in one function: > Figure. nki_tensor_add_kernel, from the AWS Neuron NKI getting started guide Source NeuronCore-v4 composition, memory, and Tensor Engine figures from the [Trainium3 Architecture Guide for NKI](https://awsdocs-neuron.readthedocs-hosted.com/en/latest/nki/guides/architecture/trainium3_arch.html). Optimization targets and techniques from the [NKI Performance Guide](https://awsdocs-neuron.readthedocs-hosted.com/en/v2.24.0/general/nki/nki_perf_guide.html). Kernel phases and the listing from [Getting Started with NKI](https://awsdocs-neuron.readthedocs-hosted.com/en/v2.24.0/general/nki/getting_started.html). > Figure. Memory on four current accelerators. Capacity and bandwidth as each vendor's own documentation states them. NVIDIA's tuning guide and architecture page give HBM3e capacity for the B200 but do not state its raw HBM bandwidth, so that cell is left blank rather than estimated.Datasheet numbers --- # NVIDIA Blackwell · Current hardware NVIDIA Blackwell-architecture GPUs pack 208 billion transistors, manufactured on a custom TSMC 4NP process, split across two reticle-limited dies. Those dies are connected by a 10 terabytes per second chip-to-chip interconnect and presented to software as a single unified GPU rather than two separate devices the driver has to stitch together. The programming model underneath does not change: the Blackwell Tuning Guide describes the architecture as retaining and extending the same CUDA model used by Ampere and Hopper, so code already following those architectures’ best practices should see speedups without any changes. At the streaming-multiprocessor level, occupancy limits move in both directions depending on which Blackwell you have. The maximum number of concurrent warps per SM is 64 for compute capability 10.0 devices and 48 for compute capability 12.0 devices; every SM carries a 64K-entry, 32-bit register file with up to 255 registers addressable per thread, and up to 32 thread blocks can be resident on one SM at once. Hopper’s **thread block clusters** (An optional grouping of thread blocks above the grid level. Thread blocks in the same cluster can read, write, and perform atomics in each other's shared memory, and are guaranteed to run concurrently on the same GPU.) carry over unchanged: a thread block can reach into another block’s shared memory within its cluster through **distributed shared memory** (Shared memory belonging to one thread block, made reachable by other thread blocks in the same cluster, so an SM can read, write, and perform atomics on memory it did not allocate.). Every Blackwell GPU supports a portable cluster size of up to 8, and the B200 additionally allows a nonportable cluster size of 16 for applications that opt in. The memory system centers on HBM3 and HBM3e: the B200 GPU supports up to 180 GB of it, and the GB200 GPU raises L2 cache capacity to 126 MB. On-chip, the combined L1 data cache, texture cache, and shared memory still tops out at 256 KB per SM, the same ceiling as Hopper, with the shared-memory carveout selectable at runtime across the same 0, 8, 16, 32, 64, 100, 132, 164, 196, and 228 KB increments on both the H100 and the B200. Scaling out relies on fifth-generation NVLink, which can connect up to 576 GPUs into a single domain. Inside a rack, the NVLink Switch Chip delivers 130 TB/s of aggregate GPU bandwidth across one 72-GPU NVL72 domain, and multiple NVL72 racks can be joined over the same 1.8 TB/s switch interconnect; NVIDIA states that a full NVL72 system supports nine times the GPU throughput of a single eight-GPU server. The link to the host side of the system is separate again: Blackwell reaches into Grace CPU memory over a dedicated connection running at 900 GB/s of bidirectional bandwidth. See [NVIDIA Blackwell documentation](https://learn-kernels.com/chapters/reading#hardware-nvidia) in the reading list for the full tuning guide and architecture materials. Source SM occupancy limits, register file size, cluster sizes, and shared memory carveout increments from the [NVIDIA Blackwell Tuning Guide](https://docs.nvidia.com/cuda/blackwell-tuning-guide/). Transistor count, die interconnect, NVLink, and Grace CPU link figures from the [NVIDIA Blackwell Architecture](https://www.nvidia.com/en-us/data-center/technologies/blackwell-architecture/) overview page. ### What CUTLASS targets on Blackwell The tuning guide describes the envelope; NVIDIA’s CUTLASS documentation describes what a matmul kernel actually executes inside it. Blackwell SM100 introduces the tcgen05.mma family: seven new matrix-multiply instructions that CUTLASS states are 2x to 4x faster than the Hopper architecture’s WGMMA instructions. They cover the legacy tensor core types (tf32, fp16, bf16, and 8-bit integers) alongside new 4-, 6-, and 8-bit floating point types, with and without scale factors, and each instruction comes in a cta\_group::1 and a cta\_group::2 variant. The 2-SM variants split one MMA tile across a pair of SMs, with tile shapes up to 256x256 in which each SM produces half of the output, and a kernel that selects a 2-SM instruction must launch with a cluster whose first dimension is a multiple of 2. The block scaled instruction kinds are where Blackwell meets the same OCP microscaling standard that AMD’s CDNA 4 adopts in the next section. The mxf8f6f4, mxf4, and mxf4nvf4 kinds compute D = C + (A x SFA) \* (B x SFB), where every 16 or 32 elements of A and B along the reduction dimension share one scale factor, so an M x K operand carries a much smaller matrix of scale factors beside it. CUTLASS’s mx\_float8, mx\_float6, and mx\_float4 types pair the data with a ue8m0 scale factor over 32-element blocks and follow the OCP specification; nv\_float4 pairs 4-bit data with a ue4m3 scale factor over 16-element blocks and is not OCP compliant. Per CUTLASS’s own throughput table, the mxf4 kinds reach 4x the throughput of Hopper’s FP8 tensor core. None of this is reached through free-form assembly. CUTLASS exposes the instructions through its collective builder interface, where a kernel names element types, layouts, alignments, an MMA tile shape, a cluster shape, and a dispatch policy such as KernelTmaWarpSpecialized2SmMxf8f6f4Sm100, and the documentation’s tables enumerate exactly which combinations are valid. The constraints tighten as the types narrow: a float4\_t operand must be 128-element aligned to target the f8f6f4 instruction kind, and the block scaled mxf4 kinds accept only the TN layout, row-major A against column-major B. Scheduling gets its own hardware assist. CUTLASS GEMMs are persistent kernels: workers stay resident on the GPU and loop over output tiles, hiding prologue and epilogue costs, but a static tile schedule balances poorly when part of the GPU is occupied by another kernel. Blackwell adds cluster launch control: the kernel launches a grid with as many thread blocks as there are output tiles, and each resident worker asks the hardware for the next unstarted block coordinate via the clusterlaunchcontrol.try\_cancel instruction. Every coordinate is guaranteed to be either launched as a worker or handed to an existing one, so tiles flow to whatever capacity actually exists, a dynamic schedule with the launch semantics of an ordinary grid. Source tcgen05.mma instruction kinds, throughput multiples versus Hopper, block scaling, narrow-precision types, alignment and layout constraints, and dispatch policies from the [CUTLASS Blackwell functionality documentation](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/blackwell_functionality.html). Persistent kernels and dynamic scheduling from the [CUTLASS Blackwell cluster launch control documentation](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/blackwell_cluster_launch_control.html). --- # Inference engines A single kernel and a single request’s KV cache only explain what happens to one sequence on one GPU. An inference engine has to spread that same memory and compute across hundreds of concurrent requests, shrink the weights it streams every step, and turn spare memory bandwidth into extra tokens per second, all without changing what the model would have said on its own. ## 6 sections 1. [4.1Scheduling and continuous batching](https://learn-kernels.com/chapters/inference-engines/scheduling-and-continuous-batching) 2. [4.2KV cache systems](https://learn-kernels.com/chapters/inference-engines/kv-cache-systems) 3. [4.3Quantization](https://learn-kernels.com/chapters/inference-engines/quantization) 4. [4.4Speculative decoding](https://learn-kernels.com/chapters/inference-engines/speculative-decoding) 5. [4.5Structured decoding and fairness](https://learn-kernels.com/chapters/inference-engines/structured-decoding-and-fairness) 6. [4.6Long context and multimodal inference](https://learn-kernels.com/chapters/inference-engines/long-context-and-multimodal) --- # KV cache systems · Inference engines [The KV cache](https://learn-kernels.com/chapters/foundations/the-kv-cache) established that every decode step has to read the cached keys and values for every token that came before, so the size of that cache is a direct tax on the memory-bound step. One multiplier on that size is the number of key-value heads the model keeps. Multi-query attention collapses them to a single key-value head shared by all query heads, which drastically speeds up decoder inference but can degrade quality. **Grouped-query attention** (A generalization of multi-query attention that uses an intermediate number of key-value heads, more than one but fewer than the number of query heads, so the KV cache shrinks without collapsing all queries onto a single shared head.) sits between the two extremes, and the GQA paper showed it does not require training a model from scratch: an existing multi-head checkpoint can be uptrained into a grouped-query one using 5 percent of the original pre-training compute, reaching quality close to multi-head attention at speed comparable to multi-query attention. Reducing the head count is not the only way to shrink what each token leaves behind. **Multi-head latent attention** (An attention variant that compresses each token's key-value state into a small latent vector, which is what gets cached, instead of storing full keys and values per head.), introduced with DeepSeek-V2, compresses the KV cache into a latent vector rather than storing full per-head keys and values. Combined with the rest of that model’s architecture, DeepSeek-V2 reduces the KV cache by 93.3 percent relative to its dense 67B predecessor while supporting a 128K token context and raising maximum generation throughput to 5.76 times. A third lever is where the cache lives at all. Mooncake, the serving platform behind the Kimi chatbot, treats the KV cache as the center of the whole system: it separates prefill and decoding onto different clusters and pools the underexploited CPU, DRAM, SSD, and NIC resources of the GPU cluster into a disaggregated KV cache, so a prefix computed once can be stored and reused far beyond a single GPU’s memory. On real traces this KV-cache-centric design increased effective request capacity by 59 to 498 percent over baseline methods while meeting latency service-level objectives, and the deployed system runs across thousands of nodes processing over 100 billion tokens a day. > Figure. What each attention variant leaves in the cache. Eight query heads in every panel; the shaded shapes underneath are what a decode step must read back per past token. MHA caches keys and values for every head, MQA shares one key-value head across all of them, GQA gives each group of queries its own, and MLA stores only a compressed latent vector that all heads reconstruct from.Illustrative numbers The cache can also be compressed in place, the same lever the next section applies to weights. KIVI starts from a study of how the cached elements of popular models are actually distributed, and the finding is asymmetric: the key cache should be quantized per-channel, grouping elements along the channel dimension, while the value cache should be quantized per-token. Built on that split, KIVI is a tuning-free 2-bit KV cache quantization algorithm, and with a hardware-friendly implementation it lets Llama-2, Falcon, and Mistral models keep almost the same quality while using 2.6 times less peak memory. The freed memory admits up to 4 times larger batches, which translates to 2.35 to 3.47 times the throughput on real inference workloads. Once a cache outlives a single machine, as in Mooncake’s disaggregated store, moving it becomes its own problem: reusing a stored prefix avoids recomputing it, but the KV cache is a large tensor, and fetching it over the network can add enough delay to undo the savings. CacheGen treats this as a transmission problem. A custom tensor encoder leverages the distributional properties of the KV cache to encode it into compact bitstreams with negligible decoding overhead, and the compression level of different parts of the cache adapts to the bandwidth available while it streams, so loading delay stays low as conditions change. Against recent systems that reuse KV caches without encoding them this way, CacheGen reduced the cache size 3.5 to 4.3 times and the total delay of fetching and processing contexts 3.2 to 3.7 times, with negligible impact on response quality. Source Multi-query and grouped-query attention and the 5 percent uptraining figure from [GQA: Training Generalized Multi-Query Transformer Models from Multi-Head Checkpoints](https://arxiv.org/abs/2305.13245) (Ainslie et al., 2023). Multi-head latent attention and the 93.3 percent and 5.76× figures from [DeepSeek-V2](https://arxiv.org/abs/2405.04434) (DeepSeek-AI, 2024). The disaggregated cache and capacity figures from [Mooncake: Trading More Storage for Less Computation](https://www.usenix.org/conference/fast25/presentation/qin) (Qin et al., FAST 2025). Per-channel key and per-token value quantization and the memory and throughput figures from [KIVI: A Tuning-Free Asymmetric 2bit Quantization for KV Cache](https://proceedings.mlr.press/v235/liu24bz.html) (Liu et al., ICML 2024). The KV cache encoding and transfer figures from [CacheGen: KV Cache Compression and Streaming for Fast Large Language Model Serving](https://cs.stanford.edu/~keithw/sigcomm2024/sigcomm24-final1571-acmpaginated.pdf) (Liu et al., SIGCOMM 2024). More in [KV cache systems](https://learn-kernels.com/chapters/reading#kv-cache-systems) in the reading list. --- # Long context and multimodal inference · Inference engines Everything so far assumed a request fits on the devices serving it. Long inputs break that assumption twice: the KV cache for a single very long sequence can outgrow any one device, and prefill compute grows quadratically with input length. The memory side is exactly what the Ring Attention paper targets, motivated by videos, actions, and other long-form sequences and modalities whose token counts overwhelm a single accelerator. **Ring attention** (Computing exact attention over a sequence split into blocks across multiple devices arranged in a ring, with each device passing key-value blocks to the next while it computes blockwise attention, so communication is fully overlapped with compute.) distributes blockwise computation of self-attention and feedforward across devices while fully overlapping the communication of key-value blocks with the computation of blockwise attention. The result is exact, not an approximation, and it scales the feasible sequence length by up to the number of devices in the ring, which the paper demonstrates at context sizes in the millions of tokens. > Figure. Ring attention. Four devices in a ring, the sequence split into blocks across them. Key-value blocks circulate device to device while each device keeps its own query block and computes blockwise attention against whichever KV block just arrived, so the communication is fully overlapped with compute and the feasible sequence length scales by up to the number of devices in the ring.Illustrative numbers The compute side of long context is a prefill problem. The MInference paper measures the cost concretely: because attention is quadratic in the prompt length, an 8B parameter model takes 30 minutes to prefill a 1 million token prompt on a single A100. Its observation is that long-context attention matrices are not densely important; they exhibit three recurring structures, called A-shape, Vertical-Slash, and Block-Sparse. **Dynamic sparse attention** (Attention that computes only a subset of the score matrix, where the subset is not fixed in advance but rebuilt per input, here by assigning each head a sparse pattern offline and building that head's sparse indices on the fly during inference.) exploits this by assigning each attention head its best-fitting pattern offline, then building the sparse indices for that pattern dynamically at inference time and running optimized sparse kernels over them. Applied to existing models with no change to pre-training and no fine-tuning, this cuts prefill latency by up to 10 times on an A100 while maintaining accuracy across long-context benchmarks. Between them, the two techniques bracket the long-context problem: one spreads exact attention over more hardware, the other spends less compute per unit of hardware, and both leave the model’s output distribution intact enough to serve the same requests. MInference recovers sparsity from a model trained with full attention; Native Sparse Attention, from DeepSeek, builds the sparsity in from the start. **Natively trainable sparse attention** (A sparse attention mechanism trained into the model end to end from pretraining onward, rather than imposed on a full-attention model afterward, so the model learns around the sparsity instead of merely tolerating it.) uses a dynamic hierarchical strategy that combines coarse-grained token compression, which preserves global context awareness, with fine-grained token selection, which preserves local precision. The design is hardware-aligned, balancing arithmetic intensity so the sparse kernels actually run fast on modern GPUs, and because it is trainable end to end it reduces pretraining computation as well as inference cost. A model pretrained with NSA maintains or exceeds its full-attention counterpart across general benchmarks, long-context tasks, and instruction-based reasoning, while reaching substantial speedups over full attention on 64k-length sequences across decoding, forward propagation, and backward propagation. Multimodal requests arrive through the same serving stack. In vLLM, a multimodal model takes its text prompt together with a separate multimodal data dictionary carrying the other modalities, and the documented input types cover images, video, and audio. Images can be passed as URLs, as image objects, or as pre-computed embeddings, and the maximum number of frames taken per video can be configured. The engine machinery this chapter has covered is what serves those requests too: a multimodal request enters the same scheduler as a text-only one, carrying extra inputs alongside its prompt rather than taking a separate path. Source Blockwise ring communication and the device-count scaling from [Ring Attention with Blockwise Transformers for Near-Infinite Context](https://arxiv.org/abs/2310.01889) (Liu et al., 2023). The 30 minute and 10× prefill figures and the three sparse patterns from [MInference 1.0: Accelerating Pre-filling for Long-Context LLMs via Dynamic Sparse Attention](https://arxiv.org/abs/2407.02490) (Jiang et al., 2024). Hierarchical compression and selection and the 64k-sequence results from [Native Sparse Attention: Hardware-Aligned and Natively Trainable Sparse Attention](https://arxiv.org/abs/2502.11089) (Yuan et al., 2025). Multimodal input types from the [vLLM multimodal inputs documentation](https://docs.vllm.ai/en/latest/features/multimodal_inputs.html). More in [Long context and multimodal](https://learn-kernels.com/chapters/reading#long-context-multimodal) in the reading list. --- # Quantization · Inference engines A decode step is memory bound because it spends most of its time streaming the model’s weights through memory rather than doing arithmetic on them, as established in [the memory-bound and compute-bound boundary](https://learn-kernels.com/chapters/foundations/compute-bound-and-memory-bound). Shrinking those weights shrinks that streaming cost directly. GPTQ does this with **weight-only quantization** (Compressing only the model's weights to a lower bit-width while leaving activations at their original precision.): it is a one-shot, post-training method based on approximate second-order information that compresses the weights of a model, including ones with 175 billion parameters, down to 3 or 4 bits each with negligible accuracy loss, and without any retraining. Because the compression touches only the weights, the arithmetic itself still runs at higher precision: the low-bit weights have to be dequantized back up before each matrix multiply, so the benefit is concentrated in the memory-bound decode step, where fewer bytes have to move, rather than in the compute-bound prefill step. SmoothQuant instead quantizes both operands of the matrix multiply, **weight and activation quantization** (Compressing both the model's weights and its activations to a low bit-width, most commonly 8 bits each (W8A8), so the matrix multiply itself can run in low precision instead of just being read from memory in low precision.). That is harder than quantizing weights alone because, as the paper puts it, weights are easy to quantize while activations are not: activation values in large language models develop large outliers concentrated in a handful of channels, and forcing those outliers into an 8-bit range destroys the precision available to every other value. SmoothQuant’s fix is a mathematically equivalent transformation, applied offline before serving, that migrates this quantization difficulty from the activations into the weights, making both sides easy enough to quantize to INT8 together. Because both operands end up in low precision, the matrix multiply itself can run on faster low-precision hardware paths, not just save memory traffic. > Figure. Weight-only versus weight-and-activation quantization. Two dataflows into the same matrix multiply. Weight-only quantization shrinks what streams from memory but dequantizes back to FP16 before the multiply, so the arithmetic stays high precision. Quantizing activations too lets the multiply itself run in INT8, at the cost of surviving activation outliers.Illustrative numbers AWQ sits in the same weight-only family as GPTQ but reads the problem through the activations. Its core observation is that not all weights in a model are equally important: protecting only 1 percent of **salient weights** (The small fraction of a model's weight channels whose quantization error most damages output quality, identified in AWQ by the distribution of the activations flowing through them rather than by the weights' own magnitudes.) can greatly reduce quantization error, and the way to find those salient channels is to look at the activation distribution, not at the weights themselves. Keeping that 1 percent at higher precision would leave a hardware-inefficient mixed-precision layout, so AWQ instead derives an equivalent transformation that scales up the salient channels before quantizing, with the scale determined by activation statistics collected offline. Because the method relies on no backpropagation and no reconstruction, it does not overfit its calibration set, and it generalizes across domains and modalities, including instruction-tuned and multimodal models. Low-bit weights only pay off if the kernels serving them are fast, which is why the AWQ paper ships with TinyChat, an inference framework built for 4-bit models. With kernel fusion and platform-aware weight packing, TinyChat runs more than 3 times faster than the Huggingface FP16 implementation on both desktop and mobile GPUs, and the compression is what lets a 70B Llama-2 model be deployed on a mobile GPU at all. Source Weight-only quantization and the 3 to 4 bit, 3.25× figures from [GPTQ: Accurate Post-Training Quantization for Generative Pre-trained Transformers](https://arxiv.org/abs/2210.17323) (Frantar et al., 2022). Weight-and-activation quantization and the outlier framing from [SmoothQuant: Accurate and Efficient Post-Training Quantization for Large Language Models](https://arxiv.org/abs/2211.10438) (Xiao et al., 2022). Salient-weight protection, activation-aware scaling, and the TinyChat figures from [AWQ: Activation-aware Weight Quantization for LLM Compression and Acceleration](https://proceedings.mlsys.org/paper_files/paper/2024/file/42a452cbafa9dd64e9ba4aa95cc1ef21-Paper-Conference.pdf) (Lin et al., MLSys 2024). More quantization methods in [Quantization](https://learn-kernels.com/chapters/reading#quantization) in the reading list. --- # Scheduling and continuous batching · Inference engines Batching multiple requests together lets an engine amortize the cost of streaming model weights through memory across all of them at once, which is exactly the decode-time bottleneck described in [the KV cache](https://learn-kernels.com/chapters/foundations/the-kv-cache). But requests do not arrive together, and they do not finish together: a naive scheme that groups requests into a fixed batch and waits for every member to finish either stalls new requests in a queue or pads short sequences out to the length of the longest one in the batch, wasting both compute and memory. Production engines instead run **continuous batching** (Scheduling requests at the level of individual model steps rather than whole batches: after every step, finished requests leave the batch and newly arrived ones join, so no request waits for an unrelated one to finish.), which the vLLM paper describes as iteration-level scheduling: after each step of the model, completed sequences are removed from the batch and new ones are added, so a fresh request only waits for a single iteration rather than for the whole batch to drain. The idea has a specific origin. Orca, the OSDI 2022 system that introduced iteration-level scheduling, framed the failure of earlier servers as an inflexible scheduling mechanism that cannot change the batch being processed: requests that finish earlier than the rest of their batch cannot return to the client, and new arrivals wait until the current batch completely finishes. Orca’s scheduler instead invokes the execution engine to run only a single iteration of the model at a time. Applying batching and iteration-level scheduling to a Transformer at the same time required a second technique, **selective batching** (Applying batching only to the operations that can still be batched when the requests sharing an iteration are at different phases and lengths, instead of requiring every operation in the model to run over a uniform batch.), which applies batching only to a selected set of operations, since the requests sharing an iteration no longer line up the way a fixed batch does. On a GPT-3 175B model, the combination gave Orca a 36.9 times throughput improvement over NVIDIA FasterTransformer at the same level of latency. Continuous batching alone does not solve the memory side of the problem. Each request’s KV cache grows one token at a time and its final length is not known when the request starts, so systems that store it as one contiguous tensor have to reserve space for the worst case up front. That reservation is why prior serving systems left most of their KV cache memory unused: profiling in the PagedAttention paper found only 20.4 to 38.2 percent of allocated KV cache memory actually held token state, the rest lost to padding for an unreached maximum length and to fragmentation between differently sized reservations. **Paged attention** (Splitting each request's KV cache into fixed-size blocks that do not need to sit in contiguous memory, the same idea operating systems use for virtual memory pages, so blocks can be allocated on demand and shared across sequences.) fixes this by dividing the KV cache into fixed-size blocks that can live anywhere in memory, addressed indirectly the way an OS page table addresses physical pages. Because every block is the same size, allocating one on demand as a sequence grows leaves no external fragmentation, and because blocks are addressed indirectly, several sequences that share a prefix, such as the beams in a beam search, can point at the same block instead of duplicating it. > Figure. Request-level batching versus continuous batching. The same four requests on three batch slots, one column per model step. A fixed batch holds every slot until its longest member finishes, so B's and A's slots sit idle and D waits for the drain. Continuous batching admits D the step after B ends and finishes the same work in eight steps instead of eleven.Illustrative numbers Mixing phases inside one batch creates a tension continuous batching does not resolve on its own. A prefill iteration processes the whole prompt in parallel, so it has high latency but saturates GPU compute; a decode iteration produces a single token per request, so it is fast but leaves compute idle. Interleaving the two means every prefill admitted into a running batch delays the decodes sharing it, which makes high throughput and low latency hard to achieve together. Sarathi-Serve resolves this with **chunked prefills** (Splitting a prefill request into near equal sized chunks so it can be fed into a running batch a piece at a time, instead of monopolizing an entire iteration with one long prompt.) and stall-free schedules: a prefill is split into near equal sized chunks, and new requests join the batch without pausing ongoing decodes, while the resulting uniform batches also reduce the iteration imbalance that causes pipeline bubbles. Under tail latency constraints this raised serving capacity 2.6 times for Mistral-7B on a single A100 and up to 5.6 times for Falcon-180B served with pipeline parallelism, both relative to vLLM. The scheduler can also be smarter about what it throws away. In existing engines the KV cache of a request is discarded once the request completes, so two calls that share a long prefix, a system prompt, a few-shot template, or the earlier turns of a chat, each pay to recompute it. SGLang’s **RadixAttention** (Keeping the KV cache of finished requests in a radix tree managed as an LRU cache, so a later request that shares a prefix with an earlier one can match it in the tree and reuse the cached key-value state instead of recomputing it.) instead maintains an LRU cache of the KV cache for all requests within a radix tree, so matching, insertion, and eviction are efficient and a cache-aware scheduling policy can steer requests toward their cached prefixes. On workloads built from multi-call programs, agents, reasoning chains, and multi-turn chat, this and the runtime’s other optimizations reach up to 6.4 times higher throughput than existing inference systems. Source Iteration-level scheduling, PagedAttention, and the throughput and memory-waste figures above are from [Efficient Memory Management for Large Language Model Serving with PagedAttention](https://arxiv.org/html/2309.06180) (Kwon et al., 2023), the vLLM paper. Iteration-level scheduling, selective batching, and the 36.9× figure from [Orca: A Distributed Serving System for Transformer-Based Generative Models](https://www.usenix.org/conference/osdi22/presentation/yu) (Yu et al., OSDI 2022). Chunked prefills, stall-free scheduling, and the capacity figures from [Taming Throughput-Latency Tradeoff in LLM Inference with Sarathi-Serve](https://www.usenix.org/system/files/osdi24-agrawal.pdf) (Agrawal et al., OSDI 2024). RadixAttention and the 6.4× figure from [SGLang: Efficient Execution of Structured Language Model Programs](https://arxiv.org/html/2312.07104) (Zheng et al., 2024). More engines and the papers they build on are in [Scheduling and continuous batching](https://learn-kernels.com/chapters/reading#scheduling-batching) and [KV cache systems](https://learn-kernels.com/chapters/reading#kv-cache-systems) in the reading list. --- # Speculative decoding · Inference engines Autoregressive decoding is serial by construction: generating K tokens takes K sequential runs of the model, each one waiting on the last token before it can start. But because a decode step is memory bound rather than compute bound, as [the previous chapter established](https://learn-kernels.com/chapters/foundations/compute-bound-and-memory-bound), the GPU is usually sitting on spare arithmetic capacity while it waits for weights to stream in. Speculative decoding spends that spare capacity on parallel verification instead of serial generation. A small, fast **draft model** (A smaller, cheaper model used to generate several candidate tokens ahead of the model whose output the system actually wants, run either autoregressively or in parallel.) proposes several tokens ahead, autoregressively, and the large **target model** (The larger, more capable model whose output distribution the system is trying to sample from; speculative decoding accelerates sampling from it without changing what it would have produced alone.) then scores all of those candidates in a single parallel forward pass, the same cost as generating just one token normally. The verification step is not a simple accept-or-reject on whether the two models agree. It uses a modified rejection sampling scheme: a draft token is kept whenever the target model would have assigned it at least as much probability as the draft model did, and when the target assigns it less probability, the token is only kept with probability equal to that ratio; a rejected token is replaced by resampling from a distribution built out of the leftover probability mass. That correction is what makes the method exact rather than an approximation: the sequence that comes out has exactly the distribution the target model would have produced sampling on its own, regardless of what draft model was used or how often it was wrong. What varies with the draft model’s quality is the **acceptance rate** (The fraction of proposed draft tokens the target model keeps during verification; a higher acceptance rate means more tokens are produced per parallel scoring pass.), since every rejection means one parallel scoring pass produced only one or a few tokens instead of many. Both papers that introduced this measured real speedups without retraining or changing the target model at all. Leviathan et al. report a 2 to 3 times latency improvement sampling from an 11 billion parameter T5-XXL model, walltime-tested against the standard T5X implementation; in one worked example, a 38-token sentence was produced from only 9 serial runs of the target model. Chen et al. report a 2 to 2.5 times decoding speedup on the 70 billion parameter Chinchilla model in a distributed serving setup, using the same draft-then-verify structure under the name speculative sampling. > Figure. One round of draft-then-verify. The draft model proposes five tokens one at a time; the target model scores all of them in a single forward pass. The first three are accepted, the fourth is rejected and replaced by a token resampled from the target's leftover probability mass, and everything after the rejection is discarded. One target-model step produced four tokens instead of one.Illustrative numbers The draft model itself is the operational weak point of this design: a separate model has to be acquired and maintained, and the Medusa paper names that as the obstacle impeding adoption. Medusa removes it. Instead of a second model, it adds extra **decoding heads** (Additional output heads attached to the backbone model's last hidden states, each a single feed-forward layer with a residual connection, trained so the k-th head predicts the token k+1 positions ahead.) on top of the backbone’s last hidden states, each a single feed-forward layer with a residual connection predicting a token several positions ahead. The heads emit multiple top predictions per position, which are assembled into several candidate continuations and verified simultaneously in one decoding step through a tree-based attention mechanism, a simple adjustment to the attention mask. Fine-tuning only the heads on a frozen backbone, called Medusa-1, keeps the acceleration lossless and reaches over 2.2 times speedup; training the heads and backbone together, Medusa-2, raises that to 2.3 to 2.8 times. EAGLE keeps a small drafting network but moves the drafting to a different level of the model. Its starting observations are that autoregression at the feature level, the second-to-top layer of the transformer, is more straightforward than at the token level, and that the inherent uncertainty in feature-level autoregression is what constrains its performance. EAGLE resolves that uncertainty by feeding the draft network a token sequence advanced by one time step alongside the features, which makes precise feature prediction possible with minimal overhead. On LLaMA2-Chat 70B this reaches latency speedups of 2.7 to 3.5 times and doubles throughput, while still maintaining the distribution of the generated text, the same exactness guarantee the original draft-and-verify schemes provide. Source Draft/target framing and the T5-XXL figures from [Fast Inference from Transformers via Speculative Decoding](https://proceedings.mlr.press/v202/leviathan23a.html) (Leviathan et al., 2023). Acceptance-rate framing and the Chinchilla figures from [Accelerating Large Language Model Decoding with Speculative Sampling](https://arxiv.org/abs/2302.01318) (Chen et al., 2023). Decoding heads, tree attention, and the Medusa-1 and Medusa-2 figures from [Medusa: Simple LLM Inference Acceleration Framework with Multiple Decoding Heads](https://arxiv.org/html/2401.10774) (Cai et al., 2024). Feature-level drafting and the LLaMA2-Chat 70B figures from [EAGLE: Speculative Sampling Requires Rethinking Feature Uncertainty](https://proceedings.mlr.press/v235/li24bt.html) (Li et al., ICML 2024). More on batched and production variants in [Speculative decoding](https://learn-kernels.com/chapters/reading#speculative-decoding) in the reading list. --- # Structured decoding and fairness · Inference engines Callers increasingly want output that is not just plausible text but valid JSON, a function call, or a string in some grammar. **Constrained decoding** (Enforcing strict formal language constraints, such as a grammar or schema, during generation, typically by restricting which tokens the model is allowed to emit at each step.) methods enforce such formal language constraints during generation. Doing this naively is costly in two ways: Beurer-Kellner et al. show that many existing methods not only add performance overhead at generation time but also significantly impair task accuracy, because the constraint is defined over text while the model emits subword tokens, and misaligning the two distorts what the model is allowed to say. Their DOMINO algorithm enforces constraints in a fully subword-aligned fashion, and by leaning on pre-computation and speculative decoding it runs with virtually no overhead, in some cases almost 2 times faster than unconstrained decoding. > Figure. Constrained decoding as a token mask. One decode step while generating JSON, with the output so far ending in a key and a colon. The constraint restricts which tokens the model is allowed to emit: candidates that would make the output invalid, here a bare word and a stray closing bracket, are masked out, and the next token is sampled only from the set the mask lets through.Illustrative numbers The other trust problem an engine has to solve is between clients rather than within one request. First-come-first-serve with a per-client request rate limit is how most services protect themselves, but Sheng et al. call that notion of fairness rudimentary: it leaves capacity idle when a heavy client is throttled against an underutilized server, and a request cap treats a 2,000 token request the same as a 200 token one. Their fix is **fair queueing** (A scheduling discipline in which each of n clients sharing a resource is guaranteed at least a 1/n share of it, and any share a client does not use is redistributed to clients with more demand.) applied at the token level, and LLM serving breaks the classic algorithms in specific ways: a request’s output length is unknown when it is scheduled, input tokens cost less to process than output tokens, and the server’s effective capacity in tokens per second changes with the mix of sequence lengths in the batch. Their Virtual Token Counter tracks the service each client has received, counted in tokens as they are actually processed, and admits requests from the least-served client first on top of continuous batching. The scheduler is work-conserving, and they prove a tight upper bound of 2 times on the service difference between two backlogged clients. Source Subword misalignment and the DOMINO figures from [Guiding LLMs The Right Way: Fast, Non-Invasive Constrained Generation](https://proceedings.mlr.press/v235/beurer-kellner24a.html) (Beurer-Kellner et al., 2024). Token-level fairness and the VTC bound from [Fairness in Serving Large Language Models](https://arxiv.org/html/2401.00588) (Sheng et al., 2024). More in [Structured decoding](https://learn-kernels.com/chapters/reading#structured-decoding) in the reading list. --- # Introduction Getting a GPU to run a computation is easy. Getting it to run that computation anywhere near what the hardware can actually do is the entire discipline. This book is about the gap between those two things, and about the vocabulary you need to reason about it. ## 3 sections 1. [0.1Why kernels matter](https://learn-kernels.com/chapters/introduction/why-kernels-matter) 2. [0.2How this book is organized](https://learn-kernels.com/chapters/introduction/how-this-book-is-organized) 3. [0.3How to read it](https://learn-kernels.com/chapters/introduction/how-to-read-it) --- # How this book is organized · Introduction The book is ordered the way the problem is: from one inference request down to a single kernel, then back out to a fleet of machines. [Foundations](https://learn-kernels.com/chapters/foundations) covers the execution model and the compute-bound versus memory-bound boundary that everything later depends on. [Kernel optimization](https://learn-kernels.com/chapters/kernel-optimization) applies it to the three computations that matter most in practice: matrix multiplication, low-precision arithmetic, and attention. [Programming models and profiling](https://learn-kernels.com/chapters/programming-models) is about the tools you actually write and measure kernels with, since very little of this work is done in raw CUDA C++ any more. [Inference engines](https://learn-kernels.com/chapters/inference-engines) and [Distributed inference](https://learn-kernels.com/chapters/distributed-inference) move up to the systems layer, where scheduling, batching, and placement decide throughput more than any single kernel does. [Current hardware](https://learn-kernels.com/chapters/hardware) closes with what the chips being deployed right now actually provide. Two appendices sit outside that order. The [glossary](https://learn-kernels.com/chapters/glossary) defines every term of art once, quoting its source verbatim. The [reading list](https://learn-kernels.com/chapters/reading) collects the primary sources: papers, specifications, and repositories, grouped the same way the chapters are. On sources Every factual claim here traces to a citable source, linked in a sidenote at the end of each section. Where a number appears, it came from a datasheet, a specification, or a measurement someone published, not from an estimate. --- # How to read it · Introduction If you are new to the field, read [Foundations](https://learn-kernels.com/chapters/foundations) first and do not skip it. The rest of the book assumes you can say what a **warp** (Within a thread block, threads are organized into groups of 32 threads called warps, executing in a Single-Instruction Multiple-Threads paradigm.) is and why a kernel might be memory bound. After that the chapters are reasonably independent: read them in order for the argument, or jump to whichever layer you are currently stuck on. If you already write kernels, the chapters will move faster than you need in places. The reading list is probably the more valuable artifact for you: it is a curated set of primary sources, and the chapters are best treated as a map of which of them to read and in what order. Terms with a dotted underline are defined inline; hover or tap them. The same definitions, with their sources, are collected in the [glossary](https://learn-kernels.com/chapters/glossary). Figures in this book are there to show structure, not decoration, and the caption always says whether the numbers in one are from a specification or merely illustrative. --- # Why kernels matter · Introduction A modern accelerator is not one fast processor. It is a very wide machine with a deep memory hierarchy, and almost every disappointing performance number comes from the same root cause: the arithmetic units finished early and spent the rest of their time waiting for data. A kernel that reads its inputs badly can be an order of magnitude slower than one that reads them well, running the exact same arithmetic on the exact same chip. > Figure. The gap this book is about. The top bar is what the hardware can do at peak; the bottom bar is what naive code gets on the same chip. The lengths are not measurements, the shape is the point: a kernel that reads its inputs badly can be an order of magnitude slower than one that reads them well.Illustrative numbers That is why this field has a distinctive shape. The interesting work is rarely inventing new mathematics. It is arranging a known computation so that the data it needs arrives in the right place, in the right order, in large enough pieces, while the arithmetic units stay busy. Tiling, coalescing, pipelining, and quantization are all answers to that one question. The same logic scales outward. A single kernel worries about the trip from global memory to registers. An inference engine worries about the trip from one request to a batch of them. A distributed system worries about the trip across an interconnect. The unit of analysis changes; the reasoning does not. --- # Kernel optimization Getting a kernel correct is the easy part. This chapter follows one matrix multiply from 1.3% of a GPU’s advertised throughput to 93.7% of it through nothing but memory-hierarchy discipline, then turns to the two techniques that multiply that hard-won bandwidth further: tensor cores that trade numeric precision for raw throughput, and the fused attention kernels that turned a quadratic memory-traffic problem into a solved one. ## 4 sections 1. [2.1Foundational kernel exercises](https://learn-kernels.com/chapters/kernel-optimization/foundational-kernel-exercises) 2. [2.2Matrix multiplication](https://learn-kernels.com/chapters/kernel-optimization/matrix-multiplication) 3. [2.3Tensor cores and low precision](https://learn-kernels.com/chapters/kernel-optimization/tensor-cores-and-low-precision) 4. [2.4Attention](https://learn-kernels.com/chapters/kernel-optimization/attention) --- # Attention · Kernel optimization Self-attention’s naive implementation is memory-bound for a specific, avoidable reason: computing softmax(QKT/√d)V requires materializing the full N×N score matrix, writing it out to **HBM** (A GPU's main off-chip memory. An A100 has 1.5 to 2.0 TB/s of HBM bandwidth, an order of magnitude less than the roughly 19 TB/s available from the on-chip SRAM local to each streaming multiprocessor.), then reading it straight back for the softmax and the second matmul. None of that traffic does useful arithmetic, and an A100 has only 192KB of on-chip SRAM per streaming multiprocessor, across 108 of them, to avoid it with. FlashAttention avoids ever writing that matrix to HBM. It loops over blocks of K and V in an outer loop, loading each block into SRAM once, then loops over blocks of Q in an inner loop, computing softmax incrementally across blocks, the same tiling idea from matrix multiplication applied to a running reduction instead of a running sum, and fusing every step of the attention computation into one kernel so intermediate results never leave the chip. The payoff is up to a 7.6× speedup on the attention computation itself for GPT-2, up to 9× fewer HBM accesses than the standard implementation, and, end to end, a 15% wall-clock speedup training BERT-large over the MLPerf 1.1 record and a 3× speedup on GPT-2. > Figure. IO-aware attention. One block of K and V is loaded into on-chip SRAM at a time while blocks of Q stream past it. The score block, the online softmax, and the multiply by V are fused into a single kernel, so the N×N score matrix is never written to HBM; a running max and sum rescale the partial output as each new block arrives. Shaded blocks mark one tile's worth of data.Illustrative numbers The original kernel still reached only 25 to 40% of a GPU’s theoretical peak FLOPs/s, and FlashAttention-2 traced the shortfall to work partitioning rather than the algorithm: for some problem shapes too few thread blocks were active at once, hurting **occupancy** (The fraction of the warps an SM could run at once that are actually resident.), and shared-memory traffic between warps within a block was higher than it needed to be. Rebalancing that work, splitting the sequence dimension across more thread blocks and dividing work between warps to cut shared-memory communication, without changing the underlying algorithm, roughly doubled throughput over the original kernel, reaching 50 to 73% of theoretical peak FLOPs/s on an A100 and, in end-to-end GPT-style training, up to 225 TFLOP/s per A100, 72% model FLOPs utilization. On Hopper the ceiling moved again: FlashAttention-2, tuned for the A100, achieves only 35% utilization on an H100, and FlashAttention-3 closes the gap with the same asynchronous machinery the previous section’s GEMM kernels use. Warp specialization overlaps computation with TMA data movement, block-wise matmuls are interleaved with the softmax so the tensor cores are not idle while exponentials run on slower units, and block quantization with incoherent processing exploits Hopper’s hardware FP8 support. The result is a 1.5 to 2.0× speedup on H100, up to 740 TFLOP/s at FP16, 75% utilization, close to 1.2 PFLOP/s at FP8, and 2.6× lower numerical error than a baseline FP8 attention. All three kernels handle the case where a full sequence is processed at once, training or a single prefill pass. Decoding, where the [KV cache](https://learn-kernels.com/chapters/foundations/the-kv-cache) grows by one token per step, needs the same IO-aware framing applied to a much smaller, much more frequent matmul, which is why inference engines run a separate, differently tuned attention kernel at decode time rather than reusing the training kernel unchanged. In serving stacks that split has a name: FlashInfer, a library and kernel generator for inference, ships separate optimized kernels for prefill, decode, and mixed batching, behind unified attention, GEMM, and mixture-of-experts APIs that select among backends including FlashAttention-2 and 3, cuDNN, CUTLASS, and TensorRT-LLM. Its attention kernels operate directly on the paged and ragged KV-cache layouts the next chapters’ engines use, add cascade attention to share the KV cache of common prefixes across requests, and stay compatible with CUDA Graphs and torch.compile for low-latency serving. Source IO-aware tiling, fusion, and the BERT-large/GPT-2/long-range-arena speedups from [FlashAttention](https://arxiv.org/abs/2205.14135). Work-partitioning fixes and the A100 utilization and end-to-end training numbers from [FlashAttention-2](https://arxiv.org/abs/2307.08691). Hopper asynchrony techniques, utilization, and FP8 error figures from [FlashAttention-3](https://arxiv.org/abs/2407.08608). FlashInfer’s kernels, backends, and KV-cache support from the [FlashInfer](https://github.com/flashinfer-ai/flashinfer) README. More on fused and distributed attention kernels in [the reading list](https://learn-kernels.com/chapters/reading#attention). --- # Foundational kernel exercises · Kernel optimization Before anyone optimizes a matrix multiply, they work through a shorter list of exercises: transpose, reduction, scan, softmax. None of them is interesting as a computation, which is the point. Each one isolates a single mechanism that the big kernels combine. Transpose is the cleanest example because it does no arithmetic at all, so its performance is purely a question of **coalescing** (The hardware combining the memory accesses of neighboring threads in a warp into fewer, wider transactions, which happens when consecutive threads touch consecutive addresses.). NVIDIA’s classic walkthrough makes the cost concrete: a naive transpose reads its input coalesced but writes its output with a stride of 1024 elements, 4096 bytes, between neighboring threads on a 1024×1024 matrix, and on a Tesla M2050 it reaches 18.8 GB/s of effective bandwidth where a plain copy of the same data reaches 105.2 GB/s. The fix is the same staging idea the next section applies to matmul: a warp reads a 32×32 tile row by row into shared memory, the block synchronizes at a barrier, and then warps write columns of the tile back out so that the global-memory writes become contiguous again. The barrier matters because threads now consume data that other threads staged, the cooperation discipline that reduction and scan exercises then make the entire kernel. Shared memory brings its own lesson: in a 32×32 tile every element of a column lands in the same memory bank, so reading a column is a worst-case 32-way **bank conflict** (Multiple threads in a warp accessing different addresses that live in the same shared-memory bank, forcing the accesses to be serialized instead of served in parallel.), and the cure is almost comically small: declare the tile 33 elements wide instead of 32 so columns spread across banks. With both fixes the transpose reaches about 95% of copy throughput. The kernel itself is short. Every kernel in the walkthrough launches blocks of 32×8 threads to move a 32×32 tile, so each thread handles four elements and the index arithmetic is amortized across them. The tiled version reads rows of the input, waits at the barrier, then writes columns of the tile out as rows of the output: > Figure. transposeCoalesced, from An Efficient Matrix Transpose in CUDA C/C++ On the Tesla M2050 this kernel reaches 51.3 GB/s, up from the naive 18.8 but still half of copy throughput, and the post rules out the obvious suspect with a control experiment: a copy kernel routed through the same shared-memory tile and barrier runs at 104.6 GB/s, essentially full speed. The staging is not the cost. What remains is the bank conflict described above, and the one-line padding fix, `tile[TILE_DIM][TILE_DIM+1]`, takes the transpose to 99.5 GB/s. > Figure. Coalescing in the transpose. Reads land on consecutive addresses, so the hardware combines a warp's loads into fewer, wider transactions. The naive transpose's writes land in a different row each, 4096 bytes between neighboring threads, and staging a 32×32 tile in shared memory with a barrier in between is what makes the writes contiguous again. Four threads stand in for a 32-thread warp.Illustrative numbers Reduction, collapsing an array to a single sum, is the exercise where that cooperation discipline gets optimized end to end. Mark Harris’s NVIDIA walkthrough takes one kernel through seven versions on a G80 GPU whose theoretical bandwidth is 86.4 GB/s, and since a reduction performs one flop per element loaded, bandwidth is the only score that matters. Each block builds a tree in shared memory, halving the number of active threads each step, and a second kernel launch reduces the per-block results, because a kernel launch is CUDA’s only global synchronization point across blocks. The first version reads: > Figure. reduce0, from Optimizing Parallel Reduction in CUDA The modulo test looks innocent and is the whole problem: within every warp, which threads pass `tid % (2*s) == 0` alternates, so the warps are highly divergent, and the kernel manages 2.083 GB/s on 4M elements. The walkthrough then removes one bottleneck at a time. A strided index makes the branch non-divergent (2.33× faster, but now bank-conflicted), sequential addressing makes shared-memory access conflict-free (4.68× cumulative), doing a first add while loading from global memory stops half the threads idling on the first pass (8.34×), unrolling the last warp drops the barrier and the branch once only 32 threads remain (15.01×), templating the block size unrolls the rest (21.16×), and giving each thread many elements in a grid-strided loop, what Harris calls algorithm cascading, lands at 62.671 GB/s, a 30× cumulative speedup, 73 GB/s on 32M elements. The deck’s closing arithmetic is the part worth memorizing: of that 30×, the algorithmic changes, addressing and cascading, contributed 11.84×, and the code-level unrolling contributed 2.54×. Fixing how threads cooperate bought almost five times more than fixing how instructions are emitted. Softmax teaches the streaming trick. The numerically safe version every major framework uses subtracts the vector’s maximum before exponentiating, which costs three passes over the input, the max, the normalizer, then the outputs, four memory accesses per element. The **online softmax** (Computing softmax's maximum and normalization term together in a single pass by rescaling the running sum with a factor of e to the (old max minus new max) whenever a larger element arrives.) of Milakov and Gimelshein folds the first two passes into one: carry a running maximum and a running sum together, and whenever a new maximum appears, multiply the sum by e raised to the old max minus the new max before adding the next term. That cuts memory accesses from four to three per element, measured at up to a 1.3× speedup alone and up to 5× fused with top-k. The deeper payoff is structural: a normalizer that can absorb one new element at a time can absorb one new block at a time, which is exactly what the attention section below needs when it walks the score matrix tile by tile without ever holding it whole. The [reading list](https://learn-kernels.com/chapters/reading#foundational-kernel-exercises) collects worked versions of all four exercises. Source Transpose kernels, bandwidth numbers, and the bank-conflict padding fix from [An Efficient Matrix Transpose in CUDA C/C++](https://developer.nvidia.com/blog/efficient-matrix-transpose-cuda-cc/). The reduction kernel, the seven-version sequence, and its bandwidth and speedup numbers from Mark Harris’s [Optimizing Parallel Reduction in CUDA](https://developer.download.nvidia.com/assets/cuda/files/reduction.pdf). The single-pass normalizer, memory-access counts, and speedups from [Online normalizer calculation for softmax](https://arxiv.org/abs/1805.02867). --- # Matrix multiplication · Kernel optimization A naive matrix-multiply kernel assigns one output element to one thread: each thread walks the corresponding row of A and column of B and accumulates a dot product straight out of global memory. On an RTX A6000 multiplying two 4092×4092 float32 matrices, that kernel manages about 309 GFLOP/s, roughly 1.3% of what cuBLAS reaches on the same GPU and the same problem. The gap is not arithmetic: neighboring threads in a warp end up reading rows of A that sit nowhere near each other in memory, so almost none of those loads can be combined into a single wide transaction. The worklog’s first kernel is the whole algorithm in fifteen lines, launched with one thread per entry of C: > Figure. sgemm_naive, from How to Optimize a CUDA Matmul Kernel Reassigning which thread owns which output element so that threads in a warp read consecutive addresses, letting the hardware coalesce those reads, already lifts throughput to about 1986.5 GFLOP/s with no other change. The next step is **tiling** (Copying a block-sized chunk of data from slow memory into fast on-chip memory once, then reusing that on-chip copy across many arithmetic operations before moving on to the next chunk.): the kernel stages a block-sized tile of A and a tile of B into shared memory once, and every thread in the block reads back out of that on-chip copy instead of returning to global memory for each partial product, which pushes throughput to about 2980.3 GFLOP/s. The heart of that shared-memory kernel is the loop every later version elaborates: stage a tile of A and a tile of B, synchronize, accumulate, synchronize, advance to the next tile along the reduction dimension: > Figure. Shared-memory tiling, from How to Optimize a CUDA Matmul Kernel The two barriers are the cooperation cost the transpose exercise introduced, now guarding both directions: the first keeps any thread from computing against a tile another thread has not finished staging, and the second keeps a fast thread from overwriting the tile with the next chunk while a slower one is still reading the current chunk. At a block size of 32 the two tiles occupy 8KB of the 48KB of shared memory a block can address on this GPU. Even with tiling, each thread is still computing exactly one output element, so most instructions in the inner loop are shared-memory loads rather than the fused multiply-adds actually doing the work; a profiler shows warps repeatedly stalling in the “Stall MIO Throttle” state, waiting on the memory pipe rather than on arithmetic. **Register blocking** (Giving each thread several output elements to accumulate in its own registers instead of one, so a single value read from shared memory gets reused across many fused multiply-adds rather than just one.) fixes that ratio directly: giving each thread a small 1D tile of outputs held in registers moves throughput to 8474.7 GFLOP/s, and a 2D tile of outputs per thread reaches 15971.7 GFLOP/s, 68.7% of cuBLAS. Vectorized memory instructions, autotuned tile sizes, and tiling at the warp level close most of what remains, reaching 21779.3 GFLOP/s, 93.7% of cuBLAS, without a tensor core in sight. 309.0GFLOP/s Naive kernel (1.3% of cuBLAS) 21779.3GFLOP/s Warptiled kernel (93.7% of cuBLAS) 23249.6GFLOP/s cuBLAS (same GPU, same problem) None of these kernels is limited by **occupancy** (The fraction of the warps an SM could run at once that are actually resident, capped by whichever per-block resource, threads, shared memory, or registers, runs out first.) in the way a first guess might suggest: the register-blocked kernel above fits only one block per SM and still reaches 66% occupancy, and pushing that number higher would not by itself close the remaining gap to cuBLAS. What separates a kernel like this from a vendor library is mostly the same tiling idea applied recursively, at the block, warp, and instruction level. Libraries such as CUTLASS describe those nested tiles as compositions of [layouts](https://learn-kernels.com/chapters/reading#matrix-multiplication), a shape paired with a stride, and talk about a matrix being “K-major” when it is stride-1 along the reduction dimension, rather than relying on the row-major and column-major vocabulary borrowed from BLAS. The case against occupancy-first reasoning is older than any GPU in this chapter. In 2008, Volkov and Demmel benchmarked dense linear algebra across four NVIDIA GPUs and built an SGEMM that sustained 58 to 60% of each chip’s peak where NVIDIA’s own CUBLAS 1.1 sustained 36 to 44%, and the design contradicted the official guidance of the era point by point: instead of many threads, shared memory as the primary storage, and long vectors, their kernel kept each 64×16 output block of C entirely in registers, staged only B’s block through shared memory, and ran short 64-element vector threads. The paper’s method is what survived: measure, then reason. Varying the thread count showed the code reaching 32, 49, 58, and 59% of peak at one to four threads per core, and on the GTX280 those four threads correspond to 25% occupancy, from which the authors conclude that one should not over-optimize for occupancy, though extremely low occupancy can also hurt. Cycle counting on the disassembled binaries located the real bound in instruction throughput: CUBLAS ran twice as many warps yet was 1.6× slower, because keeping both input blocks in shared memory forced an extra register move for every two multiply-adds, diluting the multiply-add share of its inner loop to 56% of instructions against 82% in theirs. > Figure. Where a matmul tile lives. One thread block computes one tile of C. The band of A rows and B columns it needs is staged tile by tile into shared memory, and each thread accumulates a small register tile of outputs. Shaded regions mark the data for the highlighted C tile.Illustrative numbers Source Kernel-by-kernel numbers, the optimization sequence, and both code listings from [How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance](https://siboehm.com/articles/22/CUDA-MMM). The register-blocked SGEMM, its occupancy measurements, and the CUBLAS instruction-mix analysis from Volkov and Demmel’s [Benchmarking GPUs to Tune Dense Linear Algebra](https://mc.stanford.edu/cgi-bin/images/6/65/SC08_Volkov_GPU.pdf). The layout and tiler vocabulary is from NVIDIA’s [CuTe GEMM tutorial](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cute/0x_gemm_tutorial.html). More on tile programming and layout algebra in [the reading list](https://learn-kernels.com/chapters/reading#matrix-multiplication). --- # Tensor cores and low precision · Kernel optimization Every kernel in the previous section topped out well short of the GPU’s advertised peak because none of them touched a **tensor core** (A hardware unit that performs a small matrix multiply-accumulate as a single instruction, instead of a sequence of individual fused multiply-adds.). On the same A6000 used above, switching cuBLAS from plain fp32 to TF32 or BF16 precision, so it can dispatch to tensor cores, raises its measured throughput by 2.5× and 3.5× respectively, entirely from using different hardware for the same multiply-accumulate. Going narrower than bf16 means an 8-bit floating-point format. The interchange format published by NVIDIA, Arm, and Intel defines two 8-bit encodings rather than one: **E4M3** (An 8-bit floating-point encoding with a 4-bit exponent and 3-bit mantissa, recommended for weight and activation tensors.), with a 4-bit exponent and 3-bit mantissa, and E5M2, with a 5-bit exponent and 2-bit mantissa. The two trade off differently: E4M3 gives up representing infinity and most NaN bit-patterns to extend its range, topping out at a maximum normal value of 448, while E5M2 keeps full IEEE 754 behavior for those special values and reaches a maximum normal value of 57,344. The paper’s own recommendation is E4M3 for weight and activation tensors and E5M2 for gradients, since gradients tend to need the wider dynamic range. > Figure. Two ways to spend eight bits. Both FP8 formats spend one bit on the sign and split the remaining seven differently. E4M3 gives up infinity and most NaN bit-patterns to push its maximum normal value to 448; E5M2 keeps full IEEE 754 special values and reaches 57,344, the wider dynamic range the paper recommends for gradients. Exponent bits are shaded, mantissa bits highlighted.Constants from the source FP8’s narrow range means a tensor has to be rescaled before it is cast down: a **scaling factor** (A per-tensor multiplier applied before casting a higher-precision value down to FP8, chosen so the tensor's largest magnitude lands close to the format's representable maximum.) is applied first so a tensor’s largest values sit near 448 or 57,344 instead of overflowing or clustering near zero, then removed again once the FP8 matrix multiply has produced a higher-precision result. NVIDIA’s TransformerEngine library implements this bookkeeping so it does not have to be hand-rolled per model: its DelayedScaling recipe wraps a forward pass in FP8 at a chosen format, E4M3 by default, and it targets Hopper, Ada, and Blackwell GPUs, the same generations whose tensor cores can execute FP8 matrix multiplies directly. Because the boundary between compute-bound and memory-bound work depends on flops per byte, halving a tensor’s size in memory does two things at once: it doubles the arithmetic a tensor core can push through in the same span of time, and it halves the bytes that have to move to feed it. That is why FP8, and the narrower MXFP8 and NVFP4 formats TransformerEngine has since added for Blackwell, keep showing up wherever a kernel is waiting on bytes rather than flops, the same trade-off the KV cache made in the previous chapter, applied to the weights and activations themselves. Feeding tensor cores at these rates also changed how data moves. The Hopper generation added the **Tensor Memory Accelerator** (A hardware unit introduced in NVIDIA's Hopper architecture that copies tiles of multi-dimensional arrays between global and shared memory asynchronously, driven by a descriptor rather than by per-thread address arithmetic.) (TMA), a dedicated unit that copies tiles of multi-dimensional arrays between global and shared memory asynchronously. A single thread issues the copy against a descriptor that carries the tensor’s shape and strides, so address computation and out-of-bounds predication happen in one place instead of consuming registers in every thread, and the copy’s asynchrony is what makes warp-specialized kernel schedules practical. TMA also emits the swizzled shared-memory layouts Hopper’s tensor-core instructions expect, layouts too intricate to be worth loading by hand. The tensor-core instruction itself went asynchronous in the same generation. Hopper’s `wgmma.mma_async` family is issued not by a thread or a warp but by a **warp group** (A group of four warps, 128 threads, that cooperatively executes Hopper's asynchronous wgmma tensor-core instructions and jointly holds the accumulator in its registers.) of four warps, because the accumulator no longer fits anywhere smaller: an m64n16k16 multiply accumulates a 64×16 tile of fp32 values, 1,024 registers, where a single thread may hold at most 256. Spread across the warp group’s 128 threads that is 8 registers each, and the instruction variants scale the N dimension from 8 all the way to 256. A worklog in the same spirit as the A6000 one above shows what those two units are worth on an H100. The A6000-style kernel, ported to bf16, reaches 32 TFLOP/s where cuBLAS reaches 716, because on Hopper the tensor cores are not optional. Rewriting the inner loop around TMA loads and wgmma lifts it to 317 TFLOP/s, larger 128×128 tiles to 423, and splitting each block into a producer warp group that issues TMA loads into a circular buffer of shared-memory tiles and a consumer warp group that drains it through the tensor cores, the pattern called warp specialization, to 498. The finished kernel, several optimizations later, outruns cuBLAS by 7% at N=4096. That machinery is no longer exotic: DeepSeek’s open-source DeepGEMM is a tensor-core kernel library for exactly these SM90 and SM100 GPUs, covering FP8, FP4, and BF16 GEMMs plus fused mixture-of-experts kernels, compiled at runtime by a lightweight JIT module so installation needs no CUDA compilation. Its README describes borrowing concepts from CUTLASS and CuTe while avoiding heavy reliance on their templates, keeping a small set of core kernel functions readable as a learning resource, and reaching up to 1550 TFLOP/s on an H800 while matching or exceeding expert-tuned libraries across matrix shapes. Source E4M3/E5M2 bit layout, exponent bias, and recommended usage from [FP8 Formats for Deep Learning](https://arxiv.org/abs/2209.05433). TransformerEngine’s API and supported GPU generations from the [TransformerEngine](https://github.com/NVIDIA/TransformerEngine) README. Tensor core speedup figures from the same A6000 benchmark as the matrix multiplication section. TMA’s design and its single-threaded descriptor-driven copies from Colfax’s [CUTLASS tutorial on TMA](https://research.colfax-intl.com/tutorial-hopper-tma/). The wgmma register arithmetic and the H100 kernel progression from [Outperforming cuBLAS on H100: a Worklog](https://cudaforfun.substack.com/p/outperforming-cublas-on-h100-a-worklog). DeepGEMM’s scope and performance from the [DeepGEMM](https://github.com/deepseek-ai/DeepGEMM) README. More formats and hardware detail in [the reading list](https://learn-kernels.com/chapters/reading#tensor-cores-low-precision). --- # Programming models and profiling Chapter 1 built a kernel out of threads, blocks, and warps: the programmer decides what each thread does. Triton and CUTLASS both reject that unit of work and ask the programmer to describe tiles instead, leaving the mapping onto threads to a compiler. Once the compiler is making those decisions, profiling and correctness checking stop being optional reading of the source and become separate tools in their own right. ## 4 sections 1. [3.1Triton](https://learn-kernels.com/chapters/programming-models/triton) 2. [3.2CUTLASS, CuTe, and CUDA Tile](https://learn-kernels.com/chapters/programming-models/cutlass-cute-and-cuda-tile) 3. [3.3Other hardware stacks](https://learn-kernels.com/chapters/programming-models/other-hardware-stacks) 4. [3.4Profiling, benchmarking, and correctness](https://learn-kernels.com/chapters/programming-models/profiling-and-correctness) --- # CUTLASS, CuTe, and CUDA Tile · Programming models and profiling CUTLASS is NVIDIA’s own answer to the same problem, aimed squarely at one computation: it is “a collection of abstractions for implementing high-performance matrix-matrix multiplication (GEMM) and related computations at all levels and scales within CUDA,” decomposing that work into reusable, modular components. Through CUTLASS 2.x, that decomposition mirrored the GPU’s own hierarchy (thread, warp, threadblock) directly. CUTLASS 3.0 broke that mirroring on purpose: Hopper’s warp-group-wide instructions do not correspond to any single warp or thread-level concept, so tying the library’s structure to one generation’s hardware layout kept breaking on the next generation. What replaced it is CuTe, described as “a collection of C++ CUDA template abstractions for defining and operating on hierarchically multidimensional layouts of threads and data.” A **CuTe layout** (A mapping from a logical, multidimensional coordinate to a linear memory offset, expressed as a shape and a stride so that CuTe can compose and manipulate it with ordinary function operations.) is, at bottom, a function from integers to integers, and CuTe builds a full **layout algebra** (Operations, including functional composition, product (build a larger layout by repeating a smaller one), and divide (partition one layout according to another), for combining and manipulating CuTe layouts as if they were ordinary functions.) on top of that idea: composition, product, and divide let a kernel author build a thread-to-data mapping for an entire GEMM tile out of a handful of primitive layouts, instead of hand-writing a new iterator type for every architecture-specific access pattern. CUTLASS 3.0 replaced most of its 2.x-era named iterator types with this one vocabulary type, on the argument that a mapping expressed as an algebra can be checked at compile time: “if the code compiles, it’s probably correct.” The shape:stride notation reads mechanically. In the layout algebra documentation’s worked example B = (4,3):(3,1), the shape (4,3) says coordinates run over a 4 by 3 grid, and the stride (3,1) says a step in the first coordinate advances the underlying index by 3 while a step in the second advances it by 1: the doc’s own table evaluates B(1,0) to 3 and B(0,1) to 1. Because a layout is just a function from integers to integers, the same function can often be written more than one way; the documentation notes that a column-major layout like (\_2,\_4):(\_1,\_2) “acts identically to” \_8:\_1 for 1-D coordinates, and CuTe’s `coalesce` operation is the simplifier that finds the smaller spelling without changing the function: > Figure. coalesce, from the CuTe layout algebra documentation The algebra’s central operation is functional composition, R := A ∘ B with R(c) = A(B(c)), and the documentation calls it “the core of CuTe,” used “in just about every higher-level operation.” Its worked example composes A = (6,2):(8,2) with B = (4,3):(3,1), evaluates all twelve outputs by hand, and lands on the observation the whole library rests on: the result is itself a layout, ((2,2),3):((24,2),8), and it is compatible with B, meaning every coordinate of B is also a valid coordinate of the result. Selecting every third element of one layout through another never leaves the algebra, which is what lets the product and divide operations, and ultimately a whole thread-to-data partition for a GEMM tile, be built out of compositions and checked at compile time. > Figure. The producer-consumer pipeline. CUTLASS's asynchronous pipeline. Producer threads load tiles into a circular list of shared-memory stages while consumer threads compute on tiles one slot behind, the two sides synchronized through acquire, commit, wait, and release barrier operations. In the highlighted slot the load of tile k+1 runs while tile k is being computed: the software pipelining CUTLASS calls critical to hiding the latency of global memory loads.Illustrative numbers CUDA Tile IR, introduced in CUDA 13.1, takes the tile idea a level higher than either CUTLASS or Triton do on their own. Where CUTLASS and CuTe still hand a C++ (or now Python, via CuTe DSL) programmer explicit control over the thread-to-data layout, a **tile block** (The basic unit of execution in a Tile IR program: a single logical thread that computes over a whole multidimensional tile of data, with the mapping onto actual hardware threads left entirely to the compiler.) in Tile IR is written with no thread-to-data mapping at all; the compiler decides how a tile block’s work lands on real SM threads, the memory hierarchy, and tensor cores. Tile IR is positioned as a compilation target beneath higher-level DSLs, not a replacement for CUTLASS or Triton themselves, which is exactly the role it plays for Triton’s Tile IR backend described in the previous section. The programming model documentation shows what that looks like as actual code. A Tile IR program is a module of tile kernels declared with `entry`, and every value in a kernel is a tensor whose rank, shape, and element type are statically known; rank-0 tensors are scalars, and global memory is only ever reached through tensors built from pointer parameters. Its first worked example, a 128-element vector addition, spends most of its lines constructing a tile of 128 pointers from a scalar base pointer: an `iota` builds the offset vector 0 through 127, a `reshape` and `broadcast` replicate the base pointer across the tile, and an `offset` adds the two. The arithmetic itself is then three statements: > Figure. Tile IR vector addition, from the Tile IR programming model documentation The documentation’s summary of that kernel is the model in one sentence: “this code is written from a single thread of control, but its level of parallelism will be determined by the compiler.” Scaling past one tile reuses CUDA’s launch shape rather than replacing it: tile blocks group into a 1-d, 2-d, or 3-d tile grid, the grid size set at launch determines how many tile blocks run, and each block queries its position with `get_tile_block_id` and the grid’s dimensions with `get_num_tile_blocks`, the role `blockIdx` plays in chapter 1’s kernels, one level of hierarchy up. Source GEMM hierarchy and correctness-by-default framing from [CUTLASS 3.0 Design](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cutlass_3x_design.html). Layout algebra from [CuTe Layout Algebra](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cute/02_layout_algebra.html). Producer and consumer roles, stages, and barrier operations from the [CUTLASS pipeline documentation](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/pipeline.html). Tile IR and tile blocks from the [Tile IR introduction](https://docs.nvidia.com/cuda/tile-ir/latest/sections/introduction.html) and [programming model](https://docs.nvidia.com/cuda/tile-ir/latest/sections/prog_model.html) docs. More in [further reading](https://learn-kernels.com/chapters/reading#cutlass-cute-tile). --- # Other hardware stacks · Programming models and profiling Nothing about the tile abstraction is NVIDIA-specific, and the clearest evidence comes from Pallas, JAX’s extension for “kernel programming for both GPUs and TPUs using a Triton-like model.” Its design document makes a pointed observation about Triton itself: Triton “exposes a TPU-like programming model to users, i.e. writing programs for tiles of arrays in L1-cache,” and yet is specialized enough to GPU that it cannot be compiled directly for TPU; Triton’s atomic operations, built for parallel writes, “don’t necessarily make sense on TPU.” So Pallas keeps only the tile-based programming model, abstracts the platform details behind it, and lowers the same kernel per backend: to Mosaic GPU (formerly Triton) on GPUs, and to Mosaic on TPUs. The tile is the portable part; what surrounds it is not. AMD’s stack arrived at the same shape from two directions. Composable Kernel, in AMD’s own description, “provides a programming model for writing performance-critical kernels for machine learning workloads across multiple architectures,” written in general purpose kernel languages such as HIP C++ and resting on two ideas: “a tile-based programming model” and a complexity-reduction technique it calls **Tensor Coordinate Transformation** (A technique Composable Kernel pairs with its tile-based programming model to reduce the algorithmic complexity of complex machine learning operators.). It is layered from templated tile operators up through a client API, a decomposition recognizable from CUTLASS. HipKittens, a research library of C++ tile primitives for AMD GPUs, states this section’s through-line as an experimental finding: porting from its NVIDIA sibling ThunderKittens, the core tile and bulk compute interfaces carry over, while the decisions around memory access patterns, compute and memory scheduling, and thread block ordering within the chiplet architecture differ. Its tiles are sized to the tensor core units, and its two core scheduling patterns, 8-wave ping pong and 4-wave interleave, are organized around CDNA’s waves rather than the warps of [chapter 1’s hardware](https://learn-kernels.com/chapters/hardware). Where those AMD kernels end up is its own data point. AITER, the AI Tensor Engine for ROCm, is “AMD’s high-performance AI operator library, providing optimized GPU kernels for inference and training workloads on ROCm,” positioned as “a unified collection of production-ready operators that framework developers can integrate directly into their stacks.” What makes it relevant here is its backend list: the same operator can be served by a Triton kernel, a Composable Kernel implementation, or hand-tuned assembly, behind one C++ or Python API. The programming models this chapter has covered are not competing endpoints so much as interchangeable suppliers to an operator library, and AITER’s README notes it is the default attention backend for vLLM on ROCm, which is where a kernel chosen by that machinery actually meets production traffic. AWS Trainium is the strongest test of the claim, because the accelerator underneath is not a GPU. NKI, the Neuron Kernel Interface for writing kernels that run on Trainium devices, still hands the programmer a tile: a kernel allocates tiles in on-chip SBUF memory, DMA-copies inputs into them from HBM, and checks that a tile’s first dimension fits within the on-chip tile size limit before operating on whole tiles at a time. What changes is everything around the tile. NKI kernels use a **sequential programming model** (NKI's execution contract: the logical order of operations follows the syntactic order of statements in the kernel, and the compiler may only reorder operations that have no data dependencies.) rather than a grid of parallel blocks, and its nki.isa functions are “designed to expose the underlying hardware capabilities in as direct a way as possible,” each call running one operation on one of the device’s compute engines while the compiler unrolls, inlines, and resolves everything else ahead of time. Across all four stacks the tile survives even where warps, blocks, and threads do not; what each stack builds around the tile tracks what its particular silicon makes cheap or expensive. See [further reading](https://learn-kernels.com/chapters/reading#other-hardware-stacks) for each stack’s documentation. Source Pallas framing, the Triton-on-TPU argument, and lowering targets from the [Pallas design document](https://docs.jax.dev/en/latest/pallas/design/design.html). Composable Kernel description and layering from the [ROCm/composable\_kernel](https://github.com/ROCm/composable_kernel) README. HipKittens primitives and scheduling patterns from the [HazyResearch/HipKittens](https://github.com/HazyResearch/HipKittens) README. AITER description, backends, and vLLM integration from the [ROCm/aiter](https://github.com/ROCm/aiter) README. NKI model and quotes from the [NKI language guide](https://awsdocs-neuron.readthedocs-hosted.com/en/latest/nki/programming_model.html). --- # Profiling, benchmarking, and correctness · Programming models and profiling A blocked or tiled program is, by design, no longer readable as a literal description of what one thread does. That means neither speed nor correctness can be checked by re-reading the source the way chapter 1’s per-thread CUDA code could be: both need to be measured against what actually ran on the GPU. NVIDIA splits those two jobs across separate tools. Nsight Compute answers “how fast, and where is the time going”; Compute Sanitizer answers “did it read or write anything it shouldn’t have.” Nsight Compute profiles by inserting measurement libraries into the running application process, intercepting its communication with the CUDA driver, and collecting metrics whenever it detects a kernel launch. Its most direct performance readout is **occupancy** (The ratio of the number of active warps per streaming multiprocessor to the maximum number of possible active warps on that multiprocessor.): “occupancy is the ratio of the number of active warps per multiprocessor to the maximum number of possible active warps,” and a large gap between the theoretical and the achieved value “typically indicates highly imbalanced workloads.” Its GPU Speed Of Light section reports the achieved percentage of compute and memory throughput against the theoretical maximum for each, which is the same flop-bound-versus-memory-bound question from chapter 1, now measured per kernel instead of estimated from a datasheet. Collecting all of that is not free, and not always possible in one pass: the GPU has a limited number of hardware counters it can read concurrently, and some metrics require patch-based software counters whose overhead would itself distort the measurement. When a requested set of metrics cannot be collected together, Nsight Compute uses **kernel replay** (Running a kernel launch more than once so that different subsets of requested metrics can each be collected in a separate pass, saving and restoring the memory the kernel writes between passes.) , saving all memory the kernel can reach before the first pass and restoring whatever it wrote before each subsequent one. A profiled run is therefore not the same execution as an unprofiled one, which is a reason to treat wall-clock numbers taken while profiling as approximate. > Figure. One command, four checkers. Compute Sanitizer runs the same kernel under one of four instrumented modes, each watching for a failure class a profiler never checks: bad addresses, shared-memory races, uninitialized reads, and misused barriers.Illustrative numbers Compute Sanitizer is the correctness half, shipped as part of the CUDA toolkit. Its own documentation states the reason it exists plainly: “every programmer invariably encounters memory access errors and thread ordering hazards that are hard to detect and time consuming to debug,” and “the number of such errors increases substantially when dealing with thousands of threads.” None of its four tools measure speed at all; they instrument actual memory and synchronization behavior and report a violation the moment one occurs, which is precisely the check a profiler has no reason to perform. A kernel can report high occupancy and near-peak memory throughput in Nsight Compute while still failing racecheck, because occupancy says nothing about whether two warps are racing on the same shared-memory address. Source Profiling mechanics, occupancy, and kernel replay from the [Nsight Compute Profiling Guide](https://docs.nvidia.com/nsight-compute/ProfilingGuide/). Compute Sanitizer tools and quotes from the [Compute Sanitizer documentation](https://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/). More in [further reading](https://learn-kernels.com/chapters/reading#profiling). ### Beyond a single kernel Nsight Compute answers its questions one kernel at a time, and that is also its blind spot: it cannot say whether the kernel it is dissecting is the one worth dissecting. That is Nsight Systems’s job. By default it “collects a profile over the entire run of your application,” laying out CPU threads, CUDA API calls, and GPU activity on one timeline, and its guide recommends narrowing collection to the performance-critical region with `cudaProfilerStart()` and `cudaProfilerStop()` rather than profiling a test harness’s setup and validation. NVTX markers and ranges added to the application appear in the Timeline View and are projected onto the GPU timeline, “allowing you to see what GPU activity was launched within each CPU range.” The workflow this implies runs in one direction: find the expensive kernel, or the gap where no kernel is running, on the Systems timeline first, then drill into that kernel with Compute. Even with the right kernel in hand, a number measured casually is not a number worth reporting. NVIDIA’s GEMM performance measurement methodology, published with the CUTLASS documentation, prescribes what a reproducible benchmark harness must do: separate warmup and profiling loops, delimited by `cudaProfilerStart/Stop` or CUDA events; no allocations, copies, or extra kernels between launches; and buffer rotation, cycling each tensor through duplicate buffers whose total footprint is at least twice the L2 capacity, so every iteration starts from DRAM instead of inheriting the last iteration’s cache. It distinguishes fixed frequency tests, which measure architectural and software efficiency at a locked clock, from fixed power tests, which mimic the dynamically-scaled clocks of real-world use and vary far more. The scale it recommends is sobering for anyone timing a kernel in a loop of ten: earlier GEMM studies needed more than 1,000 iterations for stable results on a 4096 by 4096 by 4096 problem, and for large GEMMs on Blackwell the document uses 10,000 warmups, 4,000 profiling iterations, and a second of cool-down between tests. AMD’s counterpart to the kernel drill-down is ROCm Compute Profiler, also known by its package name rocprofiler-compute: “a kernel-level profiling tool for machine learning and high performance computing (HPC) workloads running on AMD Instinct GPUs.” It is built on ROCprofiler-SDK to monitor hardware performance counters, acquires those counters via application replay, and runs accelerator-specific microbenchmarks to build hierarchical roofline data. Its analysis vocabulary maps almost term for term onto the NVIDIA one: system Speed-of-Light and hardware-block-level Speed-of-Light summaries, memory chart and roofline analysis, and baseline comparisons between runs, all for the CDNA GPUs whose waves and chiplets HipKittens schedules around in the previous section. Source Whole-run collection, focused profiling, and NVTX projection from the [Nsight Systems User Guide](https://docs.nvidia.com/nsight-systems/UserGuide/). Benchmark requirements, test types, buffer rotation, and iteration counts from the [GEMM performance measurement methodology guidelines](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/gemm_performance_measurement_methodology_guidelines.html). ROCm Compute Profiler description and features from its [documentation](https://rocm.docs.amd.com/projects/rocprofiler-compute/en/latest/). --- # Triton · Programming models and profiling Triton is, in its own words, “a language and compiler for writing highly efficient custom Deep-Learning primitives,” aimed at giving programmers “higher productivity than CUDA” and “higher flexibility than other existing DSLs.” It grew out of a MAPL 2019 paper by Tillet, Kung, and Cox, and the premise stated in its own documentation is that **blocked programs** (A style of parallel program in which the code the compiler works with operates on blocks of data at a time, rather than on one thread's worth of data at a time.) can produce compute kernels competitive with hand-written CUDA, while staying far more flexible than prior compiler frameworks like Halide or TVM. The distinction Triton draws against CUDA is precise. CUDA, in its own framing, is “Scalar Program, Blocked Threads”: a programmer writes one thread’s worth of scalar arithmetic, and the hardware replicates it across a block of threads, each computing one output element of a matrix multiply. Triton inverts this to “Blocked Program, Scalar Threads”: the program itself loops over tiles of the output, stepping by block sizes MB, NB, and KB, and every statement inside that loop already operates on whole blocks of A, B, and the accumulator at once. There is no per-thread code to write, because the block is the unit the language exposes. That inversion only pays off because of what the compiler does with it. Triton relies on **block-level data-flow analysis** (Scheduling iteration blocks statically, based on the control- and data-flow structure of the program, rather than at the level of individual threads.) to turn a blocked program into a scheduled one, and that analysis is what drives automatic memory coalescing, thread swizzling, prefetching, vectorization, tensor-core-aware instruction selection, shared-memory allocation and synchronization, and asynchronous copy scheduling. Those are exactly the decisions a CUDA programmer makes by hand at the thread-block level described in chapter 1; Triton moves them into the compiler and asks the programmer to reason about tiles instead. > Figure. Scalar program versus blocked program. On the left, CUDA's framing: you write one thread's worth of scalar arithmetic, the hardware replicates it across the block, and each thread computes one output element. On the right, Triton's: every statement already operates on a whole tile, and the coalescing, swizzling, shared-memory, and thread-mapping decisions a CUDA programmer makes by hand move into the compiler.Illustrative numbers Triton’s own first tutorial makes the style concrete with the same computation chapter 1 used to introduce CUDA: vector addition. Where `vecAdd` gave every thread exactly one element, the Triton kernel is written per program, and each program handles `BLOCK_SIZE` elements at once. The tutorial’s comments call out the two moves that define the model: the kernel asks `tl.program_id` which of the parallel program instances it is, and `offsets` “is a list of pointers,” a whole block of addresses that `tl.load` reads in one statement: > Figure. add_kernel, from the Triton vector addition tutorial Nothing in the kernel names a thread. The `mask` guards the loads and stores against out-of-bounds accesses when the vector length is not a multiple of the block size, which is the blocked model’s replacement for chapter 1’s per-thread bounds check. The launch side is what the tutorial calls an SPMD grid, “analogous to CUDA launch grids”: here a 1D grid of `cdiv(n_elements, BLOCK_SIZE)` program instances, with each `torch.Tensor` argument implicitly converted into a pointer to its first element. Everything below the program, the warps, the shared memory, the vector widths, is the compiler’s problem. That same tile abstraction is why Triton and CUTLASS keep meeting in the same place. NVIDIA has since built a backend, Triton-to-TileIR, that lets Triton programs compile directly to CUDA Tile IR (covered next) instead of PTX, preserving Triton’s tile semantics rather than lowering them to per-thread SIMT instructions first. See [further reading](https://learn-kernels.com/chapters/reading#triton) for tutorials and the original paper. Source Blocked-program framing and the matmul comparison from the [Triton programming guide](https://triton-lang.org/main/programming-guide/chapter-1/introduction.html), chapter 1. Project description from the [triton-lang/triton](https://github.com/triton-lang/triton) README. The add\_kernel listing, mask, and grid mechanics from the [vector addition tutorial](https://triton-lang.org/main/getting-started/tutorials/01-vector-add.html), with comments trimmed. Triton-to-TileIR from [NVIDIA’s CUDA Tile IR backend for Triton](https://developer.nvidia.com/blog/advancing-gpu-programming-with-the-cuda-tile-ir-backend-for-openai-triton/) post. --- # Further reading The primary sources. Anything worth understanding properly is better read there than summarized here, and most of these are readable in an evening. ## Start here The minimum mental model, read in order before anything else. - [*How to Scale Your Model: Inference*One request from prefill through decode, with batching, KV memory, and parallelism.](https://jax-ml.github.io/scaling-book/inference/) - [“Attention Is All You Need”The transformer computation that the rest of the list optimizes.](https://arxiv.org/abs/1706.03762) - [*CUDA C++ basics*The shortest official introduction to the CUDA execution model.](https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/intro-to-cuda-cpp.html) - [*Programming Massively Parallel Processors*The main textbook for GPU programming, memory, and kernel design.](https://www.elsevier.com/books/programming-massively-parallel-processors/hwu/978-0-323-91231-0) - [“Roofline: An Insightful Visual Performance Model”The compute, memory-bandwidth, and arithmetic-intensity model.](https://www2.eecs.berkeley.edu/Pubs/TechRpts/2008/EECS-2008-134.html) - [*Transformer Inference Arithmetic*FLOPs, parameter bytes, KV bytes, and communication for transformer inference.](https://kipply.github.io/blog/transformer-inference-arithmetic/) - [“Efficiently Scaling Transformer Inference”Latency, memory, and parallelism costs for large-model inference.](https://proceedings.mlsys.org/paper_files/paper/2023/file/c4be71ab8d24cdfb45e3d06dbfca2780-Paper-mlsys2023.pdf) - [“Etalon”TTFT, TPOT, goodput, and latency SLOs for generative-model serving.](https://arxiv.org/html/2407.07000) - [*GPU Mode lectures*A practical companion to work through alongside the start-here list.](https://github.com/gpu-mode/lectures) ## Programming model Threads, warps, memory hierarchy, and the CUDA execution model. - [*CUDA Programming Guide*The normative CUDA reference.](https://docs.nvidia.com/cuda/cuda-programming-guide/) - [*CUDA programming model*Threads, warps, blocks, grids, and the memory hierarchy.](https://docs.nvidia.com/cuda/cuda-programming-guide/01-introduction/programming-model.html) - [*CUDA C++ Best Practices Guide*Coalescing, shared memory, occupancy, synchronization, and optimization workflow.](https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/) - [*Hopper Tuning Guide*TMA, thread-block clusters, asynchronous execution, and Hopper-specific limits.](https://docs.nvidia.com/cuda/hopper-tuning-guide/) - [*Blackwell Tuning Guide*Tensor memory, Blackwell execution features, and architecture limits.](https://docs.nvidia.com/cuda/blackwell-tuning-guide/) ## Compilation and machine code What CUDA compiles down to, and how to read it. - [*NVCC Compiler Driver*The CUDA compilation trajectory and artifact controls.](https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/) - [*PTX ISA*NVIDIA's virtual instruction set and memory model.](https://docs.nvidia.com/cuda/parallel-thread-execution/) - [*CUDA Binary Utilities*cuobjdump and nvdisasm for inspecting GPU binaries.](https://docs.nvidia.com/cuda/cuda-binary-utilities/) - [*Understanding PTX*NVIDIA's introduction to the role of PTX between CUDA and machine code.](https://developer.nvidia.com/blog/understanding-ptx-the-assembly-language-of-cuda-gpu-computing/) ## Foundational kernel exercises The exercises everyone works through first: transpose, reduction, scan, softmax. - [*Efficient Matrix Transpose in CUDA C/C++*Coalescing, shared-memory tiling, and bank conflicts.](https://developer.nvidia.com/blog/efficient-matrix-transpose-cuda-cc/) - [*Optimizing Parallel Reduction in CUDA*Synchronization, divergence, occupancy, and instruction cost.](https://developer.download.nvidia.com/assets/cuda/files/reduction.pdf) - [“Single-pass Parallel Prefix Scan with Decoupled Look-back”A work-efficient scan with one pass over memory.](https://research.nvidia.com/sites/default/files/pubs/2016-03_Single-pass-Parallel-Prefix/nvr-2016-002.pdf) - [“Online Normalizer Calculation for Softmax”Numerically stable online softmax without materialized intermediates.](https://arxiv.org/abs/1805.02867) ## Matrix multiplication Tiling, layouts, and the reference GEMM implementations. - [“Benchmarking GPUs to Tune Dense Linear Algebra”The canonical case for reasoning from measured hardware behavior instead of occupancy alone.](https://mc.stanford.edu/cgi-bin/images/6/65/SC08_Volkov_GPU.pdf) - [*CuTe GEMM tutorial*Tiling, layouts, copies, and matrix-multiply atoms.](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cute/0x_gemm_tutorial.html) - [*CUTLASS 3.x design*The collective and kernel structure used by modern CUTLASS.](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cutlass_3x_design.html) - [*DeepGEMM*A compact production FP8 GEMM implementation for Hopper.](https://github.com/deepseek-ai/DeepGEMM) ## Direct implementation work Worklogs that build a fast kernel from scratch, in the open. - [*How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance*A matrix multiplication built from naive CUDA through shared-memory and register tiling.](https://siboehm.com/articles/22/CUDA-MMM) - [*Inside NVIDIA GPUs: Anatomy of High-Performance Matmul Kernels*Layouts, tiling, PTX, machine code, and roofline analysis.](https://www.aleksagordic.com/blog/matmul) - [*Outperforming cuBLAS on H100: A Worklog*A direct Hopper optimization worklog using tensor cores and asynchronous movement.](https://cudaforfun.substack.com/p/outperforming-cublas-on-h100-a-worklog) - [*CUTLASS Tutorial: Mastering TMA*Working kernels built around the Tensor Memory Accelerator.](https://research.colfax-intl.com/tutorial-hopper-tma/) ## Tensor cores and low precision FP8, FP4, and the formats and instructions tensor cores run. - [*OCP 8-bit Floating Point Specification*E4M3 and E5M2 formats.](https://www.opencompute.org/documents/ocp-8-bit-floating-point-specification-ofp8-revision-1-1-final-pdf) - [*OCP Microscaling Formats Specification*Shared-scale MX formats.](https://www.opencompute.org/documents/ocp-microscaling-formats-mx-v1-0-spec-final-pdf) - [*NVIDIA Transformer Engine*FP8 and FP4 transformer execution with scaling controls.](https://github.com/NVIDIA/TransformerEngine) - [*Blackwell matrix multiply instructions*tcgen05, tensor memory, and Blackwell MMA programming.](https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/mma_docs/tcgen05_programming.html) ## Attention FlashAttention and the kernels it led to. - [“FlashAttention”IO-aware exact attention.](https://arxiv.org/abs/2205.14135) - [“FlashAttention-2”Better work partitioning and parallelism.](https://arxiv.org/abs/2307.08691) - [“FlashAttention-3”Asynchronous movement and tensor-core overlap on Hopper.](https://arxiv.org/abs/2407.08608) - [“FlashAttention-4”The Blackwell attention schedule.](https://proceedings.mlsys.org/paper_files/paper/2026/file/ae8b0b5838ba510daff1198474e7b984-Paper-Conference.pdf) - [*FlashInfer*Attention and related kernels for serving workloads.](https://github.com/flashinfer-ai/flashinfer) ## Triton The blocked-program language and compiler. - [“Triton paper”The original blocked-program language and compiler design.](https://eecs.harvard.edu/~htk/publication/2019-mapl-tillet-kung-cox.pdf) - [*Triton programming guide*The official programming model.](https://triton-lang.org/main/programming-guide/chapter-1/introduction.html) - [*Triton repository*Compiler, examples, tests, and backend implementation.](https://github.com/triton-lang/triton) ## CUTLASS, CuTe, and CUDA Tile NVIDIA's tile-based kernel-authoring stack. - [*CuTe layout algebra*Layouts and layout composition.](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cute/02_layout_algebra.html) - [*CUTLASS GEMM tutorial*A GEMM expressed through CuTe layouts and atoms.](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cute/0x_gemm_tutorial.html) - [*CUTLASS pipeline documentation*Producer-consumer pipelines and asynchronous stages.](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/pipeline.html) - [*CUDA Tile IR programming model*NVIDIA's compiler-owned tile abstraction.](https://docs.nvidia.com/cuda/tile-ir/latest/sections/prog_model.html) - [*CUDA Tile repository*The current implementation and examples.](https://github.com/NVIDIA/cuda-tile) ## Other hardware stacks Kernel programming models for AMD, TPU, and Trainium. - [*ROCm Composable Kernel*AMD tiling, layout, and operator primitives.](https://github.com/ROCm/composable_kernel) - [*ROCm AITER*AMD inference and transformer operator implementations.](https://github.com/ROCm/aiter) - [*HipKittens*A tile abstraction for AMD GPUs.](https://github.com/HazyResearch/HipKittens) - [*Pallas design*The JAX kernel model for GPU and TPU backends.](https://docs.jax.dev/en/latest/pallas/design/design.html) - [*NKI programming model*The tile-level programming model for AWS NeuronCore hardware.](https://awsdocs-neuron.readthedocs-hosted.com/en/latest/nki/programming_model.html) ## Profiling, benchmarking, and correctness Tools for measuring and verifying what a kernel actually does. - [*Nsight Systems User Guide*System timelines, CPU-GPU interaction, and distributed traces.](https://docs.nvidia.com/nsight-systems/UserGuide/) - [*Nsight Compute Profiling Guide*Kernel metrics, sections, replay, and roofline analysis.](https://docs.nvidia.com/nsight-compute/ProfilingGuide/) - [*Compute Sanitizer*Memory, race, initialization, and synchronization checks.](https://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/) - [*CUTLASS GEMM measurement methodology*Reproducible GEMM benchmarking.](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/gemm_performance_measurement_methodology_guidelines.html) - [*ROCm Compute Profiler*AMD performance counters and roofline analysis.](https://rocm.docs.amd.com/projects/rocprofiler-compute/en/latest/) ## Scheduling and continuous batching Iteration-level scheduling and the engines that implement it. - [“Orca”Iteration-level scheduling for autoregressive serving.](https://www.usenix.org/conference/osdi22/presentation/yu) - [“PagedAttention and vLLM”Paged KV allocation and continuous batching.](https://arxiv.org/html/2309.06180) - [“Sarathi-Serve”Chunked prefills that reduce interference with decode.](https://www.usenix.org/system/files/osdi24-agrawal.pdf) - [“SGLang”Prefix reuse, structured programs, and a serving runtime.](https://arxiv.org/html/2312.07104) - [*vLLM*One of the main production engine implementations.](https://github.com/vllm-project/vllm) - [*SGLang repository*One of the main production engine implementations.](https://github.com/sgl-project/sglang) - [*TensorRT-LLM*One of the main production engine implementations.](https://github.com/NVIDIA/TensorRT-LLM) ## KV cache systems Shrinking, quantizing, and distributing the KV cache. - [“Grouped-Query Attention”Fewer key-value heads and a smaller KV cache.](https://arxiv.org/abs/2305.13245) - [“DeepSeek-V2”Multi-head latent attention and compressed KV state.](https://arxiv.org/abs/2405.04434) - [“KIVI”KV quantization with separate treatment for keys and values.](https://proceedings.mlr.press/v235/liu24bz.html) - [“CacheGen”KV compression for transfer.](https://cs.stanford.edu/~keithw/sigcomm2024/sigcomm24-final1571-acmpaginated.pdf) - [“Mooncake”A distributed KV cache and data plane.](https://www.usenix.org/conference/fast25/presentation/qin) ## Quantization Weight and activation quantization for inference. - [“GPTQ”One-shot second-order weight quantization.](https://arxiv.org/abs/2210.17323) - [“SmoothQuant”W8A8 execution by moving quantization difficulty from activations into weights.](https://proceedings.mlr.press/v202/xiao23c.html) - [“AWQ”Low-bit weight-only inference with salient-weight protection.](https://proceedings.mlsys.org/paper_files/paper/2024/file/42a452cbafa9dd64e9ba4aa95cc1ef21-Paper-Conference.pdf) ## Speculative decoding Drafting and verifying multiple tokens per step. - [“Fast Inference from Transformers via Speculative Decoding”Exact sampling with a draft model.](https://proceedings.mlr.press/v202/leviathan23a.html) - [“Accelerating Large Language Model Decoding with Speculative Sampling”The parallel formulation and analysis.](https://arxiv.org/abs/2302.01318) - [“Medusa”Multiple prediction heads on the target model.](https://arxiv.org/html/2401.10774) - [“EAGLE”Feature-level drafting.](https://proceedings.mlr.press/v235/li24bt.html) ## Structured decoding and fairness Constrained generation and scheduling fairness. - [“Guiding LLMs the Right Way”Constrained decoding without changing the intended token distribution.](https://proceedings.mlr.press/v235/beurer-kellner24a.html) - [“XGrammar”A fast grammar engine for structured generation.](https://proceedings.mlsys.org/paper_files/paper/2025/file/5c20ca4b0b20b0bd2f1d839dc605e70f-Paper-Conference.pdf) - [“Fairness in Serving Large Language Models”Fair scheduling when request sizes are different and unknown.](https://arxiv.org/html/2401.00588) ## Long context and multimodal inference Sparse attention and multimodal engine support at long context. - [“Ring Attention”Exact distributed attention by circulating KV blocks around a device ring.](https://arxiv.org/abs/2310.01889) - [“MInference 1.0”Dynamic sparse patterns for long-context prefill on existing models.](https://arxiv.org/abs/2407.02490) - [“Native Sparse Attention”A model trained with a hardware-aligned sparse attention hierarchy.](https://arxiv.org/abs/2502.11089) - [*vLLM multimodal inputs*Current engine support for text, image, audio, and video inputs.](https://docs.vllm.ai/en/latest/features/multimodal_inputs.html) ## Parallelism, collectives, and topology Tensor and pipeline parallelism, and the interconnects underneath. - [“Megatron-LM”Tensor and pipeline parallelism for transformer models.](https://arxiv.org/abs/1909.08053) - [*NCCL*NVIDIA's collective communication implementation.](https://github.com/NVIDIA/nccl) - [*Multi-node NVLink Systems Tuning Guide*NVLink and InfiniBand topology in GB200 NVL systems.](https://docs.nvidia.com/multi-node-nvlink-systems/multi-node-tuning-guide/) - [*UALink 1.0 Specification*An open scale-up interconnect.](https://ualinkconsortium.org/wp-content/uploads/2025/04/UALink200_Specification_v1.0_Evaluation_Copy.pdf) - [*Ultra Ethernet 1.0.3 Specification*The scale-out transport specification.](https://ultraethernet.org/wp-content/uploads/sites/20/2026/08/UE-Specification-1.0.3.pdf) ## Mixture-of-experts serving Routing, dispatch, and load balancing for MoE inference. - [“DeepSeek-V3”Routed experts, shared experts, and the model-system design.](https://arxiv.org/html/2412.19437) - [*DeepEP*Expert dispatch and combine kernels.](https://github.com/deepseek-ai/DeepEP) - [*EPLB*Expert placement and replication from measured load.](https://github.com/deepseek-ai/EPLB) - [“MegaScale-Infer”Large-scale MoE inference and communication overlap.](https://arxiv.org/abs/2504.02263) ## Prefill and decode disaggregation Splitting prefill and decode across workers. - [“DistServe”Separate prefill and decode workers optimized for goodput under latency constraints.](https://arxiv.org/html/2401.09670) - [“Splitwise”Phase-specific allocation and scheduling.](https://www.microsoft.com/en-us/research/publication/splitwise-efficient-generative-llm-inference-using-phase-splitting/) - [“Mooncake”KV-centric disaggregated inference.](https://www.usenix.org/conference/fast25/presentation/qin) - [*NIXL*A transport layer for moving inference state across memory and network backends.](https://github.com/ai-dynamo/nixl) - [*Dynamo disaggregated serving*A current production implementation.](https://docs.nvidia.com/dynamo/design-docs/disaggregated-serving.md) ## Production systems Scheduling and routing for serving at production scale. - [“Clockwork”Predictable model serving through centralized scheduling.](https://www.usenix.org/conference/osdi20/presentation/gujarati) - [“ServerlessLLM”Faster model startup and live migration.](https://www.usenix.org/conference/osdi24/presentation/fu) - [*Gateway API Inference Extension*Model, accelerator, and KV-aware request routing.](https://gateway-api-inference-extension.sigs.k8s.io/) - [*llm-d*Distributed routing, scheduling, and disaggregated serving on Kubernetes.](https://github.com/llm-d/llm-d) ## Serving benchmarks Workloads and metrics for measuring serving systems. - [“MLPerf Inference”Reproducible benchmark scenarios and load generation.](https://www.cs.toronto.edu/ecosystem/papers/ISCA_20/MLPerf%20Inference.pdf) - [“Etalon: goodput under latency SLOs”Goodput under per-request latency SLOs.](https://arxiv.org/html/2407.07000) - [“ServeGen”Workload generation that preserves important production-trace properties.](https://www.usenix.org/system/files/nsdi26-xiang-servegen.pdf) - [*BurstGPT*A public trace for bursty LLM workloads.](https://github.com/HPMLL/BurstGPT) - [*MLPerf Endpoints*An endpoint-level benchmark for interactive generative AI.](https://mlcommons.org/benchmarks/endpoints/) ## NVIDIA Current NVIDIA data center GPU architecture. - [*Blackwell architecture brief*Blackwell and Blackwell Ultra system architecture.](https://resources.nvidia.com/en-us-blackwell-architecture/blackwell-architecture-technical-brief) - [*Blackwell Tuning Guide*Programming and optimization guidance.](https://docs.nvidia.com/cuda/blackwell-tuning-guide/) - [*CUTLASS Blackwell documentation*Blackwell matrix multiply and data-movement support.](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/blackwell.html) ## AMD Current AMD data center GPU architecture. - [*CDNA 4 architecture whitepaper*MI350 compute, memory, and chiplet architecture.](https://www.amd.com/content/dam/amd/en/documents/instinct-tech-docs/white-papers/amd-cdna-4-architecture-whitepaper.pdf) - [*CDNA 4 instruction set*The native machine instruction reference.](https://www.amd.com/content/dam/amd/en/documents/instinct-tech-docs/instruction-set-architectures/amd-instinct-cdna4-instruction-set-architecture.pdf) - [*MI350 performance counters*Counter definitions and measurement guidance.](https://rocm.docs.amd.com/en/latest/reference/gpu-arch/mi350-performance-counters.html) ## Google TPU Current TPU architecture and kernel programming model. - [“TPU v1 analysis”The original datacenter TPU paper.](https://research.google/pubs/in-datacenter-performance-analysis-of-a-tensor-processing-unit/) - [“TPU v4”The TPU v4 chip, interconnect, and system.](https://arxiv.org/abs/2304.01433) - [*Ironwood documentation*Current TPU v7 architecture and configuration.](https://docs.cloud.google.com/tpu/docs/tpu7x) - [*Pallas TPU hardware model*The TPU execution and memory model for kernel authors.](https://docs.jax.dev/en/latest/pallas/tpu/hardware.html) ## AWS Trainium Current Trainium architecture and kernel programming model. - [*Trainium and Inferentia2 architecture*NeuronCore v2 compute and memory architecture.](https://awsdocs-neuron.readthedocs-hosted.com/en/latest/nki/guides/architecture/trainium_inferentia2_arch.html) - [*Trainium3 architecture*The current NeuronCore architecture.](https://awsdocs-neuron.readthedocs-hosted.com/en/latest/nki/guides/architecture/trainium3_arch.html) - [*NKI performance guide*Kernel optimization for Trainium and Inferentia.](https://awsdocs-neuron.readthedocs-hosted.com/en/v2.24.0/general/nki/nki_perf_guide.html) ## AI-generated kernels Benchmarks for LLM-generated GPU kernels. - [“KernelBench”The original benchmark for converting PyTorch operators into faster GPU kernels.](https://proceedings.mlr.press/v267/ouyang25a.html) - [“KernelBench-Verified”Stronger correctness tests and baseline parity.](https://arxiv.org/html/2607.16241) - [*SOL-ExecBench*Correctness and performance measured against a hardware speed-of-light model.](https://github.com/nvidia/sol-execbench) ## Watchlist Hardware and techniques worth tracking before they have reproducible evidence. - [*Inside the NVIDIA Rubin GPU architecture*Rubin and Rubin CPX, pending shipped systems and reproducible measurements.](https://developer.nvidia.com/blog/inside-nvidia-rubin-gpu-architecture-powering-the-era-of-agentic-ai/) 123 sources Links go to the primary sources; anything worth understanding properly is better read there than summarized here. --- # Developers # Learn Kernels for developers and agents The chapter and section index behind this site is also a small public JSON API: no account, no API key, and no request quota to buy. It is read-only, so the live endpoints below are already safe to try. ## Quickstart List every chapter: ``` curl https://learn-kernels.com/api/v1/chapters ``` Get one chapter and its sections: ``` curl https://learn-kernels.com/api/v1/chapters/introduction ``` The full surface is described machine-readably at [/openapi.json](https://learn-kernels.com/openapi.json) (OpenAPI 3.1), with typed request and response schemas an agent can load directly into a function-calling tool definition. ## MCP server The same index is also an MCP server at `https://learn-kernels.com/api/mcp`, for MCP-native clients like Claude Desktop or Claude Code. It exposes two tools: `list_chapters` and `get_chapter`, no authentication required. ``` { "mcpServers": { "learn-kernels": { "url": "https://learn-kernels.com/api/mcp" } } } ``` It also publishes an [MCP server card](https://learn-kernels.com/.well-known/mcp-server-card) for automatic discovery, per the draft [SEP-2127](https://github.com/modelcontextprotocol/modelcontextprotocol/pull/2127) proposal. ## Errors Every non-2xx response, including an unknown path under `/api/` and an unsupported method on a real one, is JSON in the same shape: a stable `code`, a human-readable `message`, and usually a `hint` naming the next call to make. ``` curl https://learn-kernels.com/api/v1/chapters/not-a-real-slug { "error": { "code": "chapter_not_found", "message": "No chapter matches slug \"not-a-real-slug\".", "hint": "GET /api/v1/chapters for the list of valid slugs." } } ``` ## Versioning and rate limits The API is versioned in the URL path (`/api/v1/...`). A breaking change ships under a new version prefix rather than changing this one in place; this version keeps working for at least 90 days after a new one ships, announced with a `Deprecation` response header and a `Sunset` date before removal. Every response carries `RateLimit-Limit`, `RateLimit-Remaining`, and `RateLimit-Reset` headers. Going over the limit returns `429` with a `Retry-After` header, in the same JSON error shape as above. ## Full text For reading rather than querying, every page is available as Markdown by appending `.md` to its URL, or by sending `Accept: text/markdown`. The whole book is one document at [/llms-full.txt](https://learn-kernels.com/llms-full.txt), indexed at [/llms.txt](https://learn-kernels.com/llms.txt). ## Reference - [OpenAPI 3.1 spec](https://learn-kernels.com/openapi.json) - [GET /api/v1/chapters](https://learn-kernels.com/api/v1/chapters) - [GET /api/v1/chapters/{slug}](https://learn-kernels.com/api/v1/chapters/introduction) - [MCP server](https://learn-kernels.com/api/mcp) - [llms.txt](https://learn-kernels.com/llms.txt) - [llms-full.txt](https://learn-kernels.com/llms-full.txt) - [sitemap.xml](https://learn-kernels.com/sitemap.xml) --- # Privacy # Privacy Nothing here needs your name. No account, no email, nothing following you to the next site. This page is the actual mechanism, not a lawyer’s paraphrase of it. ## Ask AI Type a question and it travels to the model provider behind Ask AI (Cloudflare Workers AI), along with whatever page you’re reading, so the answer can be specific to it. It can only read pages already on this site. Not the web, not code execution, not the server’s files. The server hashes your IP before anything else touches it. That hash stops one browser from hammering the endpoint, nothing else, and the raw address is never written down. Your side of the conversation lives in local storage on your machine. The server keeps a matching copy under the same session, so a reply that was mid-stream when you reloaded can keep going. Clear the conversation and your copy disappears. ## Everything else Vercel Analytics and Speed Insights count visits and page-load speed in aggregate, no cookie involved. The API and MCP server need no key and rate-limit by IP in memory only, gone the moment the server restarts. Dark mode or light lives in local storage too. Fonts are baked into the site at build time, so even the first load never asks Google for anything.