cuda - 如何避免连续异步内核启动时出现 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 毫秒。 block 大小是经典的 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

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