Further reading
The papers, guides, and repositories worth your time
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: InferenceOne request from prefill through decode, with batching, KV memory, and parallelism.
- “Attention Is All You Need”The transformer computation that the rest of the list optimizes.
- CUDA C++ basicsThe shortest official introduction to the CUDA execution model.
- Programming Massively Parallel ProcessorsThe main textbook for GPU programming, memory, and kernel design.
- “Roofline: An Insightful Visual Performance Model”The compute, memory-bandwidth, and arithmetic-intensity model.
- Transformer Inference ArithmeticFLOPs, parameter bytes, KV bytes, and communication for transformer inference.
- “Efficiently Scaling Transformer Inference”Latency, memory, and parallelism costs for large-model inference.
- “Etalon”TTFT, TPOT, goodput, and latency SLOs for generative-model serving.
- GPU Mode lecturesA practical companion to work through alongside the start-here list.
Programming model
Threads, warps, memory hierarchy, and the CUDA execution model.
- CUDA Programming GuideThe normative CUDA reference.
- CUDA programming modelThreads, warps, blocks, grids, and the memory hierarchy.
- CUDA C++ Best Practices GuideCoalescing, shared memory, occupancy, synchronization, and optimization workflow.
- Hopper Tuning GuideTMA, thread-block clusters, asynchronous execution, and Hopper-specific limits.
- Blackwell Tuning GuideTensor memory, Blackwell execution features, and architecture limits.
Compilation and machine code
What CUDA compiles down to, and how to read it.
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.
- Optimizing Parallel Reduction in CUDASynchronization, divergence, occupancy, and instruction cost.
- “Single-pass Parallel Prefix Scan with Decoupled Look-back”A work-efficient scan with one pass over memory.
- “Online Normalizer Calculation for Softmax”Numerically stable online softmax without materialized intermediates.
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.
- CuTe GEMM tutorialTiling, layouts, copies, and matrix-multiply atoms.
- CUTLASS 3.x designThe collective and kernel structure used by modern CUTLASS.
- DeepGEMMA compact production FP8 GEMM implementation for Hopper.
Direct implementation work
Worklogs that build a fast kernel from scratch, in the open.
- How to Optimize a CUDA Matmul Kernel for cuBLAS-like PerformanceA matrix multiplication built from naive CUDA through shared-memory and register tiling.
- Inside NVIDIA GPUs: Anatomy of High-Performance Matmul KernelsLayouts, tiling, PTX, machine code, and roofline analysis.
- Outperforming cuBLAS on H100: A WorklogA direct Hopper optimization worklog using tensor cores and asynchronous movement.
- CUTLASS Tutorial: Mastering TMAWorking kernels built around the Tensor Memory Accelerator.
Tensor cores and low precision
FP8, FP4, and the formats and instructions tensor cores run.
Attention
FlashAttention and the kernels it led to.
Triton
The blocked-program language and compiler.
CUTLASS, CuTe, and CUDA Tile
NVIDIA's tile-based kernel-authoring stack.
- CuTe layout algebraLayouts and layout composition.
- CUTLASS GEMM tutorialA GEMM expressed through CuTe layouts and atoms.
- CUTLASS pipeline documentationProducer-consumer pipelines and asynchronous stages.
- CUDA Tile IR programming modelNVIDIA's compiler-owned tile abstraction.
- CUDA Tile repositoryThe current implementation and examples.
Other hardware stacks
Kernel programming models for AMD, TPU, and Trainium.
- ROCm Composable KernelAMD tiling, layout, and operator primitives.
- ROCm AITERAMD inference and transformer operator implementations.
- HipKittensA tile abstraction for AMD GPUs.
- Pallas designThe JAX kernel model for GPU and TPU backends.
- NKI programming modelThe tile-level programming model for AWS NeuronCore hardware.
Profiling, benchmarking, and correctness
Tools for measuring and verifying what a kernel actually does.
- Nsight Systems User GuideSystem timelines, CPU-GPU interaction, and distributed traces.
- Nsight Compute Profiling GuideKernel metrics, sections, replay, and roofline analysis.
- Compute SanitizerMemory, race, initialization, and synchronization checks.
- CUTLASS GEMM measurement methodologyReproducible GEMM benchmarking.
- ROCm Compute ProfilerAMD performance counters and roofline analysis.
Scheduling and continuous batching
Iteration-level scheduling and the engines that implement it.
- “Orca”Iteration-level scheduling for autoregressive serving.
- “PagedAttention and vLLM”Paged KV allocation and continuous batching.
- “Sarathi-Serve”Chunked prefills that reduce interference with decode.
- “SGLang”Prefix reuse, structured programs, and a serving runtime.
- vLLMOne of the main production engine implementations.
- SGLang repositoryOne of the main production engine implementations.
- TensorRT-LLMOne of the main production engine implementations.
KV cache systems
Shrinking, quantizing, and distributing the KV cache.
Quantization
Weight and activation quantization for inference.
Speculative decoding
Drafting and verifying multiple tokens per step.
Structured decoding and fairness
Constrained generation and scheduling fairness.
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.
- “MInference 1.0”Dynamic sparse patterns for long-context prefill on existing models.
- “Native Sparse Attention”A model trained with a hardware-aligned sparse attention hierarchy.
- vLLM multimodal inputsCurrent engine support for text, image, audio, and video inputs.
Parallelism, collectives, and topology
Tensor and pipeline parallelism, and the interconnects underneath.
- “Megatron-LM”Tensor and pipeline parallelism for transformer models.
- NCCLNVIDIA's collective communication implementation.
- Multi-node NVLink Systems Tuning GuideNVLink and InfiniBand topology in GB200 NVL systems.
- UALink 1.0 SpecificationAn open scale-up interconnect.
- Ultra Ethernet 1.0.3 SpecificationThe scale-out transport specification.
Mixture-of-experts serving
Routing, dispatch, and load balancing for MoE inference.
Prefill and decode disaggregation
Splitting prefill and decode across workers.
- “DistServe”Separate prefill and decode workers optimized for goodput under latency constraints.
- “Splitwise”Phase-specific allocation and scheduling.
- “Mooncake”KV-centric disaggregated inference.
- NIXLA transport layer for moving inference state across memory and network backends.
- Dynamo disaggregated servingA current production implementation.
Production systems
Scheduling and routing for serving at production scale.
Serving benchmarks
Workloads and metrics for measuring serving systems.
- “MLPerf Inference”Reproducible benchmark scenarios and load generation.
- “Etalon: goodput under latency SLOs”Goodput under per-request latency SLOs.
- “ServeGen”Workload generation that preserves important production-trace properties.
- BurstGPTA public trace for bursty LLM workloads.
- MLPerf EndpointsAn endpoint-level benchmark for interactive generative AI.
NVIDIA
Current NVIDIA data center GPU architecture.
AMD
Current AMD data center GPU architecture.
Google TPU
Current TPU architecture and kernel programming model.
AWS Trainium
Current Trainium architecture and kernel programming model.
AI-generated kernels
Benchmarks for LLM-generated GPU kernels.
Watchlist
Hardware and techniques worth tracking before they have reproducible evidence.
Links go to the primary sources; anything worth understanding properly is better read there than summarized here.