AI Performance Engineering

A curated resource list for learning GPU performance engineering and production inference.

The list is ordered from a single inference request to a single GPU, optimized kernels, inference engines, and distributed systems. Read Start here first. After that, use it as a reference.

The core list uses original papers, official specifications and documentation, creator repositories, and direct implementation work.

If you work on these problems, Wafer is hiring.

Contents

• Start here: the minimum mental model
• 1. GPU fundamentals
  • Programming model
  • Compilation and machine code
• 2. Kernel optimization
  • Foundational kernel exercises
  • Matrix multiplication
  • Direct implementation work
  • Tensor cores and low precision
  • Attention
• 3. Programming models and profiling
  • Triton
  • CUTLASS, CuTe, and CUDA Tile
  • Other hardware stacks
  • Profiling, benchmarking, and correctness
• 4. Inference engines
  • Scheduling and continuous batching
  • KV cache systems
  • Quantization
  • Speculative decoding
  • Structured decoding and fairness
  • Long context and multimodal inference
• 5. Distributed inference
  • Parallelism, collectives, and topology
  • Mixture-of-experts serving
  • Prefill and decode disaggregation
  • Production systems
  • Serving benchmarks
• 6. Current hardware
  • NVIDIA
  • AMD
  • Google TPU
  • AWS Trainium
• Frontier
  • AI-generated kernels
  • Watchlist
• Source policy

Start here: the minimum mental model

Read these in order if you are new to the field.

1. How to Scale Your Model: Inference - One request from prefill through decode, with batching, KV memory, and parallelism.
2. Attention Is All You Need - The transformer computation that the rest of the list optimizes.
3. CUDA C++ basics - The shortest official introduction to the CUDA execution model.
4. Programming Massively Parallel Processors - The main textbook for GPU programming, memory, and kernel design.
5. Roofline: An Insightful Visual Performance Model - The compute, memory-bandwidth, and arithmetic-intensity model.
6. Transformer Inference Arithmetic - FLOPs, parameter bytes, KV bytes, and communication for transformer inference.
7. Efficiently Scaling Transformer Inference - Latency, memory, and parallelism costs for large-model inference.
8. Etalon - TTFT, TPOT, goodput, and latency SLOs for generative-model serving.

For a practical companion, use the GPU Mode lectures.

1. GPU fundamentals

Programming model

• CUDA Programming Guide - The normative CUDA reference.
• CUDA programming model - Threads, warps, blocks, grids, and the memory hierarchy.
• CUDA C++ Best Practices Guide - Coalescing, shared memory, occupancy, synchronization, and optimization workflow.
• Hopper Tuning Guide - TMA, thread-block clusters, asynchronous execution, and Hopper-specific limits.
• Blackwell Tuning Guide - Tensor memory, Blackwell execution features, and architecture limits.

Compilation and machine code

• NVCC Compiler Driver - The CUDA compilation trajectory and artifact controls.
• PTX ISA - NVIDIA's virtual instruction set and memory model.
• CUDA Binary Utilities - cuobjdump and nvdisasm for inspecting GPU binaries.
• Understanding PTX - NVIDIA's introduction to the role of PTX between CUDA and machine code.

2. Kernel optimization

Foundational kernel exercises

• Efficient Matrix Transpose in CUDA C/C++ - Coalescing, shared-memory tiling, and bank conflicts.
• Optimizing Parallel Reduction in CUDA - Synchronization, 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.
• Defeating Nondeterminism in LLM Inference - Floating-point non-associativity and batch-dependent reduction order, with batch-invariant RMSNorm, matmul, and attention kernels.

Matrix multiplication

• Benchmarking GPUs to Tune Dense Linear Algebra - The canonical case for reasoning from measured hardware behavior instead of occupancy alone.
• CUTLASS 3.x design - The collective and kernel structure used by modern CUTLASS.
• Stream-K - Work-centric GEMM decomposition that splits k-loop iterations evenly across persistent CTAs to avoid wave quantization.
• DeepGEMM - A compact production FP8 GEMM implementation for Hopper.

Direct implementation work

• How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance - A matrix multiplication built from naive CUDA through shared-memory and register tiling.
• Inside NVIDIA GPUs: Anatomy of High-Performance Matmul Kernels - Layouts, tiling, PTX, machine code, and roofline analysis.
• Outperforming cuBLAS on H100: A Worklog - A direct Hopper optimization worklog using tensor cores and asynchronous movement.
• CUTLASS Tutorial: Mastering TMA - Working kernels built around the Tensor Memory Accelerator.
• tcgen05 for dummies - A Blackwell GEMM built step by step from tcgen05 MMA and tensor memory to 2-SM MMA and persistent scheduling.

Tensor cores and low precision

• OCP 8-bit Floating Point Specification - E4M3 and E5M2 formats.
• OCP Microscaling Formats Specification - Shared-scale MX formats.
• NVIDIA Transformer Engine - FP8 and FP4 transformer execution with scaling controls.
• Blackwell matrix multiply instructions - tcgen05, tensor memory, and Blackwell MMA programming.

