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

Question

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

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

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

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

非常感谢.

Answer 1

线程分为三个级别,这三个级别的计划方式不同.Warps利用SIMD获得更高的计算密度.线程块利用多线程来实现延迟容忍.网格提供独立的粗粒度工作单元,用于跨SM的负载平衡.

经线中的线程

硬件一起执行warp的32个线程.它可以使用不同的数据执行单个指令的32个实例.如果线程采用不同的控制流,那么它们并非都执行相同的指令,那么这32个执行资源中的一些将在指令执行时空闲.这在CUDA参考中称为控制偏差.

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

// 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' core'基本上只是一个浮点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)基本上是对已经为游戏行业大规模生产的现有硬件设计如何用于加速其他类型的并行浮点处理应用程序的事后想法.