计算机系统结构-第4章-数据级并行

数据集并行的概念:

并行场景1:

对不同数据执行相同的操作:

串行执行:

可以同时进行:

可以尝试一个多条指令,多核执行

引入:

SISD:

单核,单线程,串行执行,这样耗时

MIMD:

多核,多线程,并行执行,一条指令多次重复,变成了MIMID

存在的问题:

在标量CPU流水线中，取址、译码等操作逻辑复杂，且开销不低

对于数据级并行任务，无论是在SISD还是MIMD（多核）器件上运行，其取址、译码操作都是有冗余的

SIMD:

更多的CPU,但是有着更少的取指令和译码,更快的执行(但是也需要更多的数据寄存器,来支持增加后的多数据的存储)

两类处理机:

阵列处理机:同一指令在同一时间处理不同空间上的不同的数据元素

向量处理机:同一指令在连续的时间内在同一空间处理不同的数据元素

阵列处理机(同一个指令,同时在不同的地方处理)

向量处理机

让我们去对比SISD和MIMD,发现处理数据的时间更多了,明显更加的优秀了

(同一个指令,会处理很多个数据)

向量体系结构和GPU

向量体系结构特点:窄而宽,指令流水线深,ALU宽度窄,单词指令流水后可以处理更多的数据,掩盖不必要的流水时间

GPU的特点:宽而浅,ALU宽度宽,流水线本身简单,直接对更多的数据进行计算,同一时刻处理更多数据

向量体系结构:

-基础知识:

科学计算领域的很多问题需要处理大批量操作重复且互相没有关联的计算

向量:一组由一维数组组成的数据

标量:单个数据

在流水线处理机中，设置向量数据表示及相应的向量指令，称为向量处理机。

不具有向量数据表示和相应的向量指令的流水线处理机，称为标量处理机。

向量的处理方式:

D=A+(B+C)为例

-横向计算

一个一个的计算 di=ai*(bi+ci)

对于这个循环:先计算:ki=bi+ci,di=ai*ki

这里会造成N次数据相关,而且每次都会涉及到2次操作的切换,总:2*N次功能切换

-纵向计算

先计算:K=B+C

然后计算:D=A*K

这里涉及1次数据相关,1次功能切换,明显更优秀了

但是这里需要每次都访问完向量的所有元素,所以这就需要:存储器-存储器结构,这里可能造成每次的存储操作过于麻烦

-纵横计算

为了能够在寄存器之间进行我们的数据交换,我们会尝试将数据进行分组,每一组的大小允许我们将数据放到寄存器上,

假设向量的长度是N,每一组的长度是n,那么组数有N/n+(N%n!=0 ) =s组

所以先分组计算k=b+c,d=a*k

一共s组,每组一次数据相关,2次功能切换

这样就可以使用寄存器-寄存器结构了

总结:

1.向量体系结构应当具有很大的顺序寄存器堆 (Register File)，可加载更多向量元素以支持纵向计算

2.向量体系结构从内存中收集散落的数据，将其放入寄存器堆中，并对寄存器堆中的数据们进行操作，然后将这些结果放回内存（一次传输一组数据，LD/ST流水化）

3.一条指令能够对一个向量的数据进行操作，也就对向量中诸多独立数据元素进行了操作（纵向计算，功能单元流水化）

功能部件流水化:

比如我们将浮点加法流水化:

我们就会得到一个更快的处理过程

向量处理机的优势:

由于向量的ld和sd是深度流水化了,所以一个大型的寄存器也就起到了一个缓冲的作用,可以掩盖访存延迟带来的带宽

乱序执行的超标量处理器往往具有复杂的设计,乱序程度越高,其功能实现也就越复杂,容易触碰功能墙

顺序指令的标量处理器可以轻松的拓展为向量处理器

举例:

我们有8个Vector Reg,流水化的向量功能部件,每周期1个操作(需要控制单元来检查结构or数据冒险),ld和sd全流水化,每个周期1个字,标量寄存器集合(32位)

特殊的寄存器:

