第2板块·第7节：性能分析工具入门（Nsight Compute）

学习目标

学完本节你将能够：

• 安装并使用 Nsight Compute 分析 CUDA Kernel 性能
• 解读关键内存性能指标，并对应到前几节的知识点
• 使用 ncu 命令行快速采集性能数据
• 理解 roofline 模型，判断 Kernel 是计算瓶颈还是内存瓶颈
• 独立分析一个存在内存问题的 Kernel 并定位瓶颈

1. Nsight Compute 简介与安装

1.1 什么是 Nsight Compute

Nsight Compute（简称 ncu）是 NVIDIA 官方推出的 GPU Kernel 级性能分析工具，可以采集单个 Kernel 执行过程中的详细硬件指标，包括：

• 执行时间、占用率
• 全局内存访问效率
• 共享内存 Bank Conflict
• 寄存器使用与溢出
• 计算吞吐量与内存吞吐量
• Warp 调度与停滞原因

与 Nsight Systems（系统级时间线分析）不同，Nsight Compute 专注于单个 Kernel 内部的微观性能分析。

1.2 安装方式

方式1：通过 NVIDIA 官网下载
访问 NVIDIA Nsight Compute 下载页面，选择对应操作系统和 CUDA 版本，下载安装包安装。

方式2：通过 conda 安装（推荐）

conda install -c nvidia nsight-compute

方式3：CUDA Toolkit 自带
较新版本的 CUDA Toolkit 中已包含 ncu 命令行工具，可直接使用：

which ncu

1.3 基本使用方式

2. 关键性能指标与知识点对应

Nsight Compute 采集的指标很多，入门阶段最需要关注以下内存相关指标，它们分别对应前几节的核心知识点：

核心对应关系：

• Global Load/Store Efficiency < 70% → 存在严重的非合并访问（第2节）
• Shared Memory Bank Conflict > 0 → 存在 Bank Conflict（第3节）
• Local Memory Traffic > 0 → 存在寄存器溢出（第5节）

3. 命令行工具 ncu 快速上手

3.1 基础用法

# 最简单的分析
ncu ./your_program

# 只分析指定 Kernel
ncu –kernel-name vectorAdd ./your_program

# 指定要采集的指标
ncu –metrics gld_efficiency,gst_efficiency ./your_program

3.2 常用指标采集命令

# 内存效率指标
ncu –metrics gld_efficiency,gst_efficiency ./your_program

# Bank Conflict 指标
ncu –metrics shared_load_bank_conflict,shared_store_bank_conflict ./your_program

# 局部内存指标
ncu –metrics local_load_transactions,local_store_transactions ./your_program

# 综合完整指标集
ncu –set full ./your_program

3.3 输出解读

示例输出：

==PROF== Profiling "vectorAdd":
==PROF== Profiling "vectorAdd" – kernel: vectorAdd, grid: 4096, block: 256
==PROF== Section: Memory Workload Analysis
gld_efficiency 100.00%
gst_efficiency 100.00%
shared_load_bank_conflict 0
shared_store_bank_conflict 0
local_load_transactions 0
local_store_transactions 0

• gld_efficiency = 100%：完全合并访问，无需优化
• shared_*_bank_conflict = 0：无 Bank Conflict
• local_*_transactions = 0：无寄存器溢出

4. Roofline 模型：判断瓶颈类型

4.1 什么是 Roofline 模型

Roofline 模型是一个可视化性能分析框架，用于判断 Kernel 是计算瓶颈还是内存瓶颈。

• 横轴：算术强度（Arithmetic Intensity）= FLOPs / Bytes Moved
• 纵轴：性能（FLOPs/s）

┌─────────────────────────────
│ 峰值计算性能（天花板）
╱
╱
╱ Ridge Point（拐点）
╱
╱
╱ 峰值内存带宽限制的斜率（天花板）
╱
└─────────────────────────────
低算术强度 高算术强度

4.2 判断规则

• 算术强度低（左侧区域）：Kernel 受内存带宽限制，优化重点应放在合并访问、减少内存事务
• 算术强度高（右侧区域）：Kernel 受计算性能限制，优化重点应放在指令优化、减少计算量

4.3 在 Nsight Compute 中查看

ncu –metrics sm__throughput.avg.pct_of_peak_sustained_elapsed,\\
gpu__compute_memory_throughput.avg.pct_of_peak_sustained_elapsed \\
./your_program

• sm__throughput：SM 计算吞吐量占比
• gpu__compute_memory_throughput：内存吞吐量占比

判断：

5. 代码演示：分析一个存在内存问题的 Kernel

以下程序包含三个 Kernel，分别模拟不同问题：

#include <cstdio>
#include <cuda_runtime.h>

// 1. 非合并访问
__global__ void badCoalescing(float *in, float *out, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx * 8 < N) {
out[idx] = in[idx * 8]; // 跨步 8，非合并
}
}

// 2. Bank Conflict
__global__ void badBankConflict(float *out) {
__shared__ float sdata[1024];
int tid = threadIdx.x;
sdata[tid] = tid * 1.0f;
__syncthreads();
float val = sdata[(tid * 32) % 1024]; // 32‑way Bank Conflict
out[tid] = val;
}

