|
CUDA 编程模型为 GPU 架构提供了一个抽象概念，是应用程序与GPU硬件之间的一个桥梁，本文稍微深入理解CUDA编程模型，体会设计的精妙之处，了解c/c++语言如何与CUDA交互。
首先了解两个概念： 为了执行一个CUDA程序，包含三个步骤： 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）上进行索引和组织。 虽然blocks和threads都是三维的，但可以根据具体的计算需求灵活地使用其中的维度，不必在所有情况下都完整使用三维索引。 为什么设计为3D？CUDA设计为三维索引的原因之一是为了提供灵活性和适应性，特别是在处理多维数据结构时。在人工智能（AI）和科学计算中，常见的数据结构包括向量（vectors）、矩阵（matrices）和体积（volumes）。三维索引使得这些多维数据结构的处理更加自然和高效 下面举个矩阵相加的操作： //Kernel-AddingtwomatricesMatAandMatB
__global__voidMatAdd(floatMatA[N][N],floatMatB[N][N],
floatMatC[N][N])
{
inti=blockIdx.x*blockDim.x+threadIdx.x;
intj=blockIdx.y*blockDim.y+threadIdx.y;
if(i<N&&j<N)
MatC[i][j]=MatA[i][j]+MatB[i][j];
}
intmain()
{
...
//Matrixadditionkernellaunchfromhostcode
dim3threadsPerBlock(16,16);
dim3numBlocks((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__voidMatAdd(floatMatA[N][N],floatMatB[N][N],floatMatC[N][N])
{
//计算一维索引
intidx=blockIdx.x*blockDim.x+threadIdx.x;
inti=idx/N;//行
intj=idx%N;//列
if(i<N&&j<N)
MatC[i][j]=MatA[i][j]+MatB[i][j];
}
intmain()
{
...
//使用一维block配置
intthreadsPerBlock=256;//一个block用256个线程
intnumBlocks=(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就有了一定理解。 | 
