GPU & Accelerator Kernels for LLMs

GPU & Accelerator Kernels for LLMs

The hardware substrate every other model-layer skill sits on. Pretraining, fine-tuning, alignment, inference serving, and compression all ultimately resolve to kernels: GPU programs that move bytes through a memory hierarchy and feed tensor cores. This reference is the implementation view: how a GPU executes work, why LLM attention and decode are bottlenecked by memory bandwidth rather than FLOPs, and how kernels (CUDA, Triton, FlashAttention, paged-KV) and compilers (torch.compile, TensorRT-LLM, XLA) are written to fight that bottleneck.

Where this sits among siblings (read the boundary, then the right file):

• Distributed-parallelism strategy (FSDP / ZeRO / TP / PP / EP, the 3D-placement decision) → distributed-training. This reference is the kernel/hardware layer beneath it: the NCCL collective primitives (all-reduce / all-gather, ring vs tree) those strategies call.
• Serving-engine configuration and policy (vLLM flags, continuous batching, speculative decoding, autoscaling) → llm-inference-serving. This reference is the kernel mechanics beneath it: the paged / quantized KV-cache kernels the engine schedules.
• The transformer architecture — why GQA/MLA shrink the KV cache, and FlashAttention’s math (online softmax derivation, IO-aware exactness) → transformer-architecture. This reference is the kernel/implementation view of FlashAttention: tiling, SRAM reuse, warp specialization, WGMMA/TMA.
• Quantization algorithms (GPTQ, AWQ, SmoothQuant — how to choose the low-precision weights) → llm-compression. This reference covers the low-precision kernel / tensor-core mechanics: the FP8 / FP4 / MX / INT8 tensor-core paths that make those quantized weights fast.

When to load this reference

Load when the task is about why GPU code is slow and how to make it fast at the kernel level, not about which model or which parallelism strategy:

• “Why is decode memory-bound but prefill compute-bound?” / roofline reasoning.
• “What is occupancy / a warp / an SM / SIMT?” GPU execution-model questions.
• “How does FlashAttention actually work on the hardware?” (tiling, online softmax in SRAM, warp specialization, ping-pong).
• “What’s the difference between FP8, MXFP8, NVFP4, and INT8 on tensor cores?”
• Writing or reading a CUDA or Triton kernel; memory coalescing; shared memory bank conflicts; @triton.autotune.
• Kernel fusion — when it helps, when it can’t (reductions).
• Paged / quantized KV-cache kernel internals.
• NCCL collective primitives — ring vs tree, why one is bandwidth-optimal and the other latency-optimal.
• Profiling: Nsight Systems vs Nsight Compute, PyTorch profiler, and computing MFU (Model FLOPs Utilization).
• The hardware landscape: Hopper → Blackwell (H100/H200/B200/GB200), AMD MI300X/MI350X, Google TPU — bandwidth/FLOPs/HBM specs and what they imply.
• Compilers: torch.compile/TorchInductor, TensorRT-LLM, XLA, Mojo.

Core concepts (MECE)

1. The GPU execution model — SMs, warps, SIMT, occupancy

A GPU is a throughput machine built to hide latency with parallelism, the opposite of a latency-optimized CPU. The unit of compute is the Streaming Multiprocessor (SM) — a B200 has ~148 SMs, an H100 ~132. Each SM contains arithmetic units (FP/INT), tensor cores (matrix-multiply accelerators), a register file, shared memory / L1, and one or more warp schedulers.

