使用 CUDA 进行可扩展并行编程 -

2008年4月28日
第 6 卷，第 2 期

使用 CUDA 进行可扩展并行编程

CUDA 是应用程序开发人员一直在等待的并行编程模型吗？

JOHN NICKOLLS, IAN BUCK, 和 MICHAEL GARLAND, NVIDIA, KEVIN SKADRON, 弗吉尼亚大学

多核 CPU 和众核 GPU 的出现意味着主流处理器芯片现在是并行系统。此外，它们的并行性随着摩尔定律持续扩展。挑战在于开发主流应用程序软件，使其并行性能够透明地扩展，以利用越来越多的处理器核心，就像 3D 图形应用程序透明地将其并行性扩展到具有数量差异很大的核心的众核 GPU 一样。

根据传统观点，并行编程是困难的。然而，使用 CUDA1,2 可扩展并行编程模型和 C 语言的早期经验表明，许多复杂的程序可以用一些易于理解的抽象概念轻松表达。自从 NVIDIA 在 2007 年发布 CUDA 以来，开发人员已经为广泛的应用快速开发了可扩展的并行程序，包括计算化学、稀疏矩阵求解器、排序、搜索和物理模型。这些应用程序可以透明地扩展到数百个处理器核心和数千个并发线程。配备新型 Tesla 统一图形和计算架构（在 GPU 侧栏中描述）的 NVIDIA GPU 可以运行 CUDA C 程序，并在笔记本电脑、PC、工作站和服务器中广泛使用。CUDA 模型也适用于其他共享内存并行处理架构，包括多核 CPU。3

CUDA 提供了三个关键抽象概念——线程组的层次结构、共享内存和屏障同步——为层次结构中单个线程的传统 C 代码提供了清晰的并行结构。线程、内存和同步的多层结构提供了细粒度数据并行和线程并行，嵌套在粗粒度数据并行和任务并行中。这些抽象概念引导程序员将问题划分为可以独立并行解决的粗子问题，然后再划分为可以协作并行解决的更精细的部分。编程模型可以透明地扩展到大量的处理器核心：编译后的 CUDA 程序可以在任意数量的处理器上执行，并且只有运行时系统需要知道物理处理器计数。

CUDA 范例

CUDA 是 C 和 C++ 编程语言的最小扩展。程序员编写串行程序，调用并行内核，内核可以是简单的函数或完整的程序。一个内核在一组并行线程中并行执行。程序员将这些线程组织成线程块网格的层次结构。线程块是一组并发线程，它们可以通过屏障同步和对块私有的内存空间的共享访问来相互协作。网格是一组线程块，每个线程块都可以独立执行，因此可以并行执行。

在调用内核时，程序员指定每个块的线程数和构成网格的块数。每个线程在其线程块中被赋予唯一的线程 ID 号 threadIdx，编号为 0, 1, 2, ..., blockDim–1，并且每个线程块在其网格中被赋予唯一的块 ID 号 blockIdx。CUDA 支持包含多达 512 个线程的线程块。为了方便起见，线程块和网格可以具有一维、二维或三维，通过 .x、.y 和 .z 索引字段访问。

作为一个非常简单的并行编程示例，假设我们给定两个向量 x 和 y，每个向量包含 n 个浮点数，并且我们希望计算 y←ax + y 的结果，其中 a 为某个标量值。这是 BLAS（基本线性代数子程序）库定义的所谓 saxpy 内核。图 1 显示了在串行处理器上和使用 CUDA 并行执行此计算的代码。

__global__ 声明说明符指示该过程是内核入口点。CUDA 程序使用扩展的函数调用语法启动并行内核

kernel<<<dimGrid, dimBlock>>>(... parameter list ...);

其中 dimGrid 和 dimBlock 是 dim3 类型的三个元素向量，分别指定网格在块中的维度和块在线程中的维度。未指定的维度默认为 1。

