背景
随着时钟频率的发展陷入停滞，集成更多的计算逻辑和计算核心成为了获取更高算力的主要途径。多核处理器可以视作多路平台的自然演化，而 GPGPU 的出现利用大规模并行架构为多核 CPU 难以解决的问题提供了颠覆性的解决方案。
现代计算机模糊了 Flynn 分类的界限，将 MIMD 与 SIMD 结合以获得更高性能。现代计算机的发展趋势主要有两种：一种是增加片上内核数目与缓存容量并结合专用 SIMD 指令集，另一种采用异构架构并分别处理不同任务。异构计算主要是指使用不同类型指令集和体系架构的计算单元组成系统的计算方式。常见的计算单元类别包含 CPU、GPU 等协处理器、DSP、ASIC、FPGA 等。而 CPU+GPU 在这些异构架构中最为常见。
相较于 CPU，GPU 有更少的片上缓存和大量能够并行执行的简单 ALU。CPU 在设计上针对由短序列计算操作和难以预测的控制流程构成的动态工作负载进行优化，而 GPU 则被设计以处理计算任务主导且带有简单控制流的工作负载。因此CPU 计算适合处理控制密集型任务，而 GPU 计算适合处理包含数据并行的计算密集型任务。二者的功能互补性促成了 CPU+GPU 的异构并行计算架构的发展。
当然，一切多核与众核芯片带来的性能提升都不是免费的，需要我们对大量传统算法进行重新设计。只有充分理解 CPU 与 GPU 架构，用并行思维思考，才能设计和编写高效的异构并行计算程序。
CUDA 是一种异构计算平台。在 CUDA 平台上，可以使用标准程序语言的扩展、API、编译器指令、CUDA 加速库等编写程序，以利用 NVIDIA GPU 高效地处理复杂的并行计算问题。
CUDA 编程模型
CUDA 采用全局串行局部并行的编程模型，GPU 作为协处理器对程序的部分进行加速。我们希望将程序中的可并发部分分解为成百上千个线程，交由 GPU 并发执行以充分发挥其性能。
事实上，CPU 与 GPU 中的线程略有区别：CPU 线程是重量级实体，操作系统交替执行线程，线程上下文切换花销很大，而 GPU 线程是轻量级的，GPU 应用一般包含成千上万的线程，多数在排队状态，线程之间切换基本没有开销。CPU 的核被设计用来尽可能减少一个或两个线程运行时间的延迟，而 GPU 核则充分为大量线程做足准备，最大幅度提高吞吐量。这些线程会被调度到数以千计的流式处理单元上执行，并通过合理的分层结构来控制。在执行模型中，我们会进一步说明。
如何从编程意义上组织成百上千的并发线程？在 CUDA 中，每次核函数调用产生一组线程，它们使用公用的函数参数来执行相同的功能，通过内部结构变量定义自己在线程结构中的位置。具体地，CUDA 将线程组织成 6 维的超立体结构，它由两层 3 维组织结构嵌套而成。线程组织成 3 维的线程块结构，线程块再组织成 3 维的线程网格结构。每个线程通过两个 3 维的内部结构变量 threadIdx, blockIdx 结合 threadDim, blockDim 来定义自己在线程结构中的位置，并实现位置信息与所分配数据子集的映射。事实上，这种结构的设计综合了软件与硬件方面的考虑。
线程在编程时以函数体现，在 GPU 上执行的函数称为核函数。带有 __global__ 修饰的函数只允许在设备（即 GPU）上执行，它们只能由主机调用。主机端调用的核函数没有返回值，而核函数的输入和输出都存储在设备内存中，需要显式地与主机内存之间进行拷贝操作，后续章节中会详细讨论。
我们常用 host 指代 CPU 及其内存，而用 device 指代 GPU 及其内存。CUDA 程序中包含 host 程序和 device 程序分别在 CPU 和 GPU 上运行。它们之间通过以内存拷贝为主的方式进行通信。
典型的 CUDA 程序执行流程如下：
- 分配 host 内存，并进行数据初始化；
- 分配 device 内存，并从 host 将数据拷贝到 device 上；
- 调用 CUDA 的核函数在 device 上完成指定的运算；
- 将 device 上的运算结果拷贝到 host 上；
- 释放 device 和 host 上分配的内存。
在 CUDA 编程中通过函数类型限定词区别 host 和 device 函数，主要的三个函数类型限定词如下：
__global__：在 device 上执行，一般从 host 中调用；__device__：在 device 上执行，单仅可以从 device 中调用；__host__：在 host 上执行，仅可以从 host 上调用，一般省略不写。
部分限定词可以同时使用，读者可查阅有关资料进一步了解。
Host 程序在调用核函数时通过特定语法 <<<g,b>>> 来指定线程的组织层次。使用 dim3 数据类型来表示包含三个元素的整型向量，如果提供数据不足 3 个，则缺省补 1。只需要使用一维时，可以简写为标量。
简举一例，用核函数进行向量加法。限于篇幅，以下程序段仅展示核函数的定义与调用，完整实例程序请参阅 NVIDIA 官方示例或查阅其它资料。
__global__ void vectorAdd(const float *A, const float *B, float *C, int numElements)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < numElements) {
C[i] = A[i] + B[i];
}
}
int main()
{
// ...
cudaMemcpy(d_A, A, length, cudaMemcpyHostToDevice);
// ...
vectorAdd <<<blocksPerGrid, threadsPerBlock >>>(d_A, d_B, d_C, numElements);
// ...
cudaMemcpy(C, d_C, length, cudaMemcpyDeviceToHost);
// ...
}
最后，需要指出的是，在编程模型中，尽管我们努力地不去谈论任何硬件的组织结构，有一个问题无法避免：计算能力（Compute Capability）。计算能力反映了架构的性能特点，但名不副实的是，它的数值并非一个绝对的性能指标，而更像是架构的版本号，有时我们也将它称为 SM version。计算能力对线程与线程块的组织方式，内存的分配限制给出了详细的界限。因此，在编程时，我们务必考虑目标平台的计算能力，或者提供较大的可扩放性。
CUDA 执行模型
CUDA 采用单指令多线程（SIMT）执行模型。SIMT 非常类似于 SIMD，而不同主要在于处理单元处理的“向量”大小是软件定义的线程块的大小。核函数在设备上运行时，相同的指令序列会被大量流处理单元（SP）部分同步地执行。
流式多处理器（SM）是在同一控制单元下执行的一组流处理单元。GPU 中通常包括数十个 SM ，每个 SM 支持数百线程并发执行。
线程以线程块为单位，调度到不同的 SM 上执行。SM 和线程块是一对多的关系：线程块被调度到一个 SM 并保持在其上直到执行结束，而单个 SM 可以同时容纳多个线程块。
线程的调度单位是线程块，而执行单位则是进一步细分的线程束。换言之，同一线程块内的所有线程未必在物理意义上并行执行。线程块被调度到 SM 后，进一步细分为若干固定大小的线程束，这一大小与硬件相关，通常为 32。
同束线程被共同的控制单元执行，它们必须同时执行相同的指令。作为 SIMT 与 SIMD 的不同之处，线程束中的线程拥有独立的寄存器（物理上是寄存器堆中的条目）、指令地址计数器等状态部件，并可以拥有独立的执行路径。
实现上，SM 获取指令，将指令广播道所有 SP 中，所有 SP 一起执行。 SM 通常拥有两套以上的线程束调度器和指令分派器。线程束调度器选择一个线程束，将指令发送到一个流处理单元组中。
CUDA 存储模型
为了取得大容量、高性能与生产成本间的平衡，必须依靠内存模型来获得最佳的延迟和带宽。
回忆在 CPU 中，我们为了使得 CPU 能全速运行，多需要为任意随机地址的快速访问提供保障，因此提供了较大的多级片上缓存。而非通用 GPU 的工作以过滤和传输图形信息为典型，在处理一次读取的大量数据后，这些数据通常不必保持在片上。因此，原先的 GPU 需要高速的数据总线，而只需要少量的片上缓存。
通用计算时代，GPGPU 正在颠覆这一点。如今的 GPU 带有越来越大的片上内存与缓存。与 CPU 不同的式，GPU 的内存层次结构对用户并不透明。用户需要部分对内存类型选择、数据移动进行选择和介入。
在 CUDA 中，我们关注的存储层次结构主要包括寄存器、本地内存、共享内存、缓存、全局内存、常量内存、纹理内存。在逻辑结构中的示意图如下：
它们的位置、访问范围与生命周期的简要对比如下表所示：
| 类型 | 位置 | 逻辑访问范围 | 生命周期 |
|---|---|---|---|
| 寄存器 | 片上 | 线程 | 线程 |
| 本地内存 | 片外 | 线程 | 线程 |
| 共享内存 | 片上 | 块 | 块 |
| 全局内存 | 片外 | 网格 | 主机控制 |
| 常量/纹理内存 | 片外 | 网格 | 主机控制 |
寄存器用于保存局部变量，从而减少对全局内存或共享内存的访问，加快操作的处理速度。CUDA 中的寄存器以寄存器文件（register file，寄存器堆）的形式存在，它是由多个寄存器组成的阵列，通常由 SRAM 实现。在现代 CPU 中由于寄存器重命名技术的使用，架构寄存器对应的寄存器堆中的物理存储条目也是动态的。
若局部变量需要的空间超过计算能力的限制，则会溢出到位于本地内存的运行时栈中。注意本地内存之所以称为本地，是因为它仅能被特定线程访问，然而，它在物理上和全局内存一样位于片外。
共享内存是可以被一个 SM 上所有 SP 共享的 RAM，它可以用于存储那些本可以放在全局内存中但需要频繁使用的数据或者全局内存部分数据的副本，以及用于在 SM 的各 SP 间共享数据。某种意义上可以将共享内存视作特殊的 L1$，事实上在 Fermi 和 Kepler 架构中共享内存与 L1$ 是同一块片上 RAM，由用户编程指定如何切分。
例如，若线程频繁对某个数据进行读写操作，可以设置操作的数据为 __shared__，使其位于共享内存中（也称其常驻缓存），且同一个线程块内的所有线程共享该内存区域。当出现多个线程对同一个内存区域进行操作时，需要对线程进行同步操作。__syncthreads() 为我们提供了简洁高效的屏障同步操作。
全局内存是设备端片外内存的主要部分，容量高但速度低。主机端与设备端的主要通过全局内存进行，因为它是主机通过 CUDA 库函数可以访问的唯一部分。
常量内存也是设备端片外内存的一部分，它可以被缓存，同时支持将单个数据广播给线程束中的所有线程。
纹理内存通过对数据的大小和格式限制来实现图像数据的快速访问与变换。正因如此，它在 GPGPU 领域的应用较为有限。
后记
CUDA 平台为异构并行计算提供了丰富的工具，但高效并行程序的设计仍然面临着诸多的问题与艰巨的挑战。对 CUDA 异构并行程序设计的学习应当从硬件平台的体系结构、并行算法设计的方法与技术、并行编程的平台与工具三个方面充分着力，并主动寻求与运用领域的深入融合，充分发掘异构并行计算的潜能与优势。
