如何获取CPU时间线中的Cuda事件时间？

Question

这是伪代码：

// the code before also calls kernels

for(int i{0}; i < n; ++i){
    auto& startEvent = startEvents[i];
    auto& stopEvent = stopEvents[i];

    cudaEventRecord(startEvent);
    kernel<<<1000, dim3{32, 32, 1}>>>();
    cudaEventRecord(stopEvent);
}

cudaDeviceSynchronize();

// get elapsed times

代码计算

n

内核启动所用的时间。但问题是它计算相对时间。我想将每个

start

和

stop

事件转换为 CPU 时间线。换句话说，我想将事件转换为自纪元以来的时间（或自程序启动以来的时间）。

在循环之前保存当前的 CPU 时间不会有帮助 - 由于 GPU 工作负载很重，我不知道循环中的第一个内核何时真正启动（假设之前有很多其他内核调用），甚至如果之前没有任何工作，内核启动本身就会产生开销。

要尝试的事情之一是添加回调

cudaLaunchHostFunc

。但正如文档所说，该函数将阻止流中的后续工作。由于我每帧有数百个内核，因此此回调可能会给程序增加明显的开销。我希望有类似

cudaLaunchHostFuncAsync

的东西。

NSight

工具如何获取内核启动和停止的精确时间戳？

Answer 1

正如 @RobertCrovella 所建议的，一种方法是记录事件 (

gpuTimeline

) 并同时拍摄 CPU 时间快照 (

cpuTimeline

)。之后，可以将另一个事件与

gpuTimeline

之间的差异添加到

cpuTimeline

中以获得 GPU 时间。

我是这样做的：

static uint64_t t1;
static uint64_t t2;

cudaLaunchHostFunc(0, [](void*){ t1 = getCurrentTime(); }, nullptr);
cudaEventRecord(gpuTimeline);
cudaLaunchHostFunc(0, [](void*){ t2 = getCurrentTime(); }, nullptr);
cudaDeviceSynchronize();
cpuTimeline = t1 + (t2 - t1) / 2;

这里我在事件记录之前和之后立即设置回调。我假设事件的实时发生在这两个回调之间。它非常适合我的需求。