Attention

• FlashAttention - IO-aware exact attention.
• FlashAttention-2 - Better work partitioning and parallelism.
• FlashAttention-3 - Asynchronous movement and tensor-core overlap on Hopper.
• FlashAttention-4 - The Blackwell attention schedule.
• FlashInfer - Block-sparse KV formats, JIT-compiled attention variants, and CUDA Graph-compatible load balancing for serving.

3. Programming models and profiling

Triton

• Triton paper - The original blocked-program language and compiler design.
• Triton programming guide - The official programming model.
• Triton tutorials - Official worked kernels for fused softmax, matmul with grouped tile ordering, fused attention, persistent matmul, and block-scaled matmul.
• Linear Layouts - Triton's layout system, which represents layouts and layout conversions as binary matrices over F2.
• Triton repository - Compiler, examples, tests, and backend implementation.

CUTLASS, CuTe, and CUDA Tile

• CuTe layout algebra - Layouts and layout composition.
• CuTe GEMM tutorial - A GEMM expressed through CuTe layouts and atoms.
• CUTLASS pipeline documentation - Producer-consumer pipelines and asynchronous stages.
• Writing tile kernels - Tile-level kernels in cuTile Python and CUDA Tile C++, with the compiler mapping tiles to threads, tensor cores, and TMA.
• CUDA Tile IR programming model - NVIDIA's compiler-owned tile abstraction.

Other hardware stacks

• ROCm Composable Kernel - AMD tiling, layout, and operator primitives.
• ROCm AITER - AMD inference and transformer operator implementations.
• HipKittens - A tile abstraction for AMD GPUs.
• Pallas design - The JAX kernel model for GPU and TPU backends.
• NKI language guide - The tile-level kernel language for AWS NeuronCore hardware.

Profiling, benchmarking, and correctness

• Nsight Systems User Guide - System timelines, CPU-GPU interaction, and distributed traces.
• Nsight Compute Profiling Guide - Kernel metrics, sections, replay, and roofline analysis.
• Compute Sanitizer - Memory, race, initialization, and synchronization checks.
• CUTLASS GEMM measurement methodology - Reproducible GEMM benchmarking.
• ROCm Compute Profiler - AMD performance counters and roofline analysis.

4. Inference engines

Scheduling and continuous batching

• 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.
• vLLM, SGLang, and TensorRT-LLM - The main production engine implementations.

KV cache systems

• Grouped-Query Attention - Fewer key-value heads and a smaller KV cache.
• DeepSeek-V2 - Multi-head latent attention and compressed KV state.
• KIVI - KV quantization with separate treatment for keys and values.
• CacheGen - KV compression for transfer.
• Jenga - KV memory allocation and per-layer caching for models that mix attention types, state-space layers, and vision embeddings.

Quantization

• GPTQ - One-shot second-order weight quantization.
• SmoothQuant - W8A8 execution by moving quantization difficulty from activations into weights.
• AWQ - Low-bit weight-only inference with salient-weight protection.
• QuaRot - Output-preserving rotations that remove outliers so weights, activations, and the KV cache can run in 4 bits.
• Introducing NVFP4 for Efficient and Accurate Low-Precision Inference - Hardware-native 4-bit micro-block scaling for Blackwell inference.

Speculative decoding

• Fast Inference from Transformers via Speculative Decoding - Exact sampling with a draft model.
• Accelerating Large Language Model Decoding with Speculative Sampling - The parallel formulation and analysis.
• Medusa - Multiple prediction heads on the target model.
• EAGLE - Feature-level drafting.
• EAGLE-3 - Drafters that predict tokens directly from fused multi-layer target features, trained on their own multi-step outputs.

Structured decoding and fairness

• Efficient Guided Generation for Large Language Models - Regex constraints compiled into a finite-state index over the vocabulary, extended to grammars with an incremental parser.
• Guiding LLMs the Right Way - Constrained decoding without changing the intended token distribution.
• XGrammar - A fast grammar engine for structured generation.
• Fairness in Serving Large Language Models - Fair scheduling when request sizes are different and unknown.

Long context and multimodal inference

• 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.
• DeepSeek-V3.2 - DeepSeek Sparse Attention, where a lightweight learned indexer selects the top-k key-value entries for each query.
• vLLM multimodal inputs - Current engine support for text, image, audio, and video inputs.

5. Distributed inference

Parallelism, collectives, and topology

• Megatron-LM - Intra-layer tensor parallelism for transformer models.
• NCCL - NVIDIA's collective communication implementation.
• NCCL device-initiated communication - Kernel-issued communication over NVLink load/store, NVLink multicast, and GPU-initiated RDMA.
• Multi-node NVLink Systems Tuning Guide - NVLink and InfiniBand topology in GB200 NVL systems.
• Ultra Ethernet 1.0.3 Specification - The scale-out transport specification.