• SIMT (Single Instruction, Multiple Thread). Threads are grouped into warps of 32. All 32 threads in a warp execute the same instruction each cycle on different data (lockstep). If threads in a warp take different branch paths (warp divergence), the paths execute serially with the inactive lanes masked off, a major performance killer.
• Latency hiding, not latency reduction. A warp scheduler holds many resident warps and, every cycle, issues from whichever warp is ready. When one warp stalls on a ~400-cycle HBM load, the SM switches to another ready warp instead of idling. This is why GPUs need thousands of threads in flight to reach peak throughput.
• Occupancy = (active warps per SM) / (max warps per SM). It is capped by the scarcest per-SM resource: registers per thread, shared memory per block, or the warp/block hardware limit. Higher occupancy gives the scheduler more warps to hide latency with — but it is a means, not a goal: a register-heavy, well-pipelined kernel can hit peak at modest occupancy, and chasing 100% occupancy by shrinking tiles can hurt (less work reuse). The Triton/CUDA tuning loop is: enough occupancy to hide memory latency, large enough tiles to keep tensor cores fed.

2. The memory hierarchy — and why attention is IO-bound

Speed and capacity trade off inversely at every level. Approximate H100/B200 figures:

• HBM3 / HBM3e is the off-chip stacked DRAM: H100 HBM3 ≈ 3.35 TB/s, H200 HBM3e ≈ 4.8 TB/s, B200 HBM3e ≈ 8 TB/s @ 192 GB. It is huge but ~100× slower per byte than SRAM.
• SRAM (shared memory) is the on-chip scratchpad a kernel explicitly manages — orders of magnitude faster than HBM but only ~hundreds of KB per SM. The whole game of a fast kernel is: stage a tile into SRAM, do all the math you can on it, then write back — minimizing HBM round-trips.
• Why attention is IO-bound. Naive attention materializes the full S = QKᵀ score matrix (size O(seq²)) in HBM, runs softmax over it (another HBM read + write), then multiplies by V (another read). The arithmetic is cheap relative to the bytes moved, so the kernel spends most of its time waiting on HBM. FlashAttention exists precisely to keep S in SRAM and never write it to HBM (see §8).

3. Arithmetic intensity and the roofline model

Arithmetic intensity (AI) = FLOPs performed ÷ bytes moved from memory (FLOP/byte). The roofline model plots attainable FLOP/s against AI:

• A sloped line (peak bandwidth × AI) on the left: the memory-bound region.
• A flat line (peak compute) on the right: the compute-bound region.
• The ridge point is where they cross — the AI at which a kernel transitions from bandwidth-limited to compute-limited. For H100 BF16 the ridge is roughly ~300 FLOP/byte; a kernel below it cannot reach peak FLOP/s no matter how fast the math units are.

The decisive LLM consequence:

• Prefill (prompt processing) is compute-bound. It is a big matrix–matrix multiply (GEMM): a long sequence × weight matrices → high AI, reuses each loaded weight across many tokens. Prefill wants high AI and is limited by tensor-core FLOP/s.
• Decode (autoregressive generation) is memory-bandwidth-bound. To emit one token you must stream the entire weight matrix (and the growing KV cache) from HBM, but you only do a matrix–vector multiply (batch=1) — almost no reuse, AI ≈ O(1). Decode latency ≈ (model bytes + KV bytes) ÷ HBM bandwidth. This is why: (a) decode throughput tracks HBM bandwidth, not FLOPs; (b) batching raises AI (reuse weights across many sequences) and is the single biggest decode-throughput lever; (c) quantization (fewer bytes per weight) directly speeds decode; (d) KV-cache size directly costs latency.

4. Precision and tensor cores — BF16 / FP8 / FP4 / MX / INT8

Tensor cores are dedicated matrix-multiply-accumulate (MMA) units: they multiply small tiles (e.g. 16×16) and accumulate, delivering ~10–30× the FLOP/s of the general FP units. Lower precision = more tensor-core throughput and fewer bytes to move (helps the memory-bound regime), so the precision frontier is the central lever for both training and inference speed.

