Некоторые дочерние сетки не выполняются с помощью динамического параллелизма 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;
}

Программа запускает ядро ​​с заданным количеством блоков (1-й параметр) с заданным количеством потоков в каждом (2-й параметр). Каждый поток в этом ядре запускает другое ядро ​​с одним потоком. Это дочернее ядро ​​запишет 1 в своей части выходного массива (который был инициализирован с 0).

В конце выполнения все значения в выходном массиве должны быть равны 1. Но странно, что для некоторых размеров блоков и сеток некоторые значения массива по-прежнему равны нулю. Это в основном означает, что некоторые из дочерних сеток не выполняются.

Это происходит только в том случае, если многие из дочерних сеток создаются одновременно. В моей тестовой системе (Tesla K20x) это случай для 10 блоков, содержащих по 210 потоков в каждом. Однако 10 блоков с 200 потоками дают правильный результат. Но также 3 блока с 1024 потоками каждый вызывают ошибку.
Как ни странно, об ошибках не сообщается во время выполнения. Планировщик просто игнорирует дочерние сетки.

Кто-нибудь еще сталкивается с такой же проблемой? Это где-то задокументировано (я ничего не нашел), или это действительно ошибка во время работы устройства?