拓展知识：了解grid、block、thread 关系
拓展知识：了解grid、block、thread 关系
如何理解呢？
实际上这个角标的计算也是虚拟地址，只不过为了底层寻址方便设计的，就跟c/c++概念中的虚拟地址一样。
一个三维的东西，grid最顶层，其次是block，最底层是thread。
1. 三者的概念关系
| 名称 | 含义 | 层次关系 | 示例 |
|---|---|---|---|
| Grid（网格） | 启动一次 kernel（GPU 函数）时的整体线程集合 | 最外层 | 一个 grid 中包含多个 block |
| Block（线程块） | 一组可以共享共享内存、可同步的线程 | grid 的组成单元 | 一个 block 中包含多个 thread |
| Thread（线程） | 最基本的执行单元 | block 的组成单元 | 每个 thread 执行相同 kernel 代码，不同数据 |
2. 三者的维度
| 名称 | 含义 | 类型 | 范围 |
|---|---|---|---|
threadIdx.{x,y,z} | 当前线程在 block 内的索引 | 局部索引 | [0, blockDim.{x,y,z}-1] |
blockIdx.{x,y,z} | 当前 block 在 grid 内的索引 | 全局块索引 | [0, gridDim.{x,y,z}-1] |
blockDim.{x,y,z} | 每个 block 的线程数 | 结构体常量 | 固定值 |
gridDim.{x,y,z} | grid 中 block 的数量 | 结构体常量 | 固定值 |
3.常见的维度
3.1 . 1D(一维)
#define M 32
#define N 64
dim3 block(N) //一个block有N个线程
dim3 grid(M) //一个grid 有 32个block
//1. blockIdx.x 范围是啥？ 0~31
//2. threadIdx.x 范围是啥 0~63
//3. blockDim.x = 64
//如何求取一维的线程标号呢？
int global_id = threadIdx.x (当前线程在这个block的偏移) + blockIdx.x(第几个block) * blockDim.x(恒定的值，为64)
3.2 . 2D(二维)
#define GM 16
#define GN 32
#define BM 64
#define BN 128
dim3 block(BM,BN);
dim3 grid(GM,GN)
//gridDim.x grid 在 X 方向上的 block 数 16
//gridDim.y grid 在 Y 方向上的 block 数 32
//blockDim.x 每个 block 在 X 方向的 thread 数 64
//blockDim.y 每个 block 在 Y 方向的 thread 数 128
//blockIdx.x 范围是啥？当前 block 在 grid 中 X 方向的索引 0~15
//blockIdx.y 当前 block 在 grid 中 Y 方向的索引 0~31
// threadIdx.x 当前线程在 block 内 X 方向索引 0~63
// threadIdx.y 当前线程在block内y方向的索引 0~127//二维索引:
int global_x = threadIdx.x + blockIdx.x * blockDim.x; //x方向的全局索引
int global_y = threadIdx.y + blockIdx.y * blockDim.y;//y方向上的索引//全局唯一索引
int g_id = global_y * (gridDim.x * blockDim.x) + global_x;3.3 三维情况
只做了解，用的少
int global_x = blockIdx.x * blockDim.x + threadIdx.x;
int global_y = blockIdx.y * blockDim.y + threadIdx.y;
int global_z = blockIdx.z * blockDim.z + threadIdx.z;int global_id = (global_z * gridDim.y * blockDim.y + global_y) * (gridDim.x * blockDim.x) + global_x;