序章：一切从得到一份报告开始说起 ⭐教你读懂 Nsight Compute 报告 系列合集⭐

序章

欢迎！这一博客合集的主旨，是帮助大家读懂 Nsight Compute（也就是 ncu）的报告，并掌握最基础的分析方法。作者默认大家已经具备 CUDA 的一些基础理解，所以像 warp、block、共享内存这些就不再复述了。作者水平有限，文中若有错漏，还请大家包涵并指出来，一起交流，一起进步。

软硬件开发环境

作者是用自己的笔记本电脑，通过 Qoder IDE 远程 SSH 到一台 Ubuntu 服务器上做开发的。服务器上装着一张 RTX 3090，代码存放、编译、运行、采集运行数据全部都在服务器上完成，采出来的 .ncu-rep 文件再传回笔记本，用笔记本本地的 ncu-ui 打开来看。作者相信这一套工作形式是很有普遍性的。为了严谨，下面把相关参数一项一项列出来，它们也是本合集全程固定不变的基准。

服务器端 RTX 3090 硬件参数如下：

项目	值
GPU	NVIDIA GeForce RTX 3090
架构 / 计算能力	Ampere，sm_86（compute capability 8.6）
显存	24 GB GDDR6X
SM 数量	82
每 SM 最大常驻线程	1536（= 48 warp）
每 block 最大线程	1024
每 SM 寄存器	65536
每 SM 共享内存	100 KB（可与 L1 按比例分配）
每 block 共享内存上限	48 KB（opt-in 可到 99 KB）
L2 缓存	6 MB
显存位宽 / 频率	384 bit / 9751 MHz
显存峰值带宽	936 GB/s
驱动	570.144

服务器端软件参数如下：

项目	版本
操作系统	Ubuntu 22.04.5 LTS
CUDA Toolkit	12.8
编译器	nvcc 12.8（V12.8.93）+ gcc 11.4.0，语言标准 C++14
Nsight Compute	2025.1.1

至于开发的前端，其实是很自由的。毕竟主要工作都在服务器上完成，无论用什么电脑，只要能 SSH 到服务器就能开发。所以前端只需要装一个能改代码的编辑器，再装一个和服务端同版本的 Nsight Compute 用来打开采集文件，就足够了，CUDA Toolkit 之类的都不必安装。

目标 kernel

本合集要读的报告，就是从下面这个 kernel 上采出来的。它是 src/reduce.cu 里的 reduce_ill，一段故意埋了很多坑的加法 reduce：

