如何优雅地测量GPU CUDA Kernel耗时？
背景
工作中经常需要评估模型性能，会涉及到GPU Kernel的耗时分析。网上关于GPU耗时的文章比较散，本文记录下日常用到的一些工具及使用过程中的坑，大多是使用经验，若有错误，还请不吝指出。
GPU耗时和CPU耗时有啥区别？
在开始之前，我们需要了解GPU耗时和CPU耗时有啥区别？
在默认情况下，CPU与GPU的执行逻辑是异步的。CPU中调用kernel时，本质上是调用的kernel launcher，会把kernel提交给GPU执行，而GPU啥时候执行，并不确定。需要在当前GPU Stream执行完已有的kernel后，才会执行提交的kernel。而CPU执行完kernel launcher后会直接执行后续的逻辑，本质上是为了达到CPU与GPU并行的目的。
看一个示例：
#include <stdio.h>
#include <iostream>
#include <chrono>
// 一个示例kernel，之所以写result，是为了避免被nvcc优化掉
__global__ void kernel_test1(float *result)
{
unsigned int idx = threadIdx.x + blockDim.x * blockIdx.x;
for (int i = 1234; i >= 0; --i)
{
if (idx == 0)
{
result[0] = i;
}
}
}
__global__ void kernel_test2(float *result)
{
unsigned int idx = threadIdx.x + blockDim.x * blockIdx.x;
for (int i = 1234 * 2; i >= 0; --i)
{
if (idx == 0)
{
result[1] = i;
}
}
}
int main()
{
// 一些调用kernel的准备工作
const dim3 block_size(5, 6, 7);
const dim3 grid_size(2, 3, 4);
float *d_result;
cudaMalloc((void **)&d_result, 2 * sizeof(float));
// 一个简单的CPU函数，CPU耗时统计示例
auto t1 = std::chrono::high_resolution_clock::now();
float localResult = 0.0f;
for (unsigned int i = 0; i < 123456789; ++i)
{
localResult += float(i) * 0.00001f;
}
auto t2 = std::chrono::high_resolution_clock::now();
// 这里预先调用几次kernel，预热下GPU，避免预热的影响
for (int i = 0; i < 10; ++i)
{
kernel_test1<<<grid_size, block_size>>>(d_result);
kernel_test2<<<grid_size, block_size>>>(d_result);
}
cudaDeviceSynchronize();
// 这里实际上是统计的kernel launcher的时间
auto t3 = std::chrono::high_resolution_clock::now();
for (int i = 0; i < 100; ++i)
{
kernel_test1<<<grid_size, block_size>>>(d_result);
kernel_test2<<<grid_size, block_size>>>(d_result);
}
auto t4 = std::chrono::high_resolution_clock::now();
// 等待kernel执行完成
cudaDeviceSynchronize();
auto t5 = std::chrono::high_resolution_clock::now();
auto d21 = std::chrono::duration_cast<std::chrono::microseconds>(t2 - t1);
auto d43 = std::chrono::duration_cast<std::chrono::microseconds>(t4 - t3);
auto d54 = std::chrono::duration_cast<std::chrono::microseconds>(t5 - t4);
std::cout << "CPU:" << d21.count() << " ms" << std::endl;
std::cout << "Kernel Launcher:" << d43.count() << " ms," << std::endl;
std::cout << "cudaDeviceSynchronize:" << d54.count() << " ms" << std::endl;
cudaFree(d_result);
return 0;
}
程序输出为：
CPU:340272 ms
Kernel Launcher:942 ms,
cudaDeviceSynchronize:7798 ms
可以看出是kernel launcher执行时间为942 ms，kernel的总耗时近似cudaDeviceSynchronize的时间，对于每次kernel的时间以上并未统计出来。
那么如何才能看到每个kernel的时间呢？由此引出以下手段。
常用测量GPU耗时手段
Nsight System
Nsight System是Nvidia官方提供的profile可视化工具，详情参考https://developer.nvidia.cn/nsight-systems。
安装完成后，可以运行nsys会出现以下：
对整个程序执行profile的命令如下：
nsys profile --trace=cuda,cudnn,cublas,nvtx --force-overwrite true -o report ./xxx
运行后可得到report.nsys-rep文件，用桌面端程序打开如下：
如图所示，上边部分是GPU kernel的执行耗时，里面可以展开，查看kernel的耗时，而CUDA API这一列则是Launcher的时间。我们可以选择其中的一个kernel：
图中会自动展示对应的launcher和kernel，图中左侧就是launcher，右侧就是kernel真实执行的耗时，可以发现二者执行时间点和耗时差别巨大，这也就符合以上提到的CPU与GPU并行的逻辑。
在分析nsys时，我们可以直接圈选kernel来测量时间，很直观。另外nsys还有其他的命令，可以参考：https://docs.nvidia.com/nsight-systems/UserGuide/index.html。
CudaEvent
nsys可以观测整体kernel的耗时，但有时候程序逻辑很多，我们只想观测指定的kernel，难免被其他kernel影响观测，而且我们很多时候也需要统计指定的kernel耗时信息，如平均时间，中位数，分布等等，这时可以使用CudaEvent。主要使用如下几个API：
cudaEventCreate()
cudaEventRecord()
cudaEventSynchronize()
cudaEventElapsedTime()
cudaEventDestroy()
API的含义可参考：https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__EVENT.html
参看以下使用示例：
#include <stdio.h>
#include <iostream>
#include <chrono>
#include <vector>
#include <cuda_runtime_api.h>
#include <algorithm>
#include <numeric>
// 示例kernel
__global__ void kernel_test1(float *result)
{
unsigned int idx = threadIdx.x + blockDim.x * blockIdx.x;
for (int i = 1234; i >= 0; --i)
{
if (idx == 0)
{
result[0] = i;
}
}
}
__global__ void kernel_test2(float *result)
{
unsigned int idx = threadIdx.x + blockDim.x * blockIdx.x;
for (int i = 1234 * 2; i >= 0; --i)
{
if (idx == 0)
{
result[1] = i;
}
}
}
std::vector<std::vector<cudaEvent_t>> initEvent(int m, int n)
{
std::vector<std::vector<cudaEvent_t>> events(m, std::vector<cudaEvent_t>(n));
for (int i = 0; i < m; i++)
{
for (int j = 0; j < n; j++)
{
cudaEventCreate(&events[i][j]);
}
}
return events;
}
void syncEvent(const std::vector<std::vector<cudaEvent_t>> &events)
{
for (auto &event_vec : events)
{
for (auto &event : event_vec)
{
cudaEventSynchronize(event);
}
}
}
std::vector<std::vector<float>> computeEventTime(const std::vector<std::vector<cudaEvent_t>> &events)
{
std::vector<std::vector<float>> times(events.size() - 1, std::vector<float>(events[0].size()));
for (int i = 0; i < events.size() - 1; i++)
{
for (int j = 0; j < events[i].size(); j++)
{
cudaEventElapsedTime(×[i][j], events[i][j], events[i + 1][j]);
}
}
return times;
}
void destoryEvent(std::vector<std::vector<cudaEvent_t>> &events)
{
for (auto &event_vec : events)
{
for (auto &event : event_vec)
{
cudaEventDestroy(event);
}
}
}
void printTimeStatInfo(std::vector<float> data)
{
std::sort(data.begin(), data.end());
std::cout << "printTimeStatInfo:" << std::endl;
std::cout << "Avg:" << std::accumulate(data.begin(), data.end(), 0.0) / (data.size()) << std::endl;
std::cout << "Mid:" << data[data.size() / 2] << std::endl;
std::cout << "Min:" << data[0] << std::endl;
std::cout << "Max:" << data[data.size() - 1] << std::endl;
}
int main()
{
const dim3 block_size(5, 6, 7);
const dim3 grid_size(2, 3, 4);
float *d_result;
cudaMalloc((void **)&d_result, 2 * sizeof(float));
// warm up
for (int i = 0; i < 10; ++i)
{
kernel_test1<<<grid_size, block_size>>>(d_result);
kernel_test2<<<grid_size, block_size>>>(d_result);
}
cudaDeviceSynchronize();
// 初始化event
auto events = initEvent(3, 100);
// 对各kernel进行计时
for (int i = 0; i < 100; ++i)
{
cudaEventRecord(events[0][i], nullptr);
kernel_test1<<<grid_size, block_size>>>(d_result);
cudaEventRecord(events[1][i], nullptr);
kernel_test2<<<grid_size, block_size>>>(d_result);
cudaEventRecord(events[2][i], nullptr);
}
// 等待各个event同步
syncEvent(events);
// 计算时间
auto times = computeEventTime(events);
// 打印时间统计信息
printTimeStatInfo(times[0]);
printTimeStatInfo(times[1]);
// 销毁event
destoryEvent(events);
cudaFree(d_result);
return 0;
}
程序输出为：
printTimeStatInfo:
Avg:0.0431002
Mid:0.043008
Min:0.041984
Max:0.044032
printTimeStatInfo:
Avg:0.0522445
Mid:0.052224
Min:0.0512
Max:0.053248
即kernel1与kernel2的每次执行的耗时信息，可以看到kernel2执行时间大于kernel1，这和程序逻辑是一致的。
注意：
1. cudaEventSynchronize的方法本质是等待event执行完成，在调用该方法时，应在整个程序的最后环节，不要在循环里面去调用sync，这样会影响程序原先的性能，导致测量数据存在问题，这也是之前踩的坑。
我们对以上程序也进行nsys观测下：
可以看到CUDA API这列存在对EventRecord和kernel的调用，顺序与我们的代码一致。
再看看kernel1的耗时：38.318us
kernel2的耗时：48.346us
可以发现，我们统计的耗时与nsys并没有严格对齐，统计的耗时都大于nsys中的kernel时间。猜测原因是event测量的是两次event之间的时间，中间可能还涉及了event本身的消耗的时间，所以会多于之间kernel时间。但整体趋势是一致的，日常也可使用event时间进行对比评估。这里如果有了解差异来源细节的同学，还望留言告知~
后记
- 除了本文提到的方法，还有NCU等工具，由于目前使用较少，暂不列出，可参考官方文档：https://developer.nvidia.com/nsight-compute。
- 以上是在C++中耗时的方式，日常可能也需要在torch等框架中统计kernel耗时，可以参考：https://pytorch.org/docs/stable/generated/torch.cuda.Event.html，本质上也是调用的C++的API，只是用python进行的包装。
参考
- https://pytorch.org/docs/stable/generated/torch.cuda.Event.html
- https://developer.nvidia.cn/nsight-systems
本文链接：
http://rainlin.top/archives/233
转载请注明转载自： Rainlin Home
