CUDA 中的嵌套内核答案

【问题标题】：Nested kernels in CUDACUDA 中的嵌套内核
【发布时间】：2011-05-24 03:36:07
【问题描述】：

CUDA 当前不允许嵌套内核。

具体来说，我有以下问题：我有 N 个 M 维数据。要处理 N 个数据点中的每一个，需要按顺序运行三个内核。由于不允许内核嵌套，我无法创建调用三个内核的内核。因此，我必须串行处理每个数据点。

一种解决方案是编写一个包含所有其他三个内核功能的大内核，但我认为它不是最佳的。

谁能建议如何使用流来并行运行 N 个数据点，同时保留三个较小的内核。

谢谢。

【问题讨论】：

大内核有什么问题？
我无法实现细粒度并行。假设我对一个数据点进行三种不同的矩阵运算。我可以为它们中的每一个编写内核。假设其中一个内核是矩阵乘法 C = A*B。乘法内核将并行查找 C(i,j) 的每个条目。当我有一个包含所有三个操作的大内核时，我无法做到这一点。大内核要做的只是并行处理数据点。
您当然可以运行多个流。相当简单，基本上内核启动的第四个参数是流。在同一流上启动的内核将按顺序执行，但跨不同流启动的内核以非同步顺序执行。如果您对实施有具体问题，我可以为您提供帮助
谢谢aaa。能否请您参考一些在内核启动中使用第四个参数的示例。

标签： arrays cuda

【解决方案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]);

并使用所有四个配置参数调用您的内核（当然，sharedMemory 可以为 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)

【讨论】：

【解决方案2】：

作为所选答案的更新，NVidia 的具有 Compute Capability 3.5 的 GPU 现在允许嵌套内核，他们称之为 Dynamic Parallelism。

【讨论】：

【解决方案3】：

如今，通过 Fermi 兼容性，可以启动并行内核

【讨论】：