CUDA 会自动为您进行负载平衡吗？

Question

我希望对 CUDA C 中负载平衡的最佳实践提供一些一般性建议和说明，特别是：

• 如果经线中的 1 个线程比其他 31 个线程花费更长的时间，它会阻止其他 31 个线程完成吗？
• 如果是这样，是否会将备用处理能力分配给另一个 warp？
• 为什么我们需要warp 和块的概念？在我看来，warp 只是 32 个线程的一小块。
• 一般来说，对于给定的内核调用，我需要什么负载平衡？
  • 每个经线中的线程？
  • 每个块中的线程？
  • 跨所有块的线程？

最后，举个例子，你会为以下功能使用哪些负载平衡技术：

1. 我有一个向量x0 的N 点：[1, 2, 3, ..., N]
2. 我随机选择5%的点和log它们（或者一些复杂的函数）
3. 我将结果向量x1（例如[1, log(2), 3, 4, 5, ..., N]）写入内存
4. 我在 x1 上重复上述 2 次操作以产生 x2（例如 [1, log(log(2)), 3, 4, log(5), ..., N]），然后再进行 8 次迭代以产生 x3 ... x10
5. 我返回x10

非常感谢。

Answer 1

线程分为三个级别，它们的调度方式不同。 Warps 利用 SIMD 来获得更高的计算密度。线程块利用多线程来实现延迟容限。网格为跨 SM 的负载平衡提供独立的、粗粒度的工作单元。

warp 中的线程

硬件一起执行 warp 的 32 个线程。它可以执行具有不同数据的单个指令的 32 个实例。如果线程采用不同的控制流，因此它们不会都执行相同的指令，那么这 32 个执行资源中的一些将在指令执行时处于空闲状态。这在 CUDA 参考文献中称为 控制分歧。

如果内核表现出很多控制分歧，那么在这个级别上重新分配工作可能是值得的。这通过使所有执行资源在一个扭曲中保持忙碌来平衡工作。您可以在线程之间重新分配工作，如下所示。

// Identify which data should be processed
if (should_do_work(threadIdx.x)) {
  int tmp_index = atomicAdd(&tmp_counter, 1); 
  tmp[tmp_index] = threadIdx.x;
}
__syncthreads();

// Assign that work to the first threads in the block
if (threadIdx.x < tmp_counter) {
  int thread_index = tmp[threadIdx.x];
  do_work(thread_index); // Thread threadIdx.x does work on behalf of thread tmp[threadIdx.x]
}

块中的扭曲

在 SM 上，硬件调度会扭曲到执行单元上。一些指令需要一段时间才能完成，因此调度程序交错执行多个 warp 以保持执行单元忙碌。如果某些 warp 尚未准备好执行，则会跳过它们而不会降低性能。

通常不需要在此级别进行负载平衡。只需确保每个线程块有足够的 warp 可用，以便调度程序总能找到准备好执行的 warp。

网格中的块

运行时系统将块调度到 SM 上。多个块可以在一个 SM 上同时运行。

通常不需要在此级别进行负载平衡。只需确保有足够的线程块可用于多次填充所有 SM。当一些 SM 处于空闲状态且没有更多线程块准备好执行时，过度供应线程块以最大程度地减少内核末端的负载不平衡非常有用。

Answer 2

正如其他人已经说过的，warp 中的线程使用称为单指令多数据 (SIMD) 的方案。SIMD 意味着硬件中有一个指令解码单元控制多个算术和逻辑单元 (ALU)。 CUDA“核心”基本上只是一个浮点 ALU，而不是与 CPU 核心相同的完整核心。虽然确切的 CUDA 内核与指令解码器的比率在不同的 CUDA Compute Capability 版本之间有所不同，但它们都使用这种方案。由于它们都使用相同的指令解码器，线程束中的每个线程将在每个时钟周期执行完全相同的指令。分配给该 warp 中不遵循当前执行代码路径的线程的内核将在该时钟周期内不执行任何操作。没有办法避免这种情况，因为这是有意的物理硬件限制。因此，如果您在一个 warp 中有 32 个线程，并且这 32 个线程中的每一个线程都遵循不同的代码路径，那么您将不会在该 warp 中从并行性中获得任何加速。它将按顺序执行这 32 个代码路径中的每一个。这就是为什么warp中的所有线程尽可能遵循相同的代码路径是理想的，因为warp中的并行性只有在多个线程遵循相同的代码路径时才有可能。

