3. 性能分析基础：性能指标、分析流程与性能基准

GPU 性能分析的第一步，不是急着打开 Nsight 或 rocProf，而是先建立一套“度量语言”。如果不知道“快”和“慢”如何量化，后续所有优化都将是盲目的。本章围绕三个核心问题展开：用什么指标衡量 GPU 是否吃饱？ 按什么流程定位瓶颈？ 拿什么标准判断优化是否有效？

3.1 三大核心性能指标：吞吐量、延迟与占用率

GPU 与 CPU 的设计哲学截然不同。CPU 追求低延迟，GPU 追求高吞吐。因此，分析 GPU 时，必须优先关注吞吐量（Throughput），其次才是延迟（Latency），而占用率（Occupancy）则是连接两者、反映硬件资源利用状态的关键桥梁。

• 吞吐量（Throughput）：单位时间内完成的工作量。在 GPU 语境下，常用 GFLOPS（每秒十亿次浮点运算）或 GB/s（显存带宽）表示。它反映的是“水管有多粗”，而非“水流多快”。
• 延迟（Latency）：从指令发出到完成的时间。对于单个线程而言，访问全局内存的延迟通常为 400~800 个时钟周期。GPU 通过大量线程并行来“隐藏”延迟，而不是像 CPU 那样依赖多级缓存和分支预测来“降低”延迟。
• 占用率（Occupancy）：每个 SM（流式多处理器）上活跃的 Warp 数量与最大可容纳 Warp 数量的比值。高占用率是隐藏延迟的前提，但并非越高越好——过高的占用率可能导致寄存器溢出或缓存命中率下降。

执行时间 ≈ max( 计算量 / 峰值计算吞吐 , 访存量 / 峰值显存带宽 )

3.2 分析流程：从“黑盒”到“白盒”的五步法

面对一个性能不佳的 Kernel，切忌直接猜测并修改代码。系统化的分析流程能大幅减少无效尝试。推荐采用以下五步迭代法：

1. 收集原始指标：使用性能分析工具（如 NVIDIA Nsight Compute、AMD rocProf）获取 Kernel 的执行时间、占用率、SM 活跃度、显存吞吐量等基础数据。
2. 识别瓶颈类型：将收集到的数据与硬件峰值对比。若 SM 计算单元利用率接近 100% 而显存带宽利用率低于 50%，则为计算瓶颈；反之则为访存瓶颈。更精确的判断需借助 Roofline 模型。
3. 定位问题代码：利用工具提供的 Source 级关联，将瓶颈指令映射到具体代码行。例如，Nsight Compute 的 “Source” 页签能直接显示是哪一行触发了高延迟的全局内存访问。
4. 实施针对性优化：根据瓶颈类型选择策略。访存瓶颈优先考虑合并访问（Coalesced Access）、使用共享内存或只读缓存；计算瓶颈则考虑使用更快的数学函数（如 __fdividef()）、减少分支发散。
5. 验证并回归：修改后必须重新测量，确认性能提升且结果正确。同时关注是否引入了新的瓶颈（例如，优化访存后计算可能成为新瓶颈）。

3.3 性能基准：Roofline 模型与实用对照表

Roofline 模型是连接算法复杂度与硬件能力的经典可视化工具。它用一条“屋顶曲线”描述硬件在不同算术强度（Arithmetic Intensity，单位 FLOP/Byte）下能达到的峰值性能。

算术强度 = 总计算量（FLOP） / 总访存量（Byte）。当算术强度低于“转折点”时，性能受限于带宽（斜线区域）；高于转折点时，受限于计算峰值（水平线区域）。

3.4 案例复盘：一个典型的访存瓶颈优化过程

假设有一个向量求和 Kernel：y[i] = a[i] + b[i]，数组长度为 1M 浮点数。初始实现如下：

__global__ void vectorAdd(const float *a, const float *b, float *y, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        y[i] = a[i] + b[i];  // 每个线程独立访问全局内存
    }
}

使用 Nsight Compute 分析后，发现显存带宽利用率仅为 35%，而 SM 计算利用率不足 5%。进一步查看 “Memory” 页签，发现全局加载效率（Global Load Efficiency）仅为 50%——原因在于线程访问地址不连续（若线程块大小与数组对齐不当）。

优化措施：

• 确保 blockDim.x 为 128 或 256，且 gridDim.x * blockDim.x 恰好覆盖 n，避免边界分支。
• 将标量访问改为向量化访问：使用 float4 一次加载 4 个浮点数，减少内存事务数量。

__global__ void vectorAddVec(const float4 *a, const float4 *b, float4 *y, int n4) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n4) {
        y[i] = make_float4(a[i].x + b[i].x, a[i].y + b[i].y,
                           a[i].z + b[i].z, a[i].w + b[i].w);
    }
}

优化后，显存带宽利用率提升至 92%，Kernel 执行时间缩短为原来的 1/3。此案例说明：先判断瓶颈类型，再选择优化手段，远比盲目调整线程块大小更有效。