CUDA中的线程与线程束

前言

本文为《CUDA与TensorRT部署学习笔记》系列中CUDA编程入门篇。本文主要从cuda的grid、block、thread的概念、线程ID的计算方式以及相关的测试脚本的解析等方面进行学习记录。

Goal:

理解在CUDA中一维、二维、三维的grid, block的写法，以及遍历thread的方法

1 CUDA 编程的基础概念

1.1 CUDA中的grid和block

1.1.1 grid、block、thread

（1）grid

grid是一个由多个block组成的网格，它表示一个核函数的执行范围。一个核函数只能包含一个grid，但是一个grid可以有任意多个block，只要不超过GPU的限制。grid的组织方式可以是一维的，二维的，或者三维的，这取决于程序的需求。

网格的大小由 gridDim 变量指定，它是一个 dim3 类型的变量，其中包含网格的 x、y 和 z 维度。

（2）block

block是一个由多个thread组成的块，它表示一个grid中的一个子任务。一个block中的thread可以通过共享内存和同步机制进行通信和协作。block的组织方式也可以是一维的，二维的，或者三维的，但是一个block中的thread的总数不能超过512或1024，这取决于GPU的架构。

块的大小由 blockDim 变量指定，它也是一个 dim3 类型的变量，其中包含块的 x、y 和 z 维度。

（3）thread

thread是一个执行单元，它表示一个block中的一个基本操作。一个thread拥有自己的程序计数器和状态寄存器，并且用自己的数据执行指令。thread的执行方式是SIMT（单指令多线程），也就是说，同一个block中的thread会以32个为一组（称为warp）执行相同的指令，但是可能有不同的行为，比如分支结构。

每个线程都有一个唯一的索引，该索引由 threadIdx 变量指定。 threadIdx 变量是一个 dim3 类型的变量，其中包含线程在块内的 x、y 和 z 维度。

（4）备注

启动一个kernel的时候需要指定grid和block
总结：
- 一个kernel对应一个grid（一个 kernel 函数可以并行执行在多个线程上，这些线程组成了一个 grid。）
- 一个grid可以有多个block，一维三维
- 一个block可以有多个thread，一维三维

grid 就像一个学校，block 就像一个班级，thread 就像一个学生。学校里有多个班级，每个班级可以容纳多个学生。每个学生都有自己的学生号，学生号是学生唯一的身份标识。

可以通过 gridDim 和 blockDim 变量来控制学校的规模和班级的数量。还可以通过 threadIdx 变量来获取每个学生的学生号。

1.1.2 线程ID的计算方式

（1）公式：

线程Id = blockId * blockSize + threadId

blockId ：当前 block 在 grid 中的坐标（可能是1维到3维）
blockSize ：block 的大小，描述其中含有多少个 thread
threadId ：当前 thread 在 block 中的坐标（同样从1维到3维）

学生的ID=当前班级的ID*每个班级的人数+当前学生在班级的坐标

（2）关键点说明：

grid 中含有若干个 blocks，其中 blocks 的数量由 gridDim.x/y/z 来描述。某个 block 在此 grid 中的坐标由 blockIdx.x/y/z 描述。
blocks 中含有若干个 threads，其中 threads 的数量由 blockDim.x/y/z 来描述。某个 thread 在此 block 中的坐标由 threadIdx.x/y/z 描述。

（3）多维的坐标用一维数据表达

想两位数和三位数，就是很好的例子。数字 = 百位数字 * 100 + 十位数字 * 10 + 个位数字。
当我们得知每个维度上的大小时，就可以利用这样的进制将三维坐标转换为1维坐标。
一般来说坐标（x, y, z）分别所在的维度大小是（Dx, Dy, Dz），一般会把 z 看成高纬度，接着是 y ，最后是 x。
高维度坐标转一维坐标公式 id = Dx * Dy * z + Dx * y + x；坐标从0开始；维度从1开始；
当我们有一个三维坐标 $(x, y, z)$ ，假设这个坐标是在一个 $Dx \times Dy \times Dz$ 的三维空间中，其中 $Dx$ 是 x 轴的大小，( $Dy$ ) 是 y 轴的大小，( $Dz$ ) 是 z 轴的大小。

现在我们想把这个三维坐标映射到一个一维坐标 ( $id$ ) 上。我们采用以下方法：
1. 考虑 z 轴的影响：当我们在 $z$ 轴上移动时，每次移动 ( $Dx \times Dy$ ) 步，因为每个 (z) 的单位步长对应了一个 ( $Dx \times Dy$ ) 的平面。所以，( $z$ ) 轴的影响是 ( $Dx \times Dy \times z$ )。
2. 考虑 y 轴的影响：在已经考虑了 ( $z$ ) 轴的情况下，每次在 $y$ 轴上移动一步，对应了 ( $Dx$ ) 个点，因为 (y) 的每个单位步长对应了一个 ( $Dx$ ) 的行。所以，( $y$ ) 轴的影响是 ( $Dx \times y$ )。
3. 考虑 x 轴的影响：最后，在已经考虑了 ( $z$ ) 和 ( $y$ ) 轴的情况下，每次在 $x$ 轴上移动一步，对应了一个点。所以，( $x$ ) 轴的影响是 ( $x$ )。
将上述三个影响相加，我们得到了将三维坐标映射到一维坐标的公式：

