话说 CUDA 的软件模型和硬件模型

CUDA 的技术体系可划分为软件模型和硬件模型两部份，充分理解这两部分的体系结构和相互关系，是掌握 CUDA 技术的关键。

一、CUDA 的软件模型

1. 两层结构的线程管理

毫无疑问， CUDA 软件模型是一个海量线程（Thread）管理工具。CUDA 中的线程按照两个层次分组：

• Block：一个 Block 最多管理 512 个Thread。Thread 对应的函数称为 Kernel 函数，本质上，CUDA 以 Block 为基本单位提交代码运行。也就是说，CUDA 不会管理单独的 Thread，逻辑上看，相当于一次运行了一个 Block 的 Thread。
• Grid：一个 Grid 可以管理任意数量的 Block。（目前一个 Kernel 函数只能在一个 Grid 范围内工作。）

Kernel 函数可以查询到自己所在的 Grid 索引、Block 索引、Thread 索引，根据这些索引，就能正确地从多维数组中访问自己需要的数据。这个机制把海量数据的并行处理分解到了每一个线程。

2. 三维索引

为方便处理向量、矩阵、张量等数据类型，CUDA 中 Thread 的索引采用了 dim3 结构，这样简化了访问多维数组时计算各个分量下标的工作量，提高了计算效率。参见下面的例子：

//Kernel 函数
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]){
    int i = threadIdx.x;
    int j = threadIdx.y;
    C[i][j] = A[i][j] + B[i][j];
}

int main(){
    ...
    dim3 dimBlock(N, N);
    MatAdd<<<1, dimBlock>>>(A, B, C);
    ...

程序中的 dim3 相当于 3 个unsigned int 类型组成的结构体，方便使用一维、二维、三维的索引来标识线程。上面程序中，dimBlock(N, N) 定义了二维标识，处理矩阵这样的二维阵列运算非常方便。

上面例子中用了 1 个 Block，N*N 个 Thread。如果 N 太大，必须使用多个 Block 解决，这不是本文重点，不再赘述。

3. 为什么线程要分层？

答案很简单，Block 之间不能共享内存，但是 Block 内部的 Thread 可以共享内存。因此两层架构使得 CUDA 既可以开发出不共享内存的并行算法，也可以编写共享内存的并行算法。

当然，一个 Block 内部最多只有 512 个 Thread，因此一块内存最多只能被 512 个 Thread 共享。尽管由此限制，能够共享内存这一特征，能在关键时候解决大问题。

4. 共享内存

直接看一段例子吧

__global__ void testShareMemory(float* g_idata, flowt* g_odata){
    //extern 标识数组大小由 host 端的 Ns 参数确定    
    extern __shared__ float sdata[]
    ...
    sdata[tid_in_block] = g_idata[tid_in_grid];
    __syncthreads();
    ...
}

5. 软件模型的好处

上面讲到的 Thread、Block、Grid 等概念，并非英伟达 GPU 硬件中的物理实体，属于纯软件的定义。在这个体系下，可以把算法设计和硬件彻底隔离开，做到编写一次，可在各种 GPU 硬件板卡上运行。

需要牢记一点，我们编写的 Kernel 函数，CUDA 是按照 Block 提交执行的，同一个 Block 内的 Thread 是可以共享内存的。接下来考虑 CUDA 的硬件调度策略，我们需要牢记硬件是按照 Block 来分配资源的。

二、CUDA 的硬件模型

1. SM 处理器

英伟达的 GPU 是由一大堆 SM 处理器构成的。细节我不想多讲，讲多了容易冲淡本文的主题，以后我会专门写文章介绍细节问题。

CUDA 提交 Block 的线程给一个 SM 处理器，如果 Block 数量大于 SM 处理器的数量，就会出现多个 Block 排队等待 SM 处理器的现象。这个很有意思，因为有些 Block 可能因为访问显存等操作时待系统回应被“挂起”，这时其他等待中的 Block 可以趁虚而入，让 SM 执行其代码，这样可以大大提高运算效率。

其实了解这么多几乎就够一般性的使用了，CUDA 的硬件模型看起来足够简单。

2. SP 核

事实上英伟达 GPU 的硬件体系底层还是很复杂的，本文简单多讨论一点，说一下 SP 核。

一个 Block 最多可以有 512 个 Thread，那么一个 SM 处理器可以同时并行执行几个 Thread 呢？答案是 32 个，因为一个 SM 处理器内部包含了 8 个 SP 核（参见下图），每个SP会重复执行4次指令。当然，这个细节在 CUDA 软件模型中被屏蔽掉了，现阶段我们可以完全不考虑这一细节。事实上，还有更多的底层细节，在今后的高级程序设计阶段，我们可以进一步研究，以便针对具体GPU板卡开发更有效率的算法。

3. 软硬件体系的协作过程

简单看一张图吧：

今天就先到这里，谢谢阅读！