大规模并行处理器编程实战笔记2

1：CUDA线程组织结构

一个grid中有多个block，可以为一维，二维，三维

一个block中有多个thread，可以为一维，二维，三维

grid中的block个数由gridDim确定（最多二维）

grid中的block定位使用blockIdx来确定

block中的thread个数由blockDim确定（最多三维）

block中的thread定位使用threadIdx来确定

例子：

grid一维（大小为N），block一维（大小为M），则每个线程利用blockIdx.x * blockDim.x + threadIdx.x来标识，二维及以上的同理。

dim3类型：

dim3实质上是一种C结构，有三个无符号整数类型的字段：x，y，z。

在启动设备函数的时候，传入的参数如果是多维的，必须使用dim3类型如：

Function<<<dimGrid, dimBlock>>>(...);

注意：每个grid中的block个数和每个block中的thread个数有数量限制，需要查看GPU型号确定。

2：使用blockIdx和threadIdx

如果grid中有多个block，每个block中有多个thread，那么，必须同时使用blockIdx和threadIdx来标识一个线程。

采用多个block执行矩阵乘法的kernel函数如下：

__global__ void MatrixMultiplication_kernel(flaot * Md, float * Nd, flaot *   Pd, int width)

{

//compute row index of Pd and Md

int Row = blockIdx.y * TILE_WIDTH + threadIdx.y;

//compute column index of Pd and Nd

int Col = blockIde.x * ITLE_WIDTH + threadIdx.x;

float pValue = 0;

//every thread compute one element of subMatrix

for(int k = 0; k < Width; k++)

pValue += Md[Row * Width + k] * Nd[k * Width + Col];

pd[Row * Width + Col] = pValue;

}

启动该kernel函数的主机代码如下：

//set execute configure args

dim3 dimGrid(Width / TILE_WIDTH, Width / TILE_WIDTH);

dim3 dimBlock(TILE_WIDTH, TILE_WIDTH);

//start device to compute

MatrixMultiplication_kernel<<<dimGrid, dimBlock>>>(Md, Nd, Pd, Width);

3：同步与透明可扩展性

CUDA中的线程同步函数：__syncthreads()

同步函数用于同步一个block中的所有线程，当一个线程执行到此处，将停下来等待其他线程，直到block中所有的线程都执行到此处，才开始继续下面的处理。

一个block中的线程要么全部执行__syncthreads()，要么全部不执行__syncthreads()，如果在if-then-else语句中，if和else都有__syncthreads()，则所有的线程必须执行同一个__syncthreads()。

由于block之间无需同步，所以，可以进行block与block之间的透明扩展（设备少的时候，先执行少数的block，设备多的时候，可以同时执行多数的block，无需更改程序代码）。

4：线程分配

分配过程：

启动kernel函数后，CUDA运行系统并生成线程网格（grid）。以block为单位把线程网格（grid）分配给SM（Streaming Multiprocessor，流多处理器）。

在GT200设计中，每个SM最对哦可以分配8个block，并保证有足够的资源满足所有block的需求，如果不能满足，系统运行时将会自动减少每个SM中block的个数（透明可扩展性的一个体现），直到能够满足为止。

GT200中有30个SM，最多可同时分配240个block，但是大多数grid都包含了超过240个block，则当前block执行完毕后，需要执行和分配新的block，系统运行时会维护这样的block列表。

透明可扩展性的另一个体现：例如由于GT200和G80中能容纳的线程总量不同，那么系统运行时会根据具体的设备特性进行调整（每次分配多少个线程），而无需修改任何代码。

5：线程调度与容许时延

线程调度的实际过程：

以GT200为例，一个block一旦分配到一个SM中，该block就被划分成了一个个的warp（大小为32个线程的单元，注意，warp的大小是在具体实现时指定的，与设备硬件特性相关），在SM中，warp才是真正的执行单元。

延时隐藏（latency hiding）：

以G80为例，一个SM只有8个SP（streaming processor，流处理器），但最多却可以驻留768/32=24个warp，为何需要这么多warp呢？答案在于CUDA处理器需要高效地执行长延时操作，如果warp中的线程执行一个条指令需要等待前面启动的长延时操作的结果，那么不会选择执行该warp，而是选择执行另一个不用等待结果的驻留的warp，这样，如果有了多个warp准备执行，则总可以选择不产生延时的线程先执行，达到所谓的延时隐藏。

零开销线程调度：

只要warp足够多，任何时候硬件都有可能及时找到合适执行的warp，因此尽管有长延时操作，仍能够充分利用执行硬件。选择执行处于就绪状态的warp不会带来多余的时间开销，这就是零开销线程调度（zero-overhead thread scheduling）。

如何选择block大小比较合适：

(每个SM中线程上限/block线程总数) <= (每个SM中可容纳block个数)

(block线程总数/每个wrap中线程的数量)尽可能大

6：小结

blockIdx和threadIdx这两个唯一的坐标变量标识网格中的线程和线程的作用域。

一旦启动一个grid，其中的block就以任意顺序被分配给SM，从而导致CUDA应用程序的透明可扩展性。

透明可扩展性的限制：不同的block之间的线程不能同步，它们唯一安全的同步机制是终止kernel函数并在同步点处为对应活动启动一个新的kernel函数。

一旦一个block分配给了SM，它将进一步被划分成warp。

在任何时候，SM都仅执行它驻留的部分warp，其他warp虽然在等待长延时操作，但是由于有大量的执行单元，所以也不会减少执行时的总吞吐量。