VL向量长度寄存器:我们的向量长度就要小于64,我们通过VL这个寄存器来控制

VM向量屏蔽寄存器:如果向量的长度小于64,或者控制语句下对某些元素单独操作的时候使用,即使maskcode有大部分的0(数据大部分无效),也可以获得很快的指令执行速度

Vmips向量指令格式:

合理使用Vmips不光可以节省代码,也可以大量节约取值和指令译码的时间

一些相关概念:

循环间相关：对一个循环来说，如果各轮迭代之间存在相关性，则称为循环间相关，否则为循环间无关

可向量化：针对一组MIPS指令描述的循环，如果满足循环间无关，则循环称为可向量化的，编译器可为其生成向量指令。

指令编队（convoy）：由一组不包含结构冒险的向量指令组成，一个编队中的所有向量指令在硬件条件允许时可以并行执行。

向量处理机的优化

-多车道技术:

在最开始的串行执行中尽可能的优化,使用多个ALU,就像是把单行道,拓展为4车道

-链接技术:**

链接技术：当两条指令出现“写后读”相关时，若它们不存在功能部件冲突和向量寄存器(源或目的) 冲突，就有可能把它们所用的功能部件头尾相接，形成一个链接(长)流水线，进行流水处理。

链接过程：无链接情况下，后面的功能需要等到前一个功能的n个结果都产生才能开始；而链接情况下，后面的功能只需要等到前一个功能的第一个结果产生就可以开始，即向量数据的生产与向量数据的消费进行延迟的重叠。

链接实质：把流水线定向的思想引入到向量执行过程，对两条流水线进行联合控制，没有改变寄存器和运算电路。

以下的讨论中假设各个部件之间传递一个结果需要一拍时间。

让我们来看一个例子:(向量寄存器冲突)
V3=V1+V2

V5=V4&V1

这里就产生了数据相关,向量寄存器冲突,不能够实现链接

功能部件冲突:

V3=V1*V5

V6=V7*V6

这里都使用了乘法部件,功能部件产生冲突

让我们来看一个链接的例子:

假设功能单元的时间开销为：浮点加减(6 cycle)，浮点乘法(7 cycle)，浮点存储操作(6 cycle)。为了同步要求，将向量元素送往功能部件，以及把结果存入向量寄存器需要一拍时间，从存储器中把数据送入访存功能部件也需要一拍时间。为以下程序画出链接示意图，并分析非链接执行和链接执行两种情况下的执行时间。假设向量长度为N，N≤64。

如果不链接的话,串行执行:

执行的时间:

[（1＋6＋1）＋N－1]＋[（1＋6＋1）＋N－1]＋[（1＋7＋1）＋N－1] = 3N ＋22

如果前面两个指令并行:

[（1＋6＋1）＋N－1]＋[（1＋7＋1）＋N－1] = 2N ＋15

如果前面两个并行,第三个指令链接执行

一个指令的时间: [（1＋6＋1）] ＋[（1＋7＋1）] = 17

N个指令: 17+N-1=N+16

总结:

-无向量寄存器使用冲突和无功能部件使用冲突

-链接的指令之间需要将前一个指令的元素送入寄存器才能开始链接

-当一条向量指令的两个源操作数分别是两条先行指令的结果寄存器时，要求先行的两条指令产生运算结果的时间必须相等，即要求有关功能部件的通过时间相等。

-要进行链接执行的向量指令的向量长度必须相等，否则无法进行链接。

-一次链接行为通常仅发生在分组内部，即不对整个N进行链接，而对个分组内的n个向量元素的计算过程进行链接

-编队技术:

定义:几条能在同一个时钟周期内一起开始执行的向量指令集合称为一个编队；

要求:同一个编队内:

不存在结构冲突；

不存在数据冲突；

存在数据冲突，但是可以链接。

例题:
假设每种向量功能部件只有一个，下面一组向量指令，在不使用链接技术和使用链接技术的情况下如何编队？

不用链接的编队:

LV->(MULTSV,LV)->ADDV->SV

