GPU 高性能算子工程

深度学习框架把矩阵乘、注意力、归一化都包装成一行 Python，但真正决定训练吞吐和推理延迟的，是这些算子在 GPU 上跑得有多快。当默认算子不够快、形状不规则、需要融合、或要支持新精度时，工程师就得自己写或调 kernel。

这个领域的资料分布很零散：CUDA 官方文档讲 API 不讲怎么调优，论文（FlashAttention、CUTLASS）讲思想不讲完整工程，博客多停在”GEMM tiling 入门”。本系列试图把硬件模型、性能方法论、核心算子实现、编程框架与算子库工程连成一条可复现的学习路径，填补 transformer 系列（讲算法）、编译器与 MLIR 系列（讲编译）、大模型基础设施系列（讲框架与服务）之间”手写快算子”这一缺环。

推荐入口

• GPU 执行模型：SM、warp、线程层次与 occupancy：一切性能直觉的地基，先搞清楚硬件到底怎么调度你的线程。
• Roofline 模型：判断算子是 compute-bound 还是 memory-bound：调优前先回答”瓶颈在哪”，否则全是瞎猜。
• FlashAttention：在线 softmax 与 IO-aware 注意力：把前面所有技巧（tiling、shared memory、recompute、访存）综合到一个真实算子里。

一、这个系列要回答的五个问题

1. 为什么同一个矩阵乘，朴素 kernel 和优化 kernel 能差几十倍？ 差距来自访存而非算力：朴素实现反复读 global memory，优化实现用 shared memory tiling 与寄存器分块把数据复用起来。详见第二、三部分（内存层次、访存优化、GEMM）。

2. 怎么知道一个算子还有多少优化空间，瓶颈到底在哪？ 靠 Roofline 模型与 profiler，而不是直觉。先用 Nsight 测出算术强度、访存吞吐、occupancy，再决定优化方向。详见第二部分（Roofline、Nsight 调优）。

3. Tensor Core 到底快在哪，为什么用上它反而经常更慢？ Tensor Core 要求特定数据布局与精度，喂数据（访存与布局转换）跟不上算力就会空转。详见第三部分（Tensor Core 与 MMA、CUTLASS）。

4. FlashAttention 为什么能在不近似的前提下省显存又提速？ 它把注意力重写成分块的在线 softmax，避免落盘 \(N\times N\) 的注意力矩阵，用重算换访存。详见第三部分（FlashAttention）。

5. 该手写 CUDA、用 Triton，还是调 CUTLASS / cuBLAS？ 取决于形状规则性、融合需求、维护成本与目标硬件。本系列给出判断框架，而不是”哪个最好”的口号。详见第四、五部分（Triton、kernel fusion、算子库工程）。

二、篇目依赖关系与推荐阅读路径

全系列共五部分、21 篇。各篇按「问题—硬件/原理—代码—profiling 验证—边界」组织；动手实验集中在第 04、10、15 篇，调优工具箱集中在第 08 篇。

强依赖

flowchart TD
  A["01 全景：算子在 AI 栈的位置"] --> B["02 执行模型"]
  B --> C["03 内存层次"]
  C --> D["04 第一个 CUDA kernel"]
  D --> E["05 访存优化"]
  D --> F["06 occupancy 与延迟隐藏"]
  E --> G["07 Roofline"]
  F --> G
  G --> H["08 Nsight 调优"]
  H --> I["09 Reduction / Scan"]
  H --> J["10 GEMM 从朴素到 tiled"]
  J --> K["11 Tensor Core 与 MMA"]
  K --> L["12 CUTLASS / CuTe"]
  I --> M["13 softmax / layernorm 与融合"]
  J --> N["14 FlashAttention"]
  M --> N
  L --> O["15 Triton"]
  N --> O
  O --> P["16 kernel fusion"]
  P --> Q["17 量化与多精度算子"]
  Q --> R["18 通信与计算重叠"]
  R --> S["19 算子库工程"]
  S --> T["20 调试与数值正确性"]
  T --> U["21 趋势：TMA / Blackwell / 编译器协同"]

推荐阅读路径

• CUDA 入门到能调优 01 → 02 → 03 → 04 → 05 → 06 → 07 → 08 → 10

• 深度学习算子开发者 02 → 03 → 07 → 09 → 10 → 11 → 13 → 14 → 16

• 推理性能工程师 07 → 08 → 11 → 12 → 14 → 16 → 17 → 18

• 想用 Triton 快速落地的工程师 02 → 03 → 05 → 07 → 15 → 16 → 13（13 放在最后：掌握 Triton 编程模型与融合套路后，softmax/layernorm 正好是一个可直接上手改写的融合算子实例）

• 算子库 / 框架维护者 08 → 12 → 15 → 16 → 19 → 20 → 21

三、目录与每篇一句话价值

第一部分：硬件模型与编程基础

• 01. 全景：算子工程在 AI 计算栈的位置：从框架一行 matmul 到 PTX/SASS，看清”框架算子 → 库 → 手写 kernel”的分层，回答什么时候才需要自己写。
• 02. GPU 执行模型：SM、warp、线程层次与 occupancy：grid/block/warp 如何映射到 SM，warp 内 32 线程的 SIMT 执行与分支发散为何昂贵。
• 03. 内存层次：global / L2 / shared / register 的带宽与延迟：每一级存储的容量、带宽、延迟量级，以及”数据放哪”如何主导算子性能。
• 04. 写第一个 CUDA kernel：索引、同步与启动配置：从向量加法到逐元素算子，讲清线程索引计算、__syncthreads() 与 launch 参数选择。

第二部分：性能方法论

