使用Nsight Profiling工具对大模型进行性能调优

1. 什么是 Nsight System分析工具？

Nsight Systems 是 NVIDIA 提供的一款系统级性能分析工具，可以帮助我们进行CUDA代码的开发优化以及模型代码的优化。

1.1. 主要功能

系统级分析：Nsight Systems 能够捕获和分析整个系统的性能数据，包括 CPU 和 GPU 的活动、内存使用情况、线程调度、I/O 操作等。

时间轴视图：提供详细的时间轴视图，展示不同时间点上系统中各个组件的活动情况。开发者可以通过时间轴视图直观地看到 CPU 和 GPU 任务的执行时间、数据传输时间等。

API 跟踪：能够跟踪各种 API 调用，如 CUDA API、NVTX（NVIDIA Tools Extension）标记、操作系统线程调度等，帮助开发者了解应用程序的执行流程和时间开销。

2. 如何进行深度学习任务分析？

下面我们首先先来思考两个问题？

为什么深度学习任务提升GPU利用率和性能是比较复杂的过程？
当我们拿到待分析的深度学习任务如何开始分析？直接nsight Profiler吗？

2.1. 深度学习任务分析的难点

对于深度学习训练来说，它既可能是计算密集（Compute intensive），又可能是数据密集（Data intensive），还可能是内存密集（Memory intensive）型的工作负载。

不同的模型面临的情况不同，其瓶颈点也就不尽相同。随着对模型的优化，其瓶颈点也可能随之变化。

此外，提升模型训练的性能与提升整体的GPU利用率又不是完全的正相关，这也使得如何提升深度学习模型训练的性能与GPU利用率成为一个比较复杂的工程。

2.2. 进行任务性能分析的步骤

当我们开始进行深度学习任务分析时，我们应该把握两个关键：

1.宏观到微观，从整体到局部。

2.综合阶段划分、单侧复现、工具Profiler和源码定位步骤。

当前绝大数的DL模型都是跑在GPU上的，而GPU 任务会交替的使用 CPU 和 GPU 进行计算，当 CPU 计算成为瓶颈时，就会出现 GPU 等待，频繁的CPU与GPU切换也会导致GPU利用率下降和性能下降的问题。

所以，从宏观上我们将DL的训练分为两个阶段，如上图所示：

CPU上的数据处理（包括模型保存）
GPU上的模型计算。

所以宏观上，我们首先要认识整个模型或者代码，将模型分为多个阶段，包括数据加载、数据预处理、模型计算和模型保存。还要注意模型的训练策略，是采用哪种分布式训练方式，在当前数据量下是否适合模型。

此外，还要注意当前模型训练的数据，例如GPU利用率，显存、内存的占用量等。

宏观上，我们可以通过GPU利用率，显存、内存的占用量快速判断模型存在哪些的问题，我们可以举几个例子：

如果 GPU 利用率较低，但磁盘或 CPU 利用率较高，则数据加载或预处理可能是潜在的瓶颈，可以在训练之前就对数据进行预处理。
如果 GPU 利用率较低，并且 CPU 和磁盘利用率持续较低但不为零，尽管数据集足够大，但这可能意味着您的代码未有效利用底层资源，可以增加数据加载器 API 中的工作线程数量。

微观上，我们可以通过固定住内存中数据使模型没有IO瓶颈，判断当前模型运行状态。下面我们总结了些checklist可以帮助快速判断模型存在的问题：

数据加载阶段：

小文件是否太多，导致文件 io 耗时太长，读取会浪费很多时间在寻道上。
存储介质是否已达到瓶颈，可以监控存储介质的繁忙度，如果达到瓶颈可以增加存储介质缓解读取性能；
是否启用多进程并行读取数据，另外可以注意线程争用问题，监控线程等待时候是否过长，可以采用私有线程池进行环境；
是否启用提前加载机制来实现 CPU 和 GPU 的并行；

数据预处理阶段：

是否设置开启共享内存 pin_memory，可以直接将数据放置在pin_memory中；
优化 I/O 和网络操作，确保数据以与其计算相匹配的速率馈送到 GPU；
如果是A100或以上的机器，可以考虑开启numa绑定，缓解争用，提升性能；

模型训练阶段：

是否存在大量的CPU运算，可以通过实现GPU实现或去除指定CPU设备，尽可能的让模型运行在GPU上；
模型是否存在GPU利用率不均的情况，尽可能得不在代码里指定GPU运行的卡；
对于比较复杂运行效率较低的模块，可以通过实现融合的大GPU算子提升训练速度；
避免指标和日志打印太频繁，CPU 和 GPU 频繁切换导致 GPU 利用率低；
是否开启AMP来提升模型的训练性能；
使用最新的高性能库和 GPU 驱动程序，cuda是否升级到最新版本；

