做性能优化做久了你会遇到一个特别迷惑的现象nvidia-smi里 GPU 利用率已经顶到 99%代码也看不出什么明显毛病可实际吞吐就是上不去。我以前调一个多头注意力的 Kernel 时就被这个问题卡了快一周后来用 Nsight Compute 打开一看真相很有意思——SM 确实在忙但它忙的是等数据、搬数据真正干活的 Tensor Core 流水线大部分时间居然闲着。那一刻我才彻底理解了一个判断GPU 利用率和 Tensor Core 利用率是两个完全不同的指标。这是AI 系统性能工程学习笔记系列的第九篇。前几篇我们聊过 GPU 架构、内存层次、线程模型、算子实现这篇进入整个系列最核心的一个话题让 Tensor Core 吃饱。我会从 Tensor Core 的工作特性、算术强度的计算方法、Kernel 改造的具体动作、实测数据、分析工具和常见伪优化几个角度把怎么判断 Kernel 有没有效率、怎么让算术强度撑得起来这件事讲透。适合正在做大模型训练/推理性能优化、手写 CUDA 算子或者被 PyTorch 算子拖慢的朋友。1. 为什么 GPU 占用率 99%Tensor Core 却在饿肚子1.1 SM 忙不等于计算单元忙先厘清一个最容易被误解的概念。GPU 在跑 Kernel 的时候nvidia-smi显示的 GPU-Util 是基于采样统计出来的某个时刻至少有线程在 SM 上执行的比例。注意它看的是 SM 上有没有活跃的 warp而不是这些 warp 在做什么。一个 warp 可能在执行算术指令也可能在等显存数据、等 shared memory 地址冲突解决、等同步栅栏。只要 warp 没被调度器换出去SM 就算忙。但在等数据的那些周期里计算单元——尤其是 Tensor Core——是完全闲置的。这就是为什么你会看到 GPU 占用率 99%、但 kernel 耗时是理论最优值好几倍的怪现象。我后来在 NCU 的报告里找到了直接证据同一个 kernelsm__throughput.avg.pct_of_peak_sustained_elapsed显示 87%但sm__pipe_tensor_op_hmma_inst_executed_pct_of_peak_sustained_active只有 12%。SM 确实没闲着但 Tensor Core 这一条专用的矩阵乘加流水线绝大部分时间都在空转。1.2 算力与带宽之间的剪刀差Tensor Core 之所以容易饿肚子根源在算力和带宽的极端不平衡。A100 80GB 这张卡FP16 Tensor Core 理论算力 312 TFLOPSHBM 显存带宽 2TB/s 左右。不算不知道一算吓一跳要让 Tensor Core 满速跑每读 1 字节数据进来就得配套执行 156 次浮点运算。再看 CUDA Core 的 FP32 算力19.5 TFLOPS同样除以带宽拐点只有大约 10 FLOP/byte。两个数字摆在一起问题很尖锐Tensor Core 对数据供给密度的要求是普通 CUDA Core 的 15 倍以上。也正因为如此从 Volta 到 HopperTensor Core 的算力翻了好几番但 HBM 带宽只涨了不到一倍这个剪刀差还在持续扩大。1.3 拐点判断 Kernel 属性的分水岭上面那个 156 FLOP/byte在 Roofline 模型里叫拐点ridge point。它的物理含义很直白如果你的 Kernel 算术强度高于拐点理论上限由算力决定就是计算密集型compute-bound低于拐点理论上限被带宽卡死就是访存密集型memory-bound。判断 Kernel 到底属于哪一类是性能优化的第一步。方向搞反了后面全是白干。比如一个 memory-bound 的 Kernel你花大力气优化计算指令排列收益趋近于零反过来一个 compute-bound 的 Kernel你把数据搬来搬去调布局也是缘木求鱼。这个判断用不着上复杂的分析工具拿算术强度的估算值和拐点一比就出来了。我在实际项目里验证过很多次手写 GEMM 的时候只要矩阵规模上不去比如单次只算 64×64 的小块算术强度就很容易跌破拐点这时候无论怎么排布指令性能都上不去必须先解决数据复用的问题。所以算术强度不是纸上谈兵的学术概念它是决定优化方向和资源分配的定盘星。2. 算术强度先算出你的 Kernel 每字节干了多少活2.1 定义与估算方法算术强度Arithmetic Intensity简称 AI的定义很简单一个 Kernel 总共执行的浮点运算次数除以它从显存中读写的字节总数单位是 FLOP/byte。[ AI \frac{总浮点运算量}{总访存字节数} ]关键是把总访存字节数算对。很多人只算了输入输出张量的大小忽略了中间结果的写回和读回这样算出来的算术强度会虚高。正确做法是把 Kernel 生命周期内所有经过显存的数据都算进去原始输入、中间结果、最终输出包括临时 buffer。以最经典的 C A × B 矩阵乘法为例假设三个矩阵都是 N×N 的 FP16。运算量是 2N³ FLOP乘加算两次浮点操作。如果把 A、B 各读一次、C 写一次访存量是 3N²×2 字节。算术强度就是[ AI \frac{2N^3}{6N^2} \frac{N}{3} ]这个结果信息量很大矩阵越大算术强度线性上升。N128 时 AI≈42远低于 A100 的 156 拐点N4096 时 AI≈1365远超拐点。换句话说大矩阵乘天然是计算密集型的小矩阵乘天然是访存密集型的。这个结论解释了大模型推理里常见的窘境你用大矩阵训练时 GPU 跑得很欢一到小 batch 推理、小形状的算子比如 GQA 里的分组矩阵乘性能就断崖式下跌。2.2 中间结果正在偷偷吃掉你的算术强度估算算术强度时最容易漏掉的是中间结果。我举个例子一个典型的 Attention 计算朴素流程是Q 和 K 做矩阵乘得到 N×N 的分数矩阵 S写回显存对 S 做 softmax读出、算完、再写回显存拿 softmax 后的 S 和 V 做矩阵乘得到输出每多一次写回再读回就额外增加 2×N×N×2 字节的访存开销而且不产生任何新运算。当序列长度 N 很大的时候这个开销会直接压垮算术强度。标准 Attention 的理论算术强度是 O(N) 级别的本来挺高但因为中间结果反复落地到显存实际的有效算术强度可能只有理论值的三分之一甚至更低。这也是 FlashAttention 这类融合 Kernel 能带来数倍加速的底层原因——它做的事本质上就是砍掉中间结果的反复落地让 S 矩阵只活在 shared memory 和寄存器里。后面第 3 节我会专门讲融合的具体思路。2.3 一组实用的快速估算参照下面这组数字是我做性能预判时常用的参照都基于 A100 80GB 的 FP16 拐点 156 FLOP/byte算子形状理论算术强度初步判断4096×4096 GEMM约 1365计算密集优先优化 Tensor Core 指令1024×1024 GEMM约 341勉强越过拐点注意 tile 和复用128×128 GEMM约 42访存密集优先优化数据搬运元素级 AddN 个元素0.25FLOP/byte极重度访存密集不可能用满 Tensor Core标准 AttentionN4096理论约 1365实际受中间结果拖累改造重点在融合FlashAttentionN4096逼近理论值融合后有效算术强度大幅回升看到 Element-wise 算子的算术强度只有 0.25你可能会问那这些算子怎么办答案很现实——它们天生就不适合单独吃 Tensor Core正确做法是跟前面的 GEMM 融合让整个 Kernel 的有效算术强度被拉高。这就是为什么在生产级推理引擎里你几乎看不到裸的 element-wise kernel。3. 真正让 Tensor Core 吃饱的四个动作3.1 数据排布把零散的字节变成整块搬运Tensor Core 一次矩阵乘操作要读取的数据量很大如果数据在显存里零零散散访存效率会非常难看。GPU 访存的最小高效单位是 32 字节一个 sector最理想的是线程以 128 字节为单位整块读取。具体到代码层面要做到两点。第一让连续线程访问连续地址这样硬件能把多次访存合并成一次大事务。第二尽量使用向量化加载指令比如float416 字节或half816 字节让每个线程一次搬更多数据。CUDA 里浮点运算量翻倍很容易但访存事务数量一旦翻倍性能就崩了。我曾经把一个 FP16 GEMM kernel 的数据布局从按行连续、但每个线程挑着读改成每个线程用half8连续读 8 个元素其他逻辑完全没动kernel 耗时直接降了 25%。这就是数据排布的力量。另外注意矩阵的主序row-major 还是 column-major和 Tensor Core 指令对 A/B 矩阵布局的要求如果不匹配通常不改指令而是先做一次 layout 转换比如用cublasLtMatmul选最优 layout成本远低于在 kernel 里别扭地逐元素换序。3.2 Tile 切分让数据在 Shared Memory 里反复使用算术强度低的最直接原因是数据只用了一次就被扔回显存。要提高算术强度核心手段就是 tiling——把大矩阵切成小块每个小块被加载到 shared memory 或寄存器后参与多次计算再丢弃。以 GEMM 为例如果每个线程算一个 8×8 的输出 tile对应 A 的 8 个元素和 B 的 8 个元素这个数据只用了 8 次。但如果把 tile 扩大到 64×64同样的数据能被复用 64 次。算术强度随 tile 维度近似线性上涨直到接近数据从 L2/显存只读一次的理论上限。tile 大小的取舍很有讲究。A100 上 GEMM 的常见配置是 128×128 的 block tile配合每个线程 8×8 甚至更大的 register tile。但 tile 不是越大越好shared memory 总量有限A100 单 SM 是 164KBblock tile 太大会降低同时驻留的 block 数量占用率掉下去反而拖慢整体吞吐。实际调优时我会先看 NCU 里的shared memory占用率和achieved occupancy两者取平衡点。3.3 Kernel 融合砍掉中间结果的运送费提升算术强度另一个立竿见影的动作是 kernel 融合。原理还是那句话减少访存。每融合一个中间步骤就少一次完整的写回显存 读回显存。FlashAttention 是最好的案例。标准 Attention 需要三个独立的 kernel 分别处理 QK^T、softmax、softmax(QK^T)V中间那个 N×N 矩阵要落地到显存两次。FlashAttention 的做法是把 N×N 分数矩阵按 block 切分每一块只存在 shared memory 里softmax 的归一化系数用 online 方式增量更新最后直接把加权结果累加到输出。这样 N×N 矩阵从头到尾不碰显存算术强度直接被拉回理论值附近。在工程上你不用每次手写这么复杂的融合。PyTorch 里可以用torch.compile自动做部分融合但更可靠的路径是先把最热的那几个算子挑出来看它们的中间张量形状如果能用torch._C._nn里的融合接口或手写一个 CUDA kernel 替代三次调用收益通常非常可观。我做过的实际优化案例里一个 LayerNorm Residual Bias 的三段融合端到端提升了 18%原因就是省掉了两次中间结果落地。3.4 精度策略FP16、BF16、TF32 怎么选Tensor Core 支持多种精度输入但很多人只知道开混合精度不知道精度的选择直接决定算术强度和数值行为。FP16 输入 FP32 累加是默认配置算术强度最高但对数值范围比较敏感容易出现上溢/下溢需要配合 loss scaling。BF16 有和 FP32 一样的指数位范围完全够用但尾数只有 7 位模型收敛速度可能受影响。TF32 则是把 FP32 输入截断成 19 位精度的特殊格式Tensor Core 能跑但算力降一半A100 上是 156 TFLOPS适合不想改模型数值逻辑、只想白嫖 Tensor Core 的场景。我自己的经验是能用 BF16 就用 BF16尤其在大模型场景输出敏感的小模型先用 FP16 试跑看 loss 曲线是否稳定只有 HPC 类负载、对数值零容忍的时候才考虑 TF32。另外累加器务必保持 FP32这是 Tensor Core 指令的硬性要求也是矩阵乘数值稳定的关键千万别把累加器也降成 FP16。4. 实测一个 GEMM Kernel 从 11 TFLOPS 到 271 TFLOPS 的全过程4.1 基线完全不用 Tensor Core 的朴素 Kernel为了说明问题我专门手写了三个版本的 4096×4096 FP16 GEMM kernel在同一张 A100 80GB 上实测。基线版本是最朴素的 CUDA 实现每个线程负责输出矩阵里的一个元素三重循环不切 tile也不用 shared memory。这个版本基本是给计算机体系结构教材当反面教材用的——它根本没有数据复用每个 A 的元素和 B 的元素都从显存重复读了很多次。实测性能11.2 TFLOPS。注意A100 的 FP16 Tensor Core 理论峰值是 312 TFLOPS这个 kernel 只跑出了 3.6%。而且它还不是 memory-bound 的死局——因为访存重复次数太多远超一次问题是既没用好带宽也没用好算力。这就是最差的情况两头都不沾。4.2 第一次改造Tile Shared Memory第二个版本加了 standard 的 shared memory tiling每个 block 负责 128×128 的输出块A 和 B 的对应 128×128 子块按 16×16 的小块逐步拷入 shared memory所有计算在 shared memory 里反复复用。这一步没有用到 Tensor Core 指令只是纯 CUDA Core 计算。性能提升非常明显直接到 41.7 TFLOPS。原因就是算术强度被数据复用拉上来了显存只做每块数据读一次的搬运。但我立刻意识到一个问题——它已经逼近 FP32 CUDA Core 的硬件上限19.5 TFLOPS了吗不对它用的是 FP16 指令在 CUDA Core 上跑。CUDA Core 跑 FP16 向量指令也有加速但和 Tensor Core 的 312 TFLOPS 比还是差着数量级。这一步的启示是tiling 解决的是数据供给问题但真正干重活的引擎还没启动。4.3 第二次改造切换到 Tensor Core 的 mma 指令第三个版本把内层计算换成 Tensor Core 指令。CUDA 里可以用两层 API上层是nvcuda::wmma定义成 16×16×16 的 fragment 操作缺点是有些灵活性损失底层是 PTX 的mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32直接控制每个线程持有的 4×8 矩阵 fragment。在这个版本里我按 m16n8k16 的形状组织线程每个 warp 负责一个大一点的输出 tile通过循环累积多个 mma 操作。寄存器怎么分配、fragment 按什么顺序加载这些细节都直接影响最终性能。第一次跑出来是 183 TFLOPS我对着 NCU 报告调了两轮 bank conflict 和 shared memory padding最后稳定在 271 TFLOPS。4.4 数据汇总与复盘三个版本的实测数据版本主要改动实测算力相对峰值利用率基线朴素三重循环11.2 TFLOPS3.6%V2Tile Shared Memory41.7 TFLOPS13.4%V3mma Tensor Core 指令271 TFLOPS86.9%参考cuBLAS官方库284 TFLOPS91.0%复盘下来我的核心体会是tiling 解决管饱的基础问题Tensor Core 指令才是吃得好的关键两者缺一不可。V2 已经证明数据复用能让带宽不再是瓶颈但算力引擎还在用小水管V3 换上了 Tensor Core 这个大水龙头配合前面打好的数据管线才真正把性能释放出来。我手写的版本离 cuBLAS 还差 4-5 个百分点差距主要在一些细粒度的调度技巧上——比如 double buffering、更深的循环展开、swizzled layout这些属于进阶打磨等后面专门写一篇展开。另外注意一个细节V3 的 271 TFLOPS 里如果我只改数据布局、不改计算指令也就是跳过 V3 直接往 V2 上叠优化是绝对到不了这个数字的。我见过很多人在 V2 阶段反复折腾访存模式收益越抠越少却从没想过换计算引擎。希望这个实测能提醒你优化要两层一起动数据供给和计算执行是同一个硬币的两面。5. 用 NCU 确认吃饱了没关键指标与实操命令5.1 哪些指标值得盯手写 kernel 最终性能如何不能靠感觉要看数据。Nsight ComputeNCU是 NVIDIA 官方最细粒度的性能分析工具我每次调 kernel 都离不开它。我这里说几个最关键的指标按重要性排序sm__pipe_tensor_op_hmma_inst_executed_pct_of_peak_sustained_activeTensor Core 流水线利用率接近 100% 说明 Tensor Core 没闲着sm__throughput.avg.pct_of_peak_sustained_elapsedSM 整体吞吐包含所有计算和访存单元gpu__compute_memory_throughput.avg.pct_of_peak_sustained_elapsed显存带宽利用率配合前者判断瓶颈类型smsp__average_warps_issue_stalled_long_scoreboard等 stall 原因说明 warp 在等什么长等待基本是显存延迟l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ldshared memory bank conflict 计数器这里有一个很实用的判断框架如果 SM throughput 已经很高、但 Tensor Core 利用率低说明计算引擎用错了或没喂到如果两个都低、但 memory throughput 高说明是访存瓶颈优先做数据复用如果三个都低说明并行度不够先查 occupancy。5.2 命令行实操流程我常用的三个命令组合# 快速看整体瓶颈类型 ncu --section SpeedOfLight --section ComputeWorkloadAnalysis ./my_kernel # 完整分析包含 stall 原因、张量核心利用率 ncu --set full --kernel-name regex:my_kernel_name -o profile_result ./my_kernel # 分析完成后生成可浏览的报告 ncu --import profile_result.ncu-rep --page details对于 PyTorch 用户第一步可以先用torch.profiler看算子级别的时间分布锁定热点算子再针对那个算子单独写成可执行程序上 NCU 做指令级分析。我通常不会直接对整个训练脚本跑 NCU那样输出会大得没法看而且噪声太多。5.3 一眼识别异常信号NCU 的 Roofline 图是最直观的判断工具。它会把你所有 kernel 画在一张图上横轴是算术强度纵轴是实测性能两条斜线分别是带宽上限和算力上限。如果某个 kernel 的点落在带宽斜率上说明它被带宽卡住了往右平移提高算术强度才是出路如果点贴在算力线上但离峰值还很远说明指令调度或占用率有问题。我记忆里最典型的一个异常信号是某个 kernel 的 memory throughput 只有 35%Tensor Core 利用率也不到 50%但 L2 命中率高达 90% 以上。这说明数据大量在 L2 里打转但真正流进寄存器喂给计算单元的效率很低。后来定位到是 shared memory 分配过多导致 occupancy 掉到 25%warp 数量太少没法用并行掩盖延迟。把 block tile 从 256×256 削到 128×128 之后occupancy 回到 50%性能反而提升了 60%。这就是用数据说话的价值——没有 NCU我大概率会在布局上白忙活。6. 伪优化避坑看起来专业实则无效的做法6.1 只切块不换引擎在 V2 里无限打转这是我自己和身边同事踩过最多的坑把 tiling 从 64×64 调到 128×128再从 128×128 调到 256×256调出几个百分点的收益就以为优化到底了。但如果你始终没把计算指令换成 Tensor Core 的 mma你就是在 CUDA Core 的小水龙头里打转。第 4 节的实测已经说明引擎切换带来的收益是数量级的比任何 tile 微调都重要。优化的正确顺序永远是先确认计算引擎正确再谈数据细节。6.2 盲目扩大 tile寄存器溢出警告曾有一次我把 register tile 从每个线程 8×8 扩大到 16×16想着数据复用率又能翻一翻。结果 kernel 不仅没变快反而慢了 30%。NCU 报告显示发生了本地内存溢出local memory spill大量寄存器被换到显存里每次计算都要经历一次写回读回。寄存器总量就那么多每线程超过 128 个寄存器之后必然开始溢出。记住一个原则tile 扩大带来的数据复用收益必须大于寄存器溢出带来的访存惩罚。建议实测时盯着local memory相关指标一旦出现溢出就立刻回退。6.3 忽略 Bank Conflictshared memory 也讲究并行Shared memory 的速度是有前提的——多个线程同时访问不同 bank 才能并行。如果两个线程访问同一个 bank 的不同地址比如访问一个 stride 为 32 的数组硬件会把访问串行化性能惩罚非常隐蔽。解决手法是对数组做 padding比如把 128 宽的共享数组声明成 129 个元素或者用 swizzle 重新映射地址。这个问题在 NCU 里一眼就能看出来l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld不为零。我调 V3 时就从 183 提到 220 TFLOPS 那一步纯靠消除 bank conflict。6.4 只看 Kernel 时间不看端到端延迟还有一个更宏观的伪优化单看某个 kernel 优化后快了多少却没注意整个 pipeline 别的地方成了新瓶颈。比如我用 FlashAttention 替代标准 Attention 后单个 attention kernel 快了 2.8 倍但整个 Transformer block 只快了 30%。原因是融合后 L2 里的数据局部性变了后面 MLP 的访存模式受影响。所以每次优化完务必要回到端到端基准测一轮。我自己现在养成的习惯是任何 kernel 级优化最终都要用同一个脚本跑一次完整 step time 做验收。6.5 架构不匹配Kernel 编译参数错误导致 Tensor Core 不可用最后提一个很多人遇到但搞不清楚的报错CUDA error: no kernel image is available for execution on the device。这个报错的本质是当前编译出的 kernel 二进制cubin/SASS不支持当前 GPU 的架构。如果你在编译时指定了-archcompute_70Volta而机器上插的是 A100Amperesm_80那 Tensor Core 相关的指令集根本没被编译进去运行时就只能报找不到可用的 kernel image。解决方法是让编译目标跟实际 GPU 匹配直接用-archnative或者在 PyTorch 场景下设TORCH_CUDA_ARCH_LIST8.0再重新安装/编译扩展。这个问题和本文主题的关系在于哪怕你算术强度算得再清楚、kernel 写得再合理架构不匹配导致 Tensor Core 代码根本没进到最终二进制里一切都白搭。最后分享一点个人体会。做性能优化做到后面你会发现绝大多数问题都可以归结到同一个核心矛盾——数据流动的速度和计算消耗的速度是否匹配。让 Tensor Core 吃饱本质上不是某一个奇技淫巧而是一整套思维方式的转变先算算术强度确定方向再选对计算引擎然后用数据排布、tiling、融合把供给密度拉上去最后用 NCU 的数据验证每一步。这套流程跑通之后你会对我的 kernel 到底跑得好不好这个问题的答案越来越有把握。如果你正准备优化一个深度模型建议先别急着改代码拿本篇文章第 2 节的方法把几个热点算子的算术强度估算一遍方向明确了再动手。这篇笔记就先到这里后面我打算接着写 double buffering 和 swizzle 这些进阶调度技巧把 tensor core 从 86% 推到 95% 以上的最后那段路走完。