定时 CUDA 内核答案

【问题标题】：Timing CUDA kernels定时 CUDA 内核
【发布时间】：2012-06-17 09:33:15
【问题描述】：

大家好，我目前正在为我的一些 CUDA 代码计时。我能够使用事件为他们计时。我的内核运行了 19 毫秒。不知何故，我觉得这很可疑，因为当我运行它的顺序实现时，它大约是 5000 毫秒。我知道代码应该运行得更快，但它应该这么快吗？

我正在使用包装函数在我的 cpp 程序中调用 cuda 内核。我应该在那里或在 .cu 文件中调用它们吗？谢谢！

【问题讨论】：

CUDA 的 100 倍加速并不令人惊讶。但是你应该发布一些代码，这样我们就可以看到你在做什么！
你使用过流吗？您是否在内核调用之后和时间测量之前添加了cudaDeviceSynchronize()，以防默认流使用？
由于 OP 正在使用事件，因此 OP 应该使用 cudaEventSynchronize()，而不是 cudaDeviceSynchronize()（后者会起作用，但对于计时来说有点重锤......）。跨度>
如何从 .cpp 文件中调用 cuda 内核？如果您没有使用 >>、CUDA 驱动程序 API 或 cudaLaunch()，那么您没有在设备上启动内核。发布一些示例代码将帮助我们回答。
另一个需要解决的问题是您如何测量顺序版本的时间？一些不公平的比较从控制台测量顺序代码的全部时间，包括malloc 和free 的时间，更糟糕的是，从文件中读取输入数据的时间。因此，与 CUDA 内核（不包括 cudaMalloc、cudaFree 甚至从 CPU--GPU 传输数据）相比，可以带来令人印象深刻的加速。

标签： cuda

【解决方案1】：

检查您的程序是否正常工作的明显方法是将输出与基于 CPU 的实现的输出进行比较。如果你得到相同的输出，它按定义工作，对吧？ :)

如果您的程序是实验性的，它不会真正产生任何可验证的输出，那么编译器很可能已经优化了您的部分（或全部）代码。编译器将删除对输出数据没有贡献的代码。例如，如果存储计算值的最终语句被注释掉，这可能会导致内核的全部内容被删除。

至于你的加速。 5000 毫秒 / 19 毫秒 = 263 倍，这不太可能增加，即使对于完美映射到 GPU 架构的算法也是如此。

【讨论】：

【解决方案2】：

好吧，如果你正确地编写了你的 CUDA 代码，是的，它可能会快得多。想想看。您将代码从单个处理器上的顺序执行移动到数百个处理器上的并行执行，具体取决于您的 GPU 型号。我 179 美元的中档显卡有 480 个内核。现在一些可用的有 1500 个内核。使用 CUDA 非常有可能获得 100 倍的性能跳跃，特别是如果您的内核的计算限制比内存限制更多。

也就是说，请确保您正在衡量您认为自己正在衡量的内容。如果您在不使用任何显式流的情况下调用 CUDA 内核，则调用与主机线程同步，并且您的时间应该是准确的。如果您使用流调用内核，则需要调用 cudaDeviceSynchronise() 或让主机代码等待内核发出的事件信号。在流上调用的内核调用与主机线程异步执行，因此主机线程中的时间测量将无法正确反映内核时间，除非您让主机线程等待内核调用完成。您还可以使用 CUDA 事件来测量给定流中 GPU 上的经过时间。请参阅 NVidia GPU Computing SDK 4.2 中 CUDA 最佳实践指南的第 5.1.2 节。

【讨论】：

【解决方案3】：

在我自己的代码中，我使用clock() 函数来获取精确的时间。为方便起见，我有宏

enum {
    tid_this = 0,
    tid_that,
    tid_count
    };
__device__ float cuda_timers[ tid_count ];
#ifdef USETIMERS
 #define TIMER_TIC clock_t tic; if ( threadIdx.x == 0 ) tic = clock();
 #define TIMER_TOC(tid) clock_t toc = clock(); if ( threadIdx.x == 0 ) atomicAdd( &cuda_timers[tid] , ( toc > tic ) ? (toc - tic) : ( toc + (0xffffffff - tic) ) );
#else
 #define TIMER_TIC
 #define TIMER_TOC(tid)
#endif

这些可用于检测设备代码，如下所示：

__global__ mykernel ( ... ) {

    /* Start the timer. */
    TIMER_TIC

    /* Do stuff. */
    ...

    /* Stop the timer and store the results to the "timer_this" counter. */
    TIMER_TOC( tid_this );

    }

然后您可以读取主机代码中的cuda_timers。

一些注意事项：

计时器在每个块的基础上工作，即如果您有 100 个块执行同一个内核，则将存储它们所有时间的总和。
计时器计算时钟滴答数。要获取毫秒数，请将其除以设备上的 GHz 数，然后乘以 1000。
计时器会稍微减慢您的代码速度，这就是为什么我将它们包装在 #ifdef USETIMERS 中以便您可以轻松关闭它们。
虽然clock() 返回clock_t 类型的整数值，但我将累积值存储为float，否则对于耗时超过几秒的内核（在所有块上累积），这些值将回绕。李>
选择( toc > tic ) ? (toc - tic) : ( toc + (0xffffffff - tic) ) ) 是必要的，以防时钟计数器回绕。

【讨论】：

这种“所有块时间的总和”方法不是计算内核执行时间的准确方法。如果使用正确，OP 正在使用的 CUDA 事件是一种准确的方法来计时内核的执行。