CUDA_C_NOTES [3]

CH03 CUDA执行模型

3.1 CUDA执行模型概述

CUDA执行模型能够提供有助于在指令吞吐量和内存访问方面编写高效代码的见解

3.1.1 GPU架构概述

GPU架构是围绕一个流式多处理器（SM） (Stream Multiprocessor)的可扩展阵列搭建的,可以通过复制这种架构的构建块来实现GPU的硬件并行

Fermi SM的关键组件：

• CUDA核心
• 共享内存/一级缓存
• 寄存器文件
• 加载/存储单元
• 特殊功能单元
• 线程束调度器

GPU中的每一个SM都能支持数百个线程并发执行， 每个GPU通常有多个SM， 所以在一个GPU上并发执行数千个线程是有可能的。 当启动一个内核网格时， 它的线程块被分布在了可用的SM上来执行。 线程块一旦被调度到一个SM上， 其中的线程只会在那个指定的SM上并发执行。 多个线程块可能会被分配到同一个SM上， 而且是根据SM资源的可用性进行调度的。同一线程中的指令利用指令级并行性进行流水线化， 另外， 在CUDA中已经介绍了线程级并行。

CUDA采用单指令多线程（SIMT）（single instruciton multi thread） 架构来管理和执行线程， 每32个线程为一组， 被称为线程束（warp） 。 线程束中的所有线程同时执行相同的指令。 每个线程都有自己的指令地址计数器和寄存器状态， 利用自身的数据执行当前的指令。 每个SM都将分配给它的线程块划分到包含32个线程的线程束中， 然后在可用的硬件资源上调度执行。

SIMT架构与SIMD（单指令多数据） 架构相似。 两者都是将相同的指令广播给多个执行单元来实现并行。 一个关键的区别是SIMD要求同一个向量中的所有元素要在一个统一的同步组中一起执行， 而SIMT允许属于同一线程束的多个线程独立执行.

SIMT确保可以编写独立的线程级并行代码、 标量线程以及用于协调线程的数据并行代码。

SIMT模型包含3个SIMD所不具备的关键特征。

• 每个线程都有自己的指令地址计数器
• 每个线程都有自己的寄存器状态
• 每个线程可以有一个独立的执行路径

一个神奇的数字： 32

从概念上讲， 它是SM用SIMD方式所同时处理的工作粒度。 优化工作负载以适应线程束（一组有32个线程） 的边界， 一般这样会更有效地利用GPU计算资源。

一个线程块只能在一个SM上被调度。 一旦线程块在一个SM上被调度， 就会保存在该SM上直到执行完成。 在同一时间， 一个SM可以容纳多个线程块.

在SM中， 共享内存和寄存器是非常重要的资源。 共享内存被分配在SM上的常驻线程块中， 寄存器在线程中被分配。

尽管线程块里的所有线程都可以逻辑地并行运行， 但是并不是所有线程都可以同时在物理层面执行。 因此， 线程块里的不同线程可能会以不同的速度前进。

在并行线程中共享数据可能会引起竞争： 多个线程使用未定义的顺序访问同一个数据， 从而导致不可预测的程序行为。 CUDA提供了一种用来同步线程块里的线程的方法，从而保证所有线程在进一步动作之前都达到执行过程中的一个特定点。 然而， 没有提供块间同步的原语。

当线程束由于任何理由闲置的时候（如等待从设备内存中读取数值） ， SM可以从同一SM上的常驻线程块中调度其他可用的线程束。 在并发的线程束间切换并没有开销， 因为硬件资源已经被分配到了SM上的所有线程和块中， 所以最新被调度的线程束的状态已经存储在SM上

SM： GPU架构的核心*

SM是GPU架构的核心。 寄存器和共享内存是SM中的稀缺资源。 CUDA将这些资源分配到SM中的所有常驻线程里。

这些有限的资源限制了在SM上活跃的线程束数量，活跃的线程束数量对应于SM上的并行量。

了解一些SM硬件组成的基本知识， 有助于组织线程和配置内核执行以获得最佳的性能

3.1.2 Fermi架构

