问题标签 [cuda-streams]

这个问题是关于 CUDA 流 (Stream) 的概念以及记录在流上的 CUDA 事件 (Event) 的明显异常。

考虑以下演示此异常的代码，

由于 Stream 上的所有工作都是顺序的，因此应该发生以下事件序列，

1. " eventStart" 记录为完整
2. kernel1运行并完成
3. " eventBetween" 记录为已完成
4. kernel2运行并完成
5. " eventEnd" 记录为已完成

将主机线程与设备同步后， 和 之间所用的时间eventStart（eventBetween由 测量cudaEventElapsedTime）应该是 的运行时间，kernel1和 之间所用的时间eventBetween应该eventEnd是 的运行时间kernel2。

但根据NVIDIA 关于 的文档cudaStreamWaitEvent，

使提交到流的所有未来工作等到事件报告完成后再开始执行。

和博文，

您还可以使用 cudaStreamWaitEvent(event) 在特定事件的单个流中同步操作

在同一流上记录事件之后安排的所有工作之前，事件不需要报告完成。这令人惊讶，因为在 Stream 上安排的所有工作都应该以顺序方式运行。

有人可以澄清 Stream 中事件的这种异步行为吗？

在 CUDA（驱动程序 API）文档中，它说

回调的开始执行与在回调之前同步记录在同一流中的事件具有相同的效果。因此，它会同步在回调之前“加入”的流。

这是否意味着如果我在每个关键点之后都有一个带有回调的管道来发送信号主机，我不需要任何 cuStreamSynchronize 来让这些点能够访问输出数组？

非常简单的代码，例如

我继承了一些基本上可以执行以下操作的代码：

看着targets/x86_64-linux/driver_types.hCUDA 8，我看到：

据我了解，演员阵容会起作用，但我担心这可能会如何面向未来，以及将代码移植到 ARM 时是否安全。上面的代码有多危险？有__device_builtin__什么影响吗？

（注意：我打算直接和开发者交谈，告诉他们在cudaStream_t整个过程中使用 and #include <cuda_runtime.h>，所以我希望在这里澄清技术问题。）

从官方 NVIDIA Multi-Process Server文档看，我不清楚它是如何与 CUDA 流交互的。

这是一个例子：

App 0：向逻辑流 0 发布内核；

应用程序 1：向（它自己的）逻辑流 0 发出内核。

在这种情况下，

1) MPS 是否/如何“劫持”这些 CUDA 调用？对于每个应用程序，它是否完全了解使用了哪些流以及哪些流中有哪些内核？

2) MPS 是否创建自己的 2 个流，并将各自的内核放入正确的流中？或者 MPS 是否可能通过流以外的机制实现内核并发？

如果有帮助，我对 MPS 在 Volta 上的工作方式很感兴趣，但是关于旧架构的信息也很受欢迎。

我在使用 CUDA 时遇到以下性能问题。当我在 Titan V 和 Titan X 卡上运行一个简单的示例代码时，运行时间符合预期。

现在，当我在代码中添加另一个内核时，它使用动态并行，但仍然不调用它或根本不使用它，Volta GPU 的性能急剧下降，但在其他卡上性能不受影响。

重要的是要强调第二个内核根本没有使用，它只是位于其余代码旁边，即它仅与其余代码一起编译。还可以在创建流的同时注释递归内核调用，并看到 Volta 的运行时间再次变好。我怀疑动态并行性的存在会对代码产生负面影响，即使它根本没有在运行时使用。关于如何解决这个问题的任何想法？

编译行：

我的程序是一个管道，它包含多个内核和 memcpys。每个任务将通过具有不同输入数据的相同管道。主机代码在处理任务时首先会选择一个 Channel，它是暂存器内存和 CUDA 对象的封装。在最后一个阶段之后，我将记录一个事件，然后去处理下一个任务。
主要流水线逻辑如下。问题是不同流中的操作不重叠。我附上了处理 10 个任务的时间表。您可以看到流中的任何操作都没有重叠。对于每个内核，一个块中有 256 个线程，一个网格中有 5 个块。用于 memcpy 的所有缓冲区都已固定，我确信我已满足这些要求用于重叠内核执行和数据传输。有人可以帮我找出原因吗？谢谢。

环境信息
GPU：Tesla K40m (GK110)
Max Warps/SM：64
Max Thread Blocks/SM：16
Max Threads/SM：2048
CUDA版本：8.0

我正在使用 OpenCV 3.4 和 cuda 库来处理视频图像。图像被抓取并使用GpuMat::upload(). 然后对图像进行两次阈值处理以创建 2 个不同的二值图像（Th1 和 Th2）。我的第一个问题是：我可以在两个阈值函数中使用 cuda 流同时运行吗？这是一个问题，因为它们都使用相同的 GpuMat 作为输入。

在阈值处理之后，我将使用两个二进制 GpuMats 使用其他 cv::cuda 函数对它们进行更多处理。第二个问题是：Stream::waitForCompletion()在使用 Th1 和 Th2 进行进一步处理之前，我是否应该等待阈值流完成？还是自动检测到这种数据依赖性？基本上我正在尝试从这里并行处理这两个二进制图像，而不是先处理 Th1 然后再处理 th2。它们将使用类似的 cuda 函数进行精确处理，但具有不同的值..

