了解CUDA网格尺寸,块尺寸和线程组织(简单说明)

Question

如何组织线程由GPU执行？

Answer 1

例如,如果GPU设备有4个多处理单元,并且它们每个可以运行768个线程:那么在给定时刻,不会有超过4*768个线程并行运行(如果您计划了更多线程,它们将等待轮到他们了.

线程按块组织.块由多处理单元执行.可以使用1维度(x),2D维度(x,y)或3Dim索引(x,y,z)来识别(索引)块的线程,但在任何情况下,对于我们的示例,x y z <= 768(其他限制适用)到x,y,z,请参阅指南和您的设备功能).

显然,如果您需要超过4*768个线程,则需要超过4个块.块也可以被索引为1D,2D或3D.有一个块队列等待进入GPU(因为,在我们的示例中,GPU有4个多处理器,并且只有4个块同时执行).

假设我们想要一个线程处理一个像素(i,j).

我们可以使用每个64个线程的块.然后我们需要512*512/64 = 4096个块(所以要有512x512个线程= 4096*64)

组织(使图像索引更容易)通常是具有blockDim = 8 x 8(每个块64个线程)的2D块中的线程.我更喜欢称它为threadsPerBlock.

dim3 threadsPerBlock(8, 8);  // 64 threads

和2D gridDim = 64 x 64块(需要4096块).我更喜欢称它为numBlocks.

dim3 numBlocks(imageWidth/threadsPerBlock.x,  /* for instance 512/8 = 64*/
              imageHeight/threadsPerBlock.y);

内核是这样启动的:

myKernel <<<numBlocks,threadsPerBlock>>>( /* params for the kernel function */ );

最后:将会出现类似"4096个块的队列"的情况,其中一个块正在等待分配GPU的多个处理器之一以执行其64个线程.

在内核中,线程处理的像素(i,j)以这种方式计算:

uint i = (blockIdx.x * blockDim.x) + threadIdx.x;
uint j = (blockIdx.y * blockDim.y) + threadIdx.y;

如果每个块可以运行768个线程,为什么只使用64个？如果您使用最大限制768,您将拥有更少的块和更好的性能. (11认同)
@Aliza:块是*逻辑*,每个*物理*处理单元的768个线程的限制.您可以根据问题的规范使用块,以便将工作分配给线程.对于您遇到的每个问题,您都不可能总是使用768个线程的块.想象一下,您必须处理64x64图像(4096像素).4096/768 = 5.333333块？ (9认同)
@thouis是的,也许吧.但情况是每个线程所需的内存量取决于应用程序.例如,在我的上一个程序中,每个线程调用最小二乘优化函数,需要"大量"内存.这么多,这些块不能超过4x4线程.即便如此,与顺序版相比,获得的加速也是戏剧性的. (5认同)
@ cibercitizen1 - 我认为Aliza的观点很好:如果可能的话,人们希望尽可能多地使用每个块的线程.如果存在需要更少线程的约束,则更好地解释为什么在第二个示例中可能是这种情况(但仍然解释更简单和更理想的情况,首先). (3认同)

Answer 2

假设一个9800GT GPU:14个多处理器,每个都有8个线程处理器,warpsize是32,这意味着每个线程处理器最多可处理32个线程.14*8*32 = 3584是实际执行线程的最大数量.

如果你用超过3584个线程执行这个内核(比如说4000个线程,你定义块和网格的方式并不重要.gpu会像对待它们一样对待它们):

func1();
__syncthreads();
func2();
__syncthreads();

那么这两个函数的执行顺序如下:

1.func1是针对前3584个线程执行的

2.func2是针对前3584个线程执行的

3.func1是为剩余的线程执行的

4.func2是为剩余的线程执行的