如何避免连续异步内核启动时出现 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();
}

我想避免这种同步，因为它会大大减慢程序的速度。

由于内核启动

是异步的，我想发生错误是因为看门狗从异步启动而不是从其执行的实际开始测量内核的执行持续时间。

我是库达的新手。这是发生错误 6 的常见情况吗？有没有办法在不改变性能的情况下避免此错误？

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

Answer 2

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

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

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

对于

命令队列的低成本"刷新"的一般（有些未记录的）建议，您可以尝试尝试一下，是使用 cudaEventQuery(0);（如此处建议的）代替cudaDeviceSynchronize();，也许每 50 个内核启动一次左右。 在某种程度上，细节可能取决于计算机配置和正在使用的 GPU。

我不确定它在你的情况下会有多有效。 我不认为它可以作为避免TDR事件的"保证"，而无需进行更多的实验。您的里程可能会有所不同。