在该示例中，我们启动一个网格，该网格为向量的每个元素分配一个线程，并在每个块中放置 256 个线程。每个线程根据其线程和块 ID 计算元素索引，然后在相应的向量元素上执行所需的计算。此代码的串行版本和并行版本非常相似。这代表了一种相当常见的模式。串行代码由一个循环组成，其中每次迭代都独立于所有其他迭代。此类循环可以机械地转换为并行内核：每次循环迭代都变成一个独立的线程。通过为每个输出元素分配一个线程，我们避免了在将结果写入内存时在线程之间进行任何同步的需要。

CUDA 内核的文本只是一个顺序线程的 C 函数。因此，通常很容易编写，并且通常比编写向量运算的并行代码更简单。并行性通过在启动内核时指定网格及其线程块的维度来明确地确定。

并行执行和线程管理是自动的。所有线程创建、调度和终止都由底层系统为程序员处理。实际上，Tesla 架构 GPU 直接在硬件中执行所有线程管理。一个块的线程并发执行，并且可以通过调用 __syncthreads() 内在函数在屏障处同步。这保证了参与屏障的任何线程在所有参与线程都到达屏障之前都无法继续。在通过屏障之后，这些线程也被保证看到参与线程在屏障之前执行的所有内存写入。因此，块中的线程可以通过在同步屏障处写入和读取每个块的共享内存来进行通信。

由于块中的线程可以共享本地内存并通过屏障同步，因此它们将驻留在同一物理处理器或多处理器上。然而，线程块的数量可以大大超过处理器的数量。这虚拟化了处理元素，并使程序员可以灵活地在最方便的粒度上进行并行化。这允许直观的问题分解，因为块的数量可以由正在处理的数据的大小而不是系统中的处理器数量来决定。这也允许相同的 CUDA 程序扩展到数量差异很大的处理器核心。

为了管理这种处理元素虚拟化并提供可扩展性，CUDA 要求线程块独立执行。必须可以按任何顺序、并行或串行地执行块。不同的块没有直接通信的方式，尽管它们可以使用对所有线程可见的全局内存上的原子内存操作来协调它们的活动——例如，通过原子地递增队列指针。

这种独立性要求允许线程块以任何顺序在任意数量的核心上调度，从而使 CUDA 模型可以跨任意数量的核心以及各种并行架构进行扩展。它还有助于避免死锁的可能性。

应用程序可以独立或依赖地执行多个网格。独立的网格可以在硬件资源充足的情况下并发执行。依赖的网格按顺序执行，它们之间存在隐式的内核间屏障，从而保证第一个网格的所有块将在启动第二个依赖网格的任何块之前完成。

线程可以在执行期间访问来自多个内存空间的数据。每个线程都有一个私有的本地内存。CUDA 将此内存用于不适合线程寄存器的线程私有变量，以及堆栈帧和寄存器溢出。每个线程块都有一个共享内存，对块的所有线程可见，并且与块具有相同的生命周期。最后，所有线程都可以访问相同的全局内存。程序使用 __shared__ 和 __device__ 类型限定符在共享内存和全局内存中声明变量。在 Tesla 架构 GPU 上，这些内存空间对应于物理上独立的内存：每个块的共享内存是低延迟的片上 RAM，而全局内存驻留在图形卡上的快速 DRAM 中。

共享内存预计是靠近每个处理器的低延迟内存，很像 L1 缓存。因此，它可以为线程块的线程之间提供高性能的通信和数据共享。由于它与其对应的线程块具有相同的生命周期，因此内核代码通常会在共享变量中初始化数据，使用共享变量进行计算，并将共享内存结果复制到全局内存。顺序依赖网格的线程块通过全局内存进行通信，使用它来读取输入和写入结果。

图 2 图解了线程、线程块和线程块网格的嵌套级别。它显示了相应的内存共享级别：每个线程、每个线程块和每个应用程序的数据共享的本地内存、共享内存和全局内存。

程序通过调用 CUDA 运行时（例如 cudaMalloc() 和 cudaFree()）来管理内核可见的全局内存空间。内核可以在物理上独立的设备上执行，就像在 GPU 上运行内核时一样。因此，应用程序必须使用 cudaMemcpy() 在分配的空间和主机系统内存之间复制数据。

