我了解CUDA的“热身”现象——首次Kernel运行有时比后续慢，但我的情况恰好相反：前几次运行速度很快，之后逐渐变慢，直至稳定在一个区间。启动配置与输入始终一致，Kernel运行时间如下：

1. 0.74ms；2. 0.94ms；3. 0.85ms；4. 3.39ms；5. 5.41ms；6. 5.76ms；7. 5.62ms；8. 5.61ms……后续运行均稳定在5.5-5.8ms区间。

代码结构如下：

for loop {
  cudaEventCreate(&start);
  cudaEventCreate(&stop);
  cudaEventRecord(start, 0);

  for (int channel = 0; channel < kNumChannels; channel++) {
    kernel_call_1(input_d, output_d);
    kernel_call_2(output_d, output_d); // In-place FFT
  }
  cudaDeviceSynchronize();

  cudaEventRecord(stop, 0);
  cudaEventSynchronize(stop);

  float elapsed_time;
  cudaEventElapsedTime(&elapsed_time, start, stop);
  cudaEventDestroy(start);
  cudaEventDestroy(stop);

  output_h = output_d;
  Verify_output();
}

原本没有cudaDeviceSynchronize()调用，因默认流认为无需添加，后参考NVIDIA文档添加但无改善，请问该现象原因是什么？

核心原因分析 #

这种先快后慢再稳定的现象，主要和GPU内存状态、硬件调度策略以及代码中的数据拷贝和验证环节有关，具体拆解如下：

• 初始内存缓存优势：前几次运行时，input_d等设备内存的数据可能还停留在GPU的L1/L2缓存或显存高速区域，Kernel访问延迟极低，因此运行速度快。但随着循环次数增加，显存中积累的中间数据、反复写入的output_d会降低缓存命中率，同时显存带宽竞争加剧，导致Kernel运行时间上升，直到内存状态达到动态平衡，时间趋于稳定。

• GPU动态频率调整（DVFS）延迟生效：GPU的动态电压频率调整机制不会立刻拉满频率。初始运行时，GPU可能处于低功耗/低频状态，此时Kernel计算量不大，反而因频率低、调度延迟小显得运行快；连续执行多次Kernel后，GPU检测到持续负载，会提升到额定频率，此时虽然计算能力更强，但频率切换的调度开销、高频率下的内存访问延迟变化，会让运行时间短暂上升，最终稳定在额定频率下的真实运行时间。

• 数据拷贝与验证的隐式干扰：每次循环末尾的output_h = output_d（实际应为cudaMemcpy类操作）和Verify_output()会占用CPU-GPU总线带宽，CPU的验证操作还可能导致GPU后续循环的调度延迟。前几次循环时，这类同步开销未显现，但多次循环后总线资源被持续占用，GPU Kernel执行不得不等待数据传输间隙，导致运行时间变长，直到总线使用达到稳定状态。

• 事件创建销毁的额外开销：每次循环重复创建销毁cudaEvent本身会带来开销，前几次运行时GPU其他操作延迟低，这部分开销被掩盖；当GPU进入稳定负载后，事件开销与Kernel运行时间叠加，使得整体时间上升并稳定。

验证与解决建议 #

• 调整事件生命周期：将cudaEventCreate和cudaEventDestroy移到外层循环外，避免重复创建销毁的开销，更准确测量Kernel运行时间。
• 增加预运行步骤：在正式计时循环前先执行1-2次空循环，让GPU内存状态和频率达到稳定，再开始计时。
• 监控显存状态：用nvidia-smi或CUDA Profiler查看显存占用和带宽使用情况，确认是否存在显存不足或带宽瓶颈。
• 优化验证环节：如果Verify_output()无需每次循环执行，可将其移到整个循环末尾，减少对GPU运行的干扰。

内容的提问来源于stack exchange，提问作者Jim Clay