cuda - CUDA 是否会自动为您进行负载平衡？

Question

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

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

如果是这样，多余的处理能力是否会分配给另一个经线？

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

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

每个经线中的线？

每个块中的线程？

跨所有块的线程？

最后，举个例子，您将使用哪些负载平衡技术来实现以下功能:

我有一个矢量 x0的 N积分:[1, 2, 3, ..., N]

我随机选择了 5% 的点和 log它们(或一些复杂的函数)

我写出结果向量 x1 (例如 [1, log(2), 3, 4, 5, ..., N] )到内存

我在 x1 上重复上述 2 个操作屈服 x2 (例如 [1, log(log(2)), 3, 4, log(5), ..., N] )，然后再进行 8 次迭代以产生 x3 ... x10

我回 x10

非常感谢。

Answer 1

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

经纱中的线

硬件一起执行一个 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 上，硬件将扭曲调度到执行单元上。有些指令需要一段时间才能完成，因此调度程序会交错执行多个扭曲以保持执行单元忙碌。如果某些经纱未准备好执行，则会跳过它们而不会降低性能。

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

网格中的块

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

通常不需要在此级别进行负载平衡。只需确保有足够的线程块可用于多次填充所有 SM。当一些 SM 空闲并且没有更多线程块准备好执行时，过度配置线程块以最小化内核末尾的负载不平衡很有用。