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 パラメーター) で、指定された数のスレッド (第 2 パラメーター) でカーネルを起動します。そのカーネルの各スレッドは、単一のスレッドで別のカーネルを起動します。この子カーネルは、出力配列 (0 で初期化された) のその部分に 1 を書き込みます。
実行の最後に、出力配列のすべての値は 1 になるはずです。しかし、奇妙なことに、一部のブロック サイズとグリッド サイズでは、配列値の一部がまだ 0 のままです。これは基本的に、子グリッドの一部が実行されないことを意味します。
これは、多数の子グリッドが同時に生成された場合にのみ発生します。私のテスト システム (Tesla K20x) では、これはそれぞれ 210 スレッドを含む 10 ブロックの場合です。ただし、200 スレッドの 10 ブロックでは正しい結果が得られます。しかし、それぞれ 1024 スレッドの 3 つのブロックでもエラーが発生します。奇妙なことに、ランタイムによってエラーが返されません。子グリッドは、スケジューラによって単に無視されるようです。
他の誰かが同じ問題に直面していますか? この動作はどこかに文書化されていますか (何も見つかりませんでした)、それとも本当にデバイス ランタイムのバグなのでしょうか?