这样设计硬件的原因是它节省了芯片空间。由于每个内核没有自己的指令解码器，内核本身占用的芯片空间更少（并且使用更少的功率）。拥有更小的内核，每个内核使用更少的功率意味着更多的内核可以封装到芯片上。拥有这样的小内核使得 GPU 可以在每个芯片上拥有数百或数千个内核，而 CPU 只有 4 或 8 个，即使同时保持相似的芯片尺寸和功耗（和散热）水平。与 SIMD 的权衡是，您可以将更多的 ALU 打包到芯片上并获得更多的并行性，但只有当这些 ALU 都执行相同的代码路径时，您才能获得加速。对于 GPU 而言，这种折衷程度如此之高的原因是 3D 图形处理中涉及的大部分计算只是浮点矩阵乘法。 SIMD 非常适合矩阵乘法，因为计算结果矩阵的每个输出值的过程是相同的，只是在不同的数据上。此外，每个输出值都可以完全独立于其他输出值进行计算，因此线程根本不需要相互通信。顺便说一句，类似的模式（通常甚至矩阵乘法本身）也恰好经常出现在科学和工程应用中。这就是 GPU 上的通用处理 (GPGPU) 诞生的原因。 CUDA（以及一般的 GPGPU）基本上是事后才想到的，即已经为游戏行业量产的现有硬件设计也可以用于加速其他类型的并行浮点处理应用程序。

Answer 3

如果经线中的 1 个线程比其他 31 个线程花费的时间更长，它会阻止其他 31 个线程完成吗？

是的。一旦您在 Warp 中出现分歧，调度程序需要获取所有分歧分支并一一处理。不在当前执行的分支中的线程的计算能力将丢失。您可以查看 CUDA 编程指南，它很好地解释了到底发生了什么。

如果是这样，空闲的处理能力会分配给另一个warp吗？

不，很遗憾，它完全丢失了。

为什么我们需要 warp 和 block 的概念？在我看来，warp 只是 32 个线程的一小部分。

因为 Warp 必须是 SIMD（单指令多数据）才能实现最佳性能，所以块内的 Warp 可以完全发散，但是它们共享一些其他资源。 （共享内存、寄存器等）

所以一般来说，对于给定的内核调用，我需要什么负载平衡？

我不认为负载平衡是正确的词。只要确保你总是有足够的线程一直在执行，并避免扭曲内部的分歧。同样，CUDA 编程指南是一个很好的读物。

现在举个例子：

您可以使用 m=0..N*0.05 执行 m 个线程，每个线程选择一个随机数并将“复杂函数”的结果放入 x1[m]。 但是，在大范围内随机读取全局内存并不是使用 GPU 可以做的最有效的事情，因此您还应该考虑是否真的需要完全随机。

Answer 4

其他人为理论问题提供了很好的答案。

对于您的示例，您可以考虑将问题重组如下：

1. 有一个向量x N 点：[1, 2, 3, ..., N]
2. 对x 的每个元素计算一些复杂的函数，得到y。
3. 随机抽样y 的子集以生成y0 到y10。

第 2 步只对每个输入元素执行一次，而不考虑是否需要该值。如果步骤 3 的采样在没有替换的情况下完成，这意味着您将计算 2 倍于实际需要的元素数量，但您将在没有控制分歧的情况下计算所有内容，并且所有内存访问都将是一致的。这些通常是 GPU 上比计算本身更重要的速度驱动因素，但这取决于复杂函数的实际作用。

第 3 步将具有不一致的内存访问模式，因此您必须决定是在 GPU 上执行此操作更好，还是将其传输回 CPU 并在那里进行采样是否更快。

根据下一个计算是什么，您可以重组第 3 步，改为在 [0,N) 中为每个元素随机绘制一个整数。如果该值在 [N/2,N) 中，则在下一次计算中忽略它。如果它在 [0,N/2) 中，则将其值与该虚拟 y* 数组（或任何适合您的计算）的累加器相关联。

Answer 5

你的例子是展示减少的一个很好的方式。

I have a vector x0 of N points: [1, 2, 3, ..., N]
I randomly pick 50% of the points and log them (or some complicated function) (1)
I write the resulting vector x1 to memory (2)
I repeat the above 2 operations on x1 to yield x2, and then do a further 8 iterations to  yield x3 ... x10 (3)
I return x10 (4)

说|x0| = 1024，然后您选择 50% 的点。

第一阶段可能是您必须从全局内存中读取的唯一阶段，我将告诉您原因。

512 个线程从内存中读取 512 个值（1），将它们存储到共享内存中（2），然后对于步骤（3），256 个线程将从共享内存中读取随机值并将它们也存储在共享内存中。你这样做，直到你最终得到一个线程，这会将它写回全局内存 (4)。

您可以在初始步骤中进一步扩展此功能，让 256 个线程读取两个值，或 128 个线程读取 4 个值，等等...