CUDA的编程模式

        本章介绍了CUDA编程模型背后的主要概念，概述了它们在C++中是如何公开的。在编程接口中对CUDA C++进行了详细的描述。

        本章和下一章中使用的矢量加法示例的完整代码可以在矢量加法CUDA示例中找到。

2.1、内核

        CUDA C++通过允许程序员定义称为内核的C++函数来扩展C++，这些函数在被调用时由N个不同的CUDA线程并行执行N次，而不是像常规C++函数那样只执行一次。

        内核是使用__global__声明说明符定义的，并且为给定内核调用执行该内核的CUDA线程数是使用新的<<…>>指定的执行配置语法（请参阅C++语言扩展）。执行内核的每个线程都有一个唯一的线程ID，可以通过内置变量在内核中访问该ID。如图所示，以下示例代码使用内置变量threadIdx，将大小为N的两个矢量A和B相加，并将结果存储到矢量C中：

// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
 int i = threadIdx.x;
 C[i] = A[i] + B[i];
}
int main()
{
 ...
 // Kernel invocation with N threads
 VecAdd<<<1, N>>>(A, B, C);
 ...
}

        这里，执行VecAdd（）的N个线程中的每一个执行一对加法。 

2.2、线程层次结构

        为了方便起见，threadIdx是一个3个组件的向量，这样线程可以使用一维、二维或三维线程索引进行标识，形成一维、二维或三维线程块，称为线程块。这为在向量、矩阵或体积等域上执行计算提供了一种自然的方式。

        线程的索引和其线程ID之间的关系非常简单：对于一维块，它们是相同的；对于大小为(Dx, Dy)的二维块，线程索引为(x, y)的线程的线程ID是(x + y Dx)；对于大小为(Dx, Dy, Dz)的三维块，线程索引为(x, y, z)的线程的线程ID是(x + y Dx + z Dx Dy)。

        作为一个例子，以下代码将两个大小为NxN的矩阵A和B相加，并将结果存储在矩阵C中：

// Kernel definition
__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()
{
 ...
 // Kernel invocation with one block of N * N * 1 threads
 int numBlocks = 1;
 dim3 threadsPerBlock(N, N);
 MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
 ...
}

        每个块中的线程数量是有限的，因为一个块的所有线程都预期驻留在同一个处理器核心上，并且必须共享该核心的有限内存资源。在当前的GPU上，一个线程块可能包含多达1024个线程。然而，一个内核可以由多个形状相同的线程块执行，因此总线程数等于每个线程块中的线程数乘以线程块数。

        块组织成1D、2D或3D的线程块网格，如图6所示。网格中的线程块数量通常由正在处理的数据的大小决定，这通常超过系统中的处理器数量。

 

图6 螺纹块网格图6螺纹块网格

        <<<…>>中指定的每个块的线程数和每个网格的块数语法可以是int或dim3类型。二维块或网格可以如上面的示例中那样指定。网格中的每个块都可以通过一个一维、二维或三维的唯一索引来识别，该索引可通过内置的blockIdx变量在内核中访问。线程块的维度可以通过内置的blockDim变量在内核中访问。

        扩展前面的MatAdd（）示例以处理多个块，代码如下。

// Kernel definition
__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()
{
 ...
 // Kernel invocation
 dim3 threadsPerBlock(16, 16);
 dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
 MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
 ...
}

        16x16（256个线程）的线程块大小虽然在这种情况下是任意的，但却是常见的选择。网格是用足够的块创建的，与以前一样，每个矩阵元素有一个线程。为了简单起见，本示例假设每个维度中每个网格的线程数可以被该维度中每个块的线程数整除，但事实并非如此。

        线程块需要独立执行：必须能够以任何顺序、并行或串行执行。如图5所示，这种独立性要求允许在任何数量的内核上以任何顺序调度线程块，使程序员能够编写随内核数量而扩展的代码。

        块内的线程可以通过一些共享内存共享数据，并通过同步它们的执行来协调内存访问来进行协作。更准确地说，可以通过调用__syncthreads（）内部函数来指定内核中的同步点__syncthreads（）充当一个屏障，块中的所有线程都必须在该屏障处等待，然后才能继续执行任何线程。共享内存提供了一个使用共享内存的示例。除了__syncthreads（）之外，协作组API还提供了一组丰富的线程同步原语。

        为了高效协作，共享内存应该是每个处理器核心附近的低延迟内存（很像一级缓存），_syncthreads（）应该是轻量级的。 

2.3、内存层次

        CUDA线程在执行过程中可以访问来自多个内存空间的数据，如图7所示。每个线程都有专用的本地内存。每个线程块都共享对该块的所有线程可见的内存，并且与该块具有相同的生存期。

        所有线程都可以访问相同的全局内存。还有两个额外的只读内存空间可供所有线程访问：常量内存空间和纹理内存空间。全局、常量和纹理内存空间针对不同的内存使用进行了优化。纹理内存还为某些特定的数据格式提供了不同的寻址模式以及数据过滤。全局、常量和纹理内存空间在同一应用程序启动的内核之间是持久的。

2.4、异构编程

        如图8所示，CUDA编程模型假设CUDA线程在物理上独立的设备上执行，该设备作为运行C++程序的主机的协处理器运行。例如，当内核在GPU上执行，而C++程序的其余部分在CPU上执行时，就是这种情况。

        CUDA编程模型还假设主机和设备在DRAM中都有各自独立的内存空间，分别称为主机内存和设备内存。因此，程序通过调用CUDA运行时来管理内核可见的全局、常量和文本内存空间（如编程接口中所述）。这包括设备内存分配和释放，以及主机和设备内存之间的数据传输。

        统一内存提供托管内存以桥接主机和设备内存空间。托管内存可以从系统中的所有CPU和GPU访问，作为一个具有公共地址空间的统一内存映像。此功能实现了对设备内存的订阅，并通过消除在主机和设备上显式镜像数据的需要，大大简化了分配应用程序的任务。

串行代码在主机上执行，而并行代码在设备上执行。 

2.5、计算能力

        设备的计算能力由版本号表示，有时也称为“SM版本”。此版本号标识GPU硬件支持的功能，并由应用程序在运行时用于确定当前GPU上可用的硬件功能和/或指令。计算能力包括主要修订号X和次要修订号Y，并由X.Y表示。

        具有相同主要修订号的设备具有相同的核心体系结构。主要修订号是基于Volta架构的设备7，基于Pascal架构的设备6，基于Maxwell架构的设备5，基于Kepler架构的设备3，基于Fermi架构的设备2，以及基于Tesla架构的设备1。

        次要修订号对应于对核心体系结构的增量改进，可能包括新功能。图灵是计算能力为7.5的设备的架构，是基于Volta架构的增量更新。

        CUDA Enabled GPU列出了所有CUDA Enabled设备及其计算能力。计算能力给出了每种计算能力的技术规范。

        特定GPU的计算能力版本不应与CUDA版本（例如，CUDA 7.5、CUDA 8、CUDA 9）混淆，CUDA版本是CUDA软件平台的版本。CUDA平台被应用程序开发人员用来创建在许多代GPU架构上运行的应用程序，包括尚未发明的未来GPU架构。虽然CUDA平台的新版本通常通过支持新GPU架构的计算能力版本来增加对该架构的本地支持，但CUDA平台新版本通常还包括独立于硬件生成的软件功能。

        分别从CUDA 7.0和CUDA 9.0开始，不再支持特斯拉和费米体系结构。