• 05. 访存优化：合并访问、bank conflict 与对齐：global memory 合并访问与 shared memory bank conflict 的成因与规避，访存模式决定带宽利用率。
• 06. Occupancy 与延迟隐藏：寄存器、shared memory 的取舍：occupancy 不是越高越好，寄存器/shared 压力与延迟隐藏之间如何权衡。
• 07. Roofline 模型：compute-bound 还是 memory-bound：用算术强度把算子定位到 Roofline 曲线，决定该优化算力还是访存。
• 08. Nsight 调优工作流：Compute 与 Systems 怎么读：Nsight Compute 的关键指标（SM/内存吞吐、stall reason）与 Nsight Systems 的时间线，定位真实瓶颈。

第三部分：核心算子实现

• 09. Reduction 与 Scan：warp shuffle、block 级与 grid 级归约：从串行归约到 __shfl warp 级归约的逐步优化，前缀和的并行实现。
• 10. GEMM：从朴素实现到 shared memory tiling 与寄存器分块：矩阵乘逐步优化的经典路径，数据复用如何把访存量从 \(O(N^3)\) 降到可接受。
• 11. Tensor Core 与 MMA：wmma、mma.sync 与数据布局：Tensor Core 的 MMA 指令、fragment 布局与 FP16/BF16/TF32/FP8 精度，喂数据为何是关键。
• 12. CUTLASS 与 CuTe：模板化 GEMM 与布局代数：CUTLASS 的分层抽象与 CuTe 的张量布局代数，如何在不手写 PTX 的前提下逼近 cuBLAS。
• 13. Softmax、LayerNorm 与逐元素融合：归约类算子的数值稳定写法（在线最大值、Welford），以及如何与逐元素操作融合减少访存。
• 14. FlashAttention：在线 softmax 与 IO-aware 注意力：分块在线 softmax 的推导、用重算换显存的设计，以及它如何把注意力从 memory-bound 拉回算力。

第四部分：编程框架与融合

• 15. Triton：tile 级编程模型与 autotune：Triton 的编程抽象、自动向量化与 autotune，与手写 CUDA 的能力边界对比。
• 16. Kernel Fusion 与 epilogue：减少 HBM 往返：为什么融合能提速、哪些算子值得融合、epilogue fusion 的工程做法与代价。
• 17. 量化与多精度算子：INT8 / FP8、反量化与 per-channel：低精度算子的布局、反量化时机与精度对齐，与推理量化的衔接。
• 18. 通信与计算重叠：NCCL collective 与 kernel overlap：多卡训练/推理中 collective 与计算的重叠技巧，与并行策略的关系。

第五部分：算子库工程与未来

• 19. 算子库工程：dispatch、autotune cache 与 JIT：面向多形状多硬件的 kernel 选择与缓存，编译期与运行期 dispatch 的取舍。
• 20. 调试与数值正确性：compute-sanitizer 与对齐测试：竞态与越界检查、数值与参考实现对齐、算子单元测试的工程方法。
• 21. 趋势：TMA、Blackwell、ThunderKittens 与编译器协同：新硬件特性（TMA、更细粒度异步）、新编程抽象与编译器自动生成算子的方向，接回编译器系列。

四、读者定位与先修要求

五、写作方法与风格

• “瓶颈优先”叙述：每篇先用 Roofline 或 profiler 指出瓶颈，再给优化手段，避免”堆技巧”。
• 代码与 profiling 配套：核心算子给出可编译的最小 kernel，并配 Nsight 关键指标的读法；性能数据只在真实跑过后给出，并标注 GPU 型号、CUDA 版本、矩阵规模等口径。
• 可视化优先：线程/内存层次、tiling 数据流、Roofline、注意力分块默认配图。
• 诚实的边界：明确每种技巧的适用形状、精度与硬件代际，不把单一结论当万能。

六、关键参考资料

官方文档与规范（A 级）

• CUDA C++ Programming Guide（NVIDIA，注明使用版本）
• CUDA C++ Best Practices Guide
• PTX ISA 与 Nsight Compute / Nsight Systems 文档
• CUTLASS 与 CuTe 官方文档

论文（A 级）

• Dao, T. et al. FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness, NeurIPS 2022.
• Dao, T. FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioning, 2023.
• Williams, S. et al. Roofline: An Insightful Visual Performance Model for Multicore Architectures, CACM 2009.
• Tillet, P. et al. Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations, MAPL 2019.

开源项目（B 级）

• CUTLASS
• Triton
• FlashAttention

七、与其他系列的关联

• 注意力、Transformer 的算法层面请看 Transformer 与注意力机制系列。
• 算子如何被编译器自动生成与降阶请看 编译器与 MLIR 系列，尤其是面向异构硬件的代码生成部分。
• GPU 在训练/推理基础设施中的位置、NCCL 互联与并行策略请看 大模型基础设施工程系列。

八、本系列的承诺与不承诺

承诺：

• 核心算子配可编译的最小 CUDA / Triton 代码；
• 所有性能数据标注 GPU 型号、CUDA 版本、问题规模与采样口径，没跑过的不写成结论；
• 解释”为什么这样写更快”而不仅是”用哪个 API”。

不承诺：

• 不替代 CUDA 官方文档，本系列是学习与调优路径，不是 API 手册；
• 不覆盖所有 GPU 厂商与所有代际特性，重点放在通用方法与主流 NVIDIA 平台；
• 不预测哪种编程框架会赢，只给出按场景选择的判断框架。

更新日期：2026-06-26
作者：ltl
系列相关：Transformer 与注意力机制 | 编译器与 MLIR | 大模型基础设施工程