# 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.