1. 线程模型概述

在CUDA编程中，线程的组织形式是多维的，主要通过网格（Grid）和线程块（Block）来管理。了解线程模型对于有效设计并行计算非常重要。以下是一些核心概念：

线程模型的重要概念：

• grid （网格）：网格是线程的整体组织结构。它由多个线程块（Block）组成，每个线程块又包含多个线程。一个Grid可以视为整个计算任务的执行单元。
• block （线程块）：线程块是网格中的一个子集，包含多个线程。每个线程块的执行可以独立于其他线程块。线程块之间没有数据共享，但可以通过共享内存进行通信。

线程分块是逻辑上的，物理上线程不分块。CUDA允许开发者通过指定grid_size和block_size来配置线程的组织结构：

• grid_size：指定一个网格（Grid）中包含多少个线程块（Block）。
• block_size：指定每个线程块中包含多少个线程。

线程配置采用如下的语法：

<<<grid_size, block_size>>>

• 最大线程块大小：每个线程块最多允许包含 1024 个线程。
• 最大网格大小：每个网格最多可以包含 $2^{31}-1$ 个线程块。

可以理解为：

• 一个Grid最多有 $2^{31}-1$ 个线程块。
• 每个 线程块（Block） 最多有 1024个线程。

2. 一维线程模型

• 每个线程在核函数中都有一个唯一的身份标识
• 每个线程的唯一标识由 <<<grid_size, block_size>>> 确定，grid_size, block_size保存在内建变量（build-in variable），目前考虑的是一维情况：
  （1）gridDim.x ：该变量的数值等于执行配置中变量grid_size的值
  （2）blockDim.x ：该变量的数值等于执行配置中变量block_size的值
• 线程索引保存成内建变量（build-in variable）
  （1）blockIdx.x：该变量指定一个线程在一个网格中的线程块索引值，范围为0~gridDim.x -1
  （2）threadIdx.x：该变量指定一个线程在一个网格中的线程块索引值，范围0~blockDim.x -1

例如： kernal_fun<<<2,4>>>():

• $gridDim.x$ ：值为 2
• $blockDim.x$ ：值为 4
• $blockIdx.x$ ：取值为 0~1
• $threadIdx.x$ ：取值为 0~3

线程唯一标识：
$$Idx=threadIdx.x+blockIdx.x\ast blockDim.x$$

3. 多维线程模型

3.1 多维线程模型概述

CUDA可以组织三维网格和线程块

blockIdx和threadIdx是类型为uint3的变量，该类型是一个结构体，具有x, y, z三个成员（3个成员都为无符号类型的成员构成）：

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

---

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

gridDim和blockDim是类型为dim3的变量，该类型是一个结构体，具有x， y， z三个成员：

• $gridDim.x$
• $gridDim.y$
• $gridDim.z$

---

• $blockDim.x$
• $blockDim.y$
• $blockDim.z$

取值范围：

• $blockIdx.x$ ：范围 $[0,gridDim.x-1]$
• $blockIdx.y$ ：范围 $[0,gridDim.y-1]$
• $blockIdx.z$ ：范围 $[0,gridDim.z-1]$

---

• $threadIdx.x$ : 范围 $[0,blockDim.x-1]$
• $threadIdx.y$ : 范围 $[0,blockDim.y-1]$
• $threadIdx.z$ : 范围 $[0,blockDim.z-1]$

注意：内建变量只在核函数有效，且无需定义

在一维线程模型中 <<<grid_size, block_size>>> :

• $grid\_size$ 对应 $gridDim.x$
• $block\_size$ 对应 $blockDim.x$

$gridDim$ 和 $blockDim$ 中没有指定的维度默认为1：

• $gridDim.x=grid\_size$
• $gridDim.y=1$
• $gridDim.z=1$
• $blockDim.x=block\_size$
• $blockDim.y=1$
• $blockDim.z=1$

3.2 多维线程模型定义

定义多维网格和线程块（C++构造函数语法）：

dim3 grid_size(Gx,Gy,Gz);
dim3 block_size(Bx,By,Bz);

举个例子：定义一个 2*2*1 的网格，5*3*1的线程块，代码定义如下：

dim3 grid_size(2,2);   //等价于dim3 grid_size(2,2,1);
dim3 block_size(5,3);  //等价于dim3 block_size(5,3,1);

多维网格和多维线程块本质上是一维，GPU物理上不分块
每个线程都有一个唯一标识：

• int tid = threadIdx.y*blockDim.x + threadIdx.x;
• int bid = blockIdx.y*gridDim.x + blockIdx.x;

3.3 多维线程模型中的索引

• 多维度线程块中的线程索引：
  int tid = threadIdx.z * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;

• 多维网格中的线程块索引：
  int bid = blockIdx.z * gridDim.x * gridDim.y + blockIdx.y * gridDim.x + blockIdx.x;

4. 网格和线程块的限制条件

网格大小限制：

• $gridDim.x$ 最大值 $2^{31}-1$
• $gridDim.y$ 最大值 $2^{16}-1$
• $gridDim.z$ 最大值 $2^{16}-1$

线程块大小限制：

• $blockDim.x$ 最大值1024
• $blockDim.y$ 最大值1024
• $blockDim.z$ 最大值64

注意：线程块总的大小最大为1024！！！， 即 blockDim.x*blockDim.y*blockDim.z 不能超过1024

如有错误欢迎指正
本文参考：https://www.bilibili.com/video/BV1sM4y1x7of/?spm_id_from=333.788.videopod.episodes&vd_source=cf0b4c9c919d381324e8f3466e714d7a&p=4