如何避免连续异步内核启动导致的Cuda错误6（启动超时）？

Question

我使用这段（简化后的）代码时，遇到了Cuda错误6（也称为cudaErrorLaunchTimeout和CUDA_ERROR_LAUNCH_TIMEOUT）：

for(int i = 0; i < 650; ++i)
{
    int param = foo(i); //some CPU computation here, but no memory copy
    MyKernel<<<dimGrid, dimBlock>>>(&data, param);
}

Cuda错误6表示内核花费的时间过长。尽管单个MyKernel持续时间仅约为60毫秒，块大小为经典的16×16。现在，当我每隔50次迭代调用cudaDeviceSynchronize()时，错误就不会发生了。

for(int i = 0; i < 650; ++i)
{
    int param = foo(i); //some CPU computation here, but no memory copy
    MyKernel<<<dimGrid, dimBlock>>>(&data, param);
    if(i % 50 == 0) cudaDeviceSynchronize();
}

我希望避免这种同步方式，因为它会严重降低程序的速度。由于内核启动是异步的，我猜测错误发生是因为看门狗测量了内核从异步启动到实际开始执行所用的时间。我对Cuda还不熟悉。这种错误6是否经常出现？有没有一种方法可以避免这种错误而不影响性能？

- floriang

---

这是在Windows平台上吗？ - talonmies

是的。我听说过注册表技巧，但我不想强迫客户应用它。 - floriang

问题可能来自于Windows驱动器将GPU操作批处理在一起以分摊延迟。 - talonmies

Answer 1

感谢talonmies和Robert Crovella（他的解决方案对我不起作用），我终于找到了一个可接受的解决方法。

为了防止CUDA驱动程序将内核启动批处理在一起，必须在每个内核启动之前或之后执行另一个操作。例如，一个虚拟复制就可以解决问题:

void* dummy;
cudaMalloc(&dummy, 1);

for(int i = 0; i < 650; ++i)
{
    int param = foo(i); //some CPU computation here, but no memory copy
    cudaMemcpyAsync(dummy, dummy, 1, cudaMemcpyDeviceToDevice);
    MyKernel<<<dimGrid, dimBlock>>>(&data, param);
}

这个解决方案比那个包含对 cudaDeviceSynchronize() 调用的方案快了8秒（从50秒降至42秒）（详见问题）。此外，它更加可靠，因为 50 是一个任意的、特定于设备的时间段。

- floriang

---

它在一定程度上有所帮助，但仍然出现了cudaErrorLaunchTimeout错误。 - Fedor

Answer 2

看门狗并不是直接测量内核的执行时间。看门狗正在跟踪发送到GPU的命令队列中的请求，并确定在超时期间内是否有任何请求未被GPU确认。

正如评论中@talonmies所指出的，如果您确定没有内核执行超过超时期限，则我最好的猜测是这种行为是由CUDA驱动程序WDDM批处理机制引起的，该机制通过将GPU命令批处理在一起并以批次发送到GPU来减少平均延迟。

您无法直接控制批处理行为，因此通常情况下，尝试在不禁用或修改Windows TDR机制的情况下解决此问题将是一个不精确的练习。

一个低成本的“刷新”命令队列的一般（有些不受文档支持）建议是尝试使用cudaEventQuery(0);（如此处所建议的）代替cudaDeviceSynchronize();，也许每50个内核启动左右。在某种程度上，具体细节可能取决于机器配置和使用的GPU。

我不确定它在您的情况下会有多大效果。我认为，如果没有进行更多的实验，它不能作为避免TDR事件的“保证”。您的结果可能会有所不同。

- Robert Crovella

---

感谢您的回复。您提供的解决方案对我来说没有起作用（出现了 cudaErrorInvalidResourceHandle和其他未知错误），但是多亏了您的建议和解释，我找到了一种可行的解决方法。 - floriang