使用链接的编队:

(LV,MULTSV)->(LV,ADDV)->SV

-分段开采技术

当向量的长度N大于向量寄存器的长度n时，必须把长向量N分成长度固定为n的段，然后循环分段处理，每一次循环只处理一个向量段。这种技术称为分段开采技术。

由系统硬件和编译软件合作完成控制，对程序员是透明的。

影响一个向量体系结构的因素:

操作数向量的长度

向量启动时间

数据相关,可否链接

结构相关,可否编队,多车道,发射限制

GPU体系结构:

介绍:

GPU:具有极高的计算吞吐率和内存带宽,专用图形图像处理器,适合纹理图像处理

GPGPU，通用图形处理器:可进行通用计算编程：例如CUDA, OpenCL适合处理SIMD程序

GPU拥有众核(上千个运算单元,支持更大量的运算)

GPU的处理模式:

流水线?

还是并发?

SPMD(单程序多数据)编程模型:

程序采用多线程的模型，而非向量指令,与SIMD编程模型不同

特点:

不同的进程/线程运行同一个程序源代码（SP），但是分别使用不同的输入数据进行各自的计算（MD）

不同进程/线程相互独立，没有执行顺序的要求

常用的并行编程模型多数采用SPMD模式

CUDA:采用SPMD编程方式的,基于C++,由NVIDIA提出的通用并行计算平台和编程模型

异构编程:

Host：

主机(CPU)

运行C++程序

Device：

物理上分离的协处理器(GPU)

运行kernel程序

运行thread

运行CUDA目标代码

我们以向量加法为例:

就可以在host上唤醒多个协处理器,,交给他们来帮助我们快速的去处理(线程)

GPU的执行方式:SIMT,

单指令多线程的方式

特点:

是SIMD和多线程的结合版,

线程按照固定的方式运行

固定个数的线程一起执行(SIMD)(Warp)

但是使用的标量指令(和传统的SIMD不同)

GPU处理单元的执行方式:SIMD

总结:
GPU的编程模型:SPMD

GPU的执行方式:SIMT

计算单元的处理方式:SIMD

优点:

编程灵活,支持任意大小的工作量,任意宽度的硬件

每个线程单独对待

wrap过程的对程序员透明

一个显卡的结构:

GPGPU微体系结构:

线程以warp为单位调度,

同一个warp的不同的线程按照simt的方式执行

大容量寄存器组用于存放所有线程的数据

SP（Streaming Processor）

进行数据的并行计算，i.e., ALU

同一个wrap中的不同线程按照SIMT的方式执行

流水线

Warp:

在NvidiaGPU中,Warp为32个标量线程构成的SIMT执行单元

,在GPUSM中,Warp是最小的调度和执行单位(也就是说GPU没法控制某个单独的线程)

Warp中不同的Thread执行保持同步(执行系统的PC对应的指令)

通过颗粒度多线程并行掩盖延迟

(右下角)

Warp调度:

通过Warp之间的交替调度掩盖长时间的方程带来的延迟

Warp执行:

通过多线程和并行处理单元来加速计算的原理。GPU 中 warp 是线程的基本调度单位，利用多个功能单元并行处理可以显著提升计算效率。

Functional Unit（功能单元）

功能单元是执行实际运算操作（如加法、乘法等）的部件。在 SIMD 架构中，它可以同时对多个数据元素执行相同的指令，提高计算效率。例如在图形处理中，对多个像素点进行色彩调整时，功能单元可并行处理。

Registers for each Thread（每个线程的寄存器）

寄存器是处理器内部高速存储单元，用于临时存储数据和指令。图中不同线程 ID 对应的寄存器组，分别存储对应线程的数据。线程在执行 SIMD 指令时，从这些寄存器中快速读取和写入数据，加快运算速度。

Lane（通道）

Lane 可理解为 SIMD 执行单元中的一条数据处理通路。每个 Lane 处理一部分数据，多个 Lane 并行工作，共同完成对多个数据的处理。例如在一个 4 - Lane 的 SIMD 单元中，一次可同时处理 4 组数据。