/*
main 函数创建一个填充随机数的长度为 2^24 的一维数组，存入 GPU 全局内存，交由 kernel 求总和：
  reduce_ill<<<gridSize, blockSize, shmemSize>>>(input, output, n);
其中
  - blockSize 取 256，作者习惯用这个值，设计每个 thread 负责 8 个数；
  - gridSize 取 8192。已设计每个 block 要负责 2048 个数
    （256 线程 × 每线程 8 个），故取 2^24 / 2048 = 8192，正好把整个
    数组不重不漏地分完；
  - shmemSize 取 8192。设计 kernel 里的每一步都建立在
    “先把 2048 个数整个搬进 smem 再慢慢折腾”的前提上，所以 smem
    需要放得下 2048 个 float，故取 2048 × sizeof(float) = 8192，即 8 KB。
  - n 取 1 << 24 即 2^24。待求和数组总长度。
  - input 取 n × sizeof(float)。
  - output 取 gridSize × sizeof(float)。每个 block 把自己负责的 2048
    个数求成一个部分和，main 函数最后再把 8192 个部分和加起来，得到最终结果。
*/
__global__ void reduce_ill(float* input, float* output, int n) {

    extern __shared__ float smem[];

    // 【病态 1】非合并的全局读
    // 循环 8 轮，每轮 block 里的 256 个 thread 各搬 1 个数到共享内存，
    // 一轮 256 个，八轮搬够 2048 个。
    for (int i = 0; i < 8; ++i) {
        // 搬到哪：故意把 smem 看成一个 8 行 256 列的矩阵，各 thread 每轮
        // 合作填满矩阵的一整行，八轮即可填满整个矩阵
        const int smem_idx = blockDim.x * i + threadIdx.x;
        // 从哪搬：又故意把 input 看成一个 2048 行 8192 列的矩阵，第 blockIdx.x 列
        // 整列归本 block 管，每轮各 thread 从这一列取 256 个数，
        // 八轮取完矩阵整个第 blockIdx.x 列
        const int globalIdx = gridDim.x * smem_idx + blockIdx.x;
        // 开始搬，对于加法 reduce，越界的元素填 0
        smem[smem_idx] = (globalIdx < n) ? input[globalIdx] : 0.0f;
    }
    __syncthreads();

    // 【病态 2】共享内存读冲突
    // 又故意把 smem 看成一个 256 行 8 列的矩阵，各 thread 都在循环中负责把
    // 第 threadIdx.x 行的 8 个数加起来
    float threadSum = 0.0f;
    for (int i = 0; i < 8; ++i) {
        threadSum += smem[8 * threadIdx.x + i];
    }
    __syncthreads();
    // 各 thread 求和完成后，再把结果存为长为 256 的一维数组排放在共享内存的低位
    smem[threadIdx.x] = threadSum;
    __syncthreads();

    // 【病态 3】朴素树形归并
    // 这里是很常见的朴素树形归并，随着步数的翻倍，各 thread 不停检查现在需不需要自己干活，
    // 有活就干没活闲着，直到 256 个数完成求和
    for (int step = 1; step < blockDim.x; step *= 2) {
        if (threadIdx.x % (2 * step) == 0) {
            smem[threadIdx.x] += smem[threadIdx.x + step];
        }
        __syncthreads();
    }

    // 【病态 4】稀疏写回
    // 终于，本 block 负责的 2048 个数的和，存入结果数组，待 main 读回CPU侧后求最终和
    if (threadIdx.x == 0) {
        output[blockIdx.x] = smem[0];
    }
}

这个 kernel 确实是故意设计的，问题重重，病得不轻，但它病得很在点上，常见的几种疑难杂症一个不落全犯了，拿它来学习 ncu 报告再好不过。

病态一：非合并的全局读

GPU 对全局内存的读写是整条存储链路里最慢的一级，要从全局内存里读数据，硬件性质定死了以 32 字节、也就是一个 sector 为最小单位来读。如果一个 warp 里 32 个 thread 各要读 1 个字节，而这 32 个字节在全局内存里恰好是连续码放在一起存储的，硬件会自动只进行一次读取操作，把 1 个 sector 读回来，各 thread 一下子就各拿到自己要的数了，这就叫内存合并访问。如果一个 warp 里 32 个 thread 要读的是同一个字节，那无论这个字节在全局内存里怎么存放，硬件也只会进行一次读取操作，把包含待读字节的那 1 个 sector 读回来，再以广播的形式发给各 thread，各 thread 同样一下子就各拿到自己要的数，多读回来的 31 个字节丢掉就是了。

可在当前 kernel 里，偏偏把 input 看成一个 2048 行 8192 列的矩阵、按列来读。C 语言的数组默认是行主序存放的，这意味着同一个 warp 里各 thread 要读的数，在全局内存里的存放位置彼此隔了整整一行的距离，即 8192 × sizeof(float) = 32 KB。于是硬件为了满足第一个 thread，得先读 1 个 sector，把里面的 4 个字节交给它，剩下的丢掉；再跳过 32 KB 去读下一个 sector，把读回来的 32 个字节里的 4 个交给第二个 thread，剩下的又丢掉……如此重复 32 次。这既不是内存合并访问，也不是广播。当前场景中，每个 warp 实际要读的数据一共 32 × 4 B = 128 B，如果按合并访问来算，128 B 恰好铺满 4 个 sector，硬件读 4 次就收工；现在却是每个数各占 1 个 sector，硬件得读 32 次，搬回 32 × 32 B = 1024 B。1024 / 128 = 8，读流量就这么放大了 8 倍。

工程中，这就是很常见的转置式访问，虽然同一个矩阵，但是按行读和按列读，天壤之别。

病态二：共享内存读冲突

