并行运行多个流（而不是线程/块）

Question

我有一个内核，我想从配置“1 块 x 32 线程”开始。为了增加并行度，我想启动多个流，而不是运行比“1 块 x 32 线程”更大的“工作包”。我想在数据来自网络的程序中使用 GPU。我不想等到更大的“工作包”可用。 代码如下：

Thread(i=0..14) {
  - copy data Host -> GPU [cudaMemcpyAsync(.., stream i)]
  - run kernel(stream i)
  - copy data GPU -> Host [cudaMemcpyAsync(.., stream i)]
}

实际代码要复杂得多，但我想保持简单（15 个 CPU 线程使用 GPU）。

代码有效，但流未按预期同时运行。 GTX 480 有 15 个 SM，每个 SM 有 32 个着色器处理器。我希望如果我启动内核 15 次，所有 15 个流并行运行，但事实并非如此。我使用了 Nvidia Visual Profiler，最多有 5 个并行运行的流。通常只运行一个流。性能真的很差。

我使用“64 块 x 1024 线程”配置获得了最佳结果。如果我改用“32 块 x 1024 线程”配置，但使用两个流，则这些流一个接一个地执行，性能下降。我正在使用 Cuda Toolkit 5.5 和 Ubuntu 12.04。

有人能解释一下为什么会这样并能给我一些背景信息吗？它应该在较新的 GPU 上更好地工作吗？在您不想缓冲数据的时间要求严格的应用程序中使用 GPU 的最佳方式是什么？可能这是不可能的，但我正在寻找能让我更接近解决方案的技术。

新闻：

我做了一些进一步的研究。问题是最后一个 cudaMemcpyAsync(..) (GPU->Host copy) 调用。如果我删除它，所有流都会同时运行。我认为问题在幻灯片 21 上的http://on-demand.gputechconf.com/gtc-express/2011/presentations/StreamsAndConcurrencyWebinar.pdf 中有所说明。他们说在 Fermi 上有两个复制队列，但这仅适用于 tesla 和 quadro 卡，对吗？我认为问题在于 GTX 480 只有一个复制队列，所有复制命令（主机->GPU 和 GPU->主机）都放在这个队列中。一切都是非阻塞的，第一个线程的 GPU->host memcopy 阻塞了其他线程的 host->GPU memcopy 调用。 这里有一些观察：

Thread(i=0..14) {
  - copy data Host -> GPU [cudaMemcpyAsync(.., stream i)]
  - run kernel(stream i)
}

-> 有效：流同时运行

Thread(i=0..14) {
  - copy data Host -> GPU [cudaMemcpyAsync(.., stream i)]
  - run kernel(stream i)
  - sleep(10)
  - copy data GPU -> Host [cudaMemcpyAsync(.., stream i)]
}

-> 有效：流同时运行

Thread(i=0..14) {
  - copy data Host -> GPU [cudaMemcpyAsync(.., stream i)]
  - run kernel(stream i)
  - cudaStreamSynchronize(stream i)
  - copy data GPU -> Host [cudaMemcpyAsync(.., stream i)]
}

-> 不工作！！！也许 cudaStreamSynchronize 被放在了复制队列中？

有人知道这个问题的解决方案吗？像阻塞内核调用这样的东西会很酷。如果内核已完成，则应调用最后一个 cudaMemcpyAsync() (GPU->device)。

编辑2： 这里有一个例子来澄清我的问题： 为简单起见，我们有 2 个流：

Stream1:
------------
HostToGPU1
kernel1
GPUToHost1

Stream2:
------------
HostToGPU2
kernel2
GPUToHost2