• FP16 / BF16. 16-bit. BF16 (8-bit exponent, 7-bit mantissa) has the same dynamic range as FP32 — the default training/inference precision; rarely overflows, so usually no loss scaling.
• TF32. A 19-bit internal tensor-core mode for FP32 inputs (10-bit mantissa); a near-free Ampere+ speedup for FP32 workloads.
• FP8 (E4M3 / E5M2). 8-bit, native on Hopper+ tensor cores. E4M3 (more mantissa) for forward/weights, E5M2 (more range) for gradients. Needs scaling (per-tensor delayed scaling, or finer) to fit values in the narrow range. H100 FP8 ≈ 2× its BF16 FLOP/s.
• MX microscaling formats (OCP standard). Instead of one scale per tensor/row, an MX block of 32 consecutive elements shares one power-of-two scale stored as UE8M0 (8-bit exponent). MXFP8 = FP8 (E4M3/E5M2) elements + 1×32 block scale; MXFP4 = FP4 (E2M1) elements + 1×32 block scale; MXFP6 also exists. Finer-grained scaling than per-tensor → better accuracy at low bit-width. Native on Blackwell (SM 10.0+).
• NVFP4 (NVIDIA’s Blackwell FP4). Compatible E2M1 elements but a smaller block of 16 with two-level scaling: a per-block FP8 E4M3 scale plus a per-tensor FP32 scale. The smaller block + higher-precision scale localizes to the data’s dynamic range better than MXFP4’s 32-block UE8M0 scale, giving lower quantization error. Blackwell’s 5th-gen tensor cores handle FP4 grouping, dynamic scaling, and the 4-bit MMA in hardware. Reported: ~3.5× smaller memory vs FP16, ~1.8× vs FP8, with <1% accuracy degradation on key LM tasks for many models. Block size is hardware-fixed — picking the wrong block produces a checkpoint the tensor cores can’t consume.
• INT8. Integer 8-bit MMA, very high throughput (H100 INT8 ≈ 2× BF16); common for weight/activation quant (W8A8) and KV-cache quant. Needs careful scale/zero-point calibration.

The algorithm for choosing which weights to quantize and how (GPTQ, AWQ, SmoothQuant) lives in llm-compression. This reference is the tensor-core path those algorithms target.

5. CUDA basics — kernels, grids/blocks, coalescing, shared memory

CUDA is the C++ dialect for writing GPU kernels. The launch hierarchy:

• A kernel is a function run by many threads. You launch it over a grid of thread blocks; each block holds up to 1024 threads (executed as warps of 32). blockIdx, threadIdx, blockDim index the data each thread owns. A block runs entirely on one SM and shares that SM’s shared memory.
• Memory coalescing is the #1 global-memory rule: when the 32 threads of a warp access consecutive, aligned addresses, the hardware merges them into one (or a few) wide HBM transactions. Strided or scattered access splits into many transactions and wastes most of the bandwidth, the dominant cause of a slow memory-bound kernel.
• Shared memory is the SRAM scratchpad a block uses to stage and reuse data. Tiled matmul is the canonical pattern: each block cooperatively loads a tile of A and B from HBM into shared memory (coalesced), does the partial dot products from SRAM, advances to the next tile. This converts repeated HBM reads into one HBM read + many SRAM reads, raising arithmetic intensity.
• Bank conflicts. Shared memory is split into 32 banks. If multiple threads in a warp hit different addresses in the same bank, the accesses serialize. The standard fix is padding (e.g. a [32][33] tile) so consecutive threads land in distinct banks.
• CUDA graphs capture a sequence of kernel launches and replay them as one unit, eliminating per-launch CPU overhead, which matters in decode, where each step is many tiny kernels (see §13 for the compiler/CUDA-graph tie-in).

6. Triton — block-level kernels and autotuning

OpenAI Triton is a Python DSL+compiler for GPU kernels at a tile (block) granularity, sitting between hand-CUDA and framework ops. You write what each program instance does to a block of data; the compiler handles intra-block thread scheduling, vectorization, shared-memory allocation, and coalescing automatically — you do not manage individual threads or banks.

