阿甘的博客 | Arganzheng's Blog

GPU Kernel 工程（11）：系列总结与通关自测

GPU Kernel Engineering: Series Recap and Final Self-Test

十篇正文回答了一个问题：一个 kernel 为什么快、为什么慢，以及如何把它写到接近硬件极限。第一篇把 GPU 拆开并建立 Roofline，第二篇写出第一个 kernel 并学会测量，第三、四篇把 memory-bound 的 elementwise 与 reduction 推到带宽墙，第五、六篇把 GEMM 从 naive 推到 Tensor Core，第七篇用 Triton 看编译器接管了哪一层，第八、九篇把这些工具用到 attention、量化与融合 kernel 上组装出一个 decoder layer，第十篇讲怎么剖析、测试、接入框架并合入一个 PR。本文不讲新内容，做三件事：把十篇压成一张表与十段回顾，把贯穿全系列的几条线拎出来，然后给一套三段式的...

Posted by arganzheng | May 20, 2026

GPU Kernel 工程（10）：剖析、测试与贡献——把 kernel 做成产品

Profiling, Testing and Contributing: Turning a Kernel into a Product

前九篇结束时，手上有一个用自己写的 kernel 跑通的 decoder layer 前向：RMSNorm、RoPE、BF16 Tensor Core GEMM、FlashAttention 前向、SiLU-mul、fused residual+RMSNorm、INT4 weight-only GEMM。它们能跑、结果和 PyTorch eager 对得上、每一个都在自己的 benchmark 里比 naive 版本快很多。但”能跑”和”能合入”之间还有一整段工程。一个 kernel 要成为别人敢用的东西，需要回答四个问题：它到底卡在哪里（剖析）；它在所有会遇到的输入上都对（测试）；它确实比原来快、而且以后不会悄悄变慢（benchmark）；它在别人的 GPU ...

Posted by arganzheng | May 20, 2026

GPU Kernel 工程（09）：量化与融合 kernel——推理系统的其余部分

Quantized and Fused Kernels: The Rest of the Inference Stack

上一篇把 attention 讨论完了。一个 decoder layer 里除了 attention 和标准 GEMM，剩下的是一堆”小 kernel”：RMSNorm、RoPE、SiLU-mul、把 KV 写进分页 cache、把权重从 INT4 解开、把激活压成 FP8、MoE 的 token 重排、采样。它们单个都不复杂，但数量多、变化快，加起来占掉推理时间的一个可观比例——而且是 vLLM、SGLang 这些项目里 PR 最活跃的区域。这一篇把它们放在同一个方法论下过一遍：先算这个 kernel 理论上要搬多少字节、做多少 FLOPs，再看实现，再解释差距。核心问题是总纲提出的那个：一个 INT4 weight-only GEMM，decode...

Posted by arganzheng | May 19, 2026

GPU Kernel 工程（08）：Attention Kernel——FlashAttention 与 PagedAttention

Attention Kernels: FlashAttention and PagedAttention from Derivation to Code

前七篇分别处理了 GPU 的硬件结构与 Roofline、CUDA 执行模型、访存合并、shared memory 与 reduction（其中包括 online softmax）、GEMM 的分块、Tensor Core 与 CUTLASS、Triton。这一篇是它们的汇合点：attention 同时包含两个 GEMM（\(QK^T\) 与 \(PV\)）、一个逐行的 reduction（softmax），以及推理时特有的内存访问模式（分页的 KV cache）。它是 Transformer 推理里最重要、也最难写好的 kernel。总纲给这一篇提的核心问题是：一个序列长度 4k、head dim 128 的 attention，标准实现和 Flas...

Posted by arganzheng | May 18, 2026

GPU Kernel 工程（07）：Triton——块级编程与编译器的边界

Triton: Block-Level Programming and Where the Compiler Stops

前六篇一直在 CUDA 的世界里：每个线程算什么、warp 怎么合并访存、shared memory 怎么分块、mma.sync 怎么喂 fragment。到了第六篇，一个能跑到 cuBLAS 七八成性能的 BF16 GEMM 已经是两三百行代码，而且每一行都有”为什么这样写”的理由——tile 尺寸、bank conflict 的 padding、cp.async 的 stage 数、寄存器分块的形状。这一篇换一种写法。Triton 把”线程”从编程模型中拿掉：程序员以 block 为单位思考，写的是”这个 block 加载哪一块数据、做什么张量运算、存到哪里”，而线程到元素的映射、shared memory 的分配、向量化、软件流水、mma 指令的选择，全部...

Posted by arganzheng | May 17, 2026

