1. GPU线程分配示意图

Figure1: GPU结构示意图：2D Grid contains 2D blocks

• 由单个kernel启动产生的所有线程成为grid。
• grid中所有线程共享相同的global memory space。

2. CUDA的两种变量：

2.1 手动定义的dim3数据类型。

2.1.1 dim3特点：

• 在host，可以使用dim3定义grid和block的尺寸，作为kernel调用的一部分。
• dim3数据类型的手动定义的grid和block变量仅在host端可见。
• dim3是基于uint3的整数矢量类型。
• 当定义dim3的变量时，任何未指定的fileds都将初始化为1。类型为dim3的变量中的每个field都可以通过其x，y和z字段进行访问，

2.1.2 grid 和block的尺寸由以下两个内置变量指定：

• blockDim（block尺寸，以thread为单位）
  • blockDim.x, blockDim.y, blockDim.z
• gridDim（grid尺寸，以block为单位）

2.1.3 Example: 在Host定义 grid 和 block尺寸

#include <cuda_runtime.h>
#include <stdio.h>

int main(int argc, char **argv) { 
    // define total data elements 
    int nElem = 1024;
    
    // define grid and block structure
    dim3 block (1024);
    dim3 grid ((nElem+block.x-1)/block.x); 
    printf("grid.x %d block.x %d \n",grid.x, block.x);

    // reset block
    block.x = 512;
    grid.x = (nElem+block.x-1)/block.x; 
    printf("grid.x %d block.x %d \n",grid.x, block.x);

    // reset block
    block.x = 256;
    grid.x = (nElem+block.x-1)/block.x; 
    printf("grid.x %d block.x %d \n",grid.x, block.x);

    // reset block
    block.x = 128;
    grid.x = (nElem+block.x-1)/block.x; 
    printf("grid.x %d block.x %d \n",grid.x, block.x);
    
    // reset device before you leave cudaDeviceReset();
    return(0);
}

2.2 预定义的uint3数据类型。

2.2.1 uint3特点：

• uint3是CUDA内置的矢量类型，它是从基本整数类型派生的。
• 当kernel执行时，CUDA运行时会生成预初始化的grid，block和thread变量，这些变量可在kernel函数中访问并且类型为uint3。
• uint3数据类型的预初始化的grid和block变量仅在device可见。
• CUDA运行时，将坐标变量 (uint3) 分配给每个线程。根据线程，可以将部分数据分配给不同的线程。

2.2.2 线程用两个唯一的坐标来区分：

• blockIdx（网格内的块索引）

  • blockIdx.x
  • blockIdx.y
  • blockIdx.z

• threadIdx（块中的线程索引）

  • threadIdx.x
  • threadIdx.y
  • threadIdx.z

• 通常，grid被组织为block的2D阵列，而block被组织为thread的3D阵列。

• grid和block都使用带有三个无符号整数字段的dim3类型。未使用的字段将被初始化为1并被忽略。