问题标签 [cuda-streams]

I have created the Scan Algorithm for CUDA from scratch and trying to use it for smaller data amounts less than 80,000 bytes.

Two separate instances were created where, one runs the kernels using streams where possible and the other runs only in the default stream.

What Ive observed is that for this range of data sizes, running with streams takes longer time to complete the task compared with the other method.

When analysed using the nvprofiler, what was observed is that for smaller amount of data sizes, running in streams will not provide parallelism for separate kernals

Without Streams

With Streams

But when the data size is increased some kind of parallelism could be obtained

With Streams for 400,000bytes

My problem is, is there some additional parameters to reduce this kernel invoking time delays or is it normal to have this kind of behavior for smaller data sizes where using streams are disadvantageous

UPDATE :

I've included the Runtime API calls timeline as well to clarify the answer

我正在尝试使用流来运行 H2D 副本和内核并行运行。为此，我创建了 2 个带有 cudaStreamNonBlocking 标志的流。在一个循环中，我执行以下伪代码：

主机内存已固定。结果是每秒 cudaMemcpyAsync 都会延迟，这会导致代码运行速度变慢。见图：

我设法通过在复制的同一流上运行假内核来欺骗复制引擎，以使其立即运行。

有什么正常的方法可以让 GPU 立即执行 cudaMemcpyAsync 吗？

我使用 GeForce GTX 1060 6GB GPU

我目前正在通过计算两个向量之间的点积来学习 CUDA 流。成分是一个核函数，它接受向量x和y并返回大小等于块数的向量结果，其中每个块贡献自己的缩减和。

我还有一个宿主函数dot_gpu，它调用内核并将向量结果减少到最终的点积值。

同步版本就是这样做的：

而异步的就像：

使用流时我的性能越来越差，并试图调查原因。我尝试管道下载、内核调用和上传，但没有结果。

我的猜测是问题出在dot_gpu()函数内部，它不仅调用内核。告诉我是否正确理解了以下流执行

主机执行所有三个指令而不会被阻塞，因为 cudaMemcpyAsync 和内核会立即返回（但是在 GPU 上，它们将按顺序执行，因为它们被分配给同一流）。所以主机继续下一个流（即使stream1知道它处于哪个阶段，但谁在乎......它在GPU上做他的工作，对吧？）并再次执行三个指令而不会被阻塞......等等等等。但是，我的代码在dot_gpu()函数内部某处处理下一个流之前阻塞了主机。是因为我正在分配和释放东西，以及将内核返回的数组减少为单个值吗？

我正在实时处理我尝试使用 GeForce GTX 960M 处理的视频流。（Windows 10、VS 2013、CUDA 8.0）

每一帧都必须被捕捉，轻微模糊，只要有可能，我需要对最新的 10 帧进行一些艰苦的计算。所以我需要以 30 fps 的速度捕获所有帧，并且我希望以 5 fps 的速度获得艰苦的工作结果。

我的问题是我无法保持捕获以正确的速度运行：无论是在 CPU 级别还是在 GPU 级别，艰苦的计算似乎都会减慢帧的捕获速度。我错过了一些帧...

我尝试了很多解决方案。没有工作：

1. 我尝试在 2 个流上设置作业（下图）：
   • 主机得到一个框架
   • 第一个流（称为 Stream2）：cudaMemcpyAsync 将帧复制到设备上。然后，第一个内核进行基本的模糊计算。（在附图中，模糊显示为 3.07 秒和 3.085 秒处的一个短槽。然后什么都没有……直到大部分完成）
   • 由于 CudaEvent，主机检查第二个流是否“可用”，并在可能的情况下启动它。实际上，流在 1/2 的尝试中可用。
   • 第二个流（称为 Stream4）：在内核（kernelCalcul_W2）中开始艰苦的计算，输出结果，并记录一个事件。

NSight 捕获

实际上，我写道：

2. 我试图只使用一个流。与上面的代码相同，但在启动 hard_work 时更改了 1 个参数。
   • 主机得到一个框架
   • 流：cudaMemcpyAsync 将帧复制到设备上。然后，内核进行基本的模糊计算。然后，如果 CudaEvent Event_1 没问题，我会努力工作，并添加一个 Event_1 以获得下一轮的状态。实际上，流始终可用：我从不属于“其他”部分。

这样，在努力工作的同时，我希望“缓冲”所有要复制的帧，而不会丢失任何帧。但我确实失去了一些：事实证明，每次我得到一个帧并复制它时，Event_1 似乎还可以，所以我开始努力工作，直到很晚才得到下一帧。

3. 我试图将两个流放在两个不同的线程中（在 C 中）。没有更好（甚至更糟）。

