CUDA10.0官方文档的翻译与学习之编程模型

背景

在文章CUDA10.0官方文档的翻译与学习之介绍中我翻译了CUDA10.0官方文档中的简介部分，这里书接前文，翻译第二章——编程模型

这一章介绍一些cuda编程模型背后的主要概念，以及勾勒出他们如何以C的形式表达，cuda的C接口的集中介绍请参见下一章编程接口。本章和下一章使用的向量加法样例源码可以在cuda的vectorAdd样例中得到

核函数

cudaC通过允许程序员定义被称之为核函数的C函数来扩展C语言，相比普通C语言，核函数被调用时会在N个不同的cuda线程中并行执行N次(一个线程执行一次)。

一个核函数使用__global__标识符声明，调用它时，要指定执行这个核的线程数量，并且用<<<...>>>包裹起来。每个执行核函数的线程被分配了一个全局唯一的线程id，这个id可以在核函数内使用内置的threadIdx遍历获取。为了阐释，下面的例子累加了两个N维向量A和B，并且把结果存到了C中

这里，N个线程(1表示用1个线程块)中的每一个都会执行一次核函数VecAdd()，来进行成对的相加

__global__ void VecAdd(float* A, float* B, float* C) {
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}

int main() {
    .....
    VecAdd<<<1, N>>>(A, B, C); // 用N个线程执行核函数
    ....

}

线程层次

为了方便，threadIdx是一个三维向量，因此线程可以通过使用一维、二维或三维的线程索引来表示，对应的就是一维、二维或三维的成块的线程，亦称之为线程块。这种表示方法让我们对元素的计算在一种类似于向量、矩阵和立方体的空间中进行。

线程索引和线程id以一种直接的方式相互关联：对于一维的块来说，这俩直接向等；对于二维的块(Dx, Dy)来说，索引为(x, y)的线程的线程id等于x + y * Dx；对于三维的块(Dx, Dy, Dz)来说，索引为(x, y, z)的线程的线程id等于x + yDx + zDxDy。下面的例子将两个N*N矩阵A和B相加，结果存到C里

__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() {
    ...
    int numBlocks = 1;
    dim3 threadsPerBlock(N, N);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

因为一个块中的线程应该存在于一个处理器上，并且必须共享其上的有限内存资源，所以每个块的线程数量应该有限。在现有的GPU上，一个线程块最多可以包含1024个线程。然而，一个核函数可以被多个形状相同的线程块执行，所以线程总数 = 每个块上的线程数 * 块数。线程块可以以一维、二维、三维的形式组成网格(如下图所示)，一个网格中的线程块数通常由正在处理的数据量或处理器数量决定，但线程块数可以远超过这两个值

每个块的线程数和每个网格的线程块数用<<<...>>>中的两个参数决定，类型可以是int或dim3，如上面两个例子所示。网格中的每个块可以由一维、二维、三维索引来指定，索引可以在核内通过内置的blockIdx变量获取，线程块的维度可以在核内通过内置的blockDim变量获取，处理多核块的加强版MatAdd()样例如下所示

__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;

    if (i < N && j < N) {
        C[i][j] = A[i][j] + B[i][j];
    }
}

int main() {
    ...
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
}

线程块的大小是16 * 16，这也是通常的选择，然后用足够多的线程块来创建网格，以让每个线程处理一个矩阵元素。为了简便，此例中单元格在每个维度上的线程数可以被每个块在那个维度上的线程数整除，但这不是必须的要求。线程块被要求独立执行，也就是必须可以以任何顺序串行或并行执行它们，这种独立性要求允许线程块可以在任何处理器上以任何顺序执行(如下图所示)，也允许程序员根据内核数写代码。

一个块内的线程可以通过共享位于一块叫做共享内存的内存区域中的数据来进行合作，并且通过同步执行来协调内存访问。准确来说，我们可以在核函数内部调用__syncthreads()函数来指定同步点，__syncthreads()函数就像一个栅栏一样，所有块内的线程必须等待直到放行信号发出。在共享内存一节会给出使用共享内存的一个例子，除了__syncthreads()之外，官方手册中的合作组部分还提供了大量的线程同步原语。

为了有效合作，共享内存应该是一个位于处理器内部、距离每个处理器核都很近的低延迟内存(就像L1缓存一样)，而且__syncthreads()函数应该是轻量级的