Memory Subsystem（内存子系统）

内存子系统负责数据的存储和传输，为 SIMD 执行单元提供运算所需的数据，并存储运算结果。它与寄存器和功能单元协同工作，确保数据在不同存储层次间高效流动。

Streaming Processor(流处理器)

GPU（图形处理单元）中用于处理数据并行计算任务的核心组件。它专注于对数据流进行高效处理，适用于需要大量并行计算的场景，如图形渲染、科学计算、深度学习等。

众多流处理器（SP）并行工作，每个流处理器又包含多个 CUDA core ，可以同时处理大量数据。像在深度学习训练中，大量神经元的计算任务就可由这些并行的流处理器高效完成，极大提升运算速度。

Warp指令级的并行执行

可以将不同指令交替执行

假设机器中有8个通道，1个warp中有32个线程

每周期完成24个操作，但是每周期流出1个warp

同一时间可以操作8个,4个周期一个warp的对应操作

SIMT的访存方式:

执行同一指令的不同线程使用线程ID来访问不同的数据

Coalescing:

将一个warp中不同thread的内存访问合并成更少的访问次数:

如果一个warp中的32个thread访问内存中连续的4B大小的位置,合并成一个128B的访存请求(coalescing),而不是发送32个4B的访存请求

有效降低SM和DRAM之间的访存次数 :减少片上网络、片上存储划分、DRAM的工作量

如图:

GPU程序层次模型:

一个程序kernel grid包含有多个block,一个block有多个thread,而thread作为标量线程执行完成的步骤代码

kernel以block为单位来分配到硬件中,多个block可以分配到同一个SM,只要空间足够

在SM内部,来自同一个block的不同线程组成一个Warp,这里的BLOCK和THREAD对程序员可见,但是warp对程序员不可见

同一kernel中的所有线程都共享相同的代码

不同线程执行进度可能不同

但是同一warp中的线程一同执行

可以在Block级别进行同步

例子:一个grid可以划分区域为多个可处理的单元*Block,里面多个线程,一次以SIMD方式处理

GPU和传统的SIMD的区别:

传统SIMD采用单线程

向量指令之间锁步执行：一个向量指令完成后后续指令才能开始执行

编程方式为SIMD，软件需要知道向量长度

ISA中有向量/SIMD指令

GPU以SIMD方式执行大量标量线程

可以非锁步执行

每个线程可以被独立看待，由硬件进行线程组合(Warp)

编程方式为SPMD

ISA是标量的

Nvidia ISA的特点

Parallel Thread Execution

(PTX)抽象特性：它是 NVIDIA 编译器的指令集目标，是对硬件指令集的一种抽象表示。不依赖于特定的 GPU 硬件版本，可实现各代 GPU 的兼容性，方便开发者编写通用代码，不同硬件都能处理基于 PTX 编写的程序。

指令格式：“opcode.type d,a,b,c;” 是其指令格式。其中 opcode 为操作码，指明要执行的操作（如加法、乘法等）；type 表示操作数的数据类型（如整数、浮点数等）；d 是目的操作数，a、b、c 是源操作数。

虚拟寄存器使用：使用虚拟寄存器，在编译和运行时由软件进行映射和管理，便于优化寄存器分配，提高代码执行效率，也增加了编程灵活性。

软件翻译：PTX 指令需要由软件（如 NVIDIA 的编译器）翻译成机器语言，才能被 GPU 硬件执行。

PTX 与 SASS 的转换

PTX（Parallel Thread Execution）：是 NVIDIA GPU 指令集的一种中间表示形式，具有硬件无关性，方便实现跨代 GPU 的代码兼容性。它是一种面向软件编程的抽象指令集，由 NVIDIA 编译器生成。

SASS（Serialized Assembler）：PTX 在运行时会被转换成 SASS。SASS 是 GPU 实际执行的机器指令形式，更贴近硬件底层，是一种二进制编码的汇编指令。这种转换过程类似于将高级语言编译为机器语言，使程序能在特定 GPU 硬件上执行。

