CUDA编程概念小记

很久没有更新博客了，最近打算快速过一下GPU-MODE这个课程，重新系统性地学习一下一些CUDA的概念，所以可能会断断续续更新一些博客来记录一些我个人认为值得记录的概念。

语言层

众所周知，CUDA编程模型中有 grid - block -thread 三层概念，而启动的时候却只要两个<<<GridDim,BlockDim>>>这样子的参数就可以了，这件事情一度比较困扰我：Grid是做什么用的？

现在看来Grid可能是设计来用于多kernel launch或者在多GPU上launch的东西？就现在的编程实践来看都是只有一个Grid的。

所以这个GridDim其实是有多少个Block的意思，BlockDim其实是每一个Block里面有多少个Thread的意思。从这个角度来理解对我个人而言就会通畅很多。

而这两个接受的都是dim3类型的变量，其实我认为也把他统一理解为offset，这样子会更利于编程理解。其实就是一个(x,y,z)的坐标，可以把它展开成z*Dim_y*Dim_x+y*Dim_x+x的一维的形式。本质上是编程语言层给你设计了一些等价的多维的表示方式，在处理矩阵/张量数据的时候可能编程上更易读一点。但对于我这个常年写c的人来说其实按照offset的方式理解会更加舒服。

那么如何理解kernel函数内的执行逻辑呢？让我们通过一个简单的向量加法示例来理解：

// 向量加法的kernel函数
__global__ void vectorAdd(float* a, float* b, float* c, int n) {
    // 计算当前线程的全局索引
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 确保不会访问越界
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    ...
    
    // 设置kernel启动参数
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    
    // 启动kernel
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);
    
    ...
}

在这个例子中，我们可以看到几个重要的概念：

1. kernel函数中通过blockIdx.x和threadIdx.x来获取当前线程的位置信息
2. 通过blockDim.x（每个block中的线程数）来计算全局索引
3. 启动参数<<<blocksPerGrid, threadsPerBlock>>>决定了并行的规模
4. 需要考虑数组边界检查，因为线程数可能会超过数组大小

可能对于初学者不太好理解CUDA编程的kernel执行逻辑，但在当前这个简单的场景下，一个我认为比较好的理解方式是（如果你学习过OpenMP）的话：其实就是一个没有for的OpenMP，然后利用kernel函数固定给你传入的线程索引（blockIdx和threadIdx）来实现下标计算，然后每个线程在自己所需要对应的那个下标执行相同的操作。

硬件层

另外一个伴随cuda编程经常提到的一个概念是：不要轻易在kernel函数里面使用if语句，会把并行执行变成串行执行。但究竟是为什么会这样子，就要从GPU的硬件结构上来理解了。

GPU的基本计算单元是流式多处理器（Streaming Multiprocessors, SMs），而SM的基本计算单元是warp。warp是SM中一个非常重要的概念，它是GPU执行并行计算的基本单位。一个warp包含32个并行执行的线程，这些线程在SIMT（Single Instruction, Multiple Threads）模式下执行相同的指令。

特别需要注意的点是warp是一个硬件概念，而不是编程概念。

一个warp中的线程会做同一件事，但是if条件意味着你会让一个warp中的线程执行不同的指令，这和硬件的实际情况不符合，但它仍然能执行，代价就是串行执行true部分的指令，再串行执行false部分的指令。（插个题外话，在帕斯卡架构前是这样的，伏特架构后每个线程有自己的PC，可以交替执行掩盖延迟，比如一些有数据依赖的load-then-compute的，就可以先load然后掩盖一下延迟，这种情况会好一点）

所以其实不是一定不能写if，如果你的if条件对于一个warp内的所有线程都是一样的，那么这个if语句就不会有什么问题。

condition = ...;
if(condition is true) {
    A;  
    B;
} else {
    C;
    D;
}

下面这个伪代码里面的A,B和C,D就会被分开执行，没办法并行执行（如果一个warp内的condition又有true又有false的话）

而前面提到的延迟掩盖，可以这么理解：比如A是从global memory载入到shm，然后B再对shm的数据读取进行计算，那你交替执行A和C就可以让不同的硬件先忙起来，通过一点小流水线掩盖掉load的延迟。（这个和CPU的流水线的概念其实是类似的）

同步

CUDA的同步其实分为很多层，给了你很细粒度的同步控制。

1. 同一个block内的线程同步：__syncthreads()
2. 同一个warp内的线程同步：shfl_sync()

然后还有一个不同kernel之间（以及host和device之间）的同步：cudaDeviceSynchronize()

所以其实可以在不同粒度上进行同步，以满足不同的需求，尽可能地减少不被重叠的时间。

• 本文作者： henry_y
• 本文链接： http://henry-y.github.io/2024/12/06/cuda编程概念小记/
• 版权声明： 本博客所有文章除特别声明外，均采用 MIT 许可协议。转载请注明出处！