[ $id = Dx \times Dy \times z + Dx \times y + x$ ]

这个公式允许唯一地将三维空间中的每个点映射到一个一维坐标上。

（4）以下是一些示例

1D grid, 1D block
- blockSize = blockDim.x
- blockId = blockIdx.x
- threadId = threadIdx.x
Id = blockIdx.x * blockDim.x + threadIdx.x （公式1）
3D grid, 1D block
- blockSize = blockDim.x（一维 block 的大小）
- blockId = Dx * Dy * z + Dx * y + x （三维 grid 中 block 的 id，用公式）= gridDim.x * gridDim.y * blockIdx.z + gridDim.x * blockIdx.y + blockIdx.x
- threadId = threadIdx.x （一维 block 中 thread 的 id）
Id = (gridDim.x * gridDim.y * blockIdx.z + gridDim.x * blockIdx.y + blockIdx.x ) * blockDim.x + threadIdx.x
1D grid, 2D block
- blockSize = blockDim.x * blockDim.y（二维 block 的大小）
- blockId = blockIdx.x（一维 grid 中 block id）
- threadId = Dx * y + x （二维 block 中 thread 的 id）
  = blockDim.x * threadIdx.y + threadIdx.x
Id = blockIdx.x * (blockDim.x * blockDim.y) + blockDim.x * threadIdx.y + threadIdx.x
3D grid, 3D block
- blockSize = blockDim.x * blockDim.y * blockDim.z（三维 block 的大小）
- blockId = Dx * Dy * z + Dx * y + x（三维 grid 中 block 的 id，用公式）
  = gridDim.x * gridDim.y * blockIdx.z + gridDim.x * blockIdx.y + blockIdx.x
- threadId = Dx * Dy * z + Dx * y + x（三维 block 中 thread 的 id，用公式）
  = blockDim.x * blockDim.y * threadIdx. z + blockDim.x * threadIdx.y + threadIdx.x
Thread ID = (gridDim.x * gridDim.y * blockIdx.z + gridDim.x * blockIdx.y + blockIdx.x) * (blockDim.x * blockDim.y * blockDim.z) + blockDim.x * blockDim.y * threadIdx. z + blockDim.x * threadIdx.y + threadIdx.x（公式2）
公式2为终极公式；坐标从0开始；维度从1开始；1D时，yz坐标为0，yz的维度为1,代入上式，即可得公式1；

1..2 host 、device

host 指的是 CPU 及其内存，device 指的是 GPU 及其内存。

host 是 CUDA 程序的执行环境。 CUDA 程序的编写和调试都在 host 上进行。host 上的内存称为 host 内存，通常用于存储程序代码、数据结构和临时数据。

device 是 CUDA 程序的执行目标。 CUDA 程序中的计算和并行操作在 device 上执行。device 上的内存称为 device 内存，通常用于存储大规模数据和计算结果。

host 和 device 之间的数据传输是 CUDA 程序的重要部分。 CUDA 程序需要将数据从 host 内存复制到 device 内存，然后在 device 上执行计算，最后将计算结果从 device 内存复制回 host 内存。

以下是 host 和 device 的一些主要区别：

属性	host	device
执行环境	CPU	GPU
内存	host 内存	device 内存
计算能力	有限	强大
数据传输速度	慢	快

在编写 CUDA 程序时，需要注意 host 和 device 之间的区别。例如，在将数据从 host 内存复制到 device 内存时，需要注意数据类型和大小，以避免数据丢失。

‍

1.3 总结

（1）主机： CPU 及其内存

（2）设备： GPU 及其内存

（3）线程(Thread)：一般通过GPU的一个核进行处理;

（4）线程块(Block)：由多个线程组成；各block是并行执行的，block间无法通信，也没有执行顺序。

（5）线程格(Grid)：由多个线程块组成。

（6）核函数(Kernel)：在GPU上执行的函数通常称为核函数；一般通过标识符__global__修饰，调用通过<<<参数1,参数2>>>，用于说明内核函数中的线程数量，以及线程是如何组织的。

2 代码示例注释

2.1 main.cu

#include <iostream>
#include <cuda_runtime.h> //用于声明 CUDA 运行时 API 的函数和宏的头文件。CUDA 运行时 API 提供了与 CUDA 硬件进行交互的接口
#include <stdio.h>


__global__ void print_idx_kernel(){
    printf("block idx: (%3d, %3d, %3d), thread idx: (%3d, %3d, %3d)\n",
           blockIdx.z, blockIdx.y, blockIdx.x,threadIdx.z, threadIdx.y, threadIdx.x);
}

__global__ void print_dim_kernel(){
    printf("grid dimension: (%3d, %3d, %3d), block dimension: (%3d, %3d, %3d)\n",
           gridDim.z, gridDim.y, gridDim.x,
           blockDim.z, blockDim.y, blockDim.x);
}

