CUDA执行模型

CUDA执行模型概述

一般来说，执行模型会提供一个操作视图，说明如何在特定的计算架构上执行指令。CUDA执行模型揭示了GPU并行架构的抽象视图，使我们能够据此分析线程的并发。

GPU架构概述

GPU架构是围绕一个流式多处理器（SM）的可扩展阵列搭建的。可以通过复制这种架构的构建块来实现GPU的硬件并行。
在这里插入图片描述

上图说明了Fermi SM的关键组件：

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

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

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

SIMT架构与SIMD（单指令多数据）架构相似。两者都是将相同的指令广播给多个执行单元来实现并行。一个关键的区别是SIMD要求同一个向量中的所有元素要在一个统一的同步组中一起执行，而SIMT允许属于同一线程束的多个线程独立执行。尽管一个线程束中的所有线程在相同的程序地址上同时开始执行，但是单独的线程仍有可能有不同的行为。SIMT确保可以编写独立的线程级并行代码、标量线程以及用于协调线程的数据并行代码。

一个线程块只能在一个SM上被调度。一旦线程块在一个SM上被调度，就会保存在该SM上直到执行完成。在同一时间，一个SM可以容纳多个线程块。
在SM中，共享内存和寄存器是非常重要的资源。共享内存被分配在SM上的常驻线程块中，寄存器在线程中被分配。线程块中的线程通过这些资源可以进行相互的合作和通信。

尽管线程块里的所有线程都可以逻辑地并行运行，但是并不是所有线程都可以同时在物理层面执行。因此，线程块里的不同线程可能会以不同的速度前进。
在这里插入图片描述
在并行线程中共享数据可能会引起竞争：多个线程使用未定义的顺序访问同一个数据，从而导致不可预测的程序行为。CUDA提供了一种用来同步线程块里的线程的方法，从而保证所有线程在进一步动作之前都达到执行过程中的一个特定点。然而，没有提供块间同步的原语。

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

SM是GPU架构的核心。寄存器和共享内存是SM中的稀缺资源。CUDA将这些资源分配到SM中的所有常驻线程里。因此，这些有限的资源限制了在SM上活跃的线程束数量，活跃的线程束数量对应于SM上的并行量。

Fermi架构

Fermi架构是第一个完整的GPU计算架构，能够为大多数高性能计算应用提供所需要
的功能。Fermi已经被广泛应用于加速生产工作负载中。
在这里插入图片描述

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

Fermi架构包含一个耦合的768 KB的二级缓存，被16个SM所共享。在图中，一个垂直矩形条表示一个SM，包含了以下内容：

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

每一个多处理器有16个加载/存储单元，允许每个时钟周期内有16个线程（线程束的一半）计算源地址和目的地址。特殊功能单元（SFU）执行固有指令，如正弦、余弦、平方根和插值。每个SFU在每个时钟周期内的每个线程上执行一个固有指令。每个SM有两个线程束调度器和两个指令调度单元。当一个线程块被指定给一个SM时，线程块中的所有线程被分成了线程束。两个线程束调度器选择两个线程束，再把一个指令从线程束中发送到一个组上，组里有16个CUDA核心、16个加载/存储单元或4个特殊功能单元。Fermi架构，计算性能2.x，可以在每个SM上同时处理48个线程束，即可在一个SM上同时常驻1536个线程。

Fermi架构的一个关键特征是有一个64KB的片内可配置存储器，它在共享内存与一级缓存之间进行分配。对于许多高性能的应用程序，共享内存是影响性能的一个关键因素。共享内存允许一个块上的线程相互合作，这有利于芯片内数据的广泛重用，并大大降低了片外的通信量。CUDA提供了一个运行时API，它可以用来调整共享内存和一级缓存的数量。根据给定的内核中共享内存或缓存的使用修改片内存储器的配置，可以提高性能

Fermi架构也支持并发内核执行：在相同的GPU上执行相同应用程序的上下文中，同时启动多个内核。并发内核执行允许执行一些小的内核程序来充分利用GPU。Fermi架构允许多达16个内核同时在设备上运行。从程序员的角度看，并发内核执行使GPU表现得更像MIMD架构。

Kepler架构

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

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

理解线程束执行的本质

线程束是SM中基本的执行单元。当一个线程块的网格被启动后，网格中的线程块分布在SM中。一旦线程块被调度到一个SM上，线程块中的线程会被进一步划分为线程束。一个线程束由32个连续的线程组成，在一个线程束中，所有的线程按照单指令多线程（SIMT）方式执行；也就是说，所有线程都执行相同的指令，每个线程在私有数据上进行操作。下图展示了线程块的逻辑视图和硬件视图之间的关系。
在这里插入图片描述
然而，从硬件的角度来看，所有的线程都被组织成了一维的，线程块可以被配置为一维、二维或三维的。在一个块中，每个线程都有一个唯一的ID。对于一维的线程块，唯一的线程ID被存储在CUDA的内置变量threadIdx.x中，并且，threadIdx.x中拥有连续值的线程被分组到线程束中。例如，一个有128个线程的一维线程块被组织到4个线程束里，如下所示：
在这里插入图片描述
从逻辑角度来看，线程块是线程的集合，它们可以被组织为一维、二维或三维布局。从硬件角度来看，线程块是一维线程束的集合。在线程块中线程被组织成一维布局，每32个连续线程组成一个线程束。

资源分配

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

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

由SM处理的每个线程束的执行上下文，在整个线程束的生存期中是保存在芯片内的。因此，从一个执行上下文切换到另一个执行上下文没有损失。每个SM都有32位的寄存器组，它存储在寄存器文件中，并且可以在线程中进行分配，同时固定数量的共享内存用来在线程块中进行分配。对于一个给定的内核，同时存在于同一个SM中的线程块和线程束的数量取决于在SM中可用的且内核所需的寄存器和共享内存的数量。

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

资源可用性通常会限制SM中常驻线程块的数量。每个SM中寄存器和共享内存的数量因设备拥有不同的计算能力而不同。如果每个SM没有足够的寄存器或共享内存去处理至少一个块，那么内核将无法启动。当计算资源（如寄存器和共享内存）已分配给线程块时，线程块被称为活跃的块。它所包含的线程束被称为活跃的线程束。活跃的线程束可以进一步被分为以下3种类型：