深入理解CUDA编程模型

深入理解CUDA编程模型

CUDA（Compute Unified Device Architecture）是NVIDIA开发的一种并行计算平台和编程模型，它允许开发者利用GPU的强大计算能力来加速计算密集型应用程序。CUDA编程模型为GPU架构提供了一个抽象概念，是应用程序与GPU硬件之间的一个桥梁。本文将深入探讨CUDA编程模型的核心概念，包括host和device的定义、CUDA程序执行的三个主要步骤、CUDA kernel和线程层次结构、以及GPU的内存层次结构。

基本概念

在CUDA编程中，有两个核心概念需要理解：

• host：host是系统中的CPU，与CPU对应的系统内存也叫host内存
• device：device是系统中的GPU，对应的内存也叫做Device内存

为了执行一个CUDA程序，通常包含三个步骤：

1. 将输入数据从host内存复制到device内存，也称为host到device传输。
2. 加载 GPU 程序并执行，在芯片上缓存数据以提高性能。
3. 将结果从device内存复制到host内存，也称为device到host传输。

CUDA Kernel和线程层次结构

CUDA kernel是一个在GPU 上执行的函数，应用程序中的并行部分是由 K 个不同的 CUDA 线程并行执行 K 次，而不是像普通的 C/C++ 函数那样只执行一次。CUDA kernel以 global 声明指定符开始，使用内置变量为每个线程提供唯一的全局 ID。

CUDA的线程层次结构可以分为以下几个层次：

• CUDA block：一组线程组成一个CUDA block
• CUDA grid：CUDA blocks组成一个grid

每个CUDA block由一个SM（Streaming Multiprocessor）处理器执行，不能由其它SM执行；一个SM可以运行多个CUDA块，具体取决于CUDA blocks需要的数量；每个CUDA kernel在一个device上运行，CUDA 支持同时在一个device上运行多个cuda kernel。

在CUDA中，threads（线程）和blocks（块）都被定义为三维变量。你可以把blocks想象成一个工厂的厂区，而threads则是这个厂区中的房间。每个block和thread都可以在三个维度（x、y、z）上进行索引和组织。

• 块索引（blockIdx）：用于在网格中索引块。它有三个分量：blockIdx.x、blockIdx.y 和 blockIdx.z。
• 线程索引（threadIdx）：用于在块中索引线程。它也有三个分量：threadIdx.x、threadIdx.y 和 threadIdx.z。

虽然blocks和threads都是三维的，但可以根据具体的计算需求灵活地使用其中的维度，不必在所有情况下都完整使用三维索引。

矩阵相加示例

下面是一个矩阵相加的CUDA kernel示例：

// Kernel - Adding two matrices MatA and MatB
__global__ void MatAdd(float MatA[N][N], float MatB[N][N],
float MatC[N][N])
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i < N && j < N)
        MatC[i][j] = MatA[i][j] + MatB[i][j];
}

int main()
{
    ...
    // Matrix addition kernel launch from host code
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks((N + threadsPerBlock.x -1) / threadsPerBlock.x, (N+threadsPerBlock.y -1) / threadsPerBlock.y);
    MatAdd<<<numBlocks, threadsPerBlock>>>(MatA, MatB, MatC);
    ...
}

在这个例子中，blocks是2D的，有16 * 16个线程，对应blocks的x和y方向。blocks的数量也是根据维度计算的。

为了更清晰的理解，下面的代码blocks就使用了1D：

__global__ void MatAdd(float MatA[N][N], float MatB[N][N], float MatC[N][N])
{
    // 计算一维索引
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    int i = idx / N;  // 行
    int j = idx % N;  // 列
    
    if (i < N && j < N)
        MatC[i][j] = MatA[i][j] + MatB[i][j];
}

int main()
{
    ...
    // 使用一维block配置
    int threadsPerBlock = 256;  // 一个block用256个线程
    int numBlocks = (N * N + threadsPerBlock - 1) / threadsPerBlock;
    MatAdd<<<numBlocks, threadsPerBlock>>>(MatA, MatB, MatC);
    ...
}

需要注意的几个关键点

• CUDA机构会限制每一个blocks占用的线程总数
• 每一个blocks的维度通过内建变量blockDim获取
• 一个blocks中的线程可以使用__syncthreads函数同步
• <<...>>>语法表示从host代码到device代码的调用，也叫kernel启动

内存层次结构

GPU的内存层次结构设计使得数据能够尽可能地靠近计算单元，从而提高计算效率。下图展示了A100 GPU从SM维度的内存层次结构：

下图是从blocks和threads的维度描述内存层次结构：

• Registers：Registers对每个线程是私有的，它也是内存层次结构中速度最快的memory，用于存储本地变量和中间结果。Registers的数量是有限的，提高Registers的重复利用率是GPU的关键所在。
• L2 cache：L2 cache用于cache数据，由GPU硬件管理，用于存储频繁访问的数据，以减少内存延迟，从而提高整体性能。所有SM共享L2 cache。
• L1/Shared memory (SMEM)：每个SM有一个片上内存区域，在一个block上的线程可以使用共享内存，线程可以通过共享内存通信和同步；而L1 cache和L2 cache类似，用于cache 一个SM上的数据；这个区域的内存也是非常快的。
• Read-only memory：每个SM还有一个只读内存，用于cache 指令、常量等，对cuda kernel只读。
• Global Memory：最大的内存池，即HBM，所有线程都能访问它，但也是GPU内存层次结构中最慢的内存，输入数据和输出数据一般都在Global Memory中。
• Local Memory：这块内存是由OS系统给每个线程分配的，当线程变量的寄存器空间不足时，本地内存用于存储临时变量或溢出数据。

理解线程层次结构和内存层次结构，对GPU和CUDA就有了一定理解。