GPU Kernel 工程（06）：Tensor Core、CUTLASS 与 CuTe

Tensor Cores, CUTLASS and CuTe: Programming the Matrix Units

上一篇用 CUDA Core 把 GEMM 的分块结构讲透了。回顾一下那个结构，因为本篇要做的事情就是把它”接”到另一种计算单元上：一个 thread block 负责输出矩阵 \(C\) 的一个 \(BM \times BN\) 的 tile，沿 \(K\) 维以 \(BK\) 为步长循环；每一步把 \(A\) 的 \(BM \times BK\) 子块和 \(B\) 的 \(BK \times BN\) 子块搬进 shared memory，全局读取量从 naive 的 \(2MNK\) 个元素降为 \(MNK \cdot (1/BM + 1/BN)\)，\(BM = BN = 128\) 时减少 128 倍（\(2 / (2/128)\)）； ...

Posted by arganzheng | May 16, 2026

GPU Kernel 工程（05）：GEMM——从 naive 到分块

GEMM from Naive to Tiled: Reaching the Compute Ceiling on CUDA Cores

前四篇讨论的 kernel——elementwise、reduction、softmax、LayerNorm——有一个共同点：它们都是 memory-bound 的。每个元素读进来、算一两次、写回去，算术强度远低于 ridge point，优化的全部目标是”把 HBM 带宽用满”。做到了带宽的 80–90%，这类 kernel 就到头了。 GEMM 是这个系列的转折点。一个 4096×4096×4096 的矩阵乘法有 1374 亿次浮点运算，但三个矩阵合计只有 192 MiB（FP32）。它的算术强度是 elementwise 的几千倍，理论上应该被算力而不是带宽限制。但一个不加思考写出来的 GEMM kernel，跑出来的却是 FP32 峰值的百分之一二——它被...

Posted by arganzheng | May 15, 2026

GPU Kernel 工程（04）：共享内存与 reduction——softmax、LayerNorm 与 online softmax

Shared Memory and Reductions: Softmax, LayerNorm and Online Softmax

上一篇的 elementwise kernel 有一个共同特征：每个线程只管自己的元素，线程之间不需要说话。把访存合并、向量化、grid-stride 做对，带宽就能推到 90% 以上。这一篇进入需要线程协作的 kernel。RMSNorm 要先知道整行的均方才能缩放每个元素；softmax 要先知道整行的最大值和指数和才能归一化每个元素。”先算出一个整行的标量，再用它处理每个元素”——这个模式叫 reduction，它是所有 norm、softmax、loss、统计类算子的核心，也是 FlashAttention 内循环的一半。总纲给这一篇的核心问题是：一个 4096 维的 RMSNorm，读一次写一次，理论上是 memory-bound 的。为...

Posted by arganzheng | May 14, 2026

GPU Kernel 工程（03）：访存合并与 elementwise kernel

Memory Coalescing and Elementwise Kernels: Hitting the Bandwidth Ceiling

上一篇写出了第一个 kernel：一个 BF16 的 y = x + b，每个线程处理一个元素。它能跑、结果正确，但没有回答”它跑得够快吗”。这一篇就回答这个问题，并把答案推到极限。先把理论下界放在最前面，后面所有讨论都对着它算。BF16 的 y = x + b 对每个元素读 2 个 BF16、写 1 个 BF16，共 6 字节，做 1 次加法（1 FLOP）： \[\text{算术强度} = \frac{1\ \text{FLOP}}{6\ \text{B}} \approx 0.17\ \text{FLOP/B}\] A100 SXM 80GB 的 HBM2e 带宽约 2.0 TB/s，BF16 Tensor Core 约 312 TFLOPS（均为公开...

Posted by arganzheng | May 13, 2026

GPU Kernel 工程（02）：CUDA 编程模型与第一个 kernel

The CUDA Programming Model and Your First Kernel, Measured

上一篇建立了本系列的分析框架，用到的结论可以压缩成三个数字。GPU 的基本执行单位是 warp：32 个线程共用一个指令流，一条指令同时作用在 32 个数据上。以 A100 SXM 80GB 为默认分析对象（标称值）：HBM2e 带宽约 2.0 TB/s，BF16 Tensor Core 算力 312 TFLOPS，两者相除得到 Roofline 的拐点（ridge point）： \[\text{ridge} = \frac{312 \times 10^{12}\ \text{FLOP/s}}{2.0 \times 10^{12}\ \text{byte/s}} \approx 156\ \text{FLOP/byte}\] 算术强度低于 156 FLOP/b...

Posted by arganzheng | May 12, 2026

Arganzheng's Blog