CUDA 编程模型在风格上类似于熟悉的 SPMD（单程序多数据）模型——它显式地表达并行性，并且每个内核都在固定数量的线程上执行。然而，CUDA 比 SPMD 的大多数实现更灵活，因为每个内核调用都会动态创建一个新的网格，其中包含适合该应用程序步骤的正确数量的线程块和线程。程序员可以为每个内核使用方便的并行度，而不必设计计算的所有阶段都使用相同数量的线程。

图 3 显示了一个类似 SPMD 的 CUDA 代码序列示例。它首先在 3×2 块的 2D 网格上实例化 kernelF，其中每个 2D 线程块由 5×3 个线程组成。然后，它在四个 1D 线程块（每个块有六个线程）的 1D 网格上实例化 kernelG。由于 kernelG 依赖于 kernelF 的结果，因此它们之间通过内核间同步屏障分隔。

线程块的并发线程表达了细粒度的数据和线程并行性。网格的独立线程块表达了粗粒度的数据并行性。独立网格表达了粗粒度的任务并行性。内核只是层次结构中一个线程的 C 代码。

限制

在开发 CUDA 程序时，重要的是要了解 CUDA 模型受到限制的方式，这主要是出于效率原因。线程和线程块只能通过调用并行内核来创建，而不能从并行内核内部创建。结合线程块的必要独立性，这使得可以使用简单的调度程序执行 CUDA 程序，从而最大限度地减少运行时开销。实际上，Tesla 架构实现了线程和线程块的硬件管理和调度。

任务并行性可以在线程块级别表达，但块范围的屏障不太适合支持块中线程之间的任务并行性。为了使 CUDA 程序能够在任意数量的处理器上运行，不允许在同一内核网格内的线程块之间进行通信——它们必须独立执行。由于 CUDA 要求线程块是独立的，并允许以任何顺序执行块，因此通常必须通过在新线程块网格上启动第二个内核来完成组合由多个块生成的结果。但是，多个线程块可以使用全局内存上的原子操作（例如，管理数据结构）来协调它们的工作。

CUDA 内核中不允许递归函数调用。递归在海量并行内核中是不可取的，因为为可能处于活动状态的数万个线程提供堆栈空间将需要大量的内存。通常使用递归表达的串行算法（例如快速排序）通常最好使用嵌套数据并行而不是显式递归来实现。

为了支持结合 CPU 和 GPU（每个都有自己的内存系统）的异构系统架构，CUDA 程序必须在主机内存和设备内存之间复制数据和结果。通过使用 DMA 块传输引擎和快速互连，可以最大限度地减少 CPU-GPU 交互和数据传输的开销。当然，足够大的问题需要 GPU 性能提升，比小问题更好地分摊开销。

CUDA 应用程序经验

CUDA 编程模型使用少量的附加并行抽象概念扩展了 C 语言。熟悉 C 语言开发的程序员可以快速开始编写 CUDA 程序。

自从 CUDA 推出以来相对较短的时间内，已经使用 CUDA 模型开发了许多真实世界的并行应用程序代码。这些包括 FHD 螺旋 MRI 重建、6 分子动力学 7 和 n 体天体物理模拟。8 这些应用程序在 Tesla 架构 GPU 上运行时，与在串行 CPU 上运行的替代实现相比，能够实现显着的加速：MRI 重建快了 263 倍；分子动力学代码快了 10-100 倍；n 体模拟快了 50-250 倍。这些巨大的加速是 Tesla 架构的高度并行性及其高内存带宽的结果。

示例：稀疏矩阵-向量乘积

即使所涉及的数据结构不是简单的规则网格，也可以以相当直接的方式用 CUDA 编写各种并行算法。SpMV（稀疏矩阵-向量乘法）是一个很好的例子，说明重要的数值构建块可以使用 CUDA 提供的抽象概念直接并行化。我们在此讨论的内核与提供的 CUBLAS 向量例程相结合，使得编写迭代求解器（例如共轭梯度 9 方法）变得非常简单。

