最近刷知乎，发现一个很有意思的现象：8月份冒出来一批"手写FlashAttention翻车复盘"的文章，作者背景各不相同，但踩的坑几乎一模一样。
先说最扎心的一份。一位做了十多年游戏引擎、最近转学AI的开发者，在RTX 3060上手写简化版FlashAttention，用Nsight Compute实测：DRAM流量从朴素实现的503MB降到21MB，整整砍了24倍——按"attention是访存瓶颈"的标准叙事，这该起飞了吧？结果端到端只快了1.9倍。知乎
24倍的访存收益，最后只兑现成1.9倍的加速，中间差出去的12倍去哪了？
这个问题不只是他一个人的困惑。我把8月份能找到的几份独立实测放在一起对照，发现大家撞上的基本是同样三堵墙。如果你也在用手写算子的方式学CUDA，或者正对着profiler报告发呆，这篇能帮你省不少弯路。
三份8月实测，撞的是同一组墙
先把三份实测的数据摆在一起，都是作者自己跑出来的原始数字：
第一份，上面提到的简化版FA（RTX 3060，sm_86，FP32，seq=4096，head_dim=64）：朴素三kernel实现20.46ms，融合后10.62ms，快1.9倍；DRAM流量503MB降到21.1MB，降了约24倍。
第二份，一位博主在Blackwell系的sm_103平台上死磕FlashAttention kernel：PyTorch原生SDPA（走cuDNN auto调度）的baseline是1156 TFLOPS，他从零手写的朴素版本只有56.83 TFLOPS——差了20倍。改了第一个swizzle配置后，升到91.46 TFLOPS，提升1.6倍。知乎这个系列的目标是986 TFLOPS，也就是说，手写版要追到库实现的85%才算达标。
第三份稍微"跑题"但同样典型：有人优化Softmax算子，v2版本用warp shuffle把"寄存器→共享内存→寄存器"的路径优化成"寄存器→寄存器"，理论上稳赚，结果跑出来反而更慢、耗时暴增。知乎
三份实测，三个不同的算子场景，问题最后都指向同样三个地方：occupancy、共享内存bank conflict，以及一个特别阴险的——你的编译配置在骗你。
第一堵墙：访存降了，GPU却在空转
先看24倍访存收益为什么只兑现1.9倍。
打开Nsight Compute的报告，两个数字非常刺眼：occupancy只有7.24%，shared memory bank conflict高达2.35亿次。知乎
occupancy的问题出在block设计上：这个实现里block size只有32，也就是一个warp独自处理一整个32×64的Q tile。再叠加共享内存的限制——sm_86每个SM大约100KB共享内存，每个block占24KB，每个SM最多驻留4个block——算下来每个SM只有4个活跃warp，而上限是48个。理论occupancy就是8%左右，实测7.24%，严丝合缝。
这意味着什么？GPU绝大部分时间在空转。只有4个warp给调度器用，一旦它们全卡在访存上，SM就没有别的工作可以切过去做，访存延迟完全暴露，没有任何东西可以隐藏。你省下来的每一字节DRAM流量，都被"干等"吃掉了大半。
这也是很多人对roofline模型的一个误解：知道"这个算子是memory-bound"只是第一步。memory-bound的反面不是"随便写都能跑满带宽"，而是你得有足够的并发在途请求去把带宽填满。warp不够，带宽就是打不满的——这时候访存优化做得再漂亮，也换不来时间。
顺带说一个这份实测里很有价值的细节：作者专门区分了"访存量"的三层含义——算法逻辑上要求的读写量（algorithmic request）、经过L1/L2的实际流量（cache traffic）、真正打到显存的字节数（DRAM bytes）。只有第三层直接对应带宽成本。比如FlashAttention里K/V按算法逻辑会被每个Q block重读一遍，光请求量就有两百多MB，但因为tile小、复用密集，绝大部分被L2接住了，DRAM上几乎不可见。手算访存量通常只能算到第一层，跟实测能差一个数量级——所以别手算了，直接向ncu要 `dram__bytes_read.sum` 和 `dram__bytes_write.sum`。知乎
第二堵墙：共享内存的bank conflict，比你想的贵得多
第二堵墙在sm_103那份实测里看得最清楚。
作者第一版kernel里，PV这段矩阵乘的共享内存swizzle选了"不做处理"（SwizzleNone），当时的判断是WarpH=64、Bc=64，“感觉不太需要额外处理”。
结果用指令数一算就露馅了：P矩阵存的是fp16，row stride=128B，而共享内存是32个bank、每bank 4B、总带宽128B/cycle。128B的row stride意味着相邻行同一列的地址永远落在同一组bank上——8-way bank conflict，一次加载被拆成8次串行。算下来PV这边ldmatrix的有效延迟大约640 cycles，而mma本身只要256 cycles：计算单元在干等数据。知乎
修复方法很直接：把SwizzleNone换成Swizzle128，让同一列的相邻行打散到不同bank。改完benchmark：56.83 → 91.46 TFLOPS，提升1.6倍，profiler里bank conflict明显下降。
顺便提一句，8月份还有一份"Softmax性能悬案"也是同一堵墙的变体：作者发现v2版本加了warp shuffle后反而更慢，ncu里看到Local Memory有流量、SASS里出现STL/LDL指令。排查方向本身没错——先怀疑函数没内联，加 `forceinline`，STL/LDL确实消失了；再怀疑寄存器溢出，用 `–ptxas-options=-v` 一看，0溢出。线索全部指向死胡同。知乎
第三堵墙：你的编译配置在骗你
这份"悬案"最后的答案让人哭笑不得：问题出在CUDA C/C++的编译选项里，Generate GPU Debug Information（也就是-G模式）开着。
-G一开，`__shfl_down_sync` 这种本应内联展开的函数变成了真实的函数调用，SASS里出现 `CALL.ABS.NOINC`，还伴随 `MOV R4, R4` 这种peephole优化失效的无效指令——性能数据自然全是假的。知乎作者折腾了几个小时，最后还是问了大模型才定位到：测性能之前，确认Visual Studio是Release还不够，还得确认CUDA设备编译的debug信息是关的。
这个坑值得所有写CUDA的人记一笔：性能实测之前，先确认编译配置。知乎-G模式、没内联的设备函数、debug构建，都会让你的profiler数据变成噪声。方向错了不可怕，可怕的是拿着假数据找方向。
往上再看一层：这不是你菜，是题目本来就难
撞完这三堵墙，可能有人会怀疑：是不是手写kernel这条路本身就不值得？
不妨看看"标准答案"们现在在跟什么搏斗。
今年3月发布的FlashAttention-4（论文arXiv:2603.05451）针对Blackwell架构做了一轮彻底的重设计。背景是一个关键的硬件变化：B200的BF16 Tensor Core算力从H100的约1 PFLOP/s翻倍到约2.25 PFLOP/s，但共享内存带宽、负责exp等超越函数的SFU、通用ALU几乎没有增长——论文里叫"不对称硬件扩展"。知乎
对attention这种"两个GEMM中间夹一个softmax"的算子，这个变化是决定性的：GEMM越来越快，但GEMM之间的搬运、同步、逐元素运算没跟上。FA4的分析结论很反直觉——在B200上，前向的瓶颈是softmax里的指数运算，反向的瓶颈是共享内存带宽，都不是矩阵乘法本身。知乎为此FA4引入了TMEM（每SM 256KB的张量内存）、tcgen05.mma、2-CTA MMA、双Q tile的ping-pong调度，实现语言也从CUDA C++换成了CuTe-DSL。
Hopper时代的FA3大概能挖到峰值的75%，而Blackwell上连Tri Dao团队都要围着"exp算得太慢"重新设计整个流水线。知乎所以你在3060上被occupancy和bank conflict卡住，真不是你菜——这个算子在任何一代硬件上都在考验"除矩阵乘以外的所有东西"。
一份可以直接抄的自查顺序
把上面这些复盘归纳一下，手写kernel性能不对劲时，按这个顺序查，比乱枪打鸟效率高得多：
先对baseline。直接跟torch SDPA、cuBLAS这类库实现比，别跟自己昨天的版本比。知道差距是20倍还是20%，策略完全不同。
先验正确性，再谈性能。结果不对的快，没有任何意义。
查编译配置。确认不是debug模式、不是-G在捣乱——这一步成本最低，但最容易忘。
查occupancy。block size、每block共享内存占用、寄存器用量，任何一个都会限制活跃warp数。warp太少，延迟就藏不住。
查bank conflict和访存模式。shared load/store指令数对比conflict次数，ldmatrix的swizzle配置，row stride是不是正好撞上bank边界。
看stall reason定位瓶颈类型。是等访存（long scoreboard）、等共享内存（short scoreboard）还是指令依赖，ncu会直接告诉你。
最后才轮到流水线设计：双缓冲、warp specialization、计算与搬运重叠——这些是大招，前面没查完之前用不上。
另外两个提醒：访存量别手算，直接读ncu的DRAM字节数；优化每一步都重新跑profiler，别靠"感觉应该会好"。
手写值不值？看你想拿走什么
最后回答那个绕不开的问题：库实现都这么强了，手写还有意义吗？
分情况。如果你的目标是业务里把模型跑得更快——大多数情况下直接用SDPA、FlashAttention官方库、cuDNN，它们对新架构的适配深度（TMEM、2-CTA MMA这些东西）不是个人几周能追上的。而且同一份kernel代码，换个架构性能差几倍是常态，独立评测里连CUDA Tile这种新范式写的attention，在B200上能到FlashAttention-2的2.5倍，换到同家族的RTX PRO 6000上只剩53%——"写一次到处快"离现实还很远。知乎
但如果你的目标是学——手写的价值恰恰在"翻车"里。上面三份实测的作者，收获最大的都不是那个快了几倍的kernel，而是：看懂ncu报告的能力、对warp/shared memory/bank的肌肉记忆、以及"性能问题如何定位"的完整闭环。这些能力在你读Triton代码、调库参数、排查线上推理瓶颈时全都用得上。有招聘观察就提到，现在招算子方向的公司，看重的正是这种"能把profiler读明白"的人，而不是会背API的人。
一句话：手写是训练手段，不是交付物。用库交付，用手写练眼力，两边都不亏。
接下来值得持续关注的信号：FA4之后，CuTe-DSL和CUDA Tile这类"声明tile、编译器生成指令"的范式会不会继续吞掉手写kernel的生态位；以及sm_100和sm_120这种同代架构之间的巨大性能差异，会不会倒逼更多项目放弃单架构深度调优、转向可移植写法。这两个方向一旦有实测数据出来，值得再回来对照一次。
你手写算子时撞过最离谱的坑是什么？评论区聊聊。
