内核中的 CUDA 分析答案

【问题标题】：CUDA profiling inside kernel内核中的 CUDA 分析
【发布时间】：2013-05-30 11:17:23
【问题描述】：

是否有任何选项来分析 CUDA 内核？不是整体，而是一部分。我有一些设备功能调用，我想测量它们的时间。我是否可以设置任何标志/事件/指令，然后它将在 NVIDIA Visual Profiler 中可见？还是我需要通过插入cudaEventCreate和类似的函数来手动完成。

【问题讨论】：

标签： cuda nvvp

【解决方案1】：

您可以使用 clock() 或 clock64() 函数手动为内核的特定部分计时：

unsigned long long* time_spent;

__global__ void kernel(...)
{
    unsigned int t1, t2;
    // ...
    t1 = clock();
    // code of interest
    t2 = clock();
    atomicAdd(&time_spent, t2 - t1);
}

'clock()` 正式返回一个clock_t，但我更喜欢显式使用 unsigned int 以明确上述代码如何正确处理时钟值的环绕（只要定时代码不超过 2^32 -1 个周期完成。

确保也用

计时代码

    t1 = clock();
    t2 = clock();

背靠背，这样您就可以减去时间开销。

【讨论】：

clock()不是每个SM都返回一个本地计数器吗？
沉迷于这个答案：docs.nvidia.com/cuda/cuda-c-programming-guide/…
是的。但这并不重要，只要我们只从同一个 SM 获取时间之间的差异。（动态并行确实会在这里产生问题，为简单起见，我只是假设感兴趣的代码不会启动任何其他内核）。