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

Question

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

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

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

这只会发生在许多子网格同时产生的情况下。在我的测试系统（特斯拉K20x）上，每个包含210个线程的10个块都是这种情况。尽管如此，有200个线程的10个块可以提供正确的结果。但是也有3个块，每个1024个线程都会导致错误。奇怪的是，运行时没有报告错误。调度器似乎忽略了子网格。

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

Answer 1

你没有做任何我能看到的error checking。您可以也应该对设备内核启动进行类似的错误检查。请参阅documentation这些错误不一定会冒泡主机：

错误是每个线程的记录，以便在每个线程都可以识别，它已经​​产生的最近的错误。

您必须将它们置于设备中。文档中有很多这类设备错误检查的例子。

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

如果扫描documentation这个错误，你会发现这一点：

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