1.计算机架构
正常程序执行在host(CPU)端,当需要到device(GPU)执行时,启动kernel(核函数,一个kernel对应一个线程)调用GPU。(K-G-B-T)
- 一个kernel对应一个grid,启动一个kernel的时候需要制定grid和block的维度
- 一个grid可以有多个block,一维~三维
- 一个block可以有多个thread,一维~三维
每个thread有独立的Register和Local Memory,每个block有独立的share memory,一个block中的所有thread共享share memory,每个grid 有独立的global、constant、texture memory,所有的block共享global memory、constant memory、texture memory。
2.线程管理
由一个内核启动所产生的所有线程统称为一个网格(grid)。同一网格中的所有线程共享相同的全局内存空间(相当于CPU中的系统内存)。
一个网格由多个线程块(block)构成,一个线程块包含一组线程,同一线程块内的线程协作通过同步、共享内存实现。不同线程块(block)内的线程不能协作。
当执行一个核函数时,CUDA运行时为每个线程分配坐标变量blockIdx和threadIdx。基于这些坐标,可以将数据分配给不同的线程。
- blockIdx(线程块(block)在线程格(grid)内的索引)
- threadIdx(块(block)内的线程thread索引)
网格(grid)和块(block)的维度由下列两个内置变量指定。
- blockDim(线程块的维度,每个块中的线程数目)
- gridDim(线程格的维度,启动块的数目)
网格和块从逻辑上代表了一个核函数的线程层次结构。CUDA的特点之一就是通过编程模型揭示了一个两层的线程层次结构。
3. thread遍历
3.1 一维block、grid
dim3 block(blockDim); //一维block
这一行创建了一个三维向量block,用来定义每个block的大小。blockDim表示每个block中包含blockDim个线程。dim3数据类型是CUDA中的一个特殊数据类型,用于表示三维向量。在这个情况下,block的其余维度将被默认设置为1。这意味着你将有一个包含blockDim个线程的一维block。
dim3 grid(gridDim); // 一维grid
这一行创建了一个三维向量grid,这意味着整个grid将包含gridDim个block。与block一样,只传递一个整数值给grid,所以其余维度将被默认设置为1,得到一个一维grid。
3.2 多维
dim3 block(blockDimx, blockDimy) ; //二维
dim3 grid(gridDimx, gridDimy); // 二维grid
dim3 block(blockDimx, blockDimy,blockDimz) ; //三维
dim3 grid(gridDimx, gridDimy,,gridDimz); // 三维grid
3.3 遍历
block方向寻找threadId:
先遍历z方向(每一个z方向都有blockDim.x * blockDim.y的thread),再遍历y方向(每一个y方向都有blockDim.x的thread),最后遍历x方向。
如需在grid方向寻找threadId:
- 首先计算blockSize:blockSize = blockDim.z * blockDim.y * blockDim.z ;(每个block中含有多少个thread)
- 在grid方向查找block坐标前有多少个block,bIndex = blockIdx.z * gridDim.x * gridDim.y + blockIdx.y * gridDim.x + blockIdx.x;
- 在block方向查找thread坐标前有多少个thread:tIndex = threadIdx.z * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;
- 最后计算整个index:index = bIndex * blockSize + tIndex ;
__global__ void print_thread_idx_kernel(){
int bSize = blockDim.z * blockDim.y * blockDim.x;
int bIndex = blockIdx.z * gridDim.y * gridDim.x + \
blockIdx.y * gridDim.x + \
blockIdx.x;
int tIndex = threadIdx.z * blockDim.y * blockDim.x + \
threadIdx.y * blockDim.x + \
threadIdx.x;
int index = bIndex * bSize + tIndex;
printf("block idx: %3d, thread idx in block: %3d, thread idx: %3d\n",
bIndex, tIndex, index);
}