cuda - 测量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? ，我知道原因是:

If they are non-display cards, it might well be the driver shutting itself down after a period of inactivity. So what you are seeing on the first run might well be initialization overhead that only happens once.

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

Answer 1

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

即时编译

内核转移到 GPU 内存

缓存内容

...

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

但是，如果您有特定的应用程序和用例，那么在相关情况下对该应用程序进行基准测试总是有意义的。但是，要准备好在较少控制的测量中应对运行时间中更大的变化。