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并被忽略。