SASS 与 GPGPU 计算能力的关系

对应关系：SASS 与 GPGPU（General - Purpose computing on Graphics Processing Units，通用图形处理器计算）计算能力相对应。不同的 GPU 计算架构（如不同代的 NVIDIA GPU 架构）具有不同的硬件特性和功能单元配置。一般来说，一种计算架构有一组对应的 SASS 指令，这些指令针对特定架构的硬件资源进行优化，以充分发挥该架构的计算性能。

跨架构复用：虽然一种架构通常有一组对应的 SASS 指令，但不同架构也可能采用相同的 SASS。这可能是因为某些指令在不同架构中功能和实现方式类似，或者是为了在新架构中保持对旧代码的兼容性，降低软件开发和移植成本。

分支处理:

分支处理原理

掩码机制：GPU 分支硬件借鉴向量处理机的向量屏蔽寄存器思路，使用掩码来标识哪些线程需要执行特定分支代码。掩码是一种二进制标识，每个位对应一个线程，通过掩码可快速筛选出符合条件的线程集合，让它们执行相应分支操作。

硬件栈处理：利用硬件栈辅助分支处理。遇到分支（分歧）时，通过 “push” 操作将当前状态压入栈中，并设置新的掩码位，确定执行不同分支的线程；当分支执行完成汇聚时，执行 “pop” 操作，从栈中恢复之前的状态，确保线程执行的正确性和连贯性。

示例代码解析

高级语言逻辑：示例中if (X[i] != 0)是条件判断，满足条件时执行X[i] = X[i] - Y[i]; ，否则执行X[i] = Z[i]; 。

汇编指令解析

ld.global.f64 RD0, [X+R8]; RD0 = X[i]

//从全局内存地址X+R8处加载 64 位浮点数到寄存器RD0 ，获取X[i]的值。

setp.neq.s32 P1, RD0, #0

比较RD0与 0 是否不相等，结果存入谓词寄存器P1 ，若不相等P1为真，否则为假。

@!P1, bra ELSE1, *Push;

//若P1为假（即X[i] == 0 ），通过bra（跳转）指令跳转到ELSE1标签处执行。跳转前执行*Push操作，将旧掩码压栈并设置新掩码位，确定执行该分支的线程。

后续指令ld.global.f64 RD2, [Y+R8]; RD2 = Y[i]等是条件为真时执行的操作，完成X[i] = X[i] - Y[i];计算并存储结果。

@P1, bra ENDIF1, *Comp;

//若P1为真，跳转到ENDIF1标签处，*Comp操作对掩码位取反。

ELSE1标签下指令

ld.global.f64 RD0, [Z+R8]; RD0 = Z[i]

//执行X[i] = Z[i];操作。

ENDIF1: <next instruction>, *Pop;

//执行*Pop操作，从栈中弹出旧掩码，恢复之前的状态，继续后续指令执行。

分支分歧:

同一个Warp的不同thread在遇到分支指令时产生的分歧

分支分歧：当同一warp中的不同线程在遇到分支指令时执行不同的路径

GPU使用简化的控制逻辑来减少控制部分所占的面积:thread无法单独控制和调度,并行的处理单元同一个时钟周期只能处理相同的操作

每一个Warp使用一个栈来处理分支分歧,分支指令中加入一个汇聚指令域

支分歧处理原理

栈处理机制：在 GPU 中，每个 warp（通常包含 32 个线程）会用一个栈（SIMT Stack ）来处理分支分歧。当遇到分支指令时，通过栈来管理不同分支路径的执行状态。栈中记录程序计数器（PC）、返回程序计数器（RPC）和活动掩码（Active Mask）等信息。活动掩码标识哪些线程执行当前分支，PC 记录当前执行指令位置，RPC 用于分支结束后恢复执行位置。

汇聚指令域：分支指令中可加入汇聚指令域，用于协调 warp 内不同线程在分支执行完后的汇聚操作，确保所有线程能正确恢复到统一执行路径，避免执行混乱。