花了大价钱请来的算力猛兽A100，被一段「逻辑上无懈可击」的CUDA代码按在地上摩擦——性能只跑出峰值的1.5%。这不是段子：有博主在一块A100上。知乎
先把尺子立起来。一组137 GFLOP的矩阵乘法，在19.5 TFLOPS的A100上「跑满」应该是7毫秒左右。而最符合C程序员直觉的第一版写法——开M×N个线程、每个线程算一个输出、`threadIdx.x`对应行下标——实测。知乎知乎
注意，这不是「算不动」。就算乘加单元一个不闲着，这活儿也只值7毫秒。时间全烧在了别的地方：数据送货的方式上。
为什么慢：GPU不按个发活儿
GPU调度线程不一个个来，而是32个打包成一个warp，同一时刻发射同一条指令、各拿各的数据。这就带来一条铁律：warp里32个线程的地址挨着，硬件就能把访存请求合并到最少的事务数。反过来，地址四散就只能一个地址一个事务，内存系统被迫做大量无用功——stride越大，带宽掉得越难看。英伟达开发者博客
回头看那段「无懈可击」的代码：`threadIdx.x`映射到row，正好让一个warp的32个线程读同一列，地址相隔16KB，一次读拆成32个独立事务。修法呢？交换两行代码的映射目标，性能直接10倍：有效带宽翻8倍，再叠加事务与排队开销的减少。知乎
修完还是只有15%：第二笔账
10倍听着爽，其实只到15%。因为每个线程算1个数，要读A的一行加B的一列。这些数据明明相邻线程都能复用，却每次都回显存重新取；解法早就写在硬件手册里：把复用搬进片上共享内存做分块。系列一路推下来，从1.5%最终干到83%，55倍跳变。知乎英伟达开发者博客
有意思的是连「代码的形状」都是账：block为什么偏取32×32？换成16×16，一个warp横跨两行，合并请求被切成两段64B。知乎
再往后是更硬核的分工：有博主把256个线程一分为二，一半专职计算一半专职搬运，搬运指令从计算流里彻底消失，靠双缓冲和同步信号接力。知乎
当搬运成为「专业岗」
这套思路在最新架构上已经是标配：Blackwell上跑BF16 GEMM，8个warp里1个专职搬数、1个专职算、4个收尾。博主的说法是「warp specialization已经成为常识」。硬件甚至干脆给了专职搬运单元：新异步执行特性里有全新的TMA单元，能高效地在显存和共享内存之间批量搬数据。知乎英伟达开发者博客
都用AI写kernel了，还学这个干嘛？
这是这周中文技术社区反复吵的点：agent开发kernel已经普遍，但手写这条路的价值没人敢说可以扔。Blackwell笔记作者的开场白很诚实——。评论区网友交出的证据也不乐观：拿公司的API花了几百上千美元、几个前沿模型轮流上，也很难说满意；还有人让agent反复否定最优路线，坚持验证到底之后，它飞快认了错。知乎知乎
结论其实不复杂：AI把「能跑的代码」变便宜了，但没有把「跑得快的代码」变便宜。峰值性能的账，最后还是要懂warp、懂访存的人来核。对想入门的人，路径就摆在这些系列里：自己手写一遍GEMM，亲眼看着1.5%→10倍→83%，GPU并行计算就理解了一大半。
至于普通玩家最该在意的那件事——同一张显卡，为什么软件换个写法就能差出10倍？答案从来不在算力，在于有没有把「货怎么送」这本账算清楚。