共享内存就在 SM 片上、速度很快，它没有全局内存那套 sector 的概念，不按 32 字节打包搬运，硬件使用 bank 这一概念来组织：以 4 字节为一个字，第 1 个字放进 bank 0，第 2 个字放进 bank 1……第 32 个字放进 bank 31，第 33 个字转回来又放进 bank 0，就这样 4 字节、4 字节地循环着往后存，所以 bank 是一个循环编号的结构，每个 bank 都是一条能独立服务 thread 的通路。一个 warp 里 32 个 thread 可能同一时刻来读共享内存，如果每个 thread 各要读 1 个字，且待读的 32 个字连续存放在共享内存中，则恰好一个 thread 对上一个 bank，硬件一拍，所有 thread 就能同时拿到需要的字，这是共享内存吞吐的理想状态。反之，如果同一时刻 warp 里有多个 thread 要读或写的地址落进了同一个 bank，那这个 bank 就只能一个一个依次服务，硬件把这些访问拆成好几拍串行执行，出现排队等待，这就是 bank 冲突，报告里叫 bank conflict。同一时刻落在同一个 bank 上的地址有几个，就得拆成几拍。

当前 kernel 里，数搬进共享内存后，偏偏把 smem 看成一个 256 行 8 列的矩阵、按行来读，各 thread 读自己那一行的 8 个数，于是 warp 内相邻 thread 要读的地址相隔了 8 个字。以第一轮循环、也就是各 thread 都处理第一列时为例，同一时刻，第 1 个 thread 要读第 0 个字，0 % 32 = 0，落在 bank 0；第 2 个 thread 要读第 8 个字，8 % 32 = 8，落在 bank 8；第 3 个要读第 16 个字，落在 bank 16；第 4 个要读第 24 个字，落在 bank 24；到第 5 个 thread 要读第 32 个字，32 % 32 = 0，又撞回了 bank 0。往后都一样，thread 每 4 个转一圈，总是循环落在 bank 0、8、16、24 这 4 个 bank 上，32 个 thread 只动用了 4 个 bank。站在 bank 0 的视角看：同一时刻，第 1、5、9、…、29 这 8 个 thread 全都找它读数，可它一次只能服务一个，其余 7 个只能排队。于是本该一拍完成的读，被拆成了 8 拍串行重放，这就是 8 路共享内存读冲突。

病态三：朴素树形归并

warp 是 GPU 调度和执行的基本单位，里面 32 个 thread 同吃一条指令，永远步调一致，遇到 if 这类分支时，硬件并不会给每个 thread 发一条独立的判断：如果整个 warp 的 thread 都走同一边，一切照旧；可一旦有的 thread 走这边、有的走那边，硬件就只能把两条路径各执行一遍，每遍只让对应的 thread 生效，其余的干看着，这就是分支分歧，warp 的执行时间会按分支的路数成倍拉长。

当前 kernel 里的朴素树形归并循环正好是这个局面。每一轮开始前，各 thread 都要检查一遍现在需不需要自己干活，随着步长翻倍，能通过检查的 thread 一轮比一轮少：第一轮 warp 里还有一半 thread 有活干，另一半陪跑；越往后越冷清，到后几轮干脆整排整排的 warp 都在陪跑，只剩 warp 0 里的一两个 thread 在动。把 8 轮加起来，256 个 thread 真正做的加法只有 128 + 64 + … + 1 = 255 次，可它们陪站了整整 8 轮、同步了 8 次，2048 个 thread 轮次里只有 255 个在干活，利用率还不到 13%。

病态四：稀疏写回

和病态一同源，全局内存读的最小单位是 32 字节的 sector，对写同样成立：哪怕只想写 4 个字节，硬件也得按 1 个 sector 的粒度去占用。

当前 kernel 里，每个 block 只有 thread 0 独自往全局内存写回 1 个 float。放到整个 warp 里看，就是 32 个 thread 只有 1 个在写，硬件仍然要为这 4 个有效字节占掉一整个 32 字节的 sector，其余 28 个字节空着，写流量就这么放大了 8 倍。

四个病态互不干扰

这里大概会有一个疑问：四个病态挤在一个 kernel 里，报告上不会互相干扰吗？

作者特意让它们各自落在不同的计数器上：读放大体现在 L1/L2 的 sector 请求数和 DRAM 吞吐上，bank conflict 体现在 L1TEX 共享内存的 Wavefronts 上，树形归并的低效体现在活跃线程和指令统计上，稀疏写体现在 store 相关的那条警告上。四个信号各走各的路，不会互相掩盖，反而能在同一份报告里被分别定位出来，这正是作者想要的效果。

