如何避免连续异步内核启动时出现 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 表示内核花费了太多时间 return。不过，单个 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 的常见情况吗？有没有办法在不改变性能的情况下避免这个错误？

Answer 1

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

正如@talonmies 在评论中指出的那样，我最好的猜测是（如果您确定没有内核执行超过超时期限）此行为是由于 CUDA 驱动程序 WDDM 批处理机制造成的，该机制旨在减少平均延迟通过将 GPU 命令批处理在一起并分批发送到 GPU。

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

关于命令队列的低成本 "flush" 的一般（有点未记录）建议是使用 cudaEventQuery(0);（建议 here ) 代替 cudaDeviceSynchronize();，也许每 50 个内核启动一次左右。在某种程度上，细节可能取决于机器配置和使用的 GPU。

我不确定它对你的情况会有多大影响。我不认为它可以作为避免 TDR 事件而不进行更多实验的 "guarantee" 来推进。您的里程可能会有所不同。

Answer 2

感谢 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 是一个任意的、特定于设备的周期。

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

How to avoid Cuda error 6 (Launch Timeout) with consecutive asynchronous kernel launches?

cuda

timeout