某些子网格未使用 CUDA 动态并行执行

我正在尝试 CUDA 5.0 (GTK 110) 中的新动态并行功能。我遇到了一个奇怪的行为，即我的程序没有返回某些配置的预期结果，不仅是意外的，而且每次启动都会出现不同的结果。

现在我想我找到了问题的根源：似乎当生成太多子网格时，某些子网格（由其他内核启动的内核）有时不会执行同时.

我编写了一个小测试程序来说明这种行为：

#include <stdio.h>

__global__ void out_kernel(char* d_out, int index)
{
    d_out[index] = 1;
}

__global__ void kernel(char* d_out)
{
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    out_kernel<<<1, 1>>>(d_out, index);
}

int main(int argc, char** argv) {

    int griddim = 10, blockdim = 210;
    // optional: read griddim and blockdim from command line
    if(argc > 1) griddim = atoi(argv[1]);
    if(argc > 2) blockdim = atoi(argv[2]);

    const int numLaunches = griddim * blockdim;
    const int memsize = numLaunches * sizeof(char);

    // allocate device memory, set to 0
    char* d_out; cudaMalloc(&d_out, memsize);
    cudaMemset(d_out, 0, memsize);

    // launch outer kernel
    kernel<<<griddim, blockdim>>>(d_out);
    cudaDeviceSynchronize();

    // dowload results
    char* h_out = new char[numLaunches];
    cudaMemcpy(h_out, d_out, memsize, cudaMemcpyDeviceToHost);

    // check results, reduce output to 10 errors
    int maxErrors = 10;
    for (int i = 0; i < numLaunches; ++i) {
        if (h_out[i] != 1) {
            printf("Value at index %d is %d, should be 1.\n", i, h_out[i]);
            if(maxErrors-- == 0) break;
        }
    }

    // clean up
    delete[] h_out;
    cudaFree(d_out);
    cudaDeviceReset();
    return maxErrors < 10 ? 1 : 0;
}

该程序在给定数量的块（第一个参数）中启动一个内核，每个块具有给定数量的线程（第二个参数）。然后，该内核中的每个线程将启动另一个具有单个线程的内核。该子内核将在输出数组（用 0 初始化）的其部分写入 1。

执行结束时，输出数组中的所有值都应为 1。但奇怪的是，对于某些块和网格大小，某些数组值仍然为零。这基本上意味着一些子网格没有被执行。

仅当同时生成许多子网格时才会发生这种情况。在我的测试系统（Tesla K20x）上，有 10 个块，每个块包含 210 个线程。不过，10 个具有 200 个线程的块可以提供正确的结果。但也有 3 个块（每个块有 1024 个线程）会导致错误。 奇怪的是，运行时没有报告任何错误。子网格似乎被调度程序忽略了。

其他人也面临同样的问题吗？此行为是否记录在某处（我没有找到任何内容），或者它确实是设备运行时中的错误？

你没有做错误检查 https://stackoverflow.com/questions/14038589/what-is-the-canonical-way-to-check-for-errors-using-the-cuda-runtime-api我能看到的任何类型。您可以而且应该在设备内核启动时执行类似的错误检查。请参阅文档 http://docs.nvidia.com/cuda/pdf/CUDA_Dynamic_Parallelism_Programming_Guide.pdf这些错误不一定会冒泡到主机：

每个线程都会记录错误，以便每个线程都可以识别它生成的最新错误。

您必须将它们困在设备中。文档中有很多此类设备错误检查的示例。

如果您进行适当的错误检查，您会发现在内核启动失败的每种情况下，cuda 设备运行时 API 都会返回错误 69，cudaErrorLaunchPendingCountExceeded.

如果您扫描文档 http://docs.nvidia.com/cuda/pdf/CUDA_Dynamic_Parallelism_Programming_Guide.pdf对于这个错误，你会发现：

cudaLimitDevRuntimePendingLaunchCount

控制为缓冲由于未解决的依赖性或缺乏执行资源而尚未开始执行的内核启动预留的内存量。当缓冲区已满时，launchs 会将线程的最后一个错误设置为cudaErrorLaunchPendingCountExceeded。默认待处理启动次数为 2048 次启动。

在 10 个块 * 200 个线程中，您将启动 2000 个内核，并且一切似乎都正常。

在 10 个块 * 210 个线程中，您将启动 2100 个内核，这超出了上面提到的 2048 个限制。

请注意，这本质上是动态的；根据您的应用程序启动子内核的方式，您可以轻松启动超过 2048 个内核，而不会达到此限制。但由于您的应用程序几乎同时启动所有内核，因此您已经达到了极限。

当您的 CUDA 代码未按您预期的方式运行时，建议进行适当的 cuda 错误检查。

如果您想对上述内容进行一些确认，您可以在代码中修改主内核，如下所示：

__global__ void kernel(char* d_out)
{
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    out_kernel<<<1, 1>>>(d_out, index);
//    cudaDeviceSynchronize();  // not necessary since error 69 is returned immediately
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) d_out[index] = (char)err;
}

待处理的启动计数限制是可以修改的。请参阅文档cudaLimitDevRuntimePendingLaunchCount