cudaEventElapsedTime和nvprof运行时间

Question

考虑以下内核

__global__ void elementAccess( int *a, int N )
{
  int i = blockDim.x * blockIdx.x + threadIdx.x;
  int z = a[ i ];
}

被称为

cudaEvent_t start, stop;
float elapsedTime;
cudaEventCreate(&start);
cudaEventRecord(start,0);

elementAccess<<< 1, 1>>>( d_A, 1 );

cudaEventCreate(&stop);
cudaEventRecord(stop,0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapsedTime, start,stop);

printf("Elapsed time : %f ms\n" ,elapsedTime);

其中d_A是一个整数数组。

当我在终端运行代码时，会看到

Elapsed time : 0.015328 ms

当我使用nvprof --metrics SOME_METRICS -o e.nvvp ./element_access对其进行分析时，我发现核函数的持续时间为2.95微秒。

那么，为什么时钟事件显示核运行时间为15us，而nvvp显示为2.95us呢？

它们是两个不同的东西吗？还是有其他东西缺失了？

- mahmood

---

cudaEventCreate 可能会影响到此处。您不应该在计时区域内执行此操作。在执行任何 cudaEventRecord 操作之前，请完成所有事件的创建。 - Robert Crovella

实际上没有影响。我的意思是 cudaEventRecord(start,0); /* kernel */; cudaEventRecord(stop,0); 大约需要相同的时间，大约15微秒。 - mahmood

Answer 1

你的第一次测量（基于经过的时间）包括内核启动开销。

第二次测量（基于CUDA事件）主要排除了启动开销。

考虑到你的内核完全没有做任何事情（单个内存加载会因缺乏后续使用而被优化掉），因此经过的时间几乎全部由启动开销构成是合理的。

- tera

---

我看到了类似的测量访问延迟的方法。例如，我想要测量内存延迟。这意味着内核启动开销始终存在。 - mahmood