3. CUDA事件与计时

3.1 为什么需要CUDA事件计时

在CUDA性能分析中，测量内核执行时间是最基础也最关键的一步。很多初学者习惯用 clock() 或 gettimeofday() 等CPU计时函数来测量GPU内核耗时，但这在异步执行模型下会产生严重误差。CUDA内核启动后，CPU并不会等待GPU执行完毕，而是立即返回继续执行后续代码。因此，CPU侧计时器记录到的往往只是“内核启动指令发出”的时间，而非GPU真正完成计算的时间。

CUDA事件（Event）是CUDA运行时提供的一种轻量级时间标记机制，它由GPU硬件直接记录时间戳，能够精确测量设备端操作的耗时。其核心优势在于：

• 设备端计时：时间戳由GPU记录，不受CPU调度、中断或负载影响。
• 流内顺序保证：同一流中的事件按提交顺序执行，可精确标记内核开始与结束。
• 支持异步测量：事件本身不阻塞CPU，可配合流实现流水线计时。
• 微秒级精度：在大多数现代GPU上，事件时间戳分辨率可达约0.5微秒。

3.2 CUDA事件API与计时流程

使用CUDA事件计时通常遵循“创建→记录→同步→计算”四个步骤。下面是一个标准的内核计时代码模板：

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

// 在流中记录开始事件
cudaEventRecord(start, stream);

// 启动要计时的内核（异步）
myKernel<<<grid, block, 0, stream>>>(args...);

// 在流中记录结束事件
cudaEventRecord(stop, stream);

// 同步等待stop事件完成
cudaEventSynchronize(stop);

// 计算耗时（毫秒）
float ms = 0.0f;
cudaEventElapsedTime(&ms, start, stop);
printf("Kernel time: %.3f ms\n", ms);

// 清理资源
cudaEventDestroy(start);
cudaEventDestroy(stop);

这段代码的关键在于 cudaEventSynchronize(stop)：它阻塞CPU直到GPU执行完stop事件，确保此时内核已经完成。如果不调用同步，cudaEventElapsedTime 可能读取到尚未更新的时间戳，导致结果不准确。

下表总结了CUDA事件计时的核心API及其作用：

3.3 同步与异步：理解CUDA执行模型

要正确使用CUDA事件，必须深入理解同步与异步的语义。CUDA的异步执行体现在两个层面：

1. CPU与GPU异步：内核启动后CPU立即返回，不等待GPU完成。这是CUDA默认行为，也是高吞吐量的基础。
2. 流内操作顺序执行：同一流中的操作（内核、拷贝、事件）按提交顺序执行，但不同流之间可以并行。

基于此，CUDA事件计时存在两种典型模式：

模式一：阻塞式计时（同步）——适用于单内核精确测量。使用 cudaEventSynchronize 强制CPU等待，确保时间戳有效。优点是简单可靠，缺点是CPU在等待期间无法执行其他工作。

模式二：异步计时（流水线）——适用于多流并发场景。不立即同步，而是先记录事件，待所有工作提交后再统一同步。例如：

for (int i = 0; i < N; i++) {
    cudaEventRecord(start[i], stream[i]);
    kernel<<<grid, block, 0, stream[i]>>>(...);
    cudaEventRecord(stop[i], stream[i]);
}
// 所有流提交后，统一等待
cudaEventSynchronize(stop[N-1]);
for (int i = 0; i < N; i++) {
    cudaEventElapsedTime(&ms[i], start[i], stop[i]);
}

这种模式允许GPU在CPU提交后续内核的同时执行前面的内核，从而测量真实的重叠执行时间。

3.4 案例：测量多流并发内核时间

下面通过一个实际案例展示如何使用CUDA事件测量两个流中内核的重叠执行时间。假设有两个独立的内核 kernelA 和 kernelB，分别放入流0和流1：

cudaStream_t s0, s1;
cudaStreamCreate(&s0);
cudaStreamCreate(&s1);

cudaEvent_t start0, stop0, start1, stop1;
cudaEventCreate(&start0); cudaEventCreate(&stop0);
cudaEventCreate(&start1); cudaEventCreate(&stop1);

// 流0：记录开始→执行内核A→记录结束
cudaEventRecord(start0, s0);
kernelA<<<grid, block, 0, s0>>>(...);
cudaEventRecord(stop0, s0);

// 流1：记录开始→执行内核B→记录结束
cudaEventRecord(start1, s1);
kernelB<<<grid, block, 0, s1>>>(...);
cudaEventRecord(stop1, s1);

// 等待两个流都完成
cudaEventSynchronize(stop0);
cudaEventSynchronize(stop1);

float ms0, ms1;
cudaEventElapsedTime(&ms0, start0, stop0);
cudaEventElapsedTime(&ms1, start1, stop1);
printf("Stream0 kernel: %.3f ms\n", ms0);
printf("Stream1 kernel: %.3f ms\n", ms1);

如果GPU支持并发执行（如具有多个SM或支持Hyper-Q），两个内核可能部分重叠，此时 ms0 + ms1 会大于实际墙钟时间。通过对比串行执行与并发执行的总耗时，可以评估并行效率。

下图展示了CUDA事件在流中记录时间戳的时序关系：

图中可见，Stream 0 和 Stream 1 的内核在时间上存在重叠区域。CUDA事件分别记录了每个流的开始和结束时间点，从而可以精确计算每个内核的独立耗时以及重叠程度。

3.5 常见误区与最佳实践

3.6 复盘：一个典型的计时错误案例

假设某开发者编写了如下代码来测量内核时间：

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start);
myKernel<<<grid, block>>>(...);
cudaEventRecord(stop);

// 忘记同步！
float ms;
cudaEventElapsedTime(&ms, start, stop);
printf("Time: %.3f ms\n", ms);

这段代码存在两个问题：

1. 未调用 cudaEventSynchronize(stop)，导致 cudaEventElapsedTime 可能在stop事件尚未被GPU处理时执行，读取到无效时间戳。
2. 未指定流，事件记录在默认流（流0）中，如果内核也在默认流中，顺序尚可保证；但若使用非默认流，则事件与内核可能不在同一流中，无法正确测量。

修正后的代码应显式指定流并同步：

cudaStream_t stream;
cudaStreamCreate(&stream);

cudaEventRecord(start, stream);
myKernel<<<grid, block, 0, stream>>>(...);
cudaEventRecord(stop, stream);
cudaEventSynchronize(stop);

float ms;
cudaEventElapsedTime(&ms, start, stop);
printf("Time: %.3f ms\n", ms);

通过这个案例可以看出，CUDA事件计时虽然API简单，但正确使用需要深刻理解异步执行模型。掌握同步与异步的平衡，是获得准确性能数据的关键。