硬件

  • SP(Streaming Processor)：流处理器，是GPU最基本的处理单元，在fermi架构开始被叫做CUDA core。

  • SM(Streaming MultiProcessor)：一个SM由多个CUDA core组成。

比如说，如果一个GPU有 $4$ 个SM，并且每个SM有 $768$ 个SP(Aka CUDA core)；那么在某一时刻，真正并行运行的线程数不会超过 $4 \times 768$ 个。

软件

threads被组织成blocks。一个block的线程可以用1Dimension(x), 2Dimensions(x, y)或者3Dim indexs(x, y, z) 索引，

显然，如果你需要 $4 \times 768$ 个以上的threads的话你需要 $4$ 个以上的blocks。blocks也可以使用1D, 2D或3D索引，这些blocks被放在等待队列上进入GPU执行。

Wrap

当一个kernel被执行时，grid中的线程块被分配到SM上。一个CUDA core可以执行一个thread，一个SM的CUDA core会分成几个wrap，由wrap scheduler负责调度。

一个wrap中的线程在同一个block中，如果block所含线程数不是wrap的大小的整数倍，那么多出来的那些thread所在的wrap中，会剩余一些inactive的thread。

一个简单的case

处理一张 $512 \times 512$ 的图片。

假设我们希望一个线程处理一个像素pixel(i, j)。

我们可以使用每 $64$ 个线程的区块。所以我们需要 $\frac{512 \times 512 }{64} = 4096$ 个区块（为了拥有 $512 \times 512 $ 个线程 ）。

通常情况下，我们将线程组织在2D区块中（为了更容易索引图像像素）。blockDim= $8 * 8$ ，我更喜欢叫它threadsPerBlock。

dim3 threadsPerBlock(8, 8);

还有2D的gridDim= $64 \times 64$ （需要 $4096$ 个区块）。我更喜欢叫它numBlocks。

dim3 numBlocks(imageWidth/threadPerBlock.x, imageHeight/threadPerBlock.y);

这个kernel是这样启动的：

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

最后，会有一个类似 $4096$ 个区块的队列的东西，其中一个块正在等待被分配到GPU的一个多处理单元中，以获得其 $64$ 个线程的执行。

在kernel中，一个线程要处理的像素是这样计算的：

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

Reference