在测量CUDA内核运行时间时,需要进行热身代码



在第85页,专业CUDA C编程:

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卡上加热?,我知道原因是:

如果它们是非播放卡,则很可能是驾驶员在一段时间后就关闭的。因此,您在第一次运行中看到的很可能是初始化开销,只有一次。

因此,如果我的GPU卡很长一段时间都没有活动,例如,我只是使用它来运行一些程序,则不需要运行任何热身代码。我的理解对吗?

除了GPU处于节能状态外,还有许多其他原因是,首次启动内核可以比进一步运行更慢:

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

由于这些原因,如果您对连续内核启动实现的持续速度感兴趣,最好在定时内核之前进行至少一次"热身跑"。

但是,如果您有一个特定的应用程序和用例,则在相关情况下对应用程序进行基准测试总是有意义的。在那个较不受控制的测量中以更大的运行时间变化做好准备。

相关内容

最新更新