• Programming model: @triton.jit; pid = tl.program_id(0) identifies the block; BLOCK_SIZE is a tl.constexpr; tl.load(ptr + offs, mask=...) / tl.store(...) move tiles with boundary masks; tl.dot(a, b) issues a tensor-core matmul on tiles.
• Autotuning: decorate with @triton.autotune(configs=[triton.Config({...}, num_warps=, num_stages=), ...], key=[...]). On first call for a new shape (the key), Triton benchmarks every config on the real tensor sizes and memoizes the winner: cuDNN-style autotuning with no C++ build. num_stages controls software pipelining (overlapping loads with compute); num_warps sets the block’s warp count.
• Why it matters for LLMs: Triton is the backend torch.compile generates fused kernels into (§13), and the language most custom LLM kernels (fused softmax, layernorm, fused-attention, MoE grouped GEMM, quant kernels) are now written in. The official tutorials walk vector-add → fused-softmax → autotuned matmul → fused-attention → block-scaled (MX) matmul.

7. Kernel fusion

Fusion combines a sequence of operations into a single kernel so intermediates stay in registers/SRAM and are never written to HBM. It removes (a) HBM round-trips of intermediate tensors and (b) per-op kernel-launch overhead.

• Best targets: chains of elementwise and tile-local ops — bias → activation → dropout, RMSNorm → matmul preambles, dequant + matmul. LLM forward/backward fire hundreds of tiny ops; fusing them is a large win on launch overhead and memory traffic.
• Hard / impossible to fuse: reductions with long-range dependencies (softmax across a long axis, large all-reduce) need cross-tile/cross-SM communication that breaks single-kernel streaming. FlashAttention is the clever exception: it fuses attention by reformulating softmax into an online/streaming recurrence (§8) so no full-row reduction is materialized.
• Memory-bound ops benefit most (they were limited by bytes, and fusion cuts bytes); compute-bound GEMMs benefit less from fusion itself but still gain from fused epilogues (bias/activation folded into the GEMM store).

8. FlashAttention — the kernel case study (implementation view)

FlashAttention is IO-aware exact attention: same result as standard attention, but it never materializes the O(seq²) score matrix in HBM. (The math/derivation and the architecture motivation live in transformer-architecture; here is how the kernel is built.)

• Tiling. Q, K, V are split into blocks. The kernel loops over K/V blocks, loading each Q/K/V tile from HBM into SRAM, computing the partial scores and partial output there, and accumulating, so the score tile lives only in SRAM and is discarded, never written to HBM.
• Online (streaming) softmax. Softmax normally needs the whole row’s max and sum first. FlashAttention keeps a running max m and running denominator ℓ and rescales the accumulated output as each new K/V block arrives, producing the exact softmax without ever holding the full row. This is what makes attention fusible into one kernel.
• Recomputation in the backward pass. Rather than store the huge intermediate S, the backward pass recomputes tiles from the saved stats — trading a little extra FLOPs for a large HBM-traffic/memory saving (a selective-recompute idea).
• FlashAttention-2 raised tensor-core utilization by reducing non-matmul FLOPs, better work partitioning across warps, and parallelizing over the sequence dimension.
• FlashAttention-3 (Hopper). Exploits Hopper asynchrony: warp specialization (producer warps issue TMA async copies HBM→SRAM while consumer warps run WGMMA tensor-core matmuls), ping-pong scheduling between two warpgroups (one does GEMM while the other does softmax — ~570→620 TFLOPS), and intra-warpgroup pipelining of softmax with GEMM (→~640–660 TFLOPS FP16). It adds FP8 attention with incoherent processing (a random-sign Hadamard transform in O(d log d) to spread outliers), cutting FP8 error ~2.6× vs baseline. Result: ~740 TFLOPS FP16 (~75% of H100 peak, up from ~35%), 1.5–2.0× over FA-2; ~1.2 PFLOPS in FP8.

9. Paged and quantized KV-cache kernels

The KV cache (cached keys/values for every past token) grows with sequence and batch and dominates decode memory. Two kernel-level techniques:

