对 CUDA 内核中的不同部分进行计时答案

【问题标题】：Timing different sections in CUDA kernel对 CUDA 内核中的不同部分进行计时
【发布时间】：2026-01-19 13:25:01
【问题描述】：

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

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

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

【问题讨论】：

【解决方案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 的回复的副本，因为所需的时间是针对整个内核的，所以没有得到多少分。

【讨论】：

谢谢。很有用。查找clock()，我发现还有一个clock64()，可能不需要溢出检查和转换为float。
@RogerDahl：感谢您指出这一点！它似乎是在 CUDA 4.2 中添加的。
Fermi 添加了 64 位时钟结果。 Clock64 是在 CUDA 4.2 之前添加的。请注意，在进行这种类型的计时时，您必须小心发散——如果不同的 warp 在您的计时中采用不同的路径，那么只有线程 0 的计时将不准确。
除此之外，还要确保反汇编编译器输出并确保没有发生重新排序。编译器和汇编器（至少是较旧的 open64 工具链）可以并且确实可以移动代码，这可能意味着时钟调用可能会变成另一个，而不是包含您想要的代码。
@harrism：我对此有点不准确。函数 clock64() 仅出现在 4.2 版的 CUDA 编程指南中。至于你的第一点，我已经相应地更新了我的答案。谢谢！