问题标签 [dynamic-parallelism]

当您从 GPU 上的主内核中启动辅助内核时，会产生一些开销。造成或影响此开销金额的因素是什么？例如，内核代码的大小、正在启动内核的 SM 的占用率、内核参数的大小等。

为了这个问题，让我们包容，并将“开销”定义为以下时间间隔的总和：

开始：SM看到启动指令
结束：SM开始执行子内核的指令

加

开始：最后一个 SM 执行子内核的任何指令（或者可能是子内核指令的最后写入被提交到相关的内存空间）
结束：在子内核启动后执行父级的下一条指令。

问题1： 如果共享内存仅由子内核使用，我是否必须指定要在启动父内核时分配的动态共享内存量？

问题2： 下面是我的子内核和父内核

父内核

子内核

在这里，子内核单独工作正常。但是，当它从另一个内核启动时，在未指定启动失败时从主机启动父内核后cudaDeviceSynchronize()会给出错误（该错误不是从内核中的 printf 打印出来的）。

父内核的启动配置是<<<1,(17 17)>>>. 如果仅允许来自父级的一个线程启动子网格，则代码可以正常工作。一个区块可以启动多少个网格有限制吗？

我试图cudaMemsetAsync从内核调用（所谓的“动态并行”）。但无论我使用什么值，它总是将内存设置为 0。

这是我的测试代码：

如果我运行它，我会得到如下输出：

当我调用内存集时，我使用 value 0x7FFFFFFF。我期待非零数字，但它总是显示为零。

这是一个错误吗？还是我做错了什么？我正在使用 CUDA 8.0

产生错误的内核代码：

我测试了下面的代码以确保 OpenCL 2.0 编译器正常工作。

扫描功能提供 0,1,3,6 作为输出，因此 OpenCL 2.0 缩减功能正在工作。

动态并行是 OpenCL 2.0 的扩展吗？如果我删除enqueue_kernel命令，结果等于预期值（省略子内核）。

设备：AMD RX550，驱动：17.6.2

是否有需要在主机端运行的特殊命令才能在get_default_queue队列上运行子内核？目前，命令队列是使用 OpenCL 1.2 方式创建的，如下所示：

是否get_default_queue()必须是调用父内核的相同命令队列？问这个是因为我使用相同的命令队列将数据上传到 GPU，然后在一次同步中下载结果。

我正在使用 OpenCL 2.0 动态并行功能，并让每个工作项将另一个内核与单个工作项排入队列。当子内核的工作完成时间较长时，父内核在子内核之前完成，并且不会保留内存一致性并返回损坏的数据（随机更新的数据项）。

由于 clFinish() 和 clEnqueueMarkerWithWaitList() 仅用于主机队列，因此我不能将它们用于此默认设备无序队列。

如何使子内核在某个同步点之前或至少在缓冲区读取命令之前完成，以实现内存一致性？

这是代码：

当父内核只有 1-2 个工作项时，它可以正常工作，但是通常有 256*224 个工作项用于父内核，并且子内核在从主机访问数据之前无法完成（在 clFinish() 之后）

这是默认队列的构造（不同于父内核的队列）

编辑：这种创建队列的方式也不会使其可同步：

设备=RX550，驱动程序=17.6.2，64 位构建。

User Parallel Highway 的解决方案也不起作用：

这没有用：

这不起作用但完成，因此没有无限循环。

我们在使用 CUDA 动态并行时遇到了性能问题。目前，CDP 的执行速度至少比传统方法慢 3 倍。我们制作了最简单的可重现代码来显示此问题，即将数组的所有元素的值增加 +1。IE，

这个简单示例的目的只是看看 CDP 是否可以像其他人一样执行，或者是否存在严重的开销。

代码在这里：

可以编译

此示例可以使用 3 种方法计算结果：

1. 简单内核：只需对阵列进行一次经典内核 +1 传递。
2. 动态并行：从 main() 调用在 [0,N/3) 范围内 +1 的父内核，同时调用两个子内核。第一个孩子在 [N/3, 2*N/3) 范围内执行 +1，第二个孩子在 [2*N/3,N) 范围内执行 +1。Childs 使用不同的流启动，因此它们可以并发。
3. 来自主机的三个流：这个只是从 main() 启动三个非阻塞流，每个数组的三分之一。

我得到了方法 0（简单内核）的以下配置文件： 方法 1（动态并行性）的以下配置文件： 以及方法 2（来自主机的三个流） 的以下配置文件 运行时间是这样的：

从图片中可以看出，主要问题是在动态并行方法中，父内核在两个子内核完成后需要花费过多的时间来关闭，这使得它需要 3 倍或 4 倍的时间。即使考虑最坏的情况，如果所有三个内核（父内核和两个子内核）都串行运行，它应该花费更少。即，每个内核有 N/3 的工作，所以整个父内核应该花费大约 3 个子内核，这要少得多。有没有办法解决这个问题？

编辑：Robert Crovella 在评论中解释了子内核以及方法 2 的序列化现象（非常感谢）。内核确实以串行方式运行的事实不会使以粗体文本描述的问题无效（至少现在不是）。

我正在 GTX 980 ti 卡上尝试动态并行。运行代码的所有尝试都返回“未知错误”。下面显示了带有编译选项的简单代码。

我可以毫无问题地执行内核depth=0。第一次调用孩子时，会给出错误。在这里查看其他问题后将其cudaDeviceSynchronize()包括在内，但没有解决问题。

有任何想法吗？这可能是驱动程序问题吗？

编辑1：

操作系统：Linux-x86_64

英伟达驱动版本：384.59

nvcc 版本 7.5.17

有两个 980 ti 与 PCIe x16 Gen3 连接。该系统还在另一个配置了 RAID 的 SSD 上安装了 windows。

编译

我正在使用以下内核测试动态并行性，该内核以分而治之的方式使用动态并行性获得整数数组的最大值：

一个被称为：getMax<<<1,1>>>(d_arr, 0, N, d_max)， d_arr 是数组， N 是它的大小， d_max 是它的最大值。虽然有时我会得到正确的输出，但这个输出具有我倾向于在错误输出中看到的属性：

正如你所看到的，有很多次父亲网格在他们的孩子完成执行之前打印，尽管cudaDeviceSynchronize()正在使用。更糟糕的是，最终输出中没有考虑某些子值，从而从 GPU 得到错误的结果。

我知道在内核中使用 malloc（使用全局内存）和动态并行本身目前还不够快，无法让这段代码在 CPU 上有很好的加速。我只是想了解为什么这段代码没有正确同步。

我有一堆使用动态并行性的 .cu 文件（a.cu、b.cu、c.cu..、e.cu、f.cu），以及一个使用 MPI 从.cu 在多个节点上。我正在尝试编写一个 make 文件来编译可执行文件，但我一直面临以下错误：

这是我的生成文件：

让我们在有父内核和子内核的地方使用以下代码。从所述父内核，我们希望threadIdx.x在不同的流中启动子内核以最大化并行吞吐量。然后我们等待那些子内核，cudaDeviceSynchronize()因为父内核需要查看对global内存所做的更改。

现在假设我们还希望n使用流启动父内核，并且在我们希望并行启动的每组n父内核之间，我们还必须等待结果使用cudaDeviceSynchronize()

这将如何表现？

从Nvidia 对动态并行的官方介绍中，我认为这parent kernel[0]只会等待其中启动的流。这个对吗？如果没有，会发生什么？

注意：我知道一次只能运行这么多流（在我的情况下是 32 个），但这更多是为了最大限度地提高占用率

编辑：一个小代码示例