如何优雅地测量GPU CUDA Kernel耗时？（一）- nsys、CudaEvent使用
|
大模型与GPU编程

1544 字
|
18 分钟

背景

工作中经常需要评估模型性能，会涉及到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(&times[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时间进行对比评估。这里如果有了解差异来源细节的同学，还望留言告知~

后记

  1. 除了本文提到的方法，还有NCU等工具，由于目前使用较少，暂不列出，可参考官方文档：https://developer.nvidia.com/nsight-compute。
  2. 以上是在C++中耗时的方式，日常可能也需要在torch等框架中统计kernel耗时，可以参考：https://pytorch.org/docs/stable/generated/torch.cuda.Event.html，本质上也是调用的C++的API，只是用python进行的包装。

参考

  1. https://pytorch.org/docs/stable/generated/torch.cuda.Event.html
  2. https://developer.nvidia.cn/nsight-systems

系列文章

  1. 如何优雅地测量GPU CUDA Kernel耗时？（一）
  2. 如何优雅地测量GPU CUDA Kernel耗时？（二）
  3. 如何优雅地测量GPU CUDA Kernel耗时？（三）
本文链接：https://rainlin.top/archives/233
转载请注明转载自：https://rainlin.top
暂无评论

发送评论 编辑评论 


				

			

						

				

					

						

							

								
							

							
						

					

				

				

					

						

							

								
							

							
						

					

				

				

					

						

							

								
							

							
							
													

					

				

			

			
			

				

					
				

			

				

											

							
							
						

																				
					
											
						

	

		
|´・ω・)ノ
ヾ(≧∇≦*)ゝ
(☆ω☆)
（╯‵□′）╯︵┴─┴
￣﹃￣
(/ω＼)
∠( ᐛ 」∠)＿
(๑•̀ㅁ•́ฅ)
→_→
୧(๑•̀⌄•́๑)૭
٩(ˊᗜˋ*)و
(ノ°ο°)ノ
(´இ皿இ｀)
⌇●﹏●⌇
(ฅ´ω`ฅ)
(╯°A°)╯︵○○○
φ(￣∇￣o)
ヾ(´･ ･｀｡)ノ"
( ง ᵒ̌皿ᵒ̌)ง⁼³₌₃
(ó﹏ò｡)
Σ(っ °Д °;)っ
( ,,´･ω･)ﾉ"(´っω･｀｡)
╮(╯▽╰)╭
o(*////▽////*)q
＞﹏＜
( ๑´•ω•) "(ㆆᴗㆆ)
😂
😀
😅
😊
🙂
🙃
😌
😍
😘
😜
😝
😏
😒
🙄
😳
😡
😔
😫
😱
😭
💩
👻
🙌
🖕
👍
👫
👬
👭
🌚
🌝
🙈
💊
😶
🙏
🍦
🍉
😣
Source: github.com/k4yt3x/flowerhd
	

	

		
颜文字
Emoji
小恐龙
花!
	


									

			

			
			
		

	






上一篇
下一篇 

                    

                    
			        推荐文章
		            

		            

							
如何优雅地测量GPU占用率？（nsys gpu-metrics-devices使用）

							
							

							
如何优雅地测量GPU CUDA Kernel耗时？（三）- nsys统计kernel耗时

							
							

							
如何优雅地测量GPU CUDA Kernel耗时？（二）- 精确测量

							
							

							
如何优雅使用torch的view? 如何理解RowMajor/ColumnMajor？