Mixture-of-experts serving

• GShard - Expert parallelism through all-to-all token dispatch and combine under a per-expert capacity.
• DeepSeek-V3 - Routed experts, shared experts, and the model-system design.
• DeepEP - Expert dispatch and combine kernels.
• EPLB - Expert placement and replication from measured load.
• MegaScale-Infer - Large-scale MoE inference and communication overlap.

Prefill and decode disaggregation

• DistServe - Separate prefill and decode workers optimized for goodput under latency constraints.
• Splitwise - Phase-specific allocation and scheduling.
• Mooncake - KV-centric disaggregated inference built on a distributed KV cache pool.
• NIXL - A transport layer for moving inference state across memory and network backends.
• Dynamo - A production implementation of disaggregated prefill and decode with KV-aware routing.

Production systems

• Clockwork - Predictable model serving through centralized scheduling.
• ServerlessLLM - Faster model startup and live migration.
• Gateway API Inference Extension - Model, accelerator, and KV-aware request routing.
• llm-d - Distributed routing, scheduling, and disaggregated serving on Kubernetes.

Serving benchmarks

• MLPerf Inference - Reproducible benchmark scenarios and load generation.
• Etalon - Goodput under per-request latency SLOs.
• ServeGen - Workload generation that preserves important production-trace properties.
• BurstGPT - A public trace for bursty LLM workloads.
• MLPerf Endpoints - An endpoint-level benchmark for interactive generative AI.

6. Current hardware

Read each architecture with its ISA or tuning guide. Vendor peak numbers are not performance measurements.

NVIDIA

• Blackwell architecture brief - Blackwell and Blackwell Ultra system architecture.
• Blackwell Tuning Guide - Programming and optimization guidance.
• CUTLASS Blackwell documentation - Blackwell matrix multiply and data-movement support.

AMD

• HIP hardware implementation - Compute units, SIMDs, the LDS, matrix cores, and caches that HIP kernels map onto.
• CDNA 4 architecture whitepaper - MI350 compute, memory, and chiplet architecture.
• CDNA 4 instruction set - The native machine instruction reference.
• MI350 performance counters - Counter definitions and measurement guidance.

Google TPU

• TPU v1 analysis - The original datacenter TPU paper.
• Ten Lessons From Three Generations Shaped Google's TPUv4i - Inference-accelerator design driven by P99 latency limits, multi-tenancy, backward ML compatibility, and cost of ownership.
• TPU v4 - The TPU v4 chip, interconnect, and system.
• Ironwood documentation - Current TPU v7 architecture and configuration.
• Pallas TPU hardware model - The TPU execution and memory model for kernel authors.

AWS Trainium

• Trainium and Inferentia2 architecture - NeuronCore v2 compute and memory architecture.
• Trainium2 architecture - NeuronCore v3 compute and memory architecture.
• Trainium3 architecture - The current NeuronCore architecture.
• NKI performance guide - Kernel optimization for Trainium and Inferentia.

Frontier

Verified on 2026-10-17. This section is kept separate from the core list because the evidence changes quickly.

AI-generated kernels

• KernelBench - The original benchmark for converting PyTorch operators into faster GPU kernels.
• KernelBench-Verified - Stronger correctness tests and baseline parity.
• FlashInfer-Bench - Kernel tasks drawn from real serving traces, with correctness-checked kernels substituted into SGLang and vLLM.
• SOL-ExecBench - Correctness and performance measured against a hardware speed-of-light model.
• KernelGuard - Layered defenses against reward-hacked kernel submissions: static rules, physics-floor checks, and an LLM judge.

Watchlist

• NVIDIA Rubin, pending a tuning guide or architecture whitepaper and reproducible measurements on generally available systems.
• AMD MI455X and Helios, with the CDNA 5 whitepaper and ISA guide published, pending reproducible measurements on shipped systems.
• Google TPU 8t and TPU 8i, pending Cloud TPU architecture documentation and reproducible measurements on generally available systems.
• UALink 2.0, pending shipped silicon and reproducible measurements.
• Gluon, Triton's lower-level GPU programming model, pending a stable API across releases.
• Session-aware and agentic scheduling against public production traces.
• Real-time voice and video serving with complete quality and latency metrics.
• Inference ASICs, processing in memory, analog compute, and photonic compute with reproducible deployments.
• Individual AI kernel agents that have not been rerun on a hardened evaluator.

Source policy

A core source must be one of the following:

• the paper that introduced the mechanism;
• the specification or official documentation that defines it;
• the repository that implements it;
• a direct implementer report with code, measurements, and enough detail to reproduce the result.

Performance claims need the hardware, workload, precision, baseline, and correctness method. Otherwise the number is omitted.

See CONTRIBUTING.md before proposing a resource.

License

MIT

Maintainer

[email protected]