第一个流开始了。执行 HostToGPU1，启动 kernel1 并调用 GPUToHost1。 GPUToHost1 阻塞，因为 kernel1 正在运行。同时 Stream2 启动。 HostToGPU2 被调用，Cuda 将它放入队列中，但它无法执行，因为 GPUToHost1 阻塞，直到内核 1 完成。目前没有数据传输。 Cuda 只是等待 GPUToHost1。所以我的想法是在 kernel1 完成后调用 GPUToHost1。这似乎是它与 sleep(..) 一起工作的原因，因为 GPUToHost1 在内核完成时被调用。自动阻止 CPU 线程的内核启动会很酷。 GPUToHost1 没有在队列中阻塞（如果当时没有其他数据传输，但在我的情况下，数据传输并不耗时）。

Answer 1

在 linux 上最容易看到并发内核执行。

有关一个很好的示例和一个简单的测试，请参阅concurrent kernels sample。

内核之间良好的并发性通常需要几件事：

• 支持并发内核的设备，因此是 cc 2.0 或更新的设备
• 内核在块数和其他资源使用（寄存器、共享内存）方面足够小，因此可以实际执行多个内核。通常会观察到具有较大资源需求的内核是串行运行的。这是预期的行为。
• 正确使用流以实现并发

此外，并发内核通常意味着复制/计算重叠。为了使复制/计算重叠工作，您必须：

• 正在使用具有足够复制引擎的 GPU。有些 GPU 有一个引擎，有些有 2 个。如果您的 GPU 有一个引擎，您可以将一个复制操作（即一个方向）与内核执行重叠。如果您有 2 个复制引擎（您的 GeForce GPU 有 1 个），您可以将复制的两个方向与内核执行重叠。
• 将固定（主机）内存用于将复制到 GPU 全局内存或从 GPU 全局内存复制的任何数据，这将是您打算重叠的任何复制操作的目标（到或从）
• 正确使用流和相关 API 调用的必要异步版本（例如 cudaMemcpyAsync

关于您观察到较小的 32x1024 内核不会同时执行，这可能是资源问题（块、寄存器、共享内存）阻止了很多重叠。如果第一个内核中有足够的块来占用 GPU 执行资源，那么期望其他内核在第一个内核完成或大部分完成之前开始执行是不明智的。

编辑：回应以下问题编辑和其他 cmets。

是的，GTX480 只有一个副本“队列”（我在回答中明确提到了这一点，但我称它为副本“引擎”）。您将只能在任何给定时间运行 one cudaMemcpy... 操作，因此在任何给定时间实际上只有一个方向（H2D 或 D2H）可以移动数据，您将只看到 one cudaMemcpy... 操作与任何给定内核重叠。并且cudaStreamSynchronize 导致流等待，直到之前向该流发出的 ALL CUDA 操作完成。

请注意，我认为您在上一个示例中使用的 cudaStreamSynchronize 不是必需的。流有 2 个执行特征：

1. 无论您是否使用Async API 或任何其他考虑因素，向同一流发出的cuda 操作（API 调用、内核调用等）将始终按顺序执行。
2. 发布到单独流的 cuda 操作，假设已满足所有必要的要求，将彼此异步执行。

由于第 1 项，在您的最后一种情况下，即使没有 @ 987654327@电话。所以我认为你可以摆脱那个电话，即你列出的第二种情况应该与最后一种情况没有什么不同，在第二种情况下你也不应该需要睡眠操作。在该流中所有先前的 cuda 活动完成之前，发布到同一流的 cudaMemcpy... 不会开始。这是流的特性。

EDIT2：我不确定我们在这方面是否取得了任何进展。您在 GTC preso here（幻灯片 21）中指出的问题是一个有效的问题，但您无法通过插入额外的同步操作来解决它，“阻塞内核”也不会帮助您，也不是具有一个复制引擎或 2 的功能。如果您想在单独的流中发布内核但在没有其他干预 cuda 操作的情况下按顺序发布，那么存在这种危险。正如下一张幻灯片所指出的，对此的解决方案是不按顺序发布内核，这与您的第二种情况大致相当。我再说一遍：

• 您已经确定您的案例 2 提供了良好的并发性
• 在这种情况下，睡眠操作不需要数据完整性

如果您想提供一个简短的示例代码来演示该问题，也许可以进行其他发现。