在进行了第一步的宏观上对于模型或代码运行状态认识，和第二步检查对应checklist后，我们再进行nsight profiler分析。

3. 如何安装使用Nsight Systems？

3.1. 下载安装nsight

可以从下面这个网址下载安装nsight工具包：

nsight-systems-2022-4

如果在安装的过程中报错can't locateEnv.pmin @INC，可以安装下面的组件

yum install perl-Env

下载好后，直接bash运行即可安装，在安装过程中选择accept其协议并注意安装地址。

3.2. 下载安装nvtx

nvtx 安装包可以从nvtx下载，并注意py的版本。

此外还可以使用nvtx-plugin进行安装，它可以跟踪更多的信息。nvtx-plugin 可以有更丰富的染色方法。

3.3. 使用nsight方法

3.3.1. nvtx&nvtx-plugin染色

下面分别是 nvtx 和nvtx-plugin的使用方式：

3.3.2. 执行profile命令

在执行profile上有两种方式一种，一种是 Non-interactive 方式，一种是 Interactive, Non-interactive 方式虽然用起来麻烦，但在多种情况下都可以使用。下面是两种用法的使用方式

例子：

nsys profile --sample=none --backtrace=none --cuda backtrace=none --cpuctxsw=none --trace-fork-before-exec=true python main.py

3.4. 使用客户端加载

如何看nsight?

整体认识，寻找关键节点。
查看统计信息，寻找低性能点。

4. CUDA流与CPU线程的同步与异步处理

4.1. 基本概念

CUDA流（Stream）: CUDA流是一个命令队列，CUDA内核和内存操作可以在流中按顺序执行。不同流中的操作可以并行执行。
异步执行: CUDA操作（如内核启动和内存拷贝）可以异步启动，不需要等待前一个操作完成。
同步: 某些操作需要等待所有先前操作完成，比如 cudaDeviceSynchronize。

4.2. 分析Nsight

查看时间轴: Nsight Systems提供了详细的时间轴视图，可以显示CUDA流、内核执行时间、内存拷贝时间以及CPU线程活动。通过时间轴视图，可以识别出哪些操作是串行执行的，哪些是并行执行的。
标记同步点: 查找CUDA API调用，如 cudaDeviceSynchronize 、 cudaStreamSynchronize 等，这些是同步点。过多的同步点会导致性能瓶颈。
分析内核执行和内存拷贝: 查看内核执行和内存拷贝的时间，识别是否有长时间的空闲期，可能是由于同步导致的。

4.3. 优化策略

1. 减少不必要的同步:

尽量减少显式的同步调用，如 cudaDeviceSynchronize。
使用 cudaStreamWaitEvent 等事件机制来实现更细粒度的同步控制。

2. 使用多个流:

将独立的CUDA操作分配到不同的流中，以实现并行执行。
确保内核启动和内存拷贝操作尽可能在不同流中并行执行。

3. 优化内存拷贝:

使用异步内存拷贝函数（如 cudaMemcpyAsync ）并将其分配到不同的流中。尽量减少Host与Device之间的内存拷贝次数，使用统一内存（Unified Memory）或零拷贝（Zero Copy）技术。

举例：使用了两个流来重叠计算和内存拷贝操作。


#include <cuda_runtime.h>
__glob
al__ void kernel1(int *data) {
// 内核1的计算
}
__glob
al__ void kernel2(int *data) {
// 内核2的计算
}

void process_data(int *data, size_t size) {
int *d_data;
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
cudaMalloc(&d_data, size * sizeof(int));

// 异步内存拷贝和内核启动
cudaMemcpyAsync(d_data, data, size * sizeof(int),
cudaMemcpyHostToDevice, stream1);
kernel1<<<grid, block, 0, stream1>>>(d_data);
kernel2<<<grid, block, 0, stream2>>>(d_data);
cudaMemcpyAsync(data, d_data, size * sizeof(int),
cudaMemcpyDeviceToHost, stream2);

// 等待所有操作完成
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);
cudaFree(d_data);
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);
}

5. 内存传输的优化

5.1. 基本概念

Pinned内存拷贝：Pinned内存（也称为Page-Locked内存）可以显著提高内存传输速度，因为它避免了分页机制带来的开销。
异步内存拷贝：允许内存传输与计算重叠，从而提升性能。使用 cudaMemcpyAsync 函数可以在指定的流中异步执行内存拷贝。
Zero Copy：对于小数据集或频繁访问的数据，可以使用零拷贝技术，使得GPU直接访问主机内存。
重叠计算与数据传输：将计算任务和数据传输任务分配到不同的流中，使得它们可以并行执行，从而提高效率。