内存层次

cuda线程可以在执行时从多个内存空间中访问数据，如下图所示

每个线程都有自己的局部内存，每个线程块都有对块内所有线程可见的共享内存，并且共享内存与线程块同生命周期，所有线程都可以访问一个全局内存。除此之外，还有两个额外的只读内存，同样可以被所有线程访问：常量内存和纹理内存。全局、常量、纹理内存可以为不同的内存访问分别做优化，请参见设备内存访问一节；纹理内存也提供不同的取址模式，还为某些特定的数据格式进行数据过滤，请参见纹理和表面内存一节。最后，全局、常量和纹理内存在同一个应用的内核启动过程中是一直存在的。

异构编程

cuda编程模型假设cuda线程执行在一个单独的物理设备上，这个设备是运行C程序的主机的合作者，如下图所示。在这种情况下，核函数执行在GPU上，剩下的C程序执行在CPU上

cuda编程模型还假设，主机和设备都在内存中拥有独立空间，分别标识为主机内存和设备内存。因此，程序管理的全局、常量和纹理内存空间对核函数是可见的，调用cuda运行时函数即可进行设备内存的分配、回收、主机和设备内存之间的数据传输。统一内存提供了管理内存来连接主机和设备内存空间，管理内存作为一个单一、连续的内存镜像，对系统中的所有CPU和GPU都是可访问的，而且提供共同的地址空间，这一功能支持了设备内存的超额使用，并且可以通过消除主机和设备对镜像数据的明确需求来极大地简化应用的移植工作。

计算能力

设备的计算能力表示为一个版本号，有时也称之为SM版本。这个版本号标识了被GPU硬件支持的特征，被用来让应用在运行时判断哪种硬件特征或者指令在当前GPU上是可用的，另外此版本号由主版本号X和副版本号Y组成，表示为X.Y。

拥有主版本号的设备有同样的内核架构，主版本号为7的设备是基于Volta架构的，6则是基于Pascal架构的，5则是Maxwell架构，3、2、1分别基于Kepler、Fermi和Tesla架构；副版本号则是跟内核架构的一些增量提升相关的，有时包含新的特征。Turing时计算能力7.5设备的架构，也是基于Volta的一个增量更新，另外，Tesla和Fermi架构分别从cuda7.0和cuda9.0开始就不再支持了。

结语

这一章介绍了核函数、GPU中的线程层次和内存层次、异构编程、计算能力等内容。

核函数就是CPU调用、GPU执行的函数，<<<>>>内可以配置要执行这个核函数的线程数量(线程网格维度和线程块维度)，以及后面章节会介绍的动态分配的共享内存数和流号。我们可以在核函数中通过内置变量获取当前的线程索引，并计算出全局索引，以进行数据的全局访问；

线程层次包括线程单元格、线程块和线程，它们都位于一个流式处理器上，其可配置数量受限于流式处理器的参数；

内存层次包括全局内存、纹理内存、常量内存、共享内存和局部内存。前三者位于全局GPU内，应用程序随时可以访问，全局内存可以读写，纹理内存要通过纹理对象进行读写，常量内存则常用于缓存被经常访问的常量，以减少请求数；共享内存是片上的，存在于每个线程块中，可以在核函数中静态分配也可以在核函数的执行配置(<<<>>>)中动态分配，生命周期和线程块相同；局部内存位于每个线程内，访问速度是最快的，生命周期和线程相同。

异构编程部分说明了主机代码和设备代码、主机内存和显存的关系。执行核函数之前，一般来说我们都要把输入数据从内存复制到显存中，调用核函数后要把结果复制回来，因为内存和显存是彼此独立且不可见的。复制可以是同步或异步，异步就需要调用同步函数来同步。但如果使用管理内存的话，就不用显式地迁移数据了。

计算能力就是流式处理器(或者显卡)的版本标识，显然不同的大版本有着不一样的硬件特征和资源限制，不同的小版本之间差别就要小一些，至少内核架构是一样的。我们也要注意不同主版本设备的架构名称(比如Volta就对应7.X的设备)，以避免别的文档说起时，我们对应不上。

从下一章开始，我们就详细介绍编程接口，这是这篇文档的主力，内容将更多。