• PagedAttention kernel. Inspired by OS virtual memory: the KV cache is stored in fixed-size blocks (pages), not one contiguous per-sequence buffer. A per-sequence block table maps logical token positions to physical blocks, so blocks can be allocated on demand and shared across sequences (e.g. a shared prompt prefix, or beams). The attention kernel gathers K/V through the block table instead of a flat stride. This cuts KV fragmentation/waste to <4% and is what lets a server pack many more concurrent sequences (2–4× throughput). Engine policy (which sequences to batch, eviction, prefix caching) is llm-inference-serving; this is the kernel that the policy schedules.
• Quantized KV-cache kernels. Storing K/V in FP8 (recommended on Hopper/Blackwell) or INT8 halves/quarters KV bytes, directly relieving the memory-bandwidth-bound decode path and extending context length. The kernel must dequantize on the fly inside the attention compute (or use low-precision MMA paths). vLLM ships FP8 KV-cache; INT8 KV-cache kernels (naive/tiled/ coarsened/vectorized variants) report up to 4× KV memory reduction with small accuracy loss.

10. NCCL collectives — the communication primitives (ring vs tree)

When a model spans many GPUs, the parallelism strategy (distributed-training) is implemented on top of NCCL collective primitives. The ones that matter:

• all-reduce — sum (or other op) a tensor across all ranks, every rank gets the result (gradient sync in data parallel; the per-block sum in tensor parallel).
• all-gather — each rank contributes a shard, every rank ends with the full concatenation (FSDP/ZeRO parameter gather).
• reduce-scatter — reduce then partition (the FSDP gradient half; reduce-scatter + all-gather = one ring all-reduce).
• all-to-all — every rank sends a distinct piece to every other rank (MoE expert dispatch/combine).

Ring vs tree (the central trade-off):

• Ring all-reduce arranges ranks in a logical ring and streams shards around it (reduce-scatter phase + all-gather phase). It is bandwidth-optimal — each link is fully utilized and per-rank traffic is independent of rank count — but its latency grows linearly with the number of ranks (~2(N−1) steps), so it is poor for tiny messages at large scale.
• Tree all-reduce reduces up a (double-)binary tree and broadcasts down. Latency is logarithmic in N, so it wins for small, latency-sensitive messages and large clusters.
• NCCL auto-selects per call: it models each algorithm×protocol’s latency and bandwidth and picks the predicted winner by message size — tree for small, ring for large — and tunes to the topology (NVLink/NVSwitch intra-node, the network inter-node). Newer PAT (Parallel Aggregated Trees) gives logarithmic all-gather/reduce-scatter at scale. Overlapping collectives with compute (so comm hides behind matmuls) is the strategy-level lever covered in distributed-training.

11. Profiling and Model FLOPs Utilization (MFU)

You cannot optimize what you cannot measure; raw “GPU utilization” (percent of time a kernel was resident) is misleading — it can read 100% while tensor cores sit mostly idle. The real efficiency metric is MFU.

• MFU (Model FLOPs Utilization) = (model’s useful FLOP/s, e.g. the 6ND training estimate or the inference FLOPs) ÷ (the hardware’s peak FLOP/s at that precision). It is hardware-agnostic and tells you how close you are to the roofline. 40–50% sustained MFU is a good real-world training target; decode is far lower because it is memory-bound (low AI), so MFU is the wrong lens for decode — there, % of peak bandwidth is the metric. (HFU, hardware FLOPs utilization, additionally counts recomputed FLOPs.)
• Nsight Systems (nsys) — system-wide timeline: CPU↔GPU overlap, kernel gaps, stream/launch behavior, NCCL. Low overhead; the first tool — find the top/longest or stalling kernels and the bubbles.
• Nsight Compute (ncu) — single-kernel deep dive: achieved occupancy, memory vs compute bound, the kernel’s roofline, warp-stall reasons, bank conflicts. The second tool, once Nsight Systems names the suspect kernel.
• PyTorch profiler (torch.profiler + TensorBoard / Holistic Trace Analysis / Chrome trace) — medium overhead, framework-aware: maps kernels back to model ops, with stack traces, shapes, and memory. Best for “which layer is slow” and for correlating Python with kernels.
• Typical loop: Nsight Systems → find bubbles / a hot kernel → Nsight Compute → classify (memory- vs compute-bound on the roofline) → fix (coalesce, fuse, raise occupancy/tile size, change precision) → re-measure MFU.

