CUDA

CUDA-理解线程ID

理解线程ID

从线程逻辑结构上讲，所有线程有三层结构：threads，blocks，grids。每一层有三个维度：x，y，z。下面小例子展示了CUDA是怎样给不同的threads编号的：

假如我配置的kernel如下：

1
2
3
int nElem = 6;
dim3 block(3);
dim3 grid((nElem + block.x - 1) / block.x);

grid中结果为2，所以在kernel中：

1
checkDeviceIndex <<< grid, block >>>();

grid处为2，block处为3，即<<<2, 3>>>. 表示有2个blocks，每个blocks中有3个threads。其结构如下图：

一个蓝色矩形表示一个block，一个曲线箭头表示一个thread，在本例中一个grid由两个blocks 组成。

解释为：
对于grid

在x方向为2，表示在x方向由2个blocks。
y方向为1，表示在y方向上有1个block。
z方向为1，表示在z方向有1个block。

对于block

在x方向为2，表示在x方向上有3个threads。
y方向为1，表示在y方向上有1个thread。
在方向为1，表示在z方向上有1个thread。

threads是构成blocks和grids的最小单位，也是执行操作的最小单位。

从执行结上检验上述：

1
2
printf("grid.x %d grid.y %d grid.z %d \n", grid.x, grid.y, grid.z);
printf("block.x %d block.y %d block.z %d \n", block.x, block.y, block.z);

结果为：

1
2
grid.x=2 grid.y=1 grid.z=1 
block.x=3 block.y=1 block.z=1

与上述描述相符。

那么在kernel中是如何编号的呢！设计kernel：

1
2
3
4
5
6
7
8
9
__global__ void checkDeviceIndex(){

	printf("threadIdx:(%d, %d, %d)\n", threadIdx.x, threadIdx.y, threadIdx.z);
	printf("blockIdx:(%d, %d, %d)\n", blockIdx.x, blockIdx.y, blockIdx.z);

	printf("blockDim:(%d, %d, %d)\n", blockDim.x, blockDim.y, blockDim.z);
	printf("gridDim:(%d, %d, %d)\n", gridDim.x, gridDim.y, gridDim.z);

}

表示每一个thread都会打印4条信息，共有2*3=6个threads。结果为：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
threadIdx:(0, 0, 0)
threadIdx:(1, 0, 0)
threadIdx:(2, 0, 0)
threadIdx:(0, 0, 0)
threadIdx:(1, 0, 0)
threadIdx:(2, 0, 0)
blockIdx:(0, 0, 0)
blockIdx:(0, 0, 0)
blockIdx:(0, 0, 0)
blockIdx:(1, 0, 0)
blockIdx:(1, 0, 0)
blockIdx:(1, 0, 0)
blockDim:(3, 1, 1)
blockDim:(3, 1, 1)
blockDim:(3, 1, 1)
blockDim:(3, 1, 1)
blockDim:(3, 1, 1)
blockDim:(3, 1, 1)
gridDim:(2, 1, 1)
gridDim:(2, 1, 1)
gridDim:(2, 1, 1)
gridDim:(2, 1, 1)
gridDim:(2, 1, 1)
gridDim:(2, 1, 1)

根据threads ID的计算公式：int tid = threadIdx.x + blockIdx.x * blockDim.x
可以得到6个threads的ID分别是：
0 + 0 × 3 = 0，
1 + 0 × 3 = 1，
2 + 0 × 3 = 2，
0 + 1 × 3 = 3，
1 + 1 × 3 = 4，
2 + 1 × 3 = 5，

可以看出，CUDA kernel是根据公式给每一个threads编号的，保证每个threads有唯一的ID。