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

Question

我使用的是 CUDA 6.5 + VS2013 + GTX Titan black。我观察到当线程总数大于 65536 时，以下打印代码会崩溃。我搜索了一下，但没有看到任何有用的东西。有没有其他人观察到相同的行为？或者任何人都可以提供一些解释吗？非常感谢！

__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();
}

我在使用 block(32,16) 和 grid(16,16) 时收到以下错误消息：

Gpu API 调用（启动超时并被终止）...

Answer 1

您的内核执行时间过长：

the launch timed out and was terminated

这是a limitation of the windows operating system, when running on WDDM devices。

有多种可能的解决方法。有些是：

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

此外，内核中的 printf 功能具有 significant limits。由于各种原因，它确实不是为大规模输出而设计的。特别是这个活动的缓冲区是有限的，当溢出时，之前的缓冲区数据将丢失（即不打印出来）。

Answer 2

感谢罗伯特的回答，我意识到问题可能是由于缓冲区的大小。我用下面的代码发现默认情况下打印缓冲区的大小是1048576字节（1M）

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

当我使用以下代码将缓冲区大小增加到 100 Mb 时，错误消失了，我得到了所有预期的输出，总共 131072 行！ （我使用块（32,16）；..网格（16,16）；...）

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

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

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