10

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

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

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

4

1 回答 1

7

在我自己的代码中，我使用该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()返回 type 的整数值clock_t，但我将累积值存储为float，否则对于花费超过几秒的内核（在所有块上累积），这些值将环绕。
  • 选择( toc > tic ) ? (toc - tic) : ( toc + (0xffffffff - tic) ) )是必要的，以防时钟计数器回绕。

PS 这是我对这个问题的回复的副本，因为所需的时间是针对整个内核的，所以没有得到很多分数。

于 2012-06-26T14:12:10.643 回答