一个 n × n 稀疏矩阵是指非零条目 m 的数量仅占总数的一小部分的矩阵。稀疏矩阵表示旨在仅存储矩阵的非零元素。由于一个稀疏 n × n 矩阵通常只包含 m=O(n) 个非零元素，因此这代表了存储空间和处理时间的显着节省。

通用非结构化稀疏矩阵最常见的表示形式之一是 CSR（压缩稀疏行）表示形式。矩阵 A 的 m 个非零元素以行优先顺序存储在数组 Av 中。第二个数组 Aj 记录 Av 的每个条目的相应列索引。最后，一个包含 n+1 个元素的数组 Ap 记录了先前数组中每一行的范围；Aj 和 Av 中第 i 行的条目从索引 Ap[i] 扩展到但不包括索引 Ap[i+1]。这意味着 Ap[0] 将始终为 0，Ap[n] 将始终为矩阵中非零元素的数量。图 4 显示了一个简单矩阵的 CSR 表示形式示例。

给定 CSR 形式的矩阵 A，我们可以使用图 5 所示的 multiply_row() 过程计算乘积 y = Ax 的单行。

然后，计算完整乘积只是简单地循环遍历所有行，并使用 multiply_row() 计算该行的结果，如图 6 所示。

该算法可以很容易地转换为并行 CUDA 内核。我们只需将 csrmul_serial() 中的循环分散到许多并行线程中。每个线程将精确计算输出向量 y 的一行。图 7 显示了此内核的代码。请注意，它看起来与 csrmul_serial() 过程中使用的串行循环非常相似。实际上只有两个不同之处。首先，行索引是从分配给每个线程的块和线程索引计算得出的。其次，我们有一个条件，仅当行索引在矩阵边界内时才评估行乘积（这是必要的，因为行数 n 不必是启动内核时使用的块大小的倍数）。

假设矩阵数据结构已经复制到 GPU 设备内存，则启动此内核将类似于图 8 中的代码。

我们在这里看到的模式是一种常见的模式。原始串行算法是一个循环，其迭代彼此独立。通过简单地将循环的一个或多个迭代分配给每个并行线程，可以非常容易地并行化此类循环。CUDA 提供的编程模型使得表达这种类型的并行性特别简单直接。

将计算分解为独立工作块的这种通用策略，更具体地说是分解独立的循环迭代，并非 CUDA 独有。这是各种并行编程系统（包括 OpenMP 和英特尔的线程构建块）以一种或另一种形式使用的常用方法。

共享内存中的缓存

此处概述的 SpMV 算法相当简单。我们可以对 CPU 和 GPU 代码进行许多优化，以提高性能，包括循环展开、矩阵重新排序和寄存器阻塞。10 并行内核也可以根据数据并行扫描操作重新实现。11

CUDA 公开的一个重要架构特征是每个块的共享内存的存在，这是一种具有非常低延迟的小型片上内存。利用此内存可以显着提高性能。一种常见的做法是将共享内存用作软件管理的缓存，以保存频繁重用的数据，如图 9 所示。

在稀疏矩阵乘法的上下文中，我们观察到矩阵 A 的多行可能使用特定的数组元素 x[i]。在许多常见情况下，尤其是在矩阵已重新排序时，使用 x[i] 的行将是靠近第 i 行的行。因此，我们可以实现一个简单的缓存方案，并期望获得一些性能优势。处理第 i 行到第 j 行的线程块将 x[i] 到 x[j] 加载到其共享内存中。我们将展开 multiply_row() 循环，并在可能的情况下从缓存中获取 x 的元素。结果代码如图 9 所示。共享内存也可用于进行其他优化，例如从相邻线程获取 Ap[row+1]，而不是从内存中重新获取它。

由于 Tesla 架构提供了显式管理的片上共享内存，而不是隐式活动的硬件缓存，因此添加此类优化是相当常见的。尽管这可能会给程序员带来一些额外的开发负担，但它相对较小，并且潜在的性能优势可能是巨大的。在图 9 所示的示例中，即使这种相当简单的共享内存使用也在从 3D 表面网格派生的代表性矩阵上返回了大约 20% 的性能提升。显式管理的内存代替隐式缓存的可用性也具有这样的优势：缓存和预取策略可以专门针对应用程序需求进行定制。

