启动大量线程和块时,CUDA printf()崩溃

Question

我正在使用CUDA 6.5 + VS2013 + GTX Titan黑色.我观察到当线程总数大于65536时,以下打印代码将崩溃.我googled了一下但没有看到任何有用的东西.有没有其他人观察到同样的行为？还是有人可以提供一些解释吗？非常感谢你!

__global__ void testKernel(int val)
{
    int X = blockDim.x * blockIdx.x + threadIdx.x;
    int Y = blockDim.y * blockIdx.y + threadIdx.y;
    printf("[%d, %d]:\t" "\tValue is:%d\n", X, Y, val);
}

void main(){

    dim3 block(16,16);
    dim3 grid(16,16);
    testKernel << <grid, block >> >(10);
    cudaDeviceSynchronize();
    cudaGetLastError();

    cudaDeviceReset();
}

当我使用块(32,16)和网格(16,16)时,我收到以下错误消息:

Gpu API调用(启动超时并终止)...

Answer 1

你的内核执行时间太长:

the launch timed out and was terminated

当在WDDM设备上运行时,这是Windows操作系统的限制.

有各种可行的解决方法.有些是:

1. 减少内核执行时间
2. 如果可能,将GPU切换到TCC模式(GeForce GPU不可用).
3. 通过Windows注册表修改延长TDR超时延迟(或删除它)

此外,内核内printf功能有很大的限制.由于各种原因,它实际上不是为大规模输出而设计的.一个特别是该活动的缓冲区是有限的,并且当溢出时,先前的缓冲区数据将丢失(即,不打印出来).

Answer 2

多亏了Robert的回答，我才意识到问题可能出在缓冲区的大小上。我使用以下代码来查找默认情况下打印缓冲区的大小为1048576字节（1M）

size_t sz;
cudaDeviceGetLimit(&sz, cudaLimitPrintfFifoSize);
std::cout << sz << std::endl;

当我使用以下代码将缓冲区大小增加到100 Mb时，错误消失，并且我具有所有预期的输出，总计131072行！（我使用block（32,16）; .. grid（16,16）; ...）

sz = 1048576 * 100;
cudaDeviceSetLimit(cudaLimitPrintfFifoSize, sz);

不知何故，打印缓冲区的溢出会导致响应时间比平常更长，并触发TDR。当我相应地增加缓冲区大小时，代码设法在超时之前完成。更重要的是，足够的缓冲区大小可确保没有数据丢失。

但是，我认为缓冲区大小和执行时间的上限取决于设备。它可以在Titan Black上很好地工作，并不一定意味着它也可以在其他NVidia卡上工作。再次，我同意Robert的观点，即使用printf从CUDA内核中导出大量数据在实践中是不可靠的。我只是用它来转储一些信息来调试内核。