使用 CUDA 事件(cudaEvent_t)来测量 GPU 代码执行时间

code

cudaEvent_t start, stop; // 声明两个 CUDA 事件变量 start 和 stop，分别用于记录代码执行的开始和结束时间。
float elapsedTime; // 声明一个浮点数变量 elapsedTime，用于存储计算得到的时间差。

// 调用 cudaEventCreate 函数分别创建 start 和 stop 事件
cudaEventCreate(&start);
cudaEventCreate(&stop);

// 使用 cudaEventRecord 记录当前时间点，将其存储在 start 事件中。此时开始记录执行时间。
cudaEventRecord(start);
// 内核调用或其他操作
// 执行 CUDA 内核函数 myKernel，将其放在两个事件之间，表示要测量其执行时间的代码段。
myKernel<<<grid, block>>>();
// 使用 cudaEventRecord 记录结束事件 stop，表示内核函数执行完成的时间点。
cudaEventRecord(stop);

// 同步，确保stop事件完成 
// 使用 cudaEventSynchronize 确保所有与 stop 事件之前的 CUDA 操作已完成。这一步是必要的，因为 CUDA 操作是异步的，主机代码执行到这行时 GPU 上的内核可能尚未完成。
// cudaEventSynchronize(stop) 会阻塞主机代码，直到 stop 事件完成为止，确保我们在调用 cudaEventElapsedTime 时可以获得准确的时间间隔。
cudaEventSynchronize(stop);

// 计算 start 和 stop 事件之间的时间差，并将结果存储在 elapsedTime 中，单位为毫秒。
cudaEventElapsedTime(&elapsedTime, start, stop);
printf("Elapsed time: %f ms\n", elapsedTime);

// 销毁 start 和 stop 事件，释放与它们相关的资源，确保资源不被占用。
cudaEventDestroy(start);
cudaEventDestroy(stop);

最开始我没有加 cudaEventSynchronize(stop); 这一句，会遇到报错：cuda device not ready。

问题关键点在于：
在调用 cudaEventElapsedTime 之前，需要确保两个事件都已完成。

在 cudaEventRecord 后调用 cudaEventSynchronize 来同步两个事件，使其在时间测量之前完全完成。

为什么start不需要cudaEventSynchronize？

1. 起始事件不需要等待其他操作完成：
   • 当你在代码中记录起始事件（如 cudaEventRecord(start);）时，它只是标记当前的时间点，目的是为后续的时间测量提供起点。
   • 没有必要等待这个起始事件完成，因为它不依赖于其他操作的完成时间。即便是异步的操作，事件的记录也不需要同步。
2. 同步通常用于结束事件或依赖事件：
   • cudaEventSynchronize 的主要作用是确保设备已经完成了指定的事件或在某个时间点之前的所有操作。因此，对于结束事件（如 stop），或者希望等待设备完成的事件才需要调用 cudaEventSynchronize。
   • 在时间测量中，结束事件（如 stop）标记了需要精确测量的时间段的结束点。在 stop 事件完成前，调用 cudaEventElapsedTime 无法得到准确的时间值，所以必须确保 stop 完成。
3. CUDA 事件的记录是异步的：
   • cudaEventRecord 是异步的，不会阻塞主机代码。如果调用 cudaEventSynchronize(start);，它会使主机等待此事件，但实际上这是不必要的。
   • cudaEventRecord(start); 记录一个时间点的操作非常轻量，不涉及任何计算或设备操作，通常不会被阻塞，因此不需要同步。

因此，只需在结束事件（stop）上使用 cudaEventSynchronize，确保 CUDA 设备已经完成该事件之前的所有操作，然后调用 cudaEventElapsedTime 获取准确的时间间隔。