AI Infra、GPU 与大模型工程笔记
这里记录我对 AI Infra 与大模型工程的学习和梳理，重点关注 GPU、训练与推理系统，以及强化学习中的工程问题。从论文和源码出发，用公式、代码与图解解释实现细节。
从主题开始阅读
GPU 架构与 CUDA
从 Attention 的实现，读到 Blackwell 的计算与低精度支持。
大模型训练与推理
从编译执行，读到分片训练与推理时的 KV Cache 搬运。
强化学习与 Coding Agent
从 RLHF 与 GRPO，读到代码任务怎样成为可训练的环境。
最新文章
视频 DiT 加速：attention 为什么慢、稀疏怎么做、推理 trick
前置阅读：稀疏注意力为什么收敛到 DSA · Fast-WAM：Wan 预训练怎样变成动作策略
资料核对于 2026-11-10。标「推断」的是个人理解，没有论文或实测支撑。
1. 视频 DiT 的 attention 为什么慢
1.1 超长上下文：一段视频就是 ~76k token
Wan2.1 的 VAE 压缩比是 4×8×8（时间×高×宽），DiT 的 patchify 再压 1×2×2。81 帧视频的 latent 帧数是 $(81-1)/4+1=21$。两档分辨率下的序列长度：
分辨率
latent
token 数 $L$
480P（832×480）
...
PTX 内存一致性
前置阅读：FlashAttention：IO 感知的精确 Attention 计算 · Blackwell TMEM 与 tcgen05 MMA 深度解析
资料核对于 2026-11-10，对应 PTX ISA 9.5、FlashMLA 提交 395cb3a。标「推断」的是个人理解，文档和代码里没有直接写。
2025 年 4 月，FlashMLA 有一个只改了几行的提交，标题是 Fix synchronization issues。[1] 核心改动是把一个 __ldg 换成了普通的指针解引用：
cudaGridDependencySynchronize();
int *tile_scheduler_metadata_ptr = params.tile_...
被 mainloop 遮住的那几个点：kernel 的 Prologue 与 Epilogue
前置阅读：理解 CUTLASS 的 CLC dynamic persistent kernel · Blackwell TMEM 与 tcgen05 MMA 深度解析 · MegaMoE：EP token 流动与通信计算重叠 · FlashAttention 的空行 LSE
1. 一个 kernel 的三段式
几乎所有 tile 化的 GEMM / attention kernel 都可以切成三段：
时间轴 →
[ prologue ] → [ mainloop（稳态） ] → [ epilogue ]
建 barrier K 维（或 KV）循环 累加器收尾
取 descriptor TMA load + MMA r...
稀疏注意力为什么收敛到 DSA：相比 NSA 和 MoBA，它换来了什么
前置阅读：注意力头设计 · Blackwell TMEM 与 tcgen05 MMA 深度解析
资料核对于 2026-11-10，对应 FlashMLA 提交 395cb3a、DeepGEMM 提交 5081f85、DeepInfer 提交 9bd9cf2
2025 年 2 月，两篇块级稀疏注意力的论文几乎同时发布：DeepSeek 的 NSA 和月之暗面的 MoBA。[2][3]
NSA 专门解释了为什么要按连续的块来选：现代 GPU 读连续块的吞吐明显高于按索引随机读，按块计算也才能把 Tensor Core 用满。这也是当时的普遍看法，后来的综述总结得很直接：token 级的动态选择带来不规则访存，严重拖累 GPU 吞吐，于是块级稀疏成了主流。[14]
选择的块和 ...
DeepGEMM 的 TMA multicast 与 peer 信号：为什么停在 2，以及“先写数据再写信号”
前置阅读：Blackwell TMEM 与 tcgen05 MMA 深度解析
源码核对于 2026-11-10，对应 DeepGEMM 提交 5081f85（2026-10-20）
前阵子 TileLang 群里有人问：为什么没人给 DeepGEMM 贡献 multicast 代码？接着话题转到了 peer semaphore。有群友说他们的硬件就是把信号写进 peer 的 semaphore 寄存器，对方要么 WFI（等中断）要么轮询。另一位接话说，真正难的是 memory model：必须先写完数据才能写信号量。
翻了一遍 DeepGEMM 源码，第一个问题的前提并不成立：DeepGEMM 早就有 multicast，SM90 和 SM100 都有，只是 cluster...
kernel 的利用率怎么算：标准 MFU、TCU 与 TCA
前置阅读：MFU估算
同一个 kernel，有人报 66.7%，有人报 83.3%，还有厂商报出超过 100% 的 MFU。差别往往不在 kernel，而在分母怎么取。本文只对比三种算法：标准 MFU、TCU、TCA。
1. 标准 MFU
PaLM 给出的定义是 [1]：
\[\text{MFU}=\frac{\text{模型 FLOPs}/\text{耗时}}{P_{\text{peak}}}\]
分子：模型本身需要的 FLOPs，不计重算。计入重算的叫 HFU（Hardware FLOPs Utilization）。
分母：规格书上对应精度的 dense 峰值，例如 B200 BF16 为 2.25 PFLOPS [2]。
问题出在分母上。
规格书峰值是...
FlashAttention 的空行 LSE：FA-2 的 +inf 约定与后来的变化
前置阅读：FlashAttention：IO 感知的精确 Attention 计算
资料核对于 2026-11-09，对应 flash-attention 提交 478b5b0、FlashMLA 提交 395cb3a、FlashInfer 提交 0232a8b
一行 query 如果看不到任何 key，它的 log-sum-exp（LSE）应该是多少？
按定义是 $\log 0=-\infty$。FlashAttention-2 给出的却是 $+\infty$，FlashMLA 的接口也是如此，两种值并存让下游代码多了不少判断。Hopper 路径在 2025 年的一次重构里改回了 $-\infty$，但旧路径和一些分支至今还保留着 $+\infty$。
0. LSE 是什...
Blackwell 多轨矩乘（multi-rail GEMM）：FlashMLA 如何用一条 MMA 算完两路 QK
前置阅读：Blackwell TMEM 与 tcgen05 MMA 深度解析
资料核对于 2026-11-05，对应 PTX ISA 9.5 和 FlashMLA 提交 395cb3a
读 FlashMLA 的 SM100 代码时，可以看到一种非标准的 tcgen05.mma.ws 用法，代码里称为多轨矩乘（multi-rail GEMM，MRGEMM），目前固定为两轨。下文沿用这个叫法。
按 PTX 和 CUTLASS 的常规写法，.ws 在 M=64 时，A 要在 TMEM 的两组 lane 里各存一份相同的副本。
FlashMLA 没有照这个写法来：它往两组 lane 写入的不是两份相同的 Q，而是 Q 的不同特征维。于是一条 MMA 同时算出两份部分结果 P0 和...
共计 222 篇文章，28 页。