Fermi的特征是多达512个加速器核心， 这被称为CUDA核心。 每个CUDA核心都有一个全流水线的整数算术逻辑单元（ALU） 和一个浮点运算单元（FPU） ， 在这里每个时钟周期执行一个整数或是浮点数指令。 CUDA核心被组织到16个SM中， 每一个SM含有32个CUDA核心。 Fermi架构有6个384位的GDDR5 DRAM存储器接口， 支持多达6GB的全局机载内存， 这是许多应用程序关键的计算资源。 主机接口通过PCIe总线将GPU与CPU相连。 GigaThread引擎（图示左侧第三部分） 是一个全局调度器， 用来分配线程块到SM线程束调度器上。

一个SM(Stream Multiprocessor)包含以下内容：

• 执行单元（CUDA核心）
• 调度线程束的调度器和调度单元
• 共享内存、 寄存器文件和一级缓存

每一个多处理器有16个加载/存储单元（如图3-1所示） ， 允许每个时钟周期内有16个线程（线程束的一半） 计算源地址和目的地址。 特殊功能单元（SFU） 执行固有指令， 如正弦、 余弦、 平方根和插值。 每个SFU每个时钟周期内的每个线程上执行一个固有指令

每个SM有两个线程束调度器和两个指令调度单元。 当一个线程块被指定给一个SM时， 线程块中的所有线程被分成了线程束。 两个线程束调度器选择两个线程束， 再把一个 指令从线程束中发送到一个组上， 组里有16个CUDA核心、 16个加载/存储单元或4个特殊功能单元（如图3-4所示） 。 Fermi架构， 计算性能2.x， 可以在每个SM上同时处理48个线程束， 即可在一个SM上同时常驻1536个线程。

3.1.3 Kepler架构

发布于2012年秋季的Kepler GPU架构是一种快速、 高效、 高性能的计算架构。 Kepler 的特点使得混合计算更容易理解。 图3-6表示了Kepler K20X芯片框图， 它包含了15个SM 和6个64位的内存控制器。 以下是Kepler架构的3个重要的创新。

• 强化的SM
• 动态并行
• Hyper-Q技术

Kepler K20X的关键部分是有一个新的SM单元， 其包括一些结构的创新， 以提高编程效率和功率效率。 每个Kepler SM单元包含192个单精度CUDA核心， 64个双精度单元， 32个特殊功能单元（SFU） 以及32个加载/存储单元（LD/ST）

3.1.4 配置文件驱动优化

配置文件驱动的发展对于CUDA编程尤为重要， 原因主要有以下几个方面。

• 一个单纯的内核应用一般不会产生最佳的性能。 性能分析工具能帮助你找到代码中影响性能的关键部分， 也就是性能瓶颈。
• CUDA将SM中的计算资源当前SM中的多个常驻线程块之间进行分配。 这种分配形式导致一些资源成为了性能限制者。 性能分析工具能帮助我们理解计算资源是如何被利用的。
• CUDA提供了一个硬件架构的抽象， 它能够让用户控制线程并发。 性能分析工具可以检测和优化， 并将优化可视化。

3.2 理解线程束执行的本质

本章已经提到了把32个线程划分到一个执行单元中的概念： 线程束（warp）。 现在从硬件的角度来介绍线程束执行， 并能够获得指导内核设计的方法。

3.2.1 线程束和线程块

线程束是SM中基本的执行单元。

当一个线程块的网格被启动后， 网格中的线程块分布在SM中。 一旦线程块被调度到一个SM上， 线程块中的线程会被进一步划分为线程束。 一个线程束由32个连续的线程组成， 在一个线程束中， 所有的线程按照单指令多线程（SIMT） 方式执行； 也就是说， 所有线程都执行相同的指令， 每个线程在私有数据上进 行操作。

一个给定的二维线程块， 在一个块中每个线程的独特标识符都可以用内置变量threadIdx和blockDim来计算： threadIdx.y * blockDim.x + threadIdx.x

对于一个三维线程块， 计算如下：

threadIdx.x * blockDim.y * block.Dim.x + threadIdx.y * blockDim.x * threadIdx.x

一个线程块的线程束的数量可以根据下式确定：

$$一个线程块中线程束的数量 = 向正无穷取整（\frac{一个线程块中线程的数量}{线程束大小}）$$

因此， 硬件总是给一个线程块分配一定数量的线程束。 线程束不会在不同的线程块之间分离。 如果线程块的大小不是线程束大小的偶数倍， 那么在最后的线程束里有些线程就不会活跃。

从逻辑角度来看， 线程块是线程的集合， 它们可以被组织为一维、 二维或三维布局。

