optimization - 对 CUDA 内核中的不同部分进行计时

Question

我有一个调用一系列设备功能的 CUDA 内核。

获取每个设备功能的执行时间的最佳方法是什么？

获取设备功能之一中一段代码的执行时间的最佳方法是什么？

Answer 1

在我自己的代码中，我使用了 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) ) )在时钟计数器环绕的情况下是必要的。

附注这是我对 this question 的回复的副本，因为所需的时间是整个内核的，所以没有得到很多积分。