本文翻译自：https://developer.nvidia.com/blog/using-nsight-compute-to-inspect-your-kernels/
使用 Nsight Compute 检查您的内核
Nvidia 为开发人员提供了新的 CUDA 工具库：Nsight Compute 和 Nsight Systems 。在更新的 GPU 架构上，这些工具将会变得更重要。在本博客的示例中，我们将会使用这些新工具来获取图灵及更高架构的计算结果。建议先阅读本系列的前两篇博客，以获取相关信息：
如前几篇博客所述，Nsight Compute 和Nsight Systems 的目标和功能是不同的，所以调优行为会由一个或者几个这种新工具组合使用。Nsight Compute 的主要用途之一是提供对 Kernel 的 GPU 性能分析指标。如果您使用过 NVIDIA Visual Profiler 或 nvprof（命令行分析器），您可能已经检查了 CUDA 内核的特定指标。本博客重点介绍如何使用 Nsight Compute 做到这一点。许多其他的分析功能（例如检查时间线、测量活动持续时间等）可以使用 Nsight Systems 执行。
开始
我们将分析一段 CUDA 代码，它是上一篇博客中向量相加的代码的变体。这段代码使用了 2D CUDA grid 配置 ，以及二维数组(即双下标)访问。
#include
const size_t size_w = 1024;
const size_t size_h = 1024;
typedef unsigned mytype;
typedef mytype arr_t[size_w];
const mytype A_val = 1;
const mytype B_val = 2;
__global__ void matrix_add_2D(const arr_t * __restrict__ A, const arr_t * __restrict__ B, arr_t * __restrict__ C, const size_t sw, const size_t sh){
size_t idx = threadIdx.x+blockDim.x*(size_t)blockIdx.x;
size_t idy = threadIdx.y+blockDim.y*(size_t)blockIdx.y;
if ((idx < sh) && (idy < sw)) C[idx][idy] = A[idx][idy] + B[idx][idy];
}
int main(){
arr_t *A,*B,*C;
size_t ds = size_w*size_h*sizeof(mytype);
cudaError_t err = cudaMallocManaged(&A, ds);
if (err != cudaSuccess) {std::cout << "CUDA error: " << cudaGetErrorString(err) << std::endl; return 0;}
cudaMallocManaged(&B, ds);
cudaMallocManaged(&C, ds);
for (int x = 0; x < size_h; x++)
for (int y = 0; y < size_w; y++){
A[x][y] = A_val;
B[x][y] = B_val;
C[x][y] = 0;}
int attr = 0;
cudaDeviceGetAttribute(&attr, cudaDevAttrConcurrentManagedAccess,0);
if (attr){
cudaMemPrefetchAsync(A, ds, 0);
cudaMemPrefetchAsync(B, ds, 0);
cudaMemPrefetchAsync(C, ds, 0);}
dim3 threads(32,32);
dim3 blocks((size_w+threads.x-1)/threads.x, (size_h+threads.y-1)/threads.y);
matrix_add_2D<<<blocks,threads>>>(A,B,C, size_w, size_h);
cudaDeviceSynchronize();
err = cudaGetLastError();
if (err != cudaSuccess) {std::cout << "CUDA error: " << cudaGetErrorString(err) << std::endl; return 0;}
for (int x = 0; x < size_h; x++)
for (int y = 0; y < size_w; y++)
if (C[x][y] != A_val+B_val) {std::cout << "mismatch at: " << x << "," << y << " was: " << C[x][y] << " should be: " << A_val+B_val << std::endl; return 0;} ;
std::cout << "Success!" << std::endl;
return 0;
}
重点：
-
Managed Memory：我们使用Managed Memory 进行数据分配。对于支持页面级别的缺页异常的 GPU （即Unified Memory 2.0）我们预取数据以避免对内核造成性能影响。
-
2D：我们使用了 2D 的 grid 和 block 形状，并且使用 typedef 来简化 2D 数据的定义，其中数据宽度在编译时是已知的（在此示例中）。这让我们可以使用双下标进行访问，而不需要指针操作。
-
内核设计：内核非常简单。每个线程使用CUDA内置变量计算一组 x、y 下标，如果计算的下标是有效的(在有效的数据区域内)，就将所选元素相加。
希望上面的代码看起来很简单。作为一名 CUDA 开发人员，你应该知道两个最重要的优化事项：给 GPU 足够多的并行任务、有效使用显存子系统。我们将着重于第二个目标。因为我们仅使用了全局存储器，因此我们的重点是如何有效的使用全局存储器，为此我们需要尽量对全局存储器进行读/写联合操作。
在 Visual Profiler (nvvp) 或 nvprof 中，可以使用 gld_efficiency （全局读取效率）和 gst_efficiency （全局写入效率）等指标，可以轻松的验证全局存储器的访问是否已经合并。
有哪些指标？
一般来说，Nsight Compute 所使用的指标与以往的工具不同。例如，目前 Nsight Compute 还没有提供与以前 gld_efficiency 和 gst_efficiency 相对应的指标。
首先，有哪些新指标？有两种方式可以查看：
- 使用 Nsight Compute：与使用 nvprof 一样，您可以查询可用的指标。新工具为开发者提供了更多的指标。这些指标将会针对您正在使用的 GPU 设备显示。在 Nsight Compute 中，同系列的设备所可用的指标都应该相同。如果您有多个不同的 GPU ，您可以选择您希望显示的设备。您也可以将输出结果保存为文件。
nv-nsight-cu-cli --devices 0 --query-metrics >my_metrics.txt
(如果您需要指定完整路径，见下文)。您也可以在命令中查询任何架构的指标，而不受限于你使用了什么设备。
- 查看文档：Nsight Compute 文档在这里。Nsight Compute 文档的另一个入口在 CUDA 文档的 Tool 部分，你可以在侧边栏中找到。Nsight Compute CLI 文档中常用的一个部分是 Nvprof Transition Guide（从 CUDA 10.1 Update 2 和 Nsight Compute 2019.4 开始，现在还提供 Visual Profiler 转换指南）。该指南中有一个指标的对照表，您可以快速找到 nvvp 或 nvprof 所对应的新指标。但其中显然没有
gld_efficiency和gst_efficiency，所以我们需要其他方法。
当为了满足该代码的访问请求，所发生的显存(或缓存)传输数量最少时，则可以使可以使全局读写的性能最大化。对于每个线程 32-bit 数量的全局读取请求，例如示例代码中从 A 和 B 进行的读取操作，我们需要 128B 来满足每个 warp 范围的每次请求。因此，当我们知道每个请求的最佳传输传输是多少时，监测每个请求的传输数量就可以得到与 gld_efficiency 和 gst_efficiency metrics 类似的效果。对于 Maxwell 及更新的 GPU ，通常满足某 warp 一次 128-byte 的请求，最少需要 4 条传输（每个 32-byte）。如果数量高于此，说明并没达到最佳性能。
不幸的是，在新工具中也没有与以前 gld_transactions_per_request 或 gst_transactions_per_request 对应的新指标。但这些指数本质上是由分子为传输总数，分母为请求总数所组成的分数。至少对于计算能力7.0及以上的架构（目前为 Volta 和 Turing ），我们可以找到可以用于表示分子和分母的指标（见上述过渡指南中的对照表）。对于全局读取传输，我们使用 l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum ，而对于全局读取请求则使用 l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum 。如果你想知道指标的命名规则，可以在此文档中查看。简而言之，句点之前名称表示了该指标所显示的是哪个数据的信息，句点之后则显示了数字的统计方式。对于 Volta 和更新版本的大多数指标，基本名称和后缀（如果 .sum、.avg ...）共同组成了指标的实际名称。一旦您了解此概念，您可以轻松的将其应用在此架构上几乎其他任何指标上。
为什么要更改指标？Nsight Compute 的设计理念是更详细地展示每个 GPU 的架构和显存系统。提供了更多性能指标，更详细地映射特定架构的特征。可自定义的 analysis section and rules 还提供了一种灵活的机制来结合多种分析数据，以构建更高级的 analyzer 。
下图显示了一个带有各种指标的 GPU 显存模型：
-
l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,.per_second, l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum -
l1tex__t_bytes_pipe_lsu_mem_global_op_st.sum, .per_second -
l1tex__t_sectors_pipe_lsu_mem_local_op_ld.sum, .per_second -
l1tex__t_sectors_pipe_lsu_mem_local_op_st.sum, .per_second -
smsp__inst_executed_op_shared_ld.sum, .per_second -
smsp__inst_executed_op_shared_st.sum, .per_second -
lts__t_sectors_srcunit_tex_op_read.sum, .per_second -
lts__t_sectors_srcunit_tex_op_write.sum, .per_second -
lts__t_sectors_aperture_sysmem_op_read.sum, .per_second -
lts__t_sectors_aperture_sysmem_op_write.sum, .per_second -
dram__bytes_read.sum, .per_second -
dram__bytes_write.sum, .per_second
上表中，每行对应了图中的一个编号。每行中的第一个指标表示其累计的总量，将 .per_second 附加到指标上，可以将其转换成吞吐量。例如：dram__bytes_write.sum 是累积指标， dram__bytes_write.sum.per_second 则是吞吐量指标。此表仅提供一些代表性的示例，不一定适用于每条路径。
熟悉 Nsight Compute CLI
如果你习惯使用 nvprof 或者打算创建自动化脚本，那么 Nsight Compute CLI（命令行界面）会是最好的选择。在此，我们会使用 Linux 环境，windows 命令行的使用应该也是类似的（安装路径和路径相关设置会有所不同）。Nsight Compute tool 将会与 CUDA toolkit 10.0 及以上版本一起安装，也可以在 https://developer.nvidia.com/nsight-compute 上下载独立的安装程序直接安装 Nsight Compute tool 。在运行时，您要么将 Nsight Compute 的二进制文件路径添加到 PATH 环境变量中，要么在执行时指定完成路径。在 CUDA 10.1 上，完整路径为：/usr/local/cuda/NsightCompute-2019.3/，因此如果要指定CLI的完整路径，则使用：/usr/local/cuda/NsightCompute-2019.3/nv-nsight-cu-cli 。此时就可以运行上文所介绍的查询指标命令。对于本博客中，我们将假设您已经将路径添加到 PATH 变量中。
虽然这不是本博客的重点，但 Nsight Compute 提供了许多功能。首先我们可以在可执行文件上以“详细信息模式”运行它。使用 nvcc -arch=sm_70 example.cu -o example 来编译以上代码，修改 -arch 以适配你的 GPU 。在此示例中使用 Volta 设备 (sm_70)，但在 Turing 设备上应该同样可以顺利运行。您无法在在较早的 GPU（例如 Kepler、Maxwell、Pascal）架构上完全遵循此示例，因为计算能力 6.x 的 GPU ，与计算能力为 7.0 及更高的 GPU 可用指标有所不同。此外，计算能力 6.0 及以下的设备不支持使用 Nsight Compute。要显示详细信息页面，请尝试以下操作：
点击查看代码
$ /usr/local/cuda/NsightCompute-2019.3/nv-nsight-cu-cli ./example
==PROF== Connected to process 30244
==PROF== Profiling "matrix_add_2D" - 1: 0%....50%....100% - 48 passes
Success!
==PROF== Disconnected from process 30244
[30244] example@127.0.0.1
matrix_add_2D, 2019-Jun-06 23:12:59, Context 1, Stream 7
Section: GPU Speed Of Light
----------------------------------------- --------------- ------------------------------
Memory Frequency cycle/usecond 866.22
SOL FB % 21.46
Elapsed Cycles cycle 73,170
SM Frequency cycle/nsecond 1.21
Memory [%] % 56.20
Duration usecond 60.16
SOL L2 % 53.58
SOL TEX % 60.21
SM Active Cycles cycle 68,202.96
SM [%] % 8.97
----------------------------------------- --------------- ------------------------------
Section: Compute Workload Analysis
----------------------------------------- --------------- ------------------------------
Executed Ipc Active inst/cycle 0.18
Executed Ipc Elapsed inst/cycle 0.17
Issue Slots Max % 5.00
Issued Ipc Active inst/cycle 0.18
Issue Slots Busy % 4.57
SM Busy % 9.61
----------------------------------------- --------------- ------------------------------
Section: Memory Workload Analysis
----------------------------------------- --------------- ------------------------------
Memory Throughput Gbyte/second 251.25
Mem Busy % 56.20
Max Bandwidth % 53.58
L2 Hit Rate % 89.99
Mem Pipes Busy % 3.36
L1 Hit Rate % 90.62
----------------------------------------- --------------- ------------------------------
Section: Scheduler Statistics
----------------------------------------- --------------- ------------------------------
Active Warps Per Scheduler warp 11.87
Eligible Warps Per Scheduler warp 0.15
No Eligible % 95.39
Instructions Per Active Issue Slot inst/cycle 1
Issued Warp Per Scheduler 0.05
One or More Eligible % 4.61
----------------------------------------- --------------- ------------------------------
Section: Warp State Statistics
----------------------------------------- --------------- ------------------------------
Avg. Not Predicated Off Threads Per Warp 29.87
Avg. Active Threads Per Warp 32
Warp Cycles Per Executed Instruction cycle 261.28
Warp Cycles Per Issued Instruction 257.51
Warp Cycles Per Issue Active 257.51
----------------------------------------- --------------- ------------------------------
Section: Instruction Statistics
----------------------------------------- --------------- ------------------------------
Avg. Executed Instructions Per Scheduler inst 3,072
Executed Instructions inst 983,040
Avg. Issued Instructions Per Scheduler inst 3,116.96
Issued Instructions inst 997,428
----------------------------------------- --------------- ------------------------------
Section: Launch Statistics
----------------------------------------- --------------- ------------------------------
Block Size 1,024
Grid Size 1,024
Registers Per Thread register/thread 16
Shared Memory Configuration Size byte 0
Dynamic Shared Memory Per Block byte/block 0
Static Shared Memory Per Block byte/block 0
Threads thread 1,048,576
Waves Per SM 6.40
----------------------------------------- --------------- ------------------------------
Section: Occupancy
----------------------------------------- --------------- ------------------------------
Block Limit SM block 32
Block Limit Registers block 4
Block Limit Shared Mem block inf
Block Limit Warps block 2
Achieved Active Warps Per SM warp 48.50
Achieved Occupancy % 75.78
Theoretical Active Warps per SM warp/cycle 64
Theoretical Occupancy % 100
----------------------------------------- --------------- ------------------------------
它输出了很多内容（如果您的代码调用了多个 kernel ，它会收集并显示每个的状态）。我们不会介绍所有细节，但是有些 SOL（使用率） 方面数据需要注意：计算分析、显存分析、调度器、warp 状态、指令、配置信息、occupancy 分析。您可以使用命令行参数来决定要显示哪些部分。命令行参数帮助可以使用 --help 来获得，也可以在文档中查看。请注意，选择输出的部分和指标的选择会影响分析的时间和输出的大小。
我们可以使用以上数据来对我们的目标（全局读取/写入效率）进行判断。我们以类似 nvprof 的方式获取信息：
$ nv-nsight-cu-cli --metrics l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum ./example
==PROF== Connected to process 30749
==PROF== Profiling "matrix_add_2D" - 1: 0%....50%....100% - 4 passes
Success!
==PROF== Disconnected from process 30749
[30749] example@127.0.0.1
matrix_add_2D, 2019-Jun-06 23:25:45, Context 1, Stream 7
Section: Command line profiler metrics
------------------------------------------------ ------------ ------------------------------
l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum request 65,536
l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum sector 2,097,152
------------------------------------------------ ------------ ------------------------------
上面第一个指标表示【请求总数】 作为分母， 第二个指标表示【传输总数】 作为分子。将他们相除，即可得出每个请求需要 32 个传输的结果。这意味着，warp 中每个线程都在进行单独的传输，并没有合并。
使用 Nsight Compute GUI
如果我们想用 GUI 来显示这些数据怎么办？在 linux 上，我们需要使用 X session 来启动 GUI 程式应用。因此你需要在支持图形功能的设备上运行。要运行 Nsight Compute GUI ，在命令行输入以下命令：
/usr/local/cuda/NsightCompute-2019.3/nv-nsight-cu
如果你已经将路径加入到 PATH 变量中，可以直接输入 nv-nsight-cu ，接下来你就会看到 Nsight Compute 的图形界面：
从最简单开始，点击 Quick Launch 下的 Continue （或者你可以点击 New Project 下的 Create New Project 来创建一个新的项目。）接下来，应该会打开一个分析配置窗口，点击窗口下方的 Additional Options ，点击 Other 分页，接下来填入 Application Executable: 、 Output File: 、 Metrics: ：
在此我们输入了要分析的可执行文件的路径和名称（示例）、导出分析结果的文件和用逗号分隔开的指标：
l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum
之后，您可以最小化 Additional Options 窗口，并点击蓝色的 Launch 按钮。然后分析器就会运行并捕获数据，结果如下：
在上图中，显示了所填指标的数据，及其他资料的数据（在本例中为显存负载）。请注意，在此状态下保存到磁盘的文件是人类不可读的，需要在 Nsight Compute GUI 中查看。对于人类可读文件，在大多报告页面中都有导出按钮，通常位于右上角。
如果您想更详细地探索 GUI 功能，文档包含介绍 GUI 的快速入门部分。
修复代码
示例代码的执行效率低（每个请求的传输数量过多）的原因是由于我们使用了 2D 索引：
... C[idx][idy] = A[idx][idy] + B[idx][idy];
使用 threadIdx.x（即 idx）构建的索引应该出现在最后一个下标中，以便跨 warp 进行合并访问；相反，它出现在第一个下标中。虽然两种方法都可以给出正确的结果，但他们的性能差距很大。这种设计会倒是 warp 中每个线程访问的是显存中 “column” 的数据，而不是 “row”（即相邻）的数据。我们可以修改 kernel 代码来优化这个问题：
__global__ void matrix_add_2D(const arr_t * __restrict__ A, const arr_t * __restrict__ B, arr_t * __restrict__ C, const size_t sw, const size_t sh){
size_t idx = threadIdx.x+blockDim.x*(size_t)blockIdx.x;
size_t idy = threadIdx.y+blockDim.y*(size_t)blockIdx.y;
if ((idy < sh) && (idx < sw)) C[idy][idx] = A[idy][idx] + B[idy][idx];
}
唯一改变的只有最后一行代码，我们交换了 idx 和 idy 的位置。当我们重新编译修改后的代码并运行上面相同的分析时，可以看到：
$ nv-nsight-cu-cli --metrics l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum ./example
==PROF== Connected to process 5779
==PROF== Profiling "matrix_add_2D" - 1: 0%....50%....100% - 4 passes
Success!
==PROF== Disconnected from process 5779
[5779] example@127.0.0.1
matrix_add_2D, 2019-Jun-11 12:01:26, Context 1, Stream 7
Section: Command line profiler metrics
----------------------------------------------- --------------- ------------
l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum request 65,536
l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum sector 262,144
----------------------------------------------- --------------- ------------
现在这个比例变为了 4:1 （每个请求的传输数量），说明达到了 32 byte 的期望值，并且读取（和写入）的效率也比以前高很多。、
由于这项工作涉及将新旧结果的比较，因此我们继续演示 GUI 的附加功能。我们可以使用 GUI 收集两种情况的分析结果，并进行比较。如上文所述，我们先收集旧代码的数据，并保持 GUI 界面打开 ，然后选择左上角的 Connect 按钮，将输出文件改为新名称，并把分析文件改成修改后的文件。然后点击 Launch 创建一个新选项卡，其中的数据则是修改后的代码的。最后回到初始结果选项卡，选择顶部的 Add Baseline 并选择新结果的选项卡，即可查看各项指标的差异。
在本例中，我们看到改进后的传输次数相比原始的减少到了1/8
提高了显存的使用效率，将会提高使用此显存的代码的性能，这意味着程序运行的速度也会更快。我们可以使用上一篇博客总介绍的 Nsight Systems profiler ，检查更改前后内核的运行时间来验证这一点。我们运行以下命令，这与我们上一篇博客中减少的第一个 CLI 命令类似：
$ nsys profile -o example.nsysprofout --stats=true ./example
但是本文的重点是 Nsight Compute ，我们可以通过 GPU SOL 报告中的 Elapsed Cycles 部分来进行类似的检测。我们继续使用上一节介绍的比较功能。在 GUI 中，首先选择左上角的 Connect 按钮打开配置设置，选择 Additional Options 下拉菜单，你可以将 Other 中的指标都清楚，然后选择 Sections 选项卡，选择 GPU Speed of Light（可以将其他选项取消，以简化输出并减少分析时间）。您可能还需要更改分析输出文件的文件名。然后点击蓝色的 launch 按钮。
单击 启动 按钮以收集新的分析数据。与先前一样，我们对原版和修改版重复执行这些步骤。然后将原版设为基准，并查看对比。
如上图所述，我们可以看到修改后的代码的执行时间减少了约 68%。包括其他各项数据，也显示此更改对性能的影响。
有哪些新功能？
与 NVIDIA Visual Profiler 和 nvprof 相比，Nsight Compute 中有许多新功能，我们在本博客中仅涉及其中的一些。
Nsight Compute GUI 与 Visual Profiler 相比的新功能：
- 在工具中比较分析结果
- 交互式分析模式（使用 API 流和参数捕获）
- 具有跨操作系统支持的远程操作
Nsight Compute GUI 和 CLI 与 Visual Profiler/nvprof 相比的新功能：
- 更详细的指标
- 可自定义指标和基于 python 的分析引导
- 更稳定的数据收集（时钟控制、缓存重置……）
- 减少内核重运行的开销（与第一次的差异）
- 支持新的 CUDA/NVTX 功能（例如图形支持、nvtx 过滤器描述）
结论
与 nvprof 和 Visual Profiler 相比，新工具旨在提供相同（并且更好）的功能，但需要一些新设置和新方法才能获得相似的结果。关于作为本博客主要关注点的指标分析，熟悉新指标很重要，将新指标进行组合来获取您想要的结果。