CUDA 编程模型为 GPU 架构提供了一个抽象概念，是应用程序与GPU硬件之间的一个桥梁，本文稍微深入理解CUDA编程模型，体会设计的精妙之处，了解c/c++语言如何与CUDA交互。

首先了解两个概念：

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

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

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

• • 将输入数据从host内存复制到device内存，也称为host到device传输。

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

1：CUDA kernel和线程层次结构

CUDA kernel是一个在GPU 上执行的函数，应用程序中的并行部分是由 K 个不同的 CUDA 线程并行执行 K 次，而不是像普通的 C/C++ 函数那样只执行一次，下图是CUDA kernel和threads的关系：

CUDA kernel以 global 声明指定符开始，使用内置变量为每个线程提供唯一的全局 ID，一组线程组成一个CUDA block，CUDA blocks组成一个grid，如下图：

CUDA和GPU也存在一一对应的关系，如下图：

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

接下去具体说说threads和blocks。

在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都是三维的，但可以根据具体的计算需求灵活地使用其中的维度，不必在所有情况下都完整使用三维索引。

为什么设计为3D？CUDA设计为三维索引的原因之一是为了提供灵活性和适应性，特别是在处理多维数据结构时。在人工智能（AI）和科学计算中，常见的数据结构包括向量（vectors）、矩阵（matrices）和体积（volumes）。三维索引使得这些多维数据结构的处理更加自然和高效

下面举个矩阵相加的操作：

// 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的数量也是根据维度计算的。

为了更清晰的理解，下面的代码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启动

2：内存层次结构

内存也是有层次结构的，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就有了一定理解。