所以问题是：如何确保第一个流捕获所有帧？我真的感觉不同的流会阻塞 CPU。

我用 OpenGL 显示图像。会不会干扰？

有什么改进方法的想法吗？非常感谢！

编辑： 根据要求，我在这里放了一个 MCVE。

您可以调整一个参数 (#define ADJUST) 以查看发生了什么。基本上，主程序以异步模式发送 CUDA 请求，但它似乎阻塞了主线程。正如您将在图像中看到的那样，我每 30 毫秒进行一次“内存访问”（即捕获的图像），除非正在运行艰苦的工作（然后，我只是没有得到图像）。

最后一个细节：我正在使用 CUDA 7.5 来运行它。我试图安装 8.0 但显然编译器仍然是 7.5

我想让两个 CUBLAS API（例如.cublasDgemm）真正在两个 cudaStreams 中同时执行。

众所周知，CUBLAS API 是异步的，像 cublasDgemm 这样的 3 级例程不会阻塞主机，这意味着以下代码（默认为 cudaStream）将同时运行：

但是，当我使用“NVIDIA Visual Profiler”分析程序时，它表明它们运行有序。

然后，我尝试让它们绑定到不同的 cudaStreams，伪代码是：

当 batch_count=5 时，“NVIDIA Visual Profiler”显示的结果是：

Multi-CublasDegmm Routines 多流执行结果

结果表明，它们仍然有序地运行。如何使多个 cublas api 在多个 cudaStreams 中真正同时运行，如下所示：

多流的多内核执行结果，它们真正并发运行

有人知道吗？谢谢。

我们在使用 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 的序列化现象（非常感谢）。内核确实以串行方式运行的事实不会使以粗体文本描述的问题无效（至少现在不是）。

考虑以下程序，用于将非阻塞 GPU 流上的一些工作排​​入队列：

缓冲区的大小和内核睡眠周期的长度足够高，当它们与 CPU 线程并行执行时，它应该在它们结束之前完成排队（复制 8 毫秒 + 8 毫秒，复制 20 毫秒）内核）。

然而，看看下面的跟踪，似乎两者cudaMemcpyAsync()实际上是同步的，即它们阻塞直到（非阻塞）流实际上结束了复制。这是预期的行为吗？它似乎收缩了CUDA Runtime API 文档的相关部分。这有什么意义？

跟踪：（编号的行，以微秒为单位的时间）：

然而，看看下面的跟踪，似乎两者cudaMemcpyAsync()实际上是同步的，即它们阻塞直到（非阻塞）流实际上结束了复制。这是预期的行为吗？这似乎与 CUDA Runtime API 文档的相关部分相矛盾。这有什么意义？

以下程序：

应该会导致第一个流和第二个流上的工作之间的 I/O 和计算重叠：当第一个流的主机到设备结束时，第一个流的内核可以启动，但第二个流的主机到设备传输也可以启动。相反，我得到以下时间线，没有重叠：

我想我已经覆盖了我的基础以确保重叠。流是非阻塞的（实际上工作的排队在第一个 HtoD 完成之前就结束了）；主机内存已固定......所以我看到重叠缺少什么？

在 GNU/Linux Mint 18.2 上使用 CUDA 8.0.61 和 NVIDIA GTX 650 Ti Boost。但是驱动是v384.59。

考虑以下程序：

我在 GTX Titan X 上运行它，CUDA 8.0.61，Fedora 25，驱动程序 375.66。我看到的时间线是这样的：

这张图有几点不对：

• 据我所知，一次只能进行一次 HtoD 传输。
• 所有的内存传输都应该花费基本相同的时间——它们具有相同的数据量；PCIe 总线对传输速率的影响如此之大，没有什么其他有趣的事情发生了。
• 一些 DtoH 条就像它们被拉长一样，直到另一个流上发生某些事情。
• 有一个巨大的差距，似乎没有计算机，也没有真正的 I/O。即使所有先前完成的内核的 D​​toH 都占据了这个空白，那仍然会留下非常大量的时间。这实际上看起来像是一个调度问题，而不是一个分析错误。

那么，我应该如何解读这个时间线呢？问题出在哪里？（希望不是程序员...）

我应该提到，使用较少的流（例如 2），时间线在相同的 SW+HW 上看起来非常好：

这个程序：

产量：

这有点奇怪，因为 API 参考cudaMemcpyAsync并未将cudaErrorNotPermitted其列为潜在错误之一。从回调调度异步副本真的有问题吗？

注意：我的机器有 GTX 650 Ti (CC 3.0)、CUDA 9.0、Linux 内核 4.8.0、驱动程序 384.59。