5.2. 分析Nsight

在Nsight Systems的时间轴视图中，可以看到内存传输和计算任务的执行情况。确保内存传输和计算任务尽可能并行执行。
查找长时间的内存传输操作或空闲期，识别出可能的性能瓶颈。
查找PagedMemoryCpy的算子执行。GPU和CPU之间通过DMA来拷贝数据，如果buffer是在可分页的内存上时，此时cuda会先将buffer从可分页的内存拷贝到锁页内存，再复制到gpu上，如下图所示：

我们想到可以通过通过将Pageable Memory拷贝转换为Pinned Memory来提升前面Data处理时间。

6. 常见编程性能优化手段

6.1. 使用异步API

使用异步API如cudaMemcpyAsync可让GPU操作与CPU操作并行，CPU忙完后调用cudaStreamSynchronize，cudaEventWait等操作等待GPU任务完成。

6.2. 优化内存与显存传输效率

使用Pinned(page-locked) Memory提高传输速度
通过在不同的Stream里同时分别执行kernel调用及数据传输，使数据传输与运算并行。（注意default stream的坑）
尽量将小的数据在GPU端合成大块数据后传输

6.3. 优化Kernel访存效率

6.3.1. 提高Global Memory访存效率

对Global Memory的访存需要注意合并访存(coalesced )。
warp的访存合并后，起始地址及访存大小对齐到32字节
尽量避免跨步访存
CUDA 8.0及以上的设备可以通过编程控制L2的访存策略提高L2命中率。

6.3.2. 提高Shared Memory的访存效率

shared memory由32个bank组成
每个bank每时钟周期的带宽为4字节
连续的4字节单元映射到连续的bank。如0-3字节在bank0，4-7字节在bank1……字节128-131字节在bank0
若warp中不同的线程访问相同的bank，则会发生bank冲突(bank conflict)，bank冲突时，warp的一条访存指令会被拆分为n条不冲突的访存请求，降低shared memory的有效带宽。所以需要尽量避免bank冲突。
CUDA 11.0以上可以使用async-copy feature

6.4. 优化线程级并行

在SMSP工作时，某些warp会由于访存依赖、寄存器依赖等原因stall。此时warp scheduler可以选中另一个eligible warp，执行其指令，以隐藏前一个warp的stall，使SMSP中的各个硬件资源尽量保持忙碌。

但假如SMSP中所有的warp都不在eligible状态，则硬件只能空转等待某个 warp从stall中恢复(如从global中请求的数据终于回来了)。

Occupancy指标用来衡量SM当前activate warp数量与理论上最多支持的activate warp数量的比值。Occupancy数量越高，代表SMSP负责的activate warp越多，当某个warp stall时，有更多的备选warp，有更大的概率可以找到一个eligible warp。

极端情况Occupancy为 1/8时，SM仅4个warp，每个SMSP 1个warp，当该warp stall时，smsp没有其它warp可以选择，硬件必然空转等待。影响Occupancy指标的包括以下因素：

Thread Block 线程块的大小。
每个线程块的Shared Memory使用量
每个线程使用的Register（寄存器数量）高的Occupancy不一定代表较高的性能，如某些算法确实需要每线程128寄存器时，保持0.5的Occupancy反而是最优选择。但过低的Occupancy会对性能带来较大的负面影响。

6.5. 使用TensorCore进一步加速矩阵运算

TensorCore可以用来快速进行D=A * B+C矩阵运算，提供 load_matrix_sync，store_matrix_sync, mma_sync 等API。

7. 总结

通过Nsight Systems的时间轴视图和API跟踪功能，我们可以详细地分析整个系统的性能数据，包括CPU和GPU的活动、内存使用情况、线程调度和I/O操作。

这使得我们能够精确地定位性能瓶颈，并采取相应的优化措施，如减少不必要的同步、使用多个CUDA流、优化内存拷贝操作、提高内存传输效率、优化Kernel访存效率、提高线程级并行性，以及利用TensorCore加速矩阵运算。

性能优化是一个复杂但至关重要的过程，它要求我们不仅要对模型的宏观运行状态有清晰的认识，还要对微观的代码实现细节有深入的了解。通过结合Nsight Systems的分析能力和我们的优化策略，我们可以显著提升深度学习模型的训练效率和GPU资源的利用率。

最后，如果这篇文章可以帮助到你，希望可以点赞和分享给更多的人。

作者：Pulsar planet
来源：Tim在路上

推荐阅读

欢迎大家点赞留言，更多Arm技术文章动态请关注极术社区嵌入式AI专栏欢迎添加极术小姐姐微信（id:aijishu20)加入技术交流群，请备注研究方向。