在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) ) )如果时钟计数器环绕,则必须进行选择.

PS这是我对这个问题的回复的副本,由于所需的时间是针对整个内核的,因此没有得到很多要点.