今天 13 点 53 分，《手写 CUDA 矩阵乘法（二）》刚更新；往前翻一周，至少两位作者在死磕同一件事：不用 cuBLAS、不借任何现成轮子，从第一行核函数开始手写 FP32 矩阵乘法。同一周，还有人把 H100 上的 FlashAttention 用 CUDA/PTX 一块一块地手搓复现。知乎
外围同样热闹——《掌握GPU并行编程》连载更到了第十九章，这一章在讲向量加法。 B站的《GPU算子开发与性能优化实战》系列，也做到了"用性能模型指导 Tiled GEMM 设计"这一讲。 那个 15.6 万浏览的"学 CUDA 有什么推荐的学习路径"问题，本周还在添新回答——最新的答案说，写 CUDA 本质上跟传统的 C++ 软件工程毫无关系。 所有人绕不开的问题只有一个：从"环境装好了、程序能跑"到"我知道每一分性能是从哪来的"，中间那段路到底怎么走？两位作者刚好用两种记账方式，把这段路的账本摊开了。知乎哔哩哔哩知乎
第一本账：两行代码挣来的 10 倍
先立尺子。4096³ 方阵乘法，计算量 137 GFLOP；A100 的 FP32 峰值 19.5 TFLOPS，"跑满"约等于 7 毫秒。
最直觉的写法——开 M×N 个线程、每线程算一个输出、threadIdx.x 对到 row——跑出来是 291 GFLOPS，峰值的 1.5%，实测 472 毫秒，差了 67 倍。这 67 倍意味着瓶颈根本不在"算"：哪怕乘加单元一个都不闲着，这个 kernel 也得花 472 毫秒。知乎
病根在 warp：GPU 不逐个调度线程，而是 32 个线程编成一束、步调一致地发射同一条指令。naive 映射让一个 warp 的 32 个线程读 A 矩阵时，地址两两相差 16KB（K=4096×4 字节），散布在近 500KB 的空间里，硬件只能老老实实发 32 个独立事务。 同一个循环里读 B 却是反例：32 个线程读同一个地址，广播 1 个事务、几乎白拿。一贵一贱对照着看，慢的从来不是"读内存"这个动作，是地址的排布方式。知乎
然后是整个旅程里性价比最高的一次修改：循环体一行不动，只把两行映射互换、grid 两维对调。291 → 2972 GFLOPS，10 倍，没有增删任何一条指令。 第二篇把账算得更细：GPU 与显存之间搬运数据的最小粒度是 32 字节的 sector，修复后读 B 与写 C 变合并、读 A 变广播，有效带宽翻 8 倍，再叠加事务与排队开销的减少，落在 10 倍——量级完全对得上。知乎知乎
但修完只有 15.2%。46.5 毫秒里，计算需求依然只有 7 毫秒——这次的问题不在效率在次数：每线程算 1 个输出，A 的一行加 B 的一列，要读 8192 个 float，数据明明能被相邻线程复用，却每次都回显存重取一遍。知乎
block 形状也有代价：1024³ 下 16×16 比 32×32 慢 16%，因为一个 warp 要横跨两行、合并请求被切成两段 64B——连 block 尺寸都在为 coalescing 服务。知乎
作者自述这条路线最终把 FP32 从峰值的 1.5% 推到 83%，换上 Tensor Core 再在新尺子下推到 60%；从 15% 往后，“每一点性能都要真金白银地跟硬件特性做交易：数据复用、寄存器、指令发射、延迟隐藏，一样一样来”。知乎
第二本账：把线程劈成两半
另一位作者走的是另一种记账路线。在 ModelScope 环境上，n 从 32 到 2048 跑了七档方阵，他的手写 GEMM 在 n=1024 站上 cuBLAS 的 89.5%。知乎
手法核心是 warp specialization：256 个线程一分为二，一半 compute 线程专职计算整块 block tile、一半 loader 线程专职搬运，中间靠双缓冲和 `__syncthreads()` 交换"这批算完了、那批搬好了"的信号——搬运指令从计算线程的指令流里彻底消失。知乎
更有价值的是他前几篇踩出来的反直觉结论：绕过寄存器的异步搬运不一定更快，需要对 A 转置时，memcpy_async 反而要走 global→shared→register→shared 绕远路。知乎
双缓冲也不必然比单缓冲快——它没有改变搬运本身，重叠发不发生，取决于编译器有没有把"搬下一块"排到"算当前块"的前面；寄存器不忙的时候，异步搬运同样不会更快。 这几条恰恰是教程最常省略的部分：技术不是"用了就快"，每个都有自己的生效条件。知乎
两把尺子：89.5% 和 83% 说的可能不是同一件事
这是两本账并排放才看得出的坑。星河的百分比除以峰值算力：FP32 几篇按 19.5 TFLOPS 算尺；Tensor Core 那几篇他自己注明，“按 FP16 算力峰值 312 TFLOPS 计算，与前几篇的 19.5 不是一把尺子，不能直接横比”。知乎
东东的百分比除以 cuBLAS 实测值。所以"cuBLAS 的 89.5%"和"峰值的 83%"大概率量的是同一件事——cuBLAS 自己也很少跑满 100% 峰值；而 Tensor Core 篇的"60%"看着掉了一大截，换成绝对值（312×60%≈187 TFLOPS，对 19.5×83%≈16 TFLOPS）反而快了十倍以上。刷这类帖子先问分母是谁，再下结论。这一步两边原文都没给，是两本账交叉对照才浮出来的。
说反面：2026 年了，为什么还要手写
社区另一头的信号同样密集："TVM、MLIR、Triton 等深度学习编译器的主要区别是什么"这种问题挂着 8.7 万浏览，tilelang、cuTile 这类新写法不断有人下场。 B站《GPU编程的混乱状态》把跨平台生态的割裂摊开数了一遍——从 OpenGL、DirectX、Vulkan、WebGPU 到 CUDA。知乎哔哩哔哩
那手写图什么？东东的原话给了方向：“今天的最难的手写模式，明天就是一条指令（layout、TMA）”——手写的产出不是生产用的 kernel，而是读得懂硬件表达的那层直觉。 另一头同样直白——“不 profile 的 CUDA 优化等于瞎猜”。 profiling 报告里的 coalescing、bank conflict、stall 这些词，只有手写过一次 1.5% 的人能秒懂它们在说什么。知乎知乎
按三种人给结论
刚装好环境、没见过性能曲线的新手：别急着追 83%。今晚就能做完的闭环是——抄 naive 版跑出你的 1.5%，再互换那两行映射看它翻 10 倍。这一跳比读十篇文章更能建立"地址排布=性能"的直觉，半天可完成。
准备面试或往推理优化走的进阶用户：跟着东东的路线把双缓冲、warp 分工补上，并学会换记法汇报结果——"达到 cuBLAS 的百分之多少"比"跑多少 GFLOPS"更能让面试官对上尺子。
只在生产里用算子的人：可以不用手写，但值得把这两本账当排障地图：coalescing 修好才 15%，说明 kernel 的天花板从来不是算力问题；编译器不兜底调度，说明 autotuning 不是可选项。
两组实验都是单人单机：星河是 A100、4096³ 方阵，东东在 ModelScope、n=1024，数字搬不到你的卡上，能信的是每一跳的倍数量级和生效条件。后续值得盯的信号：星河系列还有两篇没更完——测量篇会专门拆解方阵测试对映射类错误的盲区，Tensor Core 篇也还在路上；东东说下一篇按自己的思路探索算子。两条线都还开着，适合收藏追更。知乎
