关于推理 Kernel、GPU 架构与性能工程的近期精选。
7 篇文章第4章：Kernel 优化技巧
掌握内存访问优化和计算优化，让你的 CUDA 程序飞起来 本章负责把优化技巧串成可执行的判断流程：先测量，再判断 memory-bound、compute-bound 或资源受限，然后只改一个因素并复测。具体指标解释以 Nsight Compute NCU 分析方法与优化思路 为准。
第6章：大模型推理 Kernel
实现 GPT、LLaMA 等大模型的核心算子，掌握推理优化技巧 本章负责把基础 CUDA kernel 能力映射到大模型推理：GEMM、softmax、norm、attention、RoPE、KV cache。这里先建立数据流和优化方向，具体算法语义可回到 LLM 目录，底层 GEMM/硬件细节回到 GPU 主文档。
Blackwell Tensor Core: tcgen05.mma
tcgen05.mma typically reads A and B from SMEM and updates the C/D accumulator in TMEM. The epilogue later brings the result back into registers with tcgen05.ld. cta_group::1 uses…
Building a Tiled GEMM
Builds a correct tiled GEMM from the TIRx tile primitives, starting from a single output tile. Step 1 is a single-tile GEMM, Step 2 adds the K-loop accumulation, Step 3 tiles…
Debugging Warp-Specialized Kernels
GEMM Steps 7-9 in Scaling GEMM with Warp Specialization and Clusters overlap TMA load, tcgen05 MMA, and TMEM/SMEM writeback. The same debugging method applies to Flash Attention…
Flash Attention 4
Attention runs two MMAs with softmax wedged between them, so it cannot just repeat one MMA the way GEMM does. The kernel composes the hardware primitives from Part I (TMA,…
What Makes a Kernel Fast
The roofline model gives a kernel a performance ceiling from memory bandwidth and compute throughput, while arithmetic intensity decides which ceiling applies. Low arithmetic…