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

Some child grids not being executed with CUDA Dynamic Parallelism

本文关键字：动态 并行性 执行 CUDA 未使用 网格      更新时间：2023-10-16

我正在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;
}

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

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

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

其他人也面临同样的问题吗？这种行为是在某个地方记录的（我没有发现任何东西），还是它真的是设备运行时的一个错误？

在我看来，您没有进行任何类型的错误检查。您可以也应该在设备内核启动时进行类似的错误检查。请参阅文档。这些错误不一定会出现在主机上：

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

你必须把他们困在设备里。文档中有大量此类设备错误检查的示例。

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

如果你扫描文档以查找此错误，你会发现：

cudaLimitDevRuntimePendingLaunchCount

控制为缓冲尚未开始执行的内核启动而预留的内存量，这些内核启动可能是由于未解决的依赖关系或缺乏执行资源。当缓冲区已满时，启动会将线程的最后一个错误设置为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 的文档

相关文章：