从编译到打开报告

开发环境和代码都到位了，接下来就是怎么把它跑起来并做分析。下面的所有操作都默认在远程服务器上进行。

编译

本文没有给出完整的 reduce.cu，因为 main 函数里有很多与主旨无关的如初始化、内存分配、数据拷贝、正确性检验、结果打印等逻辑，就不占篇幅了。reduce.cu 写好后，编译：

nvcc -lineinfo -std=c++14 -arch=sm_86 src/reduce.cu -o build/rtx3090/reduce -lcublas

这里面几个参数不是随意写的。-lineinfo 一定不能省，少了它 ncu 就没法把硬件指标对应回源码行，报告里的 Source 页会是一片空白，后面想定位到具体哪一行出的问题就无从查起了。-arch 要跟采集机的架构对上，3090 是 sm_86，换别的卡得改成对应的架构码。-std=c++14 是作者这个项目里统一的语言标准。-lcublas 是链接 cuBLAS 库用的，属于这个工程本身的依赖，跟 kernel 没有关系，照写就行。

运行

编译完先别急着采集，先跑一遍确认算对了：

./build/rtx3090/reduce ill

作者的程序接受一个参数指定跑哪个 kernel，传 ill 就是跑 reduce_ill。跑完会打印期望和、实际和以及 PASS/FAIL：

version=ill n=16777216 期望和=830584179.000000 实际和=830584179.000000
结果: PASS

要是打印出来是 FAIL，那就得先回头调试 kernel，把算法改对。ncu 的所有分析都建立在“结果是对的”这个前提上，一个算错的 kernel，采出来的报告再漂亮也没有意义。

采集

确认 PASS 之后就可以采报告了。作者的命令是：

ncu -f --set full --launch-count 1 --call-stack --target-processes all --import-source yes \
    -o benchmarks/rtx3090/reduce-ill \
    ./build/rtx3090/reduce ill

跑完会在 benchmarks/rtx3090/ 下得到 reduce-ill.ncu-rep。这一串参数各有各的用意：

参数	作用
-f	覆盖同名旧报告
--set full	采集完整指标集，各细节页都有数据；代价是慢，这份报告采一次要重放 40 遍
--launch-count 1	只采第一次启动。同一 kernel 每次启动的指标基本一致，采一次足够，也避免重复重放引入偏差
--call-stack	记录调用栈，便于在 UI 里定位 kernel 来源
--target-processes all	连同子进程一起采集，避免漏采
--import-source yes	把源码内嵌进报告，阅读端不必再带原始.cu

要特别注意，采集的时候务必确保显卡是完全空闲的。作者用的 RTX 3090 是服务器里的共享卡，而 ncu 采集的硬件计数器反映的是采集窗口内整张卡的状态，如果卡上面还同时跑着别人的任务，读到的占用率、带宽、缓存命中率就都掺了别人的流量。更麻烦的是，这种状况不会有任何单独提示，所以采集前务必做检查。作者的做法是采集前先记录一次 GPU 占用，GPU-Util 是 0%、显存接近空闲才继续。

安装与打开报告

报告采出来之后，就可以把它拷到自己的电脑上打开了，其实随便哪台电脑都行，只要装了版本正确的 ncu。.ncu-rep 是个二进制文件，得用 Nsight Compute 的图形界面 ncu-ui 来看。Nsight Compute 可以从 NVIDIA 官网单独下载安装，不一定要装整个 CUDA Toolkit；如果机器上已经装过 CUDA Toolkit，里面也自带一份。

这里唯一需要注意的是版本匹配：报告可以用同版本或者更高版本的 ncu-ui 打开，反过来拿更旧的版本去打开新报告就会失败。本合集采集端是 2025.1.1，作者笔记本上的阅读端也是 2025.1.1，两边一致。如果大家本地版本更低，升级到不低于采集端的版本就行。

安装好ncu后，直接在命令行敲 ncu-ui reduce-ill.ncu-rep即可打开报告，默认停在 Summary 页：

序章到这里就结束了。下一篇开始，作者就带着大家正式进入这份报告的阅读。