12. The hardware landscape (Hopper → Blackwell, MI300X, TPU)

Bandwidth and FLOPs set the roofline; HBM capacity sets how big a model/KV cache fits. Approximate per-accelerator figures (2024–mid-2026):

Takeaways: (1) each generation’s bandwidth jump is what speeds decode; (2) Blackwell’s FP4/MXFP8 is what makes 4-bit inference fast in hardware; (3) NVLink/NVSwitch fabric (GB200 NVL72) makes large collectives intra-fabric; (4) NVIDIA’s lead is partly the kernel/software moat (CUDA, cuDNN, NCCL, TensorRT, FlashAttention) — AMD MI300X has competitive silicon (more HBM) but historically trails on ready kernels; TPU is strong but XLA-only (no CUDA).

13. Compilers — torch.compile / TorchInductor, TensorRT-LLM, XLA, Mojo

Compilers turn a high-level model graph into fused, scheduled kernels so humans don’t hand-write each one.

• torch.compile (PyTorch 2.x). Front end TorchDynamo captures the graph; back end TorchInductor lowers it to fused Triton kernels (GPU) / C++/OpenMP (CPU). Fusion of elementwise chains + reduced launch overhead is the main speedup; pairs with CUDA graphs to kill per-launch cost. It is the default acceleration path and the one vLLM now uses (-O3, piecewise CUDA graphs) for its model code.
• TensorRT-LLM. NVIDIA’s inference compiler/runtime: aggressive deep graph fusion, fused multi-head attention, FP8/FP4 paths, in-flight batching, piecewise CUDA graphs, and it now uses torch.compile for lightweight vertical fusion. Tends to win on large models where deep fusion pays off; torch.compile alone can match or beat it on smaller models.
• XLA. Google’s array compiler (JAX, TF, TPU; PyTorch/XLA). The native path for TPUs and whole-graph fusion via HLO; on NVIDIA GPUs its gains over torch.compile are usually modest.
• Mojo. Modular’s Python-superset systems language aimed at writing portable high-performance kernels (an alternative to CUDA C++/Triton, MLIR-based); emerging, not yet a default in mainstream LLM stacks — watch, don’t depend.

Practical patterns

• Diagnose with the roofline first. Before optimizing, classify the kernel: memory-bound or compute-bound (Nsight Compute draws this). Memory-bound → coalesce, fuse, raise reuse, drop precision. Compute-bound → bigger tiles, use tensor cores, lower precision.
• Decode = bandwidth problem. To speed decode: batch (raise AI), quantize weights + KV (fewer bytes), shrink the KV cache (GQA/MLA — an architecture lever in transformer-architecture), use CUDA graphs to kill launch overhead. Do not expect more FLOPs to help.
• Prefer Triton + autotune over hand-CUDA for new custom kernels unless you need an instruction the DSL can’t express; let the compiler handle banks and coalescing, and let @triton.autotune find tile sizes per shape.
• Let torch.compile fuse first. Reach for hand kernels only where the compiler leaves bandwidth on the table (profile to prove it).
• Match the precision to the hardware’s fixed block size. On Blackwell, NVFP4 wants 16-element blocks, MXFP8 wants 32 — quantize to the format the tensor cores actually consume.
• Measure MFU for training, % peak bandwidth for decode. Report the right metric for the regime; a “100% GPU utilization” claim with low MFU means the tensor cores are starved.

