有时不调用 CUDA 回调

Question

我正在编写一些 CUDA 代码。该代码是一个模拟，因此它必须运行多次迭代，并且每次迭代都取决于邻居的结果。由于数据很多，我决定使用流和平铺。

这是代码的简化方案:

sync = (int *)malloc(tiles * tiles * tiles * sizeof(*sync));
memset(sync, 0, tiles * tiles * tiles * sizeof(*sync));

// At the moment tiles = 4
for (i = 0; i < tiles * tiles * tiles; ++i) {
    cudaStreamCreate(&data[i].stream);
    data[i].sync = sync;
    data[i].tiles = tiles;
    data[i].x = i / (tiles * tiles);
    data[i].y = (i / tiles) % tiles;
    data[i].z = i % tiles;

    kernel<<<grid_size, block_size, 0, data[i].stream>>>(/* parameters */);

    cudaStreamAddCallback(data[i].stream, cudaCallback, &data[i], 0);
}

// Synchronization and respawn (now trying just 1 iteration, so no respawn)
for (i = 0; i < tiles * tiles * tiles; ++i) {
    printf("Waiting %d\n", i);
    while (sync[i] != iters) { __sync_synchronize(); }
}

回调:

void CUDART_CB cudaCallback(cudaStream_t stream, cudaError_t status, void *data)
{
    struct lifeStreamData *streamData = (struct lifeStreamData *)data;

    __sync_fetch_and_add(&streamData->sync[streamData->x * streamData->tiles * 
                    streamData->tiles + streamData->y * streamData->tiles +
                    streamData->z], 1);

    printf("Callback: done tile %d\n", streamData->x * streamData->tiles * streamData->tiles +
                streamData->y * streamData->tiles + streamData->z);
}

但这行不通。仅调用最多 55 个回调。因此，程序卡在“Waiting 56”中。有 4 个图 block ，因此应该有 64 个回调。

也许内核运行得太快而无法设置回调？但为什么它适用于 55 而不是最后 9？

内核是正确的(至少它没有挂起)，因为它在没有平铺的情况下正常运行，并且通过参数，可以更改大小和输入数据。

我知道代码不是最佳的，也不漂亮，但目前我正在尝试使其工作，所以我可以从这里进行优化。

Answer 1

在更改内核数据输入的大小时，我错过了一个极端情况。这就是它失败的原因。现在正在运行。