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个块在执行同一个内核，那么它们所有时间的总和就会被存储起来。
• 已经说过，定时器假定第0个线程是活动的，所以请确保你不会在代码的一个可能不同的部分调用这些宏。
• 定时器计算时钟节拍的数量。要获得毫秒数，请将其除以设备上的GHz数并乘以1000。
• 计时器可能会使代码变慢一些，这就是为什么我将它们封装在#ifdef USETIMERS中以便您可以关闭它们的原因。clock()返回clock_t类型的整数值，我将累加值存储为float，否则这些值将对耗时超过几秒的内核进行折回(在所有块上累加)。
• 如果时钟计数器折回，则必须使用选择( toc > tic ) ? (toc - tic) : ( toc + (0xffffffff - tic) ) )。< code >H221

附言:这是我对this question的回复，因为计时是针对整个内核的，所以没有得到太多的分数。