Anti-patterns

• Chasing 100% occupancy. Occupancy is a means to hide latency, not a target; shrinking tiles to raise it can lower reuse and hurt throughput.
• Trusting nvidia-smi “GPU-Util”. It reports time a kernel was resident, not tensor-core efficiency; use MFU / Nsight Compute instead.
• Strided / uncoalesced global access in the hot loop — the most common cause of a memory-bound kernel running at a fraction of HBM bandwidth.
• Trying to fuse a long-range reduction (e.g. naive softmax across the whole sequence) into one streaming kernel — it needs cross-tile communication; reformulate (online softmax) or keep it separate.
• Optimizing FLOPs to speed decode. Decode is bandwidth-bound; FLOP-side tuning yields little. Optimize bytes moved.
• Ignoring the fixed MX/NVFP4 block size — a wrong block size produces a checkpoint the tensor cores cannot run.
• Hand-writing CUDA before profiling. Premature kernel hacking before the roofline tells you what’s actually limiting.

Troubleshooting

Cross-references (reciprocal)

• transformer-architecture — FlashAttention’s math (online-softmax derivation, IO-aware exactness), GQA/MLA why, the KV cache concept. This file is the kernel implementation of those.
• llm-inference-serving — serving-engine policy (vLLM batching, prefix caching, speculative decoding, autoscaling) on top of the paged/quantized KV kernels here.
• distributed-training — parallelism strategy (FSDP/ZeRO/TP/PP/EP, 3D placement, compute–comm overlap) on top of the NCCL collective primitives here.
• llm-compression — quantization algorithms (GPTQ/AWQ/SmoothQuant) that target the FP8/FP4/MX/INT8 tensor-core paths here.

References

1. NVIDIA — Introducing NVFP4 for Efficient and Accurate Low-Precision Inference (developer.nvidia.com, 2025).
2. NVIDIA Transformer Engine docs — MXFP8 / Using FP8 and FP4 (OCP MX block formats, UE8M0 scaling), 2025.
3. OCP — Microscaling (MX) Data Formats for Deep Learning spec / arXiv 2310.10537.
4. Tri Dao et al. — FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-precision (arXiv 2407.08608; tridao.me blog; PyTorch blog), 2024.
5. OpenAI / Triton — Introducing Triton and the official tutorials (triton-lang.org): fused softmax, autotuned matmul, fused attention, block-scaled matmul.
6. NVIDIA NCCL — developer docs + Understanding NCCL Tuning and Massively Scale … with NCCL (ring vs tree algorithm selection); PAT algorithm (arXiv 2506.20252).
7. vLLM docs — PagedAttention design and Quantized KV Cache (FP8); INT8 KV-cache quantization (arXiv 2601.04719).
8. LLM Inference Unveiled: Survey and Roofline Model Insights (arXiv 2402.16363); A Systematic Characterization of LLM Inference on GPUs (arXiv 2512.01644).
9. PyTorch — Why Is PyTorch Compile So Fast: Kernel Fusion; Introduction to torch.compile and How It Works with vLLM (vLLM blog, 2025).
10. NVIDIA TensorRT-LLM docs — Torch Compile & Piecewise CUDA Graph; Collabora torch.compile vs TensorRT (2024).
11. Trainy — GPU Utilization Is a Misleading Metric; Using Model FLOPs Utilization (MFU); NVIDIA Profiling LLM Training Workflows on Grace Hopper (Nsight Systems/Compute).
12. Hardware comparisons — Exxact Blackwell vs Hopper; SemiAnalysis MI300X vs H100/H200; Artificial Analysis TPU v6e vs MI300X vs H100/B200 (2025).
13. Deep Kernel Fusion for Transformers (arXiv 2602.11808) — fusion targets, HBM-traffic reduction.
14. Siboehm — How to Optimize a CUDA Matmul Kernel (coalescing, tiling, shared-memory, bank conflicts worklog).