CudaEvent内核启动时间

Pont 发表于 Dev

蓬

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

不幸的是，我不再可以使用与Cuda兼容的GPU进行任何测试。

非常感谢

罗伯特·克罗维拉

根据我的实验，它肯定会捕获某种开销。

我认为应该很清楚，如果内核由其他同步活动括起来，则必须包括CPU开销，因为CPU开销将先前的活动与内核启动分开了：

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

在我看来，显然上述时序必须捕获cudaMemcpy 1和内核调用之间的CPU开销（以及其他各种时间贡献，而不是内核本身）。

因此，不太明显的情况是内核是单独的或由其他异步调用括起来的：

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包括cudaEventSynchronize()紧随stop标记之后的典型时间安排中，很明显主机线程将至少在cudaEvent系统start和stop标记之间被阻塞的时间内被阻塞，因此我认为在这种情况下没有理由不对其进行引用作为CPU开销。