Reference
线程层次结构
核函数在 device 端执行时，会启动若干线程，一个核函数所启动的所有线程被称为一个线程网格（Thread Grid），同一个线程网格上的线程共享相同的全局内存空间，每个线程网格又可分为若干线程块（Thread Block），每个线程块中又包含若干线程
通过如下语句，可以定义一个 3*2 的 grid，每个 block 中含有 5*3 的 thread 的二维线程组织，其中 gird 和 block 是定义为 dim3 类型的变量，在定义时，缺省值初始化为 1，可以灵活的定义 1-dim、2-dim、3-dim 结构
1
2
3
| dim3 grid(3, 2, 1);
dim3 block(5, 3, 1);
kernel_function<<<grid, block>>>(prams...);
|
需要注意的是，每个 grid 中最多含有 65535 个 block，而每个 block 中最多含有 1024 个线程
相应的线程层次如下图所示
线程坐标
在 CUDA 中，一个线程需要两个内置的坐标变量 (blockIdx，threadIdx) 来标识，它们都是 dim3 类型变量，其中 blockIdx 指明线程所在的 block 在 grid 中的位置，而 threaIdx 指明线程在 block 中的位置
blockIdx 与 threadIdx 可通过内置变量来获得：
threadIdx.[x y z]：执行当前核函数的线程在 block 中的索引值blockIdx.[x y z]：执行当前核函数的线程所在的 block 在 grid 中的索引值
例如，threadIdx.x 是执行当前核函数的线程在 block 中的 x 方向的序号，blockIdx.x 是执行当前核函数的线程所在 block 在 grid 中的 x 方向的序号
1
2
3
4
5
6
7
8
9
10
11
12
13
| #include <stdio.h>
__global__ void hello_from_gpu() {
const int bx = blockIdx.x;
const int tx = threadIdx.x;
printf("block: %d, thread: %d\n", bx, tx);
}
int main() {
hello_from_gpu<<<4, 4>>>();
cudaDeviceSynchronize();
return 0;
}
|
线程唯一索引
唯一索引计算公式
一个线程块上的线程是放在同一个流式多处理器（SM）上的，但是单个 SM 的资源有限，这就使得 block 中的线程数是有限制的，现代 GPUs 的 block 可支持的线程数可达 1024 个，这就使得有时若想要知道一个线程在所有线程中的全局 ID，就必须要知道相应的组织结构，通过以下两个内置变量，可获得相应的组织结构信息
blockDim.[x y z]：一个 block 中包含多少个线程gridDim.[x y z]：一个 grid 中包含多少个 block
对于线程唯一索引，需要知道以下三个信息：
- blockIndex：block 在整个 grid 中的索引（1 维到 3 维）
- blockSize ：block 的大小，描述其中含有多少个线程
- threadId ：线程在 block 中的索引（1 维到 3 维）
进而有唯一索引计算公式：index = blockIndex * blockSize + threadIndex
1D Grid, 1D Block
对于 1-dim 的 grid 与 1-dim 的 block，有：
1
2
3
4
5
6
7
8
|
int blockIndex = blockIdx.x;
int blockSize = blockDim.x;
int threadIndex = threadIdx.x;
int index = blockIndex * blockSize + threadIndex;
|
1D Grid, 2D Block
对于 1-dim 的 grid 与 2-dim 的 block，有：
1
2
3
4
5
6
7
8
|
int blockIndex = blockIdx.x;
int blockSize = blockDim.x * blockDim.y;
int threadIndex = blockDim.x * threadIdx.y + threadIdx.x;
int index = blockIndex * blockSize + threadIndex;
|
1D Grid, 3D Block
对于 1-dim 的 grid 与 3-dim 的 block，有：
1
2
3
4
5
6
7
8
|
int blockIndex = blockIdx.x;
int blockSize = blockDim.x * blockDim.y * blockDim.z;
int threadIndex = blockDim.x * blockDim.y * threadIdx.z + blockDim.x * threadIdx.y + threadIdx.x;
int index = blockIndex * blockSize + threadIndex;
|
2D Grid, 1D Block
对于 2-dim 的 grid 与 1-dim 的 block，有：
1
2
3
4
5
6
7
8
|
int blockIndex = gridDim.x * blockIdx.y + blockIdx.x
int blockSize = blockDim.x;
int threadIndex = threadIdx.x;
int index = blockIndex * blockSize + threadIndex;
|
2D Grid, 2D Block
对于 2-dim 的 grid 与 2-dim 的 block，有：
1
2
3
4
5
6
7
8
|
int blockIndex = gridDim.x * blockIdx.y + blockIdx.x
int blockSize = blockDim.x * blockDim.y;
int threadIndex = blockDim.x * threadIdx.y + threadIdx.x;
int index = blockIndex * blockSize + threadIndex;
|
2D Grid, 3D Block
对于 2-dim 的 grid 与 3-dim 的 block，有：
1
2
3
4
5
6
7
8
|
int blockIndex = gridDim.x * blockIdx.y + blockIdx.x
int blockSize = blockDim.x * blockDim.y * blockDim.z;
int threadIndex = blockDim.x * blockDim.y * threadIdx.z + blockDim.x * threadIdx.y + threadIdx.x;
int index = blockIndex * blockSize + threadIndex;
|
3D Grid, 1D Block
对于 3-dim 的 grid 与 1-dim 的 block，有：
1
2
3
4
5
6
7
8
|
int blockIndex = gridDim.x * gridDim.y * blockIdx.z + gridDim.x * blockIdx.y + blockIdx.x
int blockSize = blockDim.x;
int threadIndex = threadIdx.x;
int index = blockIndex * blockSize + threadIndex;
|
3D Grid, 2D Block
对于 3-dim 的 grid 与 2-dim 的 block，有：
1
2
3
4
5
6
7
8
|
int blockIndex = gridDim.x * gridDim.y * blockIdx.z + gridDim.x * blockIdx.y + blockIdx.x
int blockSize = blockDim.x * blockDim.y;
int threadIndex = blockDim.x * threadIdx.y + threadIdx.x;
int index = blockIndex * blockSize + threadIndex;
|
3D Grid, 3D Block
对于 3-dim 的 grid 与 3-dim 的 block，有：
1
2
3
4
5
6
7
8
|
int blockIndex = gridDim.x * gridDim.y * blockIdx.z + gridDim.x * blockIdx.y + blockIdx.x
int blockSize = blockDim.x * blockDim.y * blockDim.z;
int threadIndex = blockDim.x * blockDim.y * threadIdx.z + blockDim.x * threadIdx.y + threadIdx.x;
int index = blockIndex * blockSize + threadIndex;
|