从硬件角度来看， 线程块是一维线程束的集合。 在线程块中线程被组织成一维布局，每32个连续线程组成一个线程束。

3.2.2 线程束分化

GPU是相对简单的设备， 它没有复杂的分支预测机制。 一个线程束中的所有线程在同一周期中必须执行相同的指令， 如果一个线程执行一条指令， 那么线程束中的所有线程都必须执行该指令。 如果在同一线程束中的线程使用不同的路径通过同一个应用程序， 这可能会产生问题。

如果一个线程束中的线程产生分化， 线程束将连续执行每一个分支路径， 而禁用不执行这一路径的线程。 线程束分化会导致性能明显地下降。

重要提示:

• 当一个分化的线程采取不同的代码路径时， 会产生线程束分化
• 不同的if-then-else分支会连续执行
• 尝试调整分支粒度以适应线程束大小的倍数， 避免线程束分化
• 不同的分化可以执行不同的代码且无须以牺牲性能为代价

3.2.3 资源分配

线程束的本地执行上下文主要由以下资源组成：

• 程序计数器
• 寄存器
• 共享内存

由SM处理的每个线程束的执行上下文， 在整个线程束的生存期中是保存在芯片内的。 因此， 从一个执行上下文切换到另一个执行上下文没有损失。

每个SM都有32位的寄存器组， 它存储在寄存器文件中， 并且可以在线程中进行分配， 同时固定数量的共享内存用来在线程块中进行分配。 对于一个给定的内核， 同时存在于同一个SM中的线程块和线程束的数量取决于在SM中可用的且内核所需的寄存器和共享内存的数量。

若每个线程消耗的寄存器越多， 则可以放在一个SM中的线程束就越少。 如果可以减少内核消耗寄存器的数量， 那么就可以同时处理更多的线程束。

若一个线程块消耗的共享内存越多， 则在一个SM中可以被同时处理的线程块就会变少。 如果每个线程块使用的共享内存数量变少， 那么可以同时处理更多的线程块。

当计算资源（如寄存器和共享内存） 已分配给线程块时， 线程块被称为活跃的块。 它所包含的线程束被称为活跃的线程束。 活跃的线程束可以进一步被分为以下3种类型：

• 选定的线程束
• 阻塞的线程束
• 符合条件的线程束

一个SM上的线程束调度器在每个周期都选择活跃的线程束， 然后把它们调度到执行 单元。 活跃执行的线程束被称为选定的线程束。 如果一个活跃的线程束准备执行但尚未执 行， 它是一个符合条件的线程束。 如果一个线程束没有做好执行的准备， 它是一个阻塞的 线程束。 如果同时满足以下两个条件则线程束符合执行条件。

• 32个CUDA核心可用于执行
• 当前指令中所有的参数都已就绪

3.2.4 延迟隐藏

SM依赖线程级并行， 以最大化功能单元的利用率， 因此， 利用率与常驻线程束的数量直接相关。 在指令发出和完成之间的时钟周期被定义为指令延迟。 当每个时钟周期中所有的线程调度器都有一个符合条件的线程束时， 可以达到计算资源的完全利用。 这就可以保证， 通过在其他常驻线程束中发布其他指令， 可以隐藏每个指令的延迟。

考虑到指令延迟， 指令可以被分为两种基本类型：

• 算术指令: 一个算术操作从开始到它产生输出之间的时间；
• 内存指令: 指发送出的加载或存储操作和数据到达目的地之间的时间。

你可能想知道如何估算隐藏延迟所需要的活跃线程束的数量。 利特尔法则（Little’s Law） 可以提供一个合理的近似值。 它起源于队列理论中的一个定理， 它也可以应用于 GPU中：

$$所需线程束数量 = 延迟 \times 吞吐量$$

吞吐量和带宽

吞吐量和带宽都是用来度量性能的速度指标。

带宽通常是指理论峰值， 而吞吐量是指已达到的值

带宽通常是用来描述单位时间内最大可能的数据传输量， 而吞吐量是用来描述单位时 间内任何形式的信息或操作的执行速度， 例如， 每个周期完成多少个指令。

吞吐量由SM中每个周期内的操作数量确定， 而执行一条指令的一个线程束对应32个 操作。

这个简单的单位转换表明， 有两种方法可以提高并行：

• 指令级并行（ILP） ： 一个线程中有很多独立的指令
• 线程级并行（TLP） ： 很多并发地符合条件的线程

