某些子網格未使用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