示例：并行归约

假设我们给定一个必须以某种方式组合的 N 个整数序列（例如，总和）。这发生在各种算法中，线性代数是一个常见的例子。在串行处理器上，我们将编写一个简单的循环，其中包含单个累加器变量，以构造序列中所有元素的总和。在并行机器上，使用单个累加器变量将创建一个全局序列化点，并导致非常差的性能。解决此问题的一个众所周知的解决方案是所谓的并行归约算法。每个并行线程对输入的固定长度子序列求和。然后，我们通过并行地对成对的部分和求和来收集这些部分和。此成对求和的每个步骤都将部分和的数量减半，并最终在 log2 N 步后产生最终总和。请注意，这隐式地在初始部分和上构建了一个树结构。

在图 10 所示的示例中，每个线程仅加载输入序列的一个元素（即，它最初对长度为 1 的子序列求和）。在归约结束时，我们希望线程 0 保存其块的线程最初加载的所有元素的总和。我们可以通过在类似树的模式中对值求和来并行实现此目的。此内核中的循环隐式地在输入元素上构建求和树。图 11 说明了对于八个线程的块的简单情况，此循环的操作。循环的步骤显示为图的连续级别，边缘指示从哪里读取部分和。

在此循环结束时，线程 0 保存此块加载的所有值的总和。如果我们希望 total 指向的位置的最终值包含数组中所有元素的总和，则我们必须组合网格中所有块的部分和。一种策略是让每个块将其部分和写入第二个数组，然后再次启动归约内核，重复该过程直到我们将序列归约为单个值。Tesla 架构支持的更具吸引力的替代方案是使用 atomicAdd()，这是内存子系统支持的高效原子读取-修改-写入原语。这消除了对其他临时数组和重复内核启动的需求。

并行归约是并行编程的基本原语，并突出了每个块的共享内存和低成本屏障在使线程之间的协作高效方面的重要性。如果在片外全局内存中完成这种线程之间的数据混洗，其成本将高得令人望而却步。

并行编程的普及

CUDA 是一种并行编程模型，它提供了一些易于理解的抽象概念，使程序员能够专注于算法效率并开发可扩展的并行应用程序。实际上，CUDA 是一个优秀的并行编程教学环境。弗吉尼亚大学在一个本科计算机体系结构课程中仅用了短短三周的模块就使用了它，学生们在仅仅三节课后就能够编写正确的 k 均值聚类程序。伊利诺伊大学已成功地使用 CUDA 向计算机科学和非计算机科学专业的学生教授了一个学期的并行编程课程，学生们在各种实际应用程序（包括前面提到的 MRI 重建示例）中获得了令人印象深刻的加速。

CUDA 在 NVIDIA GPU 上得到支持，这些 GPU 具有 GeForce 8 系列、最近的 Quadro、Tesla 和未来 GPU 的 Tesla 统一图形和计算架构。CUDA 提供的编程范例使开发人员能够相对轻松地利用这些可扩展的并行处理器的强大功能，使他们能够在各种复杂的应用程序上实现 100 倍或更高的加速。

然而，CUDA 抽象概念是通用的，并为多核 CPU 芯片提供了出色的编程环境。伊利诺伊大学开发的一个原型源到源翻译框架通过将并行线程块映射到单个物理线程内的循环来编译多核 CPU 的 CUDA 程序。以这种方式编译的 CUDA 内核表现出出色的性能和可扩展性。12

尽管 CUDA 在不到一年前发布，但它已经成为大规模开发活动的目标——有数以万计的 CUDA 开发人员。巨大的加速、直观的编程环境以及经济实惠、无处不在的硬件在当今市场上是罕见的。简而言之，CUDA 代表了并行编程的普及。

参考文献

