CUDA 内核计时策略：优点和缺点？

Posted 2023-04-15

【中文标题】CUDA 内核计时策略：优点和缺点？

【英文标题】：Strategies for timing CUDA Kernels: Pros and Cons?

【发布时间】：2012-11-20 11:55:28

【问题描述】：

在为 CUDA 内核计时时，以下内容不起作用，因为内核在执行时不会阻塞 CPU 程序执行：

start timer
kernel<<<g,b>>>();
end timer

---

我已经看到了三种（成功地）为 CUDA 内核计时的基本方法：

(1) 两个 CUDA eventRecords。

float responseTime; //result will be in milliseconds
cudaEvent_t start; cudaEventCreate(&start); cudaEventRecord(start); cudaEventSynchronize(start);
cudaEvent_t stop;  cudaEventCreate(&stop);
kernel<<<g,b>>>();
cudaEventRecord(stop); cudaEventSynchronize(stop);
cudaEventElapsedTime(&responseTime, start, stop); //responseTime = elapsed time

(2) 一个 CUDA eventRecord。

float start = read_timer(); //helper function on CPU, in milliseconds
cudaEvent_t stop;  cudaEventCreate(&stop);
kernel<<<g,b>>>();
cudaEventRecord(stop); cudaEventSynchronize(stop);
float responseTime = read_timer() - start;

(3) deviceSynchronize 而不是 eventRecord。 （可能仅在在单个流中使用编程时有用。）

float start = read_timer(); //helper function on CPU, in milliseconds
kernel<<<g,b>>>();
cudaDeviceSynchronize();
float responseTime = read_timer() - start;

我通过实验验证了这三种策略产生相同的计时结果。

---

问题：

这些策略的权衡是什么？这里有任何隐藏的细节吗？

除了对多个流中的多个内核进行计时之外，使用两个事件记录和cudaEventElapsedTime() 函数有什么优势吗？

---

您可能可以发挥想象力来弄清楚read_timer() 做了什么。不过，提供一个示例实现也无妨：

double read_timer()
    struct timeval start;
    gettimeofday( &start, NULL ); //you need to include <sys/time.h>
    return (double)((start.tv_sec) + 1.0e-6 * (start.tv_usec))*1000; //milliseconds

【参考方案1】：

您似乎已经排除了大多数差异，因为您说对于您展示的相对简单的情况（可能不完全正确，但我理解您的意思），它们都产生相同的结果，并且“除了时间（复杂序列）...”，其中第一种情况显然更好。

一个可能的区别是 windows 和 linux 之间的可移植性。我相信您的示例 read_timer 函数是面向 linux 的。您可能可以制作一个“可移植”的 read_timer 函数，但 cuda 事件系统（方法 1）是可移植的。

【参考方案2】：

选项 (1) 使用 cudaEventRecord 对 CPU 进行计时。这是非常低效的，我不鼓励为此目的使用 cudaEventRecord。 cudaEventRecord 可用于计时 GPU 推送缓冲区时间以执行内核，如下所示：

float responseTime; //result will be in milliseconds
cudaEvent_t start;
cudaEvent_t stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start);
kernel<<<g,b>>>();
cudaEventRecord(stop);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&responseTime, start, stop); //responseTime = elapsed time

如果您将多个工作项提交到多个流，则需要稍微更改代码。我建议阅读Difference in time reported by NVVP and counters的答案

选项 (2) 和 (3) 与给定示例类似。选项 (2) 可以更灵活。

【讨论】：

您刚刚从选项 (1) 中删除了 cudaEventSynchronize(start);，对吧？或者，我是否忽略了其他任何变化？