CudaEvent内核启动时间

Posted 2023-03-06

【中文标题】CudaEvent内核启动时间

【英文标题】：CudaEvent timing of Kernel launch

【发布时间】：2013-10-15 18:48:36

【问题描述】：

我正在尝试解释我使用 CudaEvents 所做的一些计时，通过 CudaEvents 计时内核执行是否还包括内核启动的开销时间？

很遗憾，我无法再使用与 Cuda 兼容的 GPU 进行任何测试。

非常感谢

【参考方案1】：

根据我的实验，它肯定会捕获一些类型的开销。

我认为应该清楚的是，如果内核被其他同步活动包围，那么 CPU 开销必须包括在内，因为 CPU 开销将先前的活动与内核启动分开：

cudaEventRecord(start);
cudaMemcpy(...);         // cudaMemcpy 1
my_kernel<<<...>>>(...);
cudaMemcpy(...);
cudaEventRecord(stop);

当然，在我看来，上面描述的时间很明显必须捕获 cudaMemcpy 1 和内核调用之间的 CPU 开销（以及不是由于内核本身造成的各种其他时间贡献。） p>

所以不太明显的情况是内核单独或被其他异步调用括起来时：

cudaEventRecord(start);
my_kernel<<<...>>>(...);
cudaEventRecord(stop);
cudaEventSynchronize(stop);

基于我对上述模式的测试，并使用空内核：

__global__ void my_kernel()
  

我在 linux 上观察到至少几微秒的时间，这比在我的空内核中执行两条指令所需的时间要长很多：

            Function : _Z8mykernelv
    /*0000*/     /*0x00005de428004404*/     MOV R1, c [0x1] [0x100];
    /*0008*/     /*0x00001de780000000*/     EXIT;

因此，我声称cudaEvent 系统正在捕获某种执行设置开销。如果有人想声称这种开销不是 CPU 开销，而是其他东西，那就这样吧。

我的主张是捕获了某种形式的开销，我认为没有理由不将其称为 CPU 开销。此外，典型的cudaEvent 计时包括紧跟在stop 标记之后的cudaEventSynchronize()，很明显主机线程将至少在cudaEvent 系统在start 和stop 标记，因此我认为在这种情况下没有理由不将其称为 CPU 开销。

【讨论】：

“开销捕获”是否取决于您是否使用默认流？