NVIDIA. 2007. CUDA 技术；http://www.nvidia.com/CUDA。
NVIDIA. 2007. CUDA 编程指南 1.1；http://developer.download.nvidia.com/compute/cuda/1_1/NVIDIA_CUDA_Programming_Guide_1.1.pdf。
Stratton, J.A., Stone, S. S., Hwu, W. W. 2008. M-CUDA：多核上 CUDA 内核的高效实现。IMPACT 技术报告 08-01，伊利诺伊大学厄巴纳-香槟分校，（二月）。
请参阅参考文献 3。
Buck, I., Foley, T., Horn, D., Sugerman, J., Fatahalian, K., Houston, M., Hanrahan, P. 用于 GPU 的 Brook：图形硬件上的流计算。2004 年。SIGGRAPH 会议论文集（八月）：777-786；http://doi.acm.org/10.1145/1186562.1015800。
Stone, S.S., Yi, H., Hwu, W.W., Haldar, J.P., Sutton, B.P., Liang, Z.-P. 2007. GPU 如何提高磁共振成像的质量。图形处理器通用处理第一次研讨会（十月）。
Stone, J.E., Phillips, J.C., Freddolino, P.L., Hardy, D.J., Trabuco, L.G., Schulten, K. 2007. 使用图形处理器加速分子建模应用程序。计算化学杂志 28(16)：2618–2640；http://dx.doi.org/10.1002/jcc.20829。
Nyland, L., Harris, M., Prins, J. 2007. 使用 CUDA 的快速 n 体模拟。在 GPU Gems 3 中。H. Nguyen，编辑。Addison-Wesley。
Golub, G.H., 和 Van Loan, C.F. 1996. 矩阵计算，第 3 版。约翰·霍普金斯大学出版社。
Buatois, L., Caumon, G., Lévy, B. 2007. 并发数字运算器：GPU 上高效的稀疏线性求解器。高性能计算会议论文集 (HPCC)，Springer LNCS。
Sengupta, S., Harris, M., Zhang, Y., Owens, J.D. 2007. 用于 GPU 计算的扫描原语。在 图形硬件会议论文集（八月）：97–106。
请参阅参考文献 3。

CUDA 开发工具、文档、代码示例和用户讨论论坛的最新版本的链接可以在以下位置找到：http://www.nvidia.com/CUDA。

JOHN NICKOLLS 是 NVIDIA GPU 计算架构总监。他之前曾在 Broadcom、Silicon Spice、Sun Microsystems 工作，并且是 MasPar Computer 的联合创始人。他的兴趣包括并行处理系统、语言和架构。他拥有伊利诺伊大学电气工程和计算机科学学士学位，以及斯坦福大学电气工程硕士和博士学位。

IAN BUCK 在 NVIDIA 担任 GPU 计算软件经理。他于 2004 年在斯坦福图形实验室完成了博士学位。他的论文题目是“图形硬件上的流计算”，研究了使用图形硬件作为通用计算平台的编程模型和计算策略。他的工作包括开发 Brook 软件工具链，用于将 GPU 抽象为通用流式协处理器。

MICHAEL GARLAND 是 NVIDIA Research 的研究科学家。在加入 NVIDIA 之前，他是伊利诺伊大学厄巴纳-香槟分校计算机科学系的助理教授。他获得了卡内基梅隆大学的博士和学士学位。他的研究兴趣包括计算机图形学和可视化、几何算法以及并行算法和编程模型。

KEVIN SKADRON 是弗吉尼亚大学计算机科学系的副教授，目前在 NVIDIA Research 休假。他获得了普林斯顿大学的博士学位和莱斯大学的学士学位。他的研究兴趣包括功耗和温度感知设计，以及众核架构和编程模型。他是的高级会员。

统一图形和计算 GPU

在对实时、高清 3D 图形的永不满足的市场需求的驱动下，可编程 GPU（图形处理单元）已发展成为高度并行、多线程、众核处理器。它旨在有效地支持图形着色器编程模型，其中一个线程的程序绘制一个顶点或对一个像素片段进行着色。GPU 擅长细粒度、数据并行工作负载，这些工作负载由数千个并发执行顶点、几何和像素着色器程序线程的独立线程组成。