我正在使用 cuda 9.0。如果相同的操作使用不同的数据排入不同的流两次，这仍然是一个问题吗？

我的代码示例如下所示：

我一直认为 Hyper-Q 技术不过是 GPU 中的流。后来我发现我错了（是吗？）。所以我读了一些关于 Hyper-Q 的书，却更加困惑了。我正在浏览一篇文章，它有以下两个陈述：

A. Hyper-Q 是一种灵活的解决方案，它允许来自多个 CUDA 流、多个消息传递接口 (MPI) 进程，甚至是一个进程中的多个线程的单独连接

B. Hyper-Q 增加了主机和 GK110 GPU 之间的连接总数（工作队列），允许同时进行 32 个硬件管理连接（与 Fermi 提供的单个连接相比）

在上述几点中，B 点表示可以从主机创建多个连接到单个 GPU。这是否意味着我可以通过不同的应用程序在一个简单的 GPU 上创建多个上下文？这是否意味着我必须在不同的流上执行所有应用程序？如果我的所有连接都消耗内存和计算资源，谁来管理资源（内存/核心）调度？

我正在开发一个 CUDA 矩阵乘法，但我做了一些修改以观察它们如何影响性能。

我正在尝试观察一个简单矩阵乘法内核的行为（并且我正在测量 GPU 事件时间的变化）。但我在两种特定的不同条件下对其进行测试：

• 我有一定数量的矩阵（比如matN）A，B和C，然后我将（H2D）一个矩阵转移给A，一个矩阵转移给B，然后将它们相乘，以转移回（D2H）一个C；

• 我有A、B 和 C 的任何一个，但我在时间为 A 和 BmatN传输 >1（比如说）矩阵，执行精确的乘法，然后传回结果矩阵 C。chunkchunkchunk

在第一种情况下 ( chunk = 1) 一切都按预期工作，但在第二种情况下 ( chunk > 1) 我得到一些 C 是正确的，而另一些是错误的。

但是，如果我在得到的所有结果都是正确的cudaDeviceSynchronize()之后加上 a 。cudaMemcpyAsync

这是执行我上面刚刚描述的代码的一部分：

我尝试使用 cuda-gdb 进行调试，但没有出现任何奇怪的情况，gpuErrchk不会在 CUDA API 调用中引发任何错误。我也使用 memcheck 运行代码，在有和没有的情况下cudaDeviceSynchronize，在这两种情况下我都没有收到错误。

我想我可以说这是一个同步问题，但我不明白这背后的原因是什么。有人能发现我哪里出错了吗？其他代码风格的建议也非常感谢。

我正在开发一个多 GPU 加速流求解器。目前我正在尝试实现通信隐藏。这意味着，在交换数据时，GPU 会计算网格中不参与通信的部分，并在通信完成后计算网格的其余部分。

我试图通过一个流（computeStream）用于长期运行内核（fluxKernel）和一个（communicationStream）用于不同的通信阶段来解决这个问题。具有非常低的computeStream优先级，以便允许 上的内核communicationStream交错fluxKernel，即使它使用所有资源。

这些是我正在使用的流：

所需的并发模式如下所示：

在通过 MPI 发送数据之前，我需要同步communicationStream数据，以确保在发送之前完全下载数据。

在下面的清单中，我展示了我目前正在做的事情的结构。首先，我fluxKernel在computeStream. 然后我开始sendKernel收集应该发送到第二个 GPU 的数据，然后将其下载到主机（由于硬件限制，我不能使用 cuda-aware MPI）。然后数据以非阻塞方式发送MPI_Isend，随后使用阻塞接收（MPI_recv）。当接收到数据时，该过程向后完成。首先将数据上传到设备，然后通过recvKernel. 最后对fluxKernel上的网格的剩余部分调用communicationStream。

请注意，显示的代码内核在默认流上运行之前和之后。

我使用 nvprof 和 Visual Profiler 来查看流是否实际同时执行。这是结果：

我观察到sendKernel（紫色）、上传、MPI 通信和下载与fluxKernel. 但是，recvKernel（红色）仅在另一个流完成后才开始。开启同步并不能解决问题：

对于我的实际应用程序，我不仅有一个通信，而且还有多个。我也通过两次通信对此进行了测试。程序是：

结果类似于一次通信（上图），因为第二次内核调用（这次是 a sendKernel）被延迟到内核computeStream完成。

因此，总体观察结果是，第二次内核调用被延迟，与这是哪个内核无关。

你能解释一下，为什么 GPU 会以这种方式同步，或者我如何才能让第二个内核communicationStream同时运行到 computeStream？

非常感谢。

编辑1：问题的完整返工

最小可重现示例

我构建了一个最小的可重现示例。最后，代码将数据绘制int到终端。正确的最后一个值是 32778 (=(32*1024-1) + 1 + 10)。一开始我添加了一个选项整数来测试 3 个不同的选项：

• 0：在 CPU 修改数据之前同步的预期版本
• 1：同0，但不同步
• 2：memcpys 专用流且无同步

使用代码，我看到与上述相同的行为：

选项 0（正确结果）：

选项 1（错误的结果，缺少 CPU 的 +1）：

选项2（完全错误的结果，全部10，之前下载smallKernel_1）

解决方案：

在 Linux 下运行选项 0（根据 Roberts answere 中的建议）会带来预期的行为！