测量 CUDA 内核运行时间时是否需要预热代码？

Question

在第85页，professional CUDA C programming：

int main()
{
    ......
    // run a warmup kernel to remove overhead
    size_t iStart,iElaps;
    cudaDeviceSynchronize();
    iStart = seconds();
    warmingup<<<grid, block>>> (d_C);
    cudaDeviceSynchronize();
    iElaps = seconds() - iStart;
    printf("warmup <<< %4d %4d >>> elapsed %d sec \n",grid.x,block.x, iElaps );

    // run kernel 1
    iStart = seconds();
    mathKernel1<<<grid, block>>>(d_C);
    cudaDeviceSynchronize();
    iElaps = seconds() - iStart;
    printf("mathKernel1 <<< %4d %4d >>> elapsed %d sec \n",grid.x,block.x,iElaps );

    // run kernel 3
    iStart = seconds();
    mathKernel2<<<grid, block>>>(d_C);
    cudaDeviceSynchronize();
    iElaps = seconds () - iStart;
    printf("mathKernel2 <<< %4d %4d >>> elapsed %d sec \n",grid.x,block.x,iElaps );

    // run kernel 3
    iStart = seconds ();
    mathKernel3<<<grid, block>>>(d_C);
    cudaDeviceSynchronize();
    iElaps = seconds () - iStart;
    printf("mathKernel3 <<< %4d %4d >>> elapsed %d sec \n",grid.x,block.x,iElaps);
    ......
}

我们可以看到在测量不同内核的运行时间之前有一个预热。

来自GPU cards warming up?，我知道原因是：

如果它们是非显示卡，则很可能是驱动程序在一段时间不活动后自行关闭。因此，您在第一次运行时看到的很可能是只发生一次的初始化开销。

所以如果我的 GPU 卡没有长时间处于非活动状态，例如，我只是用它来运行一些程序，它应该不需要运行任何预热代码。我的理解对吗？

Answer 1

除了 GPU 处于省电状态之外，还有许多其他原因导致内核的首次启动可能比进一步运行慢：

• 即时编译
• 将内核转移到 GPU 内存
• 缓存内容
• ...

由于这些原因，如果您对连续启动内核所达到的持续速度感兴趣，那么在定时内核运行之前至少执行一次“预热运行”总是一种好习惯。

但是，如果您有一个特定的应用程序和用例，那么在相关情况下对该应用程序进行基准测试总是有意义的。不过，请准备好在这种不太受控的测量中出现更大的运行时间变化。