延迟隐藏取决于每个SM中活跃线程束的数量， 这一数量由执行配置和资源约束隐式 决定（一个内核中寄存器和共享内存的使用情况） 。 选择一个最优执行配置的关键是在延 迟隐藏和资源利用之间找到一种平衡。

显示充足的并行 因为GPU在线程间分配计算资源并在并发线程束之间切换的消耗（在一个或两个周期 命令上） 很小， 所以所需的状态可以在芯片内获得。 如果有足够的并发活跃线程， 那么可 以让GPU在每个周期内的每一个流水线阶段中忙碌。 在这种情况下， 一个线程束的延迟可 以被其他线程束的执行隐藏。 因此， 向SM显示足够的并行对性能是有利的

3.2.5 占用率

在每个CUDA核心里指令是顺序执行的。 当一个线程束阻塞时， SM切换执行其他符 合条件的线程束。 理想情况下， 我们想要有足够的线程束占用设备的核心。 占用率是每个 SM中活跃的线程束占最大线程束数量的比值。

$$占用率 = \frac{活跃线程束数量}{最大线程束数量}$$

极端地操纵线程块会限制资源的利用：

• 小线程块： 每个块中线程太少， 会在所有资源被充分利用之前导致硬件达到每个SM的线程束数量的限制
• 大线程块： 每个块中有太多的线程， 会导致在每个SM中每个线程可用的硬件资源较少

网格和线程块大小的准则

使用这些准则可以使应用程序适用于当前和将来的设备：

• 保持每个块中线程数量是线程束大小（32） 的倍数
• 避免块太小： 每个块至少要有128或256个线程
• 根据内核资源的需求调整块大小
• 块的数量要远远多于SM的数量， 从而在设备中可以显示有足够的并行
• 通过实验得到最佳执行配置和资源使用情况

占用率唯一注重的是在每个SM中并发线程或线 程束的数量。 然而， 充分的占用率不是性能优化的唯一目标。 内核一旦达到一定级别的占 用率， 进一步增加占用率可能不会改进性能。 为了提高性能， 可以调整很多其他因素。

3.2.6 同步

在CUDA中， 同步可以在两个级别执行：

• 系统级： 等待主机和设备完成所有的工作
• 块级： 在设备执行过程中等待一个线程块中所有线程到达同一点

对于主机来说：

• cudaError_t cudaDeviceSynchronize(void):cudaDeviceSyn-chronize函数可以用来阻塞主机应用程序， 直到所有的CUDA操作（复制、核函数等） 完成;
• __device__ void __syncthreads(void);:CUDA提供了一个使用块局部栅栏来同步它们的执行的功能。
  • 当__syncthreads被调用时， 在同一个线程块中每个线程都必须等待直至该线程块中所有其他线程都已经达到这个同步点。

线程块中的线程可以通过共享内存和寄存器来共享数据。

在不同的块之间没有线程同步。 块间同步， 唯一安全的方法是在每个内核执行结束端使用全局同步点； 也就是说， 在全局同步之后， 终止当前的核函数， 开始执行新的核函数。 不同块中的线程不允许相互同步， 因此GPU可以以任意顺序执行块。 这使得CUDA程序在大规模并行GPU上是可扩展的。

3.2.7 可扩展性

对于任何并行应用程序而言， 可扩展性是一个理想的特性。 可扩展性意味着为并行应用程序提供了额外的硬件资源， 相对于增加的资源， 并行应用程序会产生加速。 例如， 若一个CUDA程序在两个SM中是可扩展的， 则与在一个SM中运行相比， 在两个SM中运行会使运行时间减半。 一个可扩展的并行程序可以高效地使用所有的计算资源以提高性能。 可扩展性意味着增加的计算核心可以提高性能。 串行代码本身是不可扩展的， 因为在成千上万的内核上运行一个串行单线程应用程序， 对性能是没有影响的。 并行代码有可扩展的潜能， 但真正的可扩展性取决于算法设计和硬件特性。

3.3 并行性的表现

3.6 动态并行

在本书中， 到目前为止， 所有核函数都是从主机线程中被调用的。 GPU的工作负载完 全在CPU的控制下。 CUDA的动态并行允许在GPU端直接创建和同步新的GPU内核。 在一 个核函数中在任意点动态增加GPU应用程序的并行性， 是一个令人兴奋的新功能。