__global__ void print_thread_idx_per_block_kernel(){
    int index = threadIdx.z * blockDim.x * blockDim.y + \
              threadIdx.y * blockDim.x + \
              threadIdx.x;

    printf("block idx: (%3d, %3d, %3d), thread idx: %3d\n",
           blockIdx.z, blockIdx.y, blockIdx.x,
           index);
}

__global__ void print_thread_idx_per_grid_kernel(){
    int bSize  = blockDim.z * blockDim.y * blockDim.x;

    int bIndex = blockIdx.z * gridDim.x * gridDim.y + \
               blockIdx.y * gridDim.x + \
               blockIdx.x;

    int tIndex = threadIdx.z * blockDim.x * blockDim.y + \
               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);
}

__global__ void print_cord_kernel(){
    int index = threadIdx.z * blockDim.x * blockDim.y + \
              threadIdx.y * blockDim.x + \
              threadIdx.x;

    int x  = blockIdx.x * blockDim.x + threadIdx.x;
    int y  = blockIdx.y * blockDim.y + threadIdx.y;

    printf("block idx: (%3d, %3d, %3d), thread idx: %3d, cord: (%3d, %3d)\n",
           blockIdx.z, blockIdx.y, blockIdx.x,
           index, x, y);
}

void print_one_dim() {
    // 定义输入的大小
    int inputSize = 8;

    // 定义每个块的线程数
    int blockDim = 4;

    // 计算网格数
    int gridDim = inputSize / blockDim;

    // 定义块的维度
    dim3 block(blockDim);

    // 定义网格的维度
    dim3 grid(gridDim);

    // 注释：调用 `print_idx_kernel()` 函数，打印每个线程的索引
    print_idx_kernel<<<grid, block>>>();

    // 注释：调用 `print_dim_kernel()` 函数，打印块的维度
    // print_dim_kernel<<<grid, block>>>();

    // 注释：调用 `print_thread_idx_per_block_kernel()` 函数，打印每个线程在块内的索引
    // print_thread_idx_per_block_kernel<<<grid, block>>>();

    // 注释：调用 `print_thread_idx_per_grid_kernel()` 函数，打印每个线程在网格内的索引
    // print_thread_idx_per_grid_kernel<<<grid, block>>>();

    // 同步设备
    cudaDeviceSynchronize();
}


void print_two_dim(){
    int inputWidth = 4;

    int blockDim = 2;
    int gridDim = inputWidth / blockDim;

    dim3 block(blockDim, blockDim);
    dim3 grid(gridDim, gridDim);

    /* 这里建议大家吧每一函数都试一遍*/
    // print_idx_kernel<<<grid, block>>>();
    // print_dim_kernel<<<grid, block>>>();
    // print_thread_idx_per_block_kernel<<<grid, block>>>();
    print_thread_idx_per_grid_kernel<<<grid, block>>>();

    cudaDeviceSynchronize();
}

void print_cord(){
    int inputWidth = 4;

    int blockDim = 2;
    int gridDim = inputWidth / blockDim;

    dim3 block(blockDim, blockDim);
    dim3 grid(gridDim, gridDim);

    print_cord_kernel<<<grid, block>>>();

    cudaDeviceSynchronize();
}

int main() {
    /*
    synchronize是同步的意思，有几种synchronize

    cudaDeviceSynchronize: CPU与GPU端完成同步，CPU不执行之后的语句，知道这个语句以前的所有cuda操作结束
    cudaStreamSynchronize: 跟cudaDeviceSynchronize很像，但是这个是针对某一个stream的。只同步指定的stream中的cpu/gpu操作，其他的不管
    cudaThreadSynchronize: 现在已经不被推荐使用的方法
    __syncthreads:         线程块内同步
    */
     print_one_dim();
//     print_two_dim();
//    print_cord();
    return 0;
}

（1）#include <cuda_runtime.h>

cuda_runtime.h是一个头文件，它包含了运行时API和其参数的定义。运行时API是一组函数，它们可以在主机（CPU）和设备（GPU）之间进行数据传输、内存管理、核函数调用等操作。其参数的定义是一些常量、类型、结构体、枚举等，它们可以帮助您指定运行时API的输入和输出。
cuda_runtime.h头文件的作用是可以使用CUDA的运行时API来编写基于GPU的并行计算程序。

（2）global void print_cord_kernel() 与print_cord_kernel<<<grid, block>>>() 声明和调用

global void print_cord_kernel() 是 CUDA 内核函数的声明，它告诉编译器该函数将在 GPU 上执行。
print_cord_kernel<<<grid, block>>>(); 是 CUDA 内核函数的调用，它告诉 CUDA 运行时系统将 print_cord_kernel() 函数调用发送到 GPU 上执行。
print_cord_kernel()函数的命名方式符合 CUDA 内核函数的命名规范。CUDA 内核函数的名称必须以 global 关键字开头，并以 kernel 结尾。

‍