// 3. 寄存器溢出
__global__ void badSpill(float *out, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float a0=0,a1=1,a2=2,a3=3,a4=4,a5=5,a6=6,a7=7;
float a8=8,a9=9,a10=10,a11=11,a12=12,a13=13,a14=14,a15=15;
float b0=0,b1=1,b2=2,b3=3,b4=4,b5=5,b6=6,b7=7;
float b8=8,b9=9,b10=10,b11=11,b12=12,b13=13,b14=14,b15=15;
float c0=0,c1=1,c2=2,c3=3,c4=4,c5=5,c6=6,c7=7;
float c8=8,c9=9,c10=10,c11=11,c12=12,c13=13,c14=14,c15=15;
float d0=0,d1=1,d2=2,d3=3,d4=4,d5=5,d6=6,d7=7;
float d8=8,d9=9,d10=10,d11=11,d12=12,d13=13,d14=14,d15=15;
// … 更多变量
float sum = a0+a1+a2+a3+a4+a5+a6+a7+a8+a9+a10+a11+a12+a13+a14+a15
+ b0+b1+b2+b3+b4+b5+b6+b7+b8+b9+b10+b11+b12+b13+b14+b15
+ c0+c1+c2+c3+c4+c5+c6+c7+c8+c9+c10+c11+c12+c13+c14+c15
+ d0+d1+d2+d3+d4+d5+d6+d7+d8+d9+d10+d11+d12+d13+d14+d15;
if (idx < N) out[idx] = sum + idx;
}

int main() {
const int N = 1 << 20;
float *d_in, *d_out;
cudaMalloc(&d_in, N * sizeof(float));
cudaMalloc(&d_out, N * sizeof(float));

int blockSize = 256;
int gridSize = (N + blockSize – 1) / blockSize;

badCoalescing<<<gridSize, blockSize>>>(d_in, d_out, N);
badBankConflict<<<1, 256>>>(d_out);
badSpill<<<gridSize, blockSize>>>(d_out, N);

cudaDeviceSynchronize();
cudaFree(d_in);
cudaFree(d_out);
return 0;
}

编译与分析：

# 编译
nvcc –ptxas-options=-v memory_problems.cu -o memory_problems

# 分析合并访问效率
ncu –metrics gld_efficiency ./memory_problems

# 分析 Bank Conflict
ncu –metrics shared_load_bank_conflict ./memory_problems

# 分析局部内存（寄存器溢出）
ncu –metrics local_load_transactions,local_store_transactions ./memory_problems

预期分析结果：

6. GUI 使用入门

6.1 启动

ncu-ui

6.2 基本工作流

6.3 快速定位问题

在 Details 页面中：

• 查看 Memory Workload Analysis 分组，检查 Global Load/Store Efficiency
• 查看 Shared Memory 分组，检查 Bank Conflict
• 查看 Registers 分组，检查寄存器使用和溢出

7. 实操案例：从定位到修复

7.1 案例描述

一个矩阵乘法 Kernel 性能远低于预期，使用 Nsight Compute 定位并修复。

7.2 分析步骤

Step 1：运行 ncu 采集基础指标

ncu –set basic ./matmul

查看 Summary 中的关键指标：

• Achieved Occupancy：远低于理论值 → 可能寄存器或共享内存占用过多
• gld_efficiency：远低于 100% → 非合并访问
• shared_load_bank_conflict：非 0 → Bank Conflict

Step 2：定位具体问题
假设分析结果显示：

gld_efficiency: 32%
shared_load_bank_conflict: 16
achieved_occupancy: 45%

对应问题：

Step 3：依次修复

Step 4：验证

ncu –metrics gld_efficiency,shared_load_bank_conflict,achieved_occupancy ./matmul_fixed

确认指标改善后，用 cudaEvent 计时验证整体性能提升。

8. 课后练习

练习1：基础指标采集
使用 ncu 分析第2节中的四个 Kernel（合并、错位、跨步、分散），记录各 Kernel 的 gld_efficiency，与理论值对比。

练习2：Bank Conflict 检测
使用 ncu 分析第3节中转置的三个版本（无共享内存、共享内存无 padding、共享内存有 padding），记录 shared_load_bank_conflict 的变化。

练习3：寄存器溢出检测
使用 ncu 分析第5节中存在溢出的 Kernel，查看 local_load_transactions 和 local_store_transactions，计算溢出导致的局部内存访问量。

练习4：Roofline 判断
选择一个你写过的 Kernel，使用 ncu 采集 sm__throughput 和 memory_throughput，判断它是计算瓶颈还是内存瓶颈，并据此决定优化方向。

练习5：完整分析流程
分析第6节中的归约 Kernel（共享内存版本），使用 ncu 采集完整指标，找出所有内存相关问题，尝试修复并验证。

9. 第二板块总结

恭喜你完成第二板块「GPU 存储层级与内存访问规则」的全部 7 节内容！

回顾本板块核心知识：

下一步

下一板块将进入 线程同步、通信、原子操作，你将学习：

• 原子操作的硬件机制与性能特性
• 多 Block 间通信与全局同步
• 锁、信号量等高级同步原语
• 归约、扫描等算子中的通信模式
• 原子操作的性能优化策略