CUDA中的嵌套内核

Question

CUDA目前不支持嵌套内核。

具体来说，我有N个M维数据。要处理每个N数据点，需要按顺序运行三个内核。由于不允许嵌套内核，因此我不能创建一个调用这三个内核的内核。因此，我必须逐个数据点进行串行处理。

一种解决方案是编写一个包含所有其他三个内核功能的大内核，但我认为这将是次优的。

有人可以建议如何使用流并保留三个较小内核来并行运行N个数据点。

谢谢。

- Prasanna

---

一个大内核有什么问题吗？ - Anycorn

我无法实现细粒度并行。假设我在一个数据点上执行三个不同的矩阵操作。我可以为每个操作编写内核。假设其中一个内核是矩阵乘法C = A * B。乘法内核将并行查找C（i，j）的每个条目。但当我有一个包含所有三个操作的大内核时，我无法这样做。大内核所做的只是并行处理数据点。 - Prasanna

你肯定可以运行多个流。这是相当简单的，基本上内核启动的第四个参数就是流。在同一流上启动的内核将按顺序执行，但在不同流上启动的内核将以非同步顺序执行。如果您对实现有具体问题，我可以帮助您解决。 - Anycorn

谢谢aaa。你能否给我提供一些使用内核启动中第四个参数的示例？ - Prasanna

Answer 1

如果你想使用流，你需要创建N个流：

cudaStream_t streams;
streams = malloc(N * sizeof(cudaStream_t));
for(i=0; i<N; i++)
{
    cudaStreamCreate(&streams[i]);
}

对于第i个数据点，您需要使用cudaMemcpyAsync进行传输：

cudaMemcpyAsync(dst, src, kind, count, streams[i]);

并使用所有四个配置参数调用您的内核（共享内存当然可以为0）：

kernel_1 <<< nBlocks, nThreads, sharedMemory, streams[i] >>> ( args );
kernel_2 <<< nBlocks, nThreads, sharedMemory, streams[i] >>> ( args );

当然，还需要进行清理：

for(i=0; i<N; i++)
{
    cudaStreamDestroy(streams[i]);
}
free(streams)

- jmilloy

Answer 2

作为所选答案的更新，NVidia的具有计算能力3.5的GPU现在允许嵌套内核，他们称之为动态并行。

- Tae-Sung Shin

Answer 3

如今，由于Fermi的兼容性，可以启动并行内核

- Dimitri