现代 GPU 惊人的原始性能促使研究人员探索将更通用的非图形计算映射到它们之上。这些 GPGPU（GPU 上的通用计算）系统已经产生了一些令人印象深刻的结果，但是通过图形 API 执行此操作的局限性和困难是众所周知的。这种将 GPU 用作更通用的并行计算设备的愿望促使 NVIDIA 开发了一种新的统一图形和计算 GPU 架构以及 CUDA 编程模型。

GPU 计算架构

NVIDIA 于 2006 年 11 月推出的 Tesla 统一图形和计算架构^1,2 显著地将 GPU 的应用扩展到图形之外——其大规模多线程处理器阵列成为一个高效的统一平台，可用于图形和通用并行计算应用。通过扩展处理器和内存分区的数量，Tesla 架构涵盖了广泛的市场范围——从高性能发烧友级的 GeForce 8800 GPU 和专业的 Quadro 和 Tesla 计算产品，到各种廉价的主流 GeForce GPU。其计算特性使得可以使用 CUDA 以 C 语言直接对 GPU 内核进行编程。Tesla 架构在笔记本电脑、台式机、工作站和服务器中的广泛可用性，以及 C 语言可编程性和 CUDA 软件，使其成为首个无处不在的超级计算平台。

Tesla 架构围绕可扩展的多线程 SM（流式多处理器）阵列构建。当前的 GPU 实现范围从 768 到 12,288 个并发执行的线程。在此广泛的可用并行性范围内实现透明扩展是 GPU 架构和 CUDA 编程模型的关键设计目标。图 A 显示了一个具有 14 个 SM 的 GPU——总共 112 个 SP（流式处理器）核心——与四个外部 DRAM 分区互连。当主机 CPU 上的 CUDA 程序调用内核网格时，CWD（计算工作分发）单元会枚举网格的块，并开始将它们分发到具有可用执行容量的 SM。线程块的线程在一个 SM 上并发执行。当线程块终止时，CWD 单元会在空出的多处理器上启动新的块。

一个 SM 由八个标量 SP 核心、两个用于超越函数的 SFU（特殊功能单元）、一个 MT IU（多线程指令单元）和片上共享内存组成。SM 在硬件中创建、管理和执行多达 768 个并发线程，且调度开销为零。它可以并发执行多达八个 CUDA 线程块，但这受线程和内存资源的限制。SM 使用单个指令实现 CUDA __syncthreads() 屏障同步内在函数。快速屏障同步以及轻量级的线程创建和零开销的线程调度有效地支持了非常细粒度的并行性，从而允许创建一个新线程来计算每个顶点、像素和数据点。

为了管理运行多个不同程序的数百个线程，Tesla SM 采用了一种我们称之为 SIMT（单指令多线程）的新架构³。SM 将每个线程映射到一个 SP 标量核心，并且每个标量线程都使用自己的指令地址和寄存器状态独立执行。SM SIMT 单元以称为 Warp 的 32 个并行线程组为单位创建、管理、调度和执行线程。（这个术语起源于编织，第一种并行线程技术。）组成 SIMT Warp 的各个线程在相同的程序地址处一起启动，但在其他方面可以自由分支和独立执行。每个 SM 管理一个由 24 个 Warp 组成的池，每个 Warp 包含 32 个线程，总共 768 个线程。

在每个指令发布时间，SIMT 单元选择一个准备好执行的 Warp，并将下一条指令发布到 Warp 的活动线程。一个 Warp 一次执行一条公共指令，因此当一个 Warp 的所有 32 个线程都同意它们的执行路径时，才能实现完全效率。如果一个 Warp 的线程通过数据相关的条件分支发散，则 Warp 会串行执行每个被采用的分支路径，禁用不在该路径上的线程，并且当所有路径完成时，线程会收敛回相同的执行路径。分支发散仅在 Warp 内发生；不同的 Warp 独立执行，而不管它们是否正在执行公共或不相交的代码路径。因此，与上一代 GPU 相比，Tesla 架构的 GPU 在分支代码上效率更高、更灵活，因为它们的 32 线程 Warp 比先前 GPU 的 SIMD（单指令多数据）宽度窄得多。

