4. 性能指标详解（上）：理论占用率与实际占用率、计算吞吐量与内存吞吐量、算术强度、L1/L2缓存命中率
各位同学，欢迎来到性能剖析的核心章节。
说实话，很多同学写CUDA代码，跑通了就以为万事大吉。但真正到了项目里，你会发现“能跑”和“跑得快”之间，隔着一道天堑。我见过太多人，代码逻辑完全正确，但性能就是上不去，卡在某个瓶颈上死活找不到原因。
为什么会这样？因为你没有用对指标去衡量你的核函数。
今天这一讲，我们就来聊聊那些最基础、也最重要的性能指标。我个人习惯，拿到一个核函数，第一件事不是看代码，而是先看这几个数字。它们就像医生的听诊器，能帮你快速定位病灶。
4.1 理论占用率 vs 实际占用率
先说说占用率。很多人一上来就问：“我的占用率怎么才50%？是不是太低了？”
嗯，这里要注意，占用率分两种：理论占用率和实际占用率。
理论占用率，是你在编译阶段就能算出来的。它取决于你的核函数用到了多少寄存器、多少共享内存，以及你的线程块大小。说白了，就是“GPU理论上最多能同时跑多少个warp”。
举个例子，我写过一段代码，每个线程用了64个寄存器。在Turing架构的GPU上，每个SM只有65536个寄存器。你算算，65536 ÷ 64 = 1024个线程。但一个SM最多能同时驻留2048个线程。所以理论占用率只有50%。
我曾经在项目中遇到过这种情况：一个看起来很简单的kernel，占用率死活上不去。后来一查，是编译器给每个线程分配了太多寄存器。解决办法？用 __launch_bounds__ 限制寄存器数量，或者手动调低 -maxrregcount。
关键点：理论占用率只是上限，不是实际值。
实际占用率，是运行时真正达到的占用率。它受限于指令发射、内存延迟、同步操作等动态因素。你想想看，即使理论占用率是100%，如果所有线程都在等内存，那实际占用率可能连10%都不到。
怎么查实际占用率？用 ncu（NVIDIA Nsight Compute）或者 nvprof。我个人习惯用 ncu --metrics sm__occupancy 直接看。
我的建议：不要盲目追求100%占用率。有时候为了降低寄存器压力，适当降低占用率反而能提升性能。我在一个矩阵乘法的优化中，把占用率从75%降到50%，性能反而提升了20%。因为每个线程有了更多寄存器，减少了全局内存访问。
4.2 计算吞吐量与内存吞吐量
这两个指标，是判断你的核函数是“计算密集型”还是“内存密集型”的关键。
计算吞吐量，就是GPU每秒能执行多少次浮点运算。单位通常是TFLOPS。你的核函数实际达到的计算吞吐量，除以GPU的理论峰值，就是计算利用率。
内存吞吐量，就是GPU每秒能读写多少字节。单位是GB/s。同样，实际值除以理论峰值，就是内存利用率。
我见过一个典型的案例：一个同学写了一个向量加法，跑在A100上。A100的理论计算峰值是19.5 TFLOPS（FP32），理论内存带宽是2 TB/s。结果他测出来，计算吞吐量只有0.5 TFLOPS，内存吞吐量却达到了1.8 TB/s。
这说明什么？说明这个kernel是内存瓶颈。你算得再快，数据送不过来也是白搭。
反过来，如果你写了一个复杂的数学计算，计算吞吐量接近峰值，但内存吞吐量很低，那就是计算瓶颈。
避坑指南：我曾经犯过一个错误，只盯着计算吞吐量看，觉得数值很高就满意了。结果发现内存吞吐量低得可怜，说明数据根本没怎么复用。后来加了共享内存做缓存，性能直接翻倍。
怎么判断你的kernel属于哪一类？看下一个指标——算术强度。
4.3 算术强度
算术强度，英文叫Arithmetic Intensity，简称AI。它的定义很简单：
算术强度 = 总浮点运算次数 / 总内存访问字节数单位是FLOP/Byte。这个值越大，说明你的计算越密集，越有可能是计算瓶颈。值越小，说明你的计算越稀疏，越有可能是内存瓶颈。
举个例子：
- 向量加法：每个元素做1次加法，读2个float（8字节），写1个float（4字节）。总内存访问12字节。算术强度 = 1 / 12 ≈ 0.08 FLOP/Byte。典型的内存密集型。
- 矩阵乘法（大矩阵）：每个输出元素需要做N次乘加，读2N个float。算术强度 ≈ N/3。当N很大时，算术强度很高，属于计算密集型。
我个人习惯，拿到一个kernel，先估算它的算术强度。然后对照GPU的“平衡点”——也就是理论计算峰值除以理论内存带宽。比如A100的平衡点大约是19.5 TFLOPS / 2 TB/s ≈ 9.75 FLOP/Byte。
如果你的算术强度远低于这个平衡点，那基本可以断定是内存瓶颈。如果远高于，那就是计算瓶颈。
记住：算术强度是指导你优化方向的罗盘。低算术强度，优先优化内存访问模式。高算术强度，优先优化计算指令和流水线。
4.4 L1/L2缓存命中率
最后聊聊缓存。GPU的缓存层次和CPU类似，但更强调吞吐量。
L1缓存：每个SM私有，延迟低，带宽高。在Turing/Ampere架构上，L1和共享内存共用一块物理空间（可配置）。
L2缓存：所有SM共享，延迟比L1高，但容量更大。
缓存命中率，就是你的数据请求在缓存中找到的比例。命中率越高，访问延迟越低，性能越好。
我记得有一次，一个同事的kernel性能很差，全局内存读取占了大部分时间。我用 ncu 一看，L1命中率只有20%，L2命中率也只有40%。这意味着大部分数据都要从显存读，延迟高达几百个周期。
怎么提高命中率？
- 空间局部性：让相邻线程访问相邻地址。比如用
float4一次读16字节，而不是一个float一个float地读。 - 时间局部性：重复使用的数据，尽量放在共享内存或寄存器里。别每次都去全局内存读。
- 对齐访问：确保起始地址是128字节对齐的，这样能充分利用缓存行。
一个小技巧：如果你发现L1命中率低，可以试试把L1缓存调大。在NVIDIA GPU上，可以通过 cudaFuncSetAttribute 设置 cudaFuncAttributeMaxDynamicSharedMemorySize 来调整L1/共享内存的分配比例。但要注意，共享内存和L1是此消彼长的关系，别顾此失彼。
好了，这一讲的内容就到这里。这些指标是性能分析的基石。下一讲，我们会继续深入，聊聊如何用这些指标去定位具体的瓶颈。