SIMT 架构类似于 SIMD 向量组织，因为单条指令控制多个处理单元。一个关键的区别是 SIMD 向量组织向软件公开 SIMD 宽度，而 SIMT 指令则指定单个线程的执行和分支行为。与 SIMD 向量机器相比，SIMT 使程序员能够为独立的标量线程编写线程级并行代码，以及为协调的线程编写数据并行代码。为了正确性，程序员基本上可以忽略 SIMT 行为；但是，通过注意代码很少需要 Warp 中的线程发散，可以实现显着的性能提升。实际上，这类似于传统代码中缓存行的作用：在设计正确性时可以安全地忽略缓存行大小，但在设计峰值性能时必须在代码结构中考虑缓存行大小。另一方面，向量架构要求软件将加载合并到向量中并手动管理发散。

线程的变量通常驻留在活动寄存器中。16KB 的 SM 共享内存具有非常低的访问延迟和类似于 L1 缓存的高带宽；它为活动线程块保存 CUDA 每个块的 __shared__ 变量。SM 提供加载/存储指令来访问 GPU 外部 DRAM 中的 CUDA __device__ 变量。当地址落在同一块中并满足对齐标准时，它会将同一 Warp 中并行线程的单个访问合并为更少的内存块访问。由于全局内存延迟可能达到数百个处理器时钟周期，因此当线程块必须多次访问数据时，CUDA 程序会将数据复制到共享内存。Tesla 加载/存储内存指令使用整数字节寻址，以方便传统的编译器代码优化。每个 SM 中的大量线程数，以及对许多未完成的加载请求的支持，有助于覆盖到外部 DRAM 的加载到使用延迟。最新的 Tesla 架构 GPU 还提供原子读-修改-写内存指令，从而方便了并行规约和并行数据结构管理。

CUDA 应用程序在 Tesla 架构 GPU 上表现良好，因为 CUDA 的并行性、同步、共享内存和线程组的层次结构有效地映射到 GPU 架构的特性，并且因为 CUDA 很好地表达了应用程序的并行性。

参考文献

Lindholm, E., Nickolls, J., Oberman, S., Montrym, J. 2008. NVIDIA Tesla: A unified graphics and computing architecture. IEEE Micro 28(2).
Nickolls, J. 2007. NVIDIA GPU parallel computing architecture. In IEEE Hot Chips 19 (August 20), Stanford, CA; http://www.hotchips.org/archives/hc19/.
参见参考文献 1。

最初发表于 Queue vol. 6, no. 2—
在数字图书馆中评论本文

更多相关文章

David Crandall, Noah Snavely - 使用互联网照片集建模人和地点
本文介绍了我们如何使用在线照片集来重建关于世界及其居民在全球和本地范围内的信息。这项工作受到了社交内容共享网站的显着增长的推动，这些网站创建了大量的用户生成视觉数据的在线集合。仅 Flickr.com 目前就托管了超过 60 亿张由超过 4000 万独立用户拍摄的图像，而 Facebook.com 表示其每天增长近 2.5 亿张照片。

Jeffrey Heer, Ben Shneiderman - 用于可视化分析的交互式动态
数字数据规模的扩大和可用性为公共政策、科学发现、商业策略甚至我们的个人生活提供了非凡的资源。然而，为了充分利用这些数据，用户必须能够理解它：提出问题，发现感兴趣的模式，并识别（并可能纠正）错误。与数据管理系统和统计算法相结合，分析需要对领域特定的集群、趋势和数据中发现的异常值的重要性进行情境化的人工判断。

Robert DeLine, Gina Venolia, Kael Rowan - 使用代码地图进行软件开发
为了更好地理解专业软件开发人员如何使用其代码的可视化表示，我们采访了微软的九位开发人员，以确定常见的场景，然后调查了 400 多位开发人员，以更深入地了解这些场景。

Brendan Gregg - 可视化系统延迟
当 I/O 延迟以可视化热图的形式呈现时，可能会出现一些有趣而美丽的模式。这些模式提供了关于系统实际性能以及最终用户应用程序体验到的延迟类型的见解。在这些模式中看到的许多特征仍然不被理解，但到目前为止，它们的分析揭示了以前未知的系统行为。