67 KiB
67 KiB
第4章 数据级并行 - 章节关键原文
曲冠南老师（qgn）
Slide 4-3: 数据级并行——任务场景
科学计算：矩阵运算
三维渲染
图像处理
Slide 4-4: 数据级并行与SPMD
无论是矩阵运算还是图形图像处理，其共性是对大量的数据施加同种变换——数据级并行（DLP）
从软件的角度，在编程模型上，我们期待SPMD
SPMD：Single Program Multiple Data
Slide 4-5: 数据级并行与SPMD
思考：一个SPMD程序如何运行在SISD（或MIMD）上?
Slide 4-6: 数据级并行与SPMD
单核，单线程，串行化执行：
把MD拆成多次SD，变为SISD
多核，多线程，并行化执行：
把SP的每条I多次重复，变为MIMD
排队、耗时、耗电
Slide 4-8: 数据级并行——传统器件的问题
分析传统的标量CPU流水线可知，取址、译码等操作逻辑复杂，且开销不低；
对于SPMD任务，无论是在SISD还是MIMD（多核）器件上运行，其取址、译码操作都是有冗余的；
对于SPMD场景，如何设计一种新型的器件？
Slide 4-9: 数据级并行——SIMD
SIMD器件：更多的ALU (Execute)；更少的Fetch和Decode
Slide 4-10: 数据级并行——SIMD
更少的Fetch和Decode（甚至其他流水部件）意味着什么？——更少的器件，更低的能耗和时间开销
更多的ALU意味着什么？——一次流水能处理更多数据，速度更快
增加数据寄存器的数量来一次存储更多数据，以减少存储器访问延迟
Slide 4-12: 数据级并行——向量体系结构和GPU
向量体系结构
"窄而深"
指令流水线深，ALU宽度窄
单次指令流水后能处理更多数据，掩盖不必要的流水线时间
GPU
"宽而浅"
指令流水线浅，ALU宽度宽
流水本身比较简单，直接对更多的数据进行并行计算，同一时刻处理更多数据
Slide 4-15: 向量 vector
向量这种数据结构，以及向量的运算，和我们对SIMD的期待不谋而合
Slide 4-17: 向量的计算方式 – 1. 横向计算
向量计算是按行的方式从左到右横向地进行。
先计算： d1←a1×(b1＋c1)
再计算： d2←a2×(b2＋c2)
……
最后计算： dN←aN×(bN＋cN)
组成循环程序进行处理。
ki←bi＋ci
di←ki×ai
数据相关：N 次 功能切换：2N 次
这种计算方式是在标量处理器上对向量的一般计算方式，不适合于向量处理机的并行处理。
Slide 4-18: 向量的计算方式 – 2. 纵向计算
刚才的计算方式中，数据相关和功能切换随着向量长度增长而增长，硬件开销过大，并不是最优的向量计算方式。
下面介绍的纵向的计算方式更适合于进行向量计算。
Slide 4-19: 向量的计算方式 – 2. 纵向计算
向量计算是按列的方式从上到下纵向地进行。
k1←b1＋c1 d1←k1×a1
先计算 …… 再计算 ……
kN←bN＋cN dN←kN×aN
表示成向量指令：
K＝B＋C
D＝K×A
两条向量指令之间：
数据相关：1次 功能切换：1次
Slide 4-20 ~ 4-22: 向量的计算方式 – 3. 纵横（分组）计算
刚刚的纵向计算方式优化了向量计算的硬件开销，但是每次计算都需要访问到向量中的全部元素；
考虑到当前计算机体系结构的存储结构往往是层次化的，指令操作数一般都会加载到寄存器中，而寄存器的数量一般不会太多（相比于可以无限增长的向量长度N来说）；
结合前面两种计算方式，我们可以使用分组计算的方法。
把向量分成若干组，组内按纵向方式处理，依次处理各组。
对于上述的例子，设：
N=S×n＋r
其中N为向量长度，S为组数，n为每组的长度，r为余数。
若余下的r个数也作为一组处理，则共有S+1组。
先算第1组：
k1～n←b1～n＋c1～n
d1～n←k1～n×a1～n
再算第2组：
k(n+1)～2n←b(n+1)～2n＋c(n+1)～2n
d(n+1)～2n←k(n+1)～2n×a(n+1)～2n
依次进行下去，直到最后一组：第S+1组。
每组内各用两条向量指令。
数据相关：1次 功能切换：2次
Slide 4-23: 向量体系结构 Vector Architecture
根据刚刚探讨的串行化向量计算优化方式，我们得到一种更适用于向量处理的体系结构
向量体系结构应当具有很大的顺序寄存器堆 (Register File)，可加载更多向量元素以支持纵向计算
向量体系结构从内存中收集散落的数据，将其放入寄存器堆中，并对寄存器堆中的数据们进行操作，然后将这些结果放回内存（一次传输一组数据，LD/ST流水化）
一条指令能够对一个向量的数据进行操作，也就对向量中诸多独立数据元素进行了操作（纵向计算，功能单元流水化）
Slide 4-24: 向量体系结构的一些优势
由于向量的Load与Store是深度流水线化的，大型寄存器堆充当了Buffer的作用，因此其能够掩盖访存延迟并充分利用内存带宽；
乱序的超标量处理器往往具有复杂的设计，且乱序程度越高，其复杂性和功耗也会越高，在此方向发展很容易触及Power Wall；
将顺序的标量处理器扩展为向量处理器则不会带来复杂度和功耗的大幅升高，且开发者也能很容易适应和转换到向量指令。
Slide 4-25: 实例：VMIPS=标量MIPS+逻辑向量扩展
向量寄存器：
64bit*64元素*8个VR，足够多读写口
向量功能单元：
全流水化，每周期1操作
需要一个控制单元检测结构冒险和数据冒险
向量load/store单元：
全流水化，每周期1个字，可操作标量
标量寄存器集合：
32通用，32浮点，可存地址和数据
可以为向量功能单元提供数据，也可以为向量load/store单元提供地址
Slide 4-26: 向量专用特殊寄存器
向量长度寄存器VL
64位，每一位对应于向量寄存器的一个单元。
VL控制所有向量运算的长度，包括ld/st
作用：将软件层程序中实际向量长度N与硬件层向量寄存器中的元素数目64相适配
向量屏蔽寄存器VM
当向量长度小于64时，或者条件语句控制下对向量某些元素进行单独运算时使用
即使maskcode中有大量的0，使用VM的向量指令速度依然远远快于标量计算模式。
Slide 4-27: Vmips向量指令格式
VMIPS指令=MIPS指令+Op1类型+Op2类型.精度
misp：add.d
vmips：addvv.d(V1,V2,V3)，addvs.d(V1,V2,F0)
Slide 4-28: 代码：VMIPS vs mips
MIPS代码：
L.D F0,a ；载入标量
DADDIU R4,Rx,#512 ；载入最后地址
LOOP: L.D F2,0(Rx) ；载入X[i]
MUL.D F2,F2,F0 ；a x X[i]
L.D F4,0(Ry) ；载入Y[i]
ADD.D F4,F4,F2 ；载入a x X[i] + Y[i]
S.D F4,0(Ry) ；存入Y[i]
DADDIU Rx,Rx,#8 ；递增X
DADDIU Ry,Ry,#8 ；递增Y
DSUBU R20,R4,Rx ；计算范围
BNEZ R20,LOOP ；检查是否完成
VMIPS代码：
L.D F0,a ；载入a
LV V1,Rx ；载入X
MULVS.D V2,V1,F0 ；a x X
LV V3,Ry ；载入Y
ADDVV.D V4,V2,V2 ；+Y
SV V4,Ry ；存入Y
减少的仅仅是代码的行数么？
Slide 4-29: 向量体系结构的相关概念
循环间相关：对一个循环来说，如果各轮迭代之间存在相关性，则称为循环间相关，否则为循环间无关
可向量化：针对一组Mips指令描述的循环，如果满足循环间无关，则循环称为可向量化的，编译器可为其生成向量指令。
指令编队（convoy）：由一组不包含结构冒险的向量指令组成，一个编队中的所有向量指令在硬件条件允许时可以并行执行。
Slide 4-30: 向量体系结构的性能优化
多车道技术
链接技术
编队技术
分段开采技术
Slide 4-31: 多车道技术
刚刚讨论的是向量计算在串行化中尽可能进行优化的结果
从并行的角度去考虑，增加功能单元（ALU）的数量也能大大提升向量的计算速度（多车道）
这就好像，将一条单车道的窄巷扩宽成四车道的公路，吞吐量自然会提升
Slide 4-32: 链接技术
链接技术：当两条指令出现"写后读"相关时，若它们不存在功能部件冲突和向量寄存器(源或目的) 冲突，就有可能把它们所用的功能部件头尾相接，形成一个链接(长)流水线，进行流水处理。
链接过程：无链接情况下，后面的功能需要等到前一个功能的n个结果都产生才能开始；而链接情况下，后面的功能只需要等到前一个功能的第一个结果产生就可以开始，即向量数据的生产与向量数据的消费进行延迟的重叠。
链接实质：把流水线定向的思想引入到向量执行过程，对两条流水线进行联合控制，没有改变寄存器和运算电路。
以下的讨论中假设各个部件之间传递一个结果需要一拍时间。
Slide 4-33: 链接技术例题
假设功能单元的时间开销为：浮点加减(6 cycle)，浮点乘法(7 cycle)，浮点存储操作(6 cycle)。为了同步要求，将向量元素送往功能部件，以及把结果存入向量寄存器需要一拍时间，从存储器中把数据送入访存功能部件也需要一拍时间。为以下程序画出链接示意图，并分析非链接执行和链接执行两种情况下的执行时间。
V3 ← 存储器 // 访存取向量A
V2 ← V0 ＋ V1 // 向量B和向量C进行浮点加
V4 ← V2 × V3 // 浮点乘，结果存入V4
Slide 4-34: 链接技术例题解答
1. 三条指令全部用串行方法执行，则执行时间为：
[（1＋6＋1）＋N－1]＋[（1＋6＋1）＋N－1]
＋[（1＋7＋1）＋N－1] = 3N ＋22 （拍）
2. 前两条指令并行执行，然后再串行执行第3条指令，则执行时间为：
[（1＋6＋1）＋N－1]＋[（1＋7＋1）＋N－1] = 2N ＋15 （拍）
3. 前两条指令并行执行，并与第3条指令链接执行。
从访存开始到把第一个结果元素存入V4所需的拍数
（亦称为链接流水线的建立时间）为：
[（1＋6＋1）] ＋[（1＋7＋1）] = 17 （拍）
生成所有结果的执行时间为：
17 ＋（N－1） = N＋16 （拍） 解毕。
Slide 4-35: 进行向量链接的要求
保证：无向量寄存器使用冲突和无功能部件使用冲突
只有在前一条指令的第一个结果元素送入结果向量寄存器的那一个时钟周期才可以进行链接。
当一条向量指令的两个源操作数分别是两条先行指令的结果寄存器时，要求先行的两条指令产生运算结果的时间必须相等，即要求有关功能部件的通过时间相等。
要进行链接执行的向量指令的向量长度必须相等，否则无法进行链接。
一次链接行为通常仅发生在分组内部，即不对整个N进行链接，而对个分组内的n个向量元素的计算过程进行链接
Slide 4-36: 编队技术
几条能在同一个时钟周期内一起开始执行的向量指令集合称为一个编队；
同一个编队中的向量指令之间
不存在结构冲突；
不存在数据冲突；
存在数据冲突，但是可以链接。
Slide 4-37: 编队技术例题
假设每种向量功能部件只有一个，下面一组向量指令，在不使用链接技术和使用链接技术的情况下如何编队？
使用链接技术分为3个编队
第一编队：LV；MULTSV；
第二编队：LV；ADDV；
第三编队：SV
解：不使用链接技术分为4个编队
第一编队：LV
第二编队：MULTSV； LV
第三编队：ADDV
第四编队：SV
Slide 4-38: 分段开采技术
当向量的长度N大于向量寄存器的长度n时，必须把长向量N分成长度固定为n的段，然后循环分段处理，每一次循环只处理一个向量段。这种技术称为分段开采技术。
由系统硬件和编译软件合作完成控制，对程序员是透明的。
Slide 4-39: 向量体系结构的性能影响因素
操作数向量的长度
向量启动时间
操作之间的数据相关，是否采用链接
操作之间的结构性相关，发射限制，车道数量，是否采用编队
Slide 4-41: GPU: 为并行而生
pixel parallelism: accelerate the building of images
Computing: Graphics oriented -> General Purpose
Programmable: Graphic API (e.g. OpenGL)-> Computing API (e.g. OpenCL)
Problems are not needed to be expressed in terms of vertex coordinates, textures and shader programs, greatly decreasing program complexity.
Slide 4-42: Transistors Dedication
Divergent Design Choices
control logics v.s. arithmetic logics
Slide 4-43: Using CPU+GPU Architecture (Heterogeneous)
Use the right processor and memory for each task
CPU excels at executing a few serial threads
Fast sequential execution
Low latency cached memory access
GPU excels at executing many parallel threads
Scalable parallel execution
High bandwidth parallel memory access
Slide 4-44: GPGPU's benefits
It frees the programmer from possessing intimate knowledge of graphics APIs and GPU architecture.
Problems are not needed to be expressed in terms of vertex coordinates, textures and shader programs, greatly decreasing program complexity.
Basic programming features such as random reads and writes to memory are supported, greatly develop the programming model.
The import of double precision support scientific applications on the GPU with high performance.
Slide 4-46: Structure of GPGPU Chips
Slide 4-52: Scheduling Hierarchy
Slide 4-55: Code Pattern
Host Code Pattern
Malloc Host Memory
Malloc Device Memory
CUDAMemoryCopy(H2D)
Invoke Kernel Function
CUDAMemoryCopy(D2H)
Device Code Pattern
Locate Thread Index (TIDX)
Compute(TIDX)
CPU端
GPU端
Slide 4-60: Computing Hierarchy
Thread
Block
Grid
Core
Software
Hardware
SM
GPU
…
Slide 4-61: SIMD/SIMT：Micro Parallelism
Slide 4-62: SIMT Implementation
为每个warp维护一个PC
取指令
指令进入不同指令队列
进入不同的车道并行执行
Slide 4-64: Warps/Wavefront
Slide 4-72: 笔记
Grid是扔到GPU上的单元。
Grid中有若干block， 每个block扔到一个SM上
Block在扔到SM上的时候，被分成warp。
比如Block中是64个thread，这64个thread是 一个程序（100条指令）的实例，64个thread分成两组，每组32个，放入两个warp。
为了节省资源，因为32个thread的代码是相同的，所以在warp中使用的是100条instruction, 放在warp的指令队列中，每个instruction在执行的时候按照32个thread，操纵32个数据，用threadID对应各自的数据。
李宏图老师（lht）
（以下内容合并自20220526和20230331两个版本的PPT）
Slide（两版本共有）: 数据级并行——任务场景
科学计算：矩阵运算
三维渲染
图像处理
Slide: 数据级并行与SPMD
无论是矩阵运算还是图形图像处理，其共性是对大量的数据施加同种变换——数据级并行（DLP）
从软件的角度，在编程模型上，我们期待SPMD
SPMD：Single Program Multiple Data
Slide: 数据级并行
对大量的数据施加同种变换——数据级并行（DLP）
迭代1
迭代2
串行执行：
for (i=0; i < N; i++)
C[i] = A[i] + B[i];
Slide: 数据级并行
对大量的数据施加同种变换——数据级并行（DLP）
for (i=0; i < N; i++)
C[i] = A[i] + B[i];
迭代1
迭代2
并行执行：
Slide: ARM A53 架构
ARM A53 架构：观察并找出前端（取址、译码）部分
参考：https://zhuanlan.zhihu.com/p/469582393
Slide: 数据级并行——传统器件的问题
分析传统的标量CPU流水线可知，取址、译码等操作逻辑复杂，且开销不低
对于数据级并行任务，无论是在SISD还是MIMD（多核）器件上运行，其取址、译码操作都是有冗余的
对于数据级并行场景，如何设计一种新型的器件？
Slide: 数据级并行——SIMD
更多的ALU (Execute)；更少的Fetch和Decode
一种设计
更少的Fetch和Decode（甚至其他流水部件）意味着什么？——更少的器件，更低的能耗和时间开销
更多的ALU意味着什么？——一次流水能处理更多数据，速度更快
增加数据寄存器的数量，一次存储更多数据，以减少存储器访问延迟
Slide（20230331版本特有）: 单指令多数据（SIMD）处理
一条指令处理多个不同的数据
有多个处理单元PE（Processing Elements）
阵列处理机：同一指令在同一时间处理不同空间上的不同数据元素
向量处理机：同一指令在连续的时间内在同一空间上处理不同的数据元素
Slide: 数据级并行两种执行方式
阵列处理机vs向量处理机
ARRAY PROCESSOR
VECTOR PROCESSOR
LD0 LD1 LD2 LD3
AD0 AD1 AD2 AD3
MU0 MU1 MU2 MU3
ST0 ST1 ST2 ST3
LD0 LD1
AD0 AD1
MU0 MU1
LD2 LD3
AD2 AD3
ST0 ST1
MU2 MU3
ST2 ST3
Space
Space
Same op @ same time
Different ops @ same space
Different ops @ time
Same op @ space
LD VR A[3:0]
ADD VR VR, 1
MUL VR VR, 2
ST A[3:0] VR
Instruction Stream
Slide: SIMD vs. VLIW
阵列处理机：对不同数据进行同一操作
Slide: SIMD vs. VLIW
VLIW：多个独立操作由编译器捆绑在一起
Slide: 数据级并行——向量体系结构和GPU
向量体系结构
"窄而深"
指令流水线深，ALU宽度窄
单次指令流水后能处理更多数据，掩盖不必要的流水线时间
GPU
"宽而浅"
指令流水线浅，ALU宽度宽
流水本身比较简单，直接对更多的数据进行并行计算，同一时刻处理更多数据
Slide: 向量 vector
向量这种数据结构，以及向量的运算，和我们对SIMD的期待不谋而合
Slide: 向量的串行化计算方式
Slide: 向量的计算方式 – 1. 横向计算
这种计算方式是在标量处理器上对向量的一般计算方式
Slide: 向量的计算方式 – 1. 横向计算
向量计算是按行的方式从左到右横向地进行。
先计算： d1←a1×(b1＋c1)
再计算： d2←a2×(b2＋c2)
……
最后计算： dN←aN×(bN＋cN)
组成循环程序进行处理。
ki←bi＋ci
di←ki×ai
数据相关：N次 功能切换：2N次
不适合于向量处理机的并行处理。
Slide: 向量的计算方式 – 2. 纵向计算
刚才的计算方式中，数据相关和功能切换随着向量长度增长而增长，硬件开销过大，并不是最优的向量计算方式
下面介绍的纵向的计算方式更适合于进行向量计算
Slide: 向量的计算方式 – 2. 纵向计算
向量计算是按列的方式从上到下纵向地进行。
k1←b1＋c1 d1←k1×a1
先计算 …… 再计算 ……
kN←bN＋cN dN←kN×aN
表示成向量指令：
K＝B＋C
D＝K×A
两条向量指令之间：
数据相关：1次 功能切换：1次
Slide: 对处理机结构的要求：存储器－存储器结构
向量指令的源向量和目的向量都存放在存储器中，运算的中间结果需要送回存储器。
存储器－存储器型操作的运算流水线
例如：STAR-100、CYBER-205
Slide: 向量的计算方式 – 3. 纵横（分组）计算
刚刚的纵向计算方式优化了向量计算的硬件开销，但是每次计算都需要访问到向量中的全部元素
考虑到当前计算机体系结构的存储结构往往是层次化的，指令操作数一般都会加载到寄存器中，而寄存器的数量一般不会太多（相比于可以无限增长的向量长度N来说）
结合前面两种计算方式，我们可以使用分组计算的方法
Slide: 向量的计算方式 – 3. 纵横（分组）计算
把向量分成若干组，组内按纵向方式处理，依次处理各组。
对于上述的例子，设：
N=S×n＋r
其中N为向量长度，S为组数，n为每组的长度，r为余数。
若余下的r个数也作为一组处理，则共有S+1组。
Slide: 向量的计算方式 – 3. 纵横（分组）计算
先算第1组：
k1～n←b1～n＋c1～n
d1～n←k1～n×a1～n
再算第2组：
k(n+1)～2n←b(n+1)～2n＋c(n+1)～2n
d(n+1)～2n←k(n+1)～2n×a(n+1)～2n
依次进行下去，直到最后一组：第S+1组。
每组内各用两条向量指令。
数据相关：1次 功能切换：2次
Slide: 对处理机结构的要求：寄存器－寄存器结构
设置能快速访问的向量寄存器，用于存放源向量、目的向量及中间结果，让运算部件的输入、输出端都与向量寄存器相联，构成寄存器－寄存器型操作的运算流水线。
典型的寄存器－寄存器结构的向量处理机
美国的CRAY-1、我国的YH-1巨型机
Slide: 向量体系结构 Vector Architecture
根据刚刚探讨的串行化向量计算优化方式，我们得到一种更适用于向量处理的体系结构
向量体系结构应当具有很大的顺序寄存器堆 (Register File)，可加载更多向量元素以支持纵向计算
向量体系结构从内存中收集散落的数据，将其放入寄存器堆中，并对寄存器堆中的数据们进行操作，然后将这些结果放回内存（一次传输一组数据，LD/ST流水化）
一条指令能够对一个向量的数据进行操作，也就对向量中诸多独立数据元素进行了操作（纵向计算，功能单元流水化）
Slide: 向量体系结构的一些优势
由于向量的Load与Store是深度流水线化的，大型寄存器堆充当了Buffer的作用，因此其能够掩盖访存延迟并充分利用内存带宽
乱序的超标量处理器往往具有复杂的设计，且乱序程度越高，其复杂性和功耗也会越高，在此方向发展很容易触及Power Wall
将顺序的标量处理器扩展为向量处理器则不会带来复杂度和功耗的大幅升高，且开发者也能很容易适应和转换到向量指令
Slide: 实例：VMIPS=标量MIPS+逻辑向量扩展
向量寄存器：
64bit64元素*8个VR，足够多读写口
向量功能单元
全流水化，每周期1操作
需要一个控制单元检测结构冒险和数据冒险
向量load/store单元
全流水化，每周期1个字，可操作标量
标量寄存器集合
32通用，32浮点，可存地址和数据
可以为向量功能单元提供数据，也可以为向量load/store单元提供地址
Slide: 向量Vi冲突和功能部件冲突
只要不出现Vi冲突和功能部件冲突，各Vi之间和各功能部件之间都能并行工作，大大加快了向量指令的处理。
Vi冲突：并行工作的各向量指令的源向量或结果向量使用了相同的Vi。
例如：源向量相同
V3←V1＋V2
V5←V4∧V1
功能部件冲突：并行工作的各向量指令要使用同一个功能部件。
例如：都需使用乘法功能部件
V3←V1×V2
V5←V4×V6
Slide: 向量专用特殊寄存器
向量长度寄存器VL
64位，每一位对应于向量寄存器的一个单元。
VL控制所有向量运算的长度，包括ld/st
作用：将软件层程序中实际向量长度N与硬件层向量寄存器中的元素数目64相适配
向量屏蔽寄存器VM
当向量长度小于64时，或者条件语句控制下对向量某些元素进行单独运算时使用
即使maskcode中有大量的0，使用VM的向量指令速度依然远远快于标量计算模式。
Slide: 向量屏蔽寄存器VM
Slide（20230331版本）: CRAY-1的基本结构
共有12条可并行工作的单功能流水线，可分别流水地进行地址、向量、标量的各种运算。
其中6个单功能流水部件：进行向量运算
整数加（3）
逻辑运算（2）
移位（4）
浮点加（6）
浮点乘（7）
浮点迭代求倒数（14）
括号中的数字为其流水经过的时钟周期。
Slide: Vmips向量指令格式
VMIPS指令=MIPS指令+Op1类型+Op2类型.精度
misp：add.d
vmips：addvv.d(V1,V2,V3)，addvs.d(V1,V2,F0)
Slide: 代码：VMIPS vs mips
MIPS代码：
L.D FO,a ；载入标量
DADDIU R4,Rx,#512 ；载入最后地址
LOOP: L.D F2,0(Rx) ；载入X[i]
MUL.D F2,F2,F0 ；a × X[i]
L.D F4,0(Ry) ；载入Y[i]
ADD.D F4,F4,F2 ；载入a x X[i] + Y[i]
S.D F4,9(Ry) ；存入Y[i]
DADDIURx,Rx,#8 ；递增X
DADDIU Ry,Ry,#8 ；递增Y
DSUBU R20,R4,Rx ；计算范围
BNEZ R20,LOOP ；检查是否完成
VMIPS代码：
L.D F0,a ；载入a
LV V1,Rx ；载入X
MULVS.D V2,V1,F0 ；a × X
LV V3,Ry ；载入Y
ADDVV.D V4,V2,V2 ；+Y
SV V4,Ry ；存入Y
减少的仅仅是代码的行数么？
Slide: 向量体系结构的相关概念
循环间相关：对一个循环来说，如果各轮迭代之间存在相关性，则称为循环间相关，否则为循环间无关
可向量化：针对一组Mips指令描述的循环，如果满足循环间无关，则循环称为可向量化的，编译器可为其生成向量指令。
指令编队（convoy）：由一组不包含结构冒险的向量指令组成，一个编队中的所有向量指令在硬件条件允许时可以并行执行。
Slide: 向量体系结构的性能优化
多车道技术
链接技术
编队技术
分段开采技术
Slide: 多车道技术
刚刚讨论的是向量计算在串行化中尽可能进行优化的结果
从并行的角度去考虑，增加功能单元（ALU）的数量也能大大提升向量的计算速度（多车道）
这就好像，将一条单车道的窄巷扩宽成四车道的公路，吞吐量自然会提升
Slide: 链接技术
链接技术：当两条指令出现"写后读"相关时，若它们不存在功能部件冲突和向量寄存器(源或目的) 冲突，就有可能把它们所用的功能部件头尾相接，形成一个链接(长)流水线，进行流水处理。
链接过程：无链接情况下，后面的功能需要等到前一个功能的n个结果都产生才能开始；而链接情况下，后面的功能只需要等到前一个功能的第一个结果产生就可以开始，即向量数据的生产与向量数据的消费进行延迟的重叠。
链接实质：把流水线定向的思想引入到向量执行过程，对两条流水线进行联合控制，没有改变寄存器和运算电路。
以下的讨论中假设各个部件之间传递一个结果需要一拍时间。
Slide: 链接技术例题
假设功能单元的时间开销为：浮点加减(6 cycle)，浮点乘法(7 cycle)，浮点存储操作(6 cycle)。为以下程序画出链接示意图，并分析非链接执行和链接执行两种情况下的执行时间。
V3 ← 存储器 // 访存取向量A
V2 ← V0 ＋ V1 // 向量B和向量C进行浮点加
V4 ← V2 × V3 // 浮点乘，结果存入V4
Slide: 链接技术例题解答
1. 三条指令全部用串行方法执行，则执行时间为：
[（1＋6＋1）＋N－1]＋[（1＋6＋1）＋N－1]
＋[（1＋7＋1）＋N－1] = 3N ＋22 （拍）
2. 前两条指令并行执行，然后再串行执行第3条指令，则执行时间为：
[（1＋6＋1）＋N－1]＋[（1＋7＋1）＋N－1] = 2N ＋15 （拍）
3. 前两条指令并行执行，并与第3条指令链接执行。
从访存开始到把第一个结果元素存入V4所需的拍数
（亦称为链接流水线的建立时间）为：
[（1＋6＋1）] ＋[（1＋7＋1）] = 17 （拍）
生成所有结果的执行时间为：
17 ＋（N－1） = N＋16 （拍） 解毕。
Slide: 进行向量链接的要求
保证：无向量寄存器使用冲突和无功能部件使用冲突
只有在前一条指令的第一个结果元素送入结果向量寄存器的那一个时钟周期才可以进行链接。
当一条向量指令的两个源操作数分别是两条先行指令的结果寄存器时，要求先行的两条指令产生运算结果的时间必须相等，即要求有关功能部件的通过时间相等。
要进行链接执行的向量指令的向量长度必须相等，否则无法进行链接。
一次链接行为通常仅发生在分组内部，即不对整个N进行链接，而对个分组内的n个向量元素的计算过程进行链接
Slide: 编队技术
几条能在同一个时钟周期内一起开始执行的向量指令集合称为一个编队；
同一个编队中的向量指令之间
不存在结构冲突；
不存在数据冲突；
存在数据冲突，但是可以链接。
Slide: 编队技术例题
假设每种向量功能部件只有一个，下面一组向量指令，在不使用链接技术和使用链接技术的情况下如何编队？
LV V1，Rx
MULTSV V2，F0，V1
LV V3，Ry
ADDV V4，V2，V3
SV Ry，V4
使用链接技术分为3个编队
第一编队：LV；MULTSV；
第二编队：LV；ADDV；
第三编队：SV
解：不使用链接技术分为4个编队
第一编队：LV
第二编队：MULTSV； LV
第三编队：ADDV
第四编队：SV
Slide: 分段开采技术
当向量的长度N大于向量寄存器的长度n时，必须把长向量N分成长度固定为n的段，然后循环分段处理，每一次循环只处理一个向量段。这种技术称为分段开采技术。
由系统硬件和编译软件合作完成控制，对程序员是透明的。
Slide: 例3.4
设A和B是长度为N的向量，考虑在Cray-1向量处理器上实现以下的循环操作：
DO I = 1:N
A（I）= 5.0 * B（I） + C ;c=1.0
当N ≤64时，可以用以下指令序列：
S1 ← 5.0 ；将常数5.0送入标量寄存器S1
S2 ← 1.0 ；将常数1.0送入标量寄存器S2
VL ← N ；在向量长度寄存器VL中设置向量长度N
V0 ← B ；从存储器中将向量B读入向量寄存器V0
V1 ← S1 × V0 ；向量B中的每个元素分别和常数S1相乘
V2 ← S2 + V1 ；向量V1中的每个元素分别和常数S2相加
A ← V2 ；将计算结果从向量寄存器V2存入存储器的向量A
Slide: 当N＞64时的分段开采
当N ＞64时，就需要进行分段开采。
循环次数K ：
余数L：
S1 ← 5.0 ；将常数5.0送入标量寄存器S1
S2 ← 1.0 ；将常数1.0送入标量寄存器S2
VL ← L ；在向量长度寄存器VL中设置向量长度L
V0 ← B ；从存储器中将向量B[0..L-1]读入向量
；寄存器V0
V1 ← S1 * V0 ；向量B中的每个元素分别和常数S1相乘
V2 ← S2 + V1 ；向量V1中的每个元素分别和常数S2相加
A ← V2 ；将计算结果从向量寄存器V2存入存储器
；的向量A[0..L-1]
处理余
数部分,
计算L
个元素
For (I=0 to K-1) {
V0 ← B ；从存储器中将向量B[L+I*64..L+I*64+63]
；读入向量寄存器V0
V1 ← S1 * V0 ；向量B中的每个元素分别和常数S1相乘；
V2 ← S2 + V1 ；向量V1中的每个元素分别和常数S2相加；
A ← V2 ；将计算结果V2存入存储器的向量
； A[L+I*64… L+I*64+63]
}
循环
K次,
分段
处理
Slide: 向量体系结构的性能影响因素
操作数向量的长度
向量启动时间
操作之间的数据相关，是否采用链接
操作之间的结构性相关，发射限制，车道数量，是否采用编队
Slide: GPU: 为并行而生
pixel parallelism: accelerate the building of images
Computing: Graphics oriented -> General Purpose
Programmable: Graphic API (e.g. OpenGL)-> Computing API (e.g. OpenCL)
Problems are not needed to be expressed in terms of vertex coordinates, textures and shader programs, greatly decreasing program complexity.
Slide: Transistors Dedication
Divergent Design Choices
control logics v.s. arithmetic logics
Slide: Using CPU+GPU Architecture (Heterogeneous)
Use the right processor and memory for each task
CPU excels at executing a few serial threads
Fast sequential execution
Low latency cached memory access
GPU excels at executing many parallel threads
Scalable parallel execution
High bandwidth parallel memory access
Slide: GPGPU's benefits
It frees the programmer from possessing intimate knowledge of graphics APIs and GPU architecture.
Problems are not needed to be expressed in terms of vertex coordinates, textures and shader programs, greatly decreasing program complexity.
Basic programming features such as random reads and writes to memory are supported, greatly develop the programming model.
The import of double precision support scientific applications on the GPU with high performance.
Slide: GPU
GPU=Graphics Processing Unit，图形处理器
具有极高的计算吞吐率和内存带宽
专用图形图像处理器
适合纹理图像处理
GPGPU=General-purpose computing on GPUs，通用图形处理器
可进行通用计算编程：例如CUDA, OpenCL
适合处理SIMD程序
Slide: CPU vs GPU
多核 vs 众核：几个核 vs 上千个核
GPU将芯片中更多的晶体管用于计算
Slide: GPU处理模式
流水线 vs 并发执行
Slide: SPMD编程模型
Single Program Multiple Data (SPMD)单程序多数据
程序采用多线程的模型，而非向量指令
与SIMD编程模型不同
不同的进程/线程运行同一个程序源代码（SP），但是分别使用不同的输入数据进行各自的计算（MD）
不同进程/线程相互独立，没有执行顺序的要求
常用的并行编程模型多数采用SPMD模式
例如：CUDA，MPI，openMP
Slide: GPU编程模型
CUDA=Compute Unified Device Architecture,统一计算设备架构
由NVIDIA提出的通用并行计算平台和编程模型
CUDA软件环境使用C++作为高级编程语言
采用SPMD编程方式
Slide: 异构编程
Host：
主机(CPU)
运行C++程序
Device：
物理上分离的协处理器(GPU)
运行kernel程序
运行thread
运行CUDA目标代码
Slide: CUDA程序示例
向量加法
// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
int main()
{
...
// Kernel invocation with N threads
VecAdd<<<1, N>>>(A, B, C);
...
}
Slide: Host-Device数据传输
Host代码
… Allocate h_A, h_B, h_C …
void vecAdd(float *h_A, float *h_B, float *h_C, int n)
{
int size = n * sizeof(float); float *d_A, *d_B, *d_C;
cudaMalloc((void **) &d_A, size);
cudaMalloc((void **) &d_B, size);
cudaMalloc((void **) &d_C, size);
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
// Kernel invocation code – to be shown later
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
cudaFree(d_A); cudaFree(d_B); cudaFree (d_C);
}
… Free h_A, h_B, h_C …
Slide: GPU执行方式
Single Instruction Multiple Thread单指令多线程
SIMD与多线程相结合
线程按照固定的方式运行（与传统SPMD不同）
固定个数的线程一起执行（SIMD）
使用标量指令（与传统SIMD不同）
Slide: SPMD, SIMT, 和SIMD
SPMD
GPU的编程模型
SIMT
GPU的执行方式
SIMD
GPU计算单元的处理方式
Slide: SIMT的优点
编程灵活
任意大小的工作量vs任意硬件SIMD宽度
每个线程可以被单独对待
分支控制
Warp细节对程序员透明
Warp大小
哪些线程由同一个warp执行
Slide: NVIDIA H100 GPU
16896 FP32 CUDA cores
66.9 TFLOPS (FP64 Tensor core)
33.5 TFLOPS (FP64)
3352 GB/s memory bandwidth
700W TDP
TSMC's 4N customized for NVIDIA
Slide: GPGPU微体系结构
线程以warp为单位进行调度
同一warp中的不同线程按照SIMT方式执行
大容量寄存器组用于存放所有线程的数据
SP（Streaming Processor）进行数据的并行计算，i.e., ALU
标量指令处理
SIMD数据访问及计算
Slide: Warp
在Nvidia GPU中，warp为32个标量线程构成的SIMT执行单位
GPU SM中以warp为最小单位进行调度和执行
GPU无法控制单独执行某一个thread
Warp中不同的Thread在执行时保持同步
所有Thread永远执行相同PC所对应的指令
Slide: 通过细粒度多线程并行掩盖延迟
细粒度多线程执行
Slide: Warp调度
LD R1, 0x00027 PC=1
LD R2, 0x00034 PC=2
ADD R3, R1, R2 PC=3
ST R3, 0x00045 PC=4
Cycle 100: 调度 warp0 PC1
Cycle 101: 调度 warp1 PC2
Cycle 102: 调度 warp2 PC1
…
如果访存时间长，则PC3对应的指令需要等待，GPU将大量warp交替调度执行，可以一定程度上掩盖长时间访存
Slide: Warp执行
32-thread warp executing ADD A[tid],B[tid] → C[tid]
Slide: Streaming Processor
CUDA core
SIMD（Single Instruction Multiple Data）执行
大量SP提供极高并行算力
Slide: Warp指令级的并行执行
可以将不同指令交替执行
假设机器中有8个通道，1个warp中有32个线程
每周期完成24个操作，但是每周期流出1个warp
Slide: SIMT的访存模式
执行同一指令的不同线程使用线程ID来访问不同的数据
假设N=16, 每个warp中有4个线程 → 共有4个warp
Slide: Coalescing
将一个warp中不同thread的内存访问合并成更少的访问次数
如果一个warp中的32个thread访问内存中连续的4B大小的位置
合并成一个128B的访存请求(coalescing)
而不是发送32个4B的访存请求
有效降低SM和DRAM之间的访存次数
减少片上网络、片上存储划分、DRAM的工作量
Slide: GPU程序层次模型
一个程序kernel grid由多个block组成
例如，256个block
一个block由多个thread组成
例如，1024个thread
Thread为标量线程，执行完整的程序源代码
Slide: GPU程序层次模型
一个kernel以block为单位分配到硬件中
多个block可以分配到同一个SM中，只要有足够的硬件资源
Slide: GPU程序层次模型
在SM内部，来自同一个block中的不同线程（编号相邻）组成一个warp
例如，1 warp = 32 thread
Block和thread对程序员可见，warp对程序员不可见，由硬件进行控制
同一kernel中的所有线程都共享相同的代码
不同线程执行进度可能不同
但是同一warp中的线程一同执行
可以在Block级别进行同步
Slide: GPU程序层次例子
要处理所有元素的代码在一个grid中
将一个grid分割成可处理的大小（Block）
例如：一个block中有512个线程
一次以SIMD方式处理32个元素
Block与向量处理循环中的向量长度为32的分段开采类似
Slide: GPU与传统SIMD的区别
传统SIMD采用单线程
向量指令之间锁步执行：一个向量指令完成后后续指令才能开始执行
编程方式为SIMD，软件需要知道向量长度
ISA中有向量/SIMD指令
GPU以SIMD方式执行大量标量线程
可以非锁步执行
每个线程可以被独立看待，由硬件进行线程组合(Warp)
编程方式为SPMD
ISA是标量的
Slide: Nvidia ISA
"Parallel Thread Execution (PTX)"
NVIDIA编译器的指令集目标时硬件指令集的一种抽象，可实现各代GPU的兼容性
opcode.type d,a,b,c;
使用虚拟寄存器
由软件将其翻译成机器语言
shl.s32 R8, blockIdx, 9 ; Thread Block ID * Block size (512 or 2^9)
add.s32 R8, R8, threadIdx ; R8 = i = my CUDA thread ID
ld.global.f64 RD0, [X+R8] ; RD0 = X[i]
ld.global.f64 RD2, [Y+R8] ; RD2 = Y[i]
mul.f64 R0D, RD0, RD4 ; Product in RD0 = RD0 * RD4 (scalar a)
add.f64 R0D, RD0, RD2 ; Sum in RD0 = RD0 + RD2 (Y[i])
st.global.f64 [Y+R8], RD0 ; Y[i] = sum (X[i]*a + Y[i])
Slide: Nvidia机器指令
PTX在运行时转换成SASS
SASS与GPGPU计算能力对应，一般一种计算架构有一组对应的SASS指令
不同架构也可能采用相同的SASS
/*0030*/ MOV R6, 0x4;
/*0040*/ S2R R3, SR_TID.X
Slide: 分支处理
GPU分支硬件也使用掩码，类似向量处理机中的向量屏蔽寄存器
此外，还使用硬件栈处理分支
分歧后push
汇聚后pop
Slide: 例子（分支处理）
if (X[i] != 0)
X[i] = X[i] – Y[i];
else X[i] = Z[i];
ld.global.f64 RD0, [X+R8] ; RD0 = X[i]
setp.neq.s32 P1, RD0, #0 ; P1 is predicate register 1
@!P1, bra ELSE1, *Push ; Push old mask, set new mask bits
; if P1 false, go to ELSE1
ld.global.f64 RD2, [Y+R8] ; RD2 = Y[i]
sub.f64 RD0, RD0, RD2 ; Difference in RD0
st.global.f64 [X+R8], RD0 ; X[i] = RD0
@P1, bra ENDIF1, *Comp ; complement mask bits
; if P1 true, go to ENDIF1
ELSE1: ld.global.f64 RD0, [Z+R8] ; RD0 = Z[i]
st.global.f64 [X+R8], RD0 ; X[i] = RD0
ENDIF1: <next instruction>, *Pop ; pop to restore old mask
*Push， *Comp， *Pop为由PTX汇编程序插入的分支同步标记
Slide: 分支分歧(branch divergence)
同一warp中的不同thread在遇到分支指令时可能会有分歧
例如，if-else
Slide: 分支分歧(branch divergence)
分支分歧：当同一warp中的不同线程在遇到分支指令时执行不同的路径
GPU使用简化的控制逻辑来减少控制部分所占的面积
thread无法单独控制和调度
并行的处理单元同一个时钟周期只能处理相同的操作
Slide: 分支分歧(branch divergence)
每个warp用一个栈来处理分支分歧
分支指令中可以加入一个汇聚指令域
A: v = foo[tid.x];
B: if (v < 10)
C: v = 0;
else
D: v = 10;
E: w = bar[tid.x]+v;
foo[] = {4,8,12,16};
One stack per warp
Slide: 分支分歧处理
每个warp使用一个栈来处理分支分歧的执行
Slide: 存储层次
Local memory
一个thread
Shared memory
一个block中的所有thread
Global memory
所有thread
Texture memory
只读，所有thread
适合特定地址访问模型
Constant memory
只读，所有thread
Slide: AMD GPU
AMD TeraScale Architecture: VLIW
CU = SM, Work group = block (CTA), work item = thread, wavefront = warp (64)
Slide: AMD GPU
AMD GCN Architecture
Wavefront size = 64
每个SIMD Unit 中有 16 ALU
Slide: AMD GPU
AMD RDNA Architecture
Wavefront size = 32
每个SIMD unit中有32个ALU
Slide: GPGPU Virtualization虚拟化
企业级做云计算时常常采用虚拟化技术：
一个GPU的SM很多，但是对于单用户只能用很少的SM,这样利用率很低；
硬件上把多个SM进行分割成多个切片；
软件上为每个用户分配一个SM切片；
使得每个用户都能满负荷使用SM;
每个用户感觉自己拥有了一个GPU。
Slide: 笔记
Grid是扔到GPU上的单元。
Grid中有若干block， 每个block扔到一个SM上
Block在扔到SM上的时候，被分成warp。
比如Block中是64个thread，这64个thread是 一个程序（100条指令）的实例，64个thread分成两组，每组32个，放入两个warp。
为了节省资源，因为32个thread的代码是相同的，所以在warp中使用的是100条instruction, 放在warp的指令队列中，每个instruction在执行的时候按照32个thread，操纵32个数据，用threadID对应各自的数据。
谭婧炜佳老师（tjwj）
Slide 4-1: 计算机系统结构第4章 数据级并行
授课教师：谭婧炜佳
jtan@jlu.edu.cn
Slide 4-2 ~ 4-3: 本章内容
数据级并行的概念
向量体系结构
GPU体系结构
Slide 4-4: 数据级并行——任务场景
对大量不同数据执行相同操作
科学计算：矩阵运算
三维渲染
图像处理
Slide 4-5: 数据级并行——计算场景
对大量不同数据执行相同操作
例如：矩阵/向量加法
串行执行：
for (i=0; i < N; i++)
C[i] = A[i] + B[i];
迭代1
迭代2
Slide 4-6: 数据级并行——计算场景
对大量不同数据执行相同操作
例如：矩阵/向量加法
并行执行：
for (i=0; i < N; i++)
C[i] = A[i] + B[i];
迭代1
迭代2
Slide 4-7: 数据级并行程序执行
Slide 4-8: 存在的问题
单核，单线程，串行化执行：
把MD拆成多次SD，变为SISD
多核，多线程，并行化执行：
把SI的每条I多次重复，变为MIMD
排队、耗时、耗电
Slide 4-9: 存在的问题
在标量CPU流水线中，取址、译码等操作逻辑复杂，且开销不低
对于数据级并行任务，无论是在SISD还是MIMD（多核）器件上运行，其取址、译码操作都是有冗余的
对于数据级并行场景，如何设计一种新型的处理器？
Slide 4-10: SIMD处理器
更多的ALU (Execute)；更少的Fetch和Decode
更少的Fetch和Decode（甚至其他流水部件）意味着什么？——更少的器件，更低的能耗和时间开销
更多的ALU意味着什么？——同时能处理更多数据，速度更快
增加数据寄存器的数量来一次存储更多数据，以减少存储器访问延迟
Slide 4-11: 单指令多数据（SIMD）处理
一条指令处理多个不同的数据
有多个处理单元PE（Processing Elements）
阵列处理机：同一指令在同一时间处理不同空间上的不同数据元素
向量处理机：同一指令在连续的时间内在同一空间上处理不同的数据元素
Slide 4-12: 数据级并行两种执行方式
阵列处理机vs向量处理机
ARRAY PROCESSOR
VECTOR PROCESSOR
LD0 LD1 LD2 LD3
AD0 AD1 AD2 AD3
MU0 MU1 MU2 MU3
ST0 ST1 ST2 ST3
LD0 LD1
AD0 AD1
MU0 MU1
LD2 LD3
AD2 AD3
ST0 ST1
MU2 MU3
ST2 ST3
Space
Space
Same op @ same time
Different ops @ same space
Different ops @ time
Same op @ space
LD VR A[3:0]
ADD VR VR, 1
MUL VR VR, 2
ST A[3:0] VR
Instruction Stream
Slide 4-13: 向量体系结构和GPU
向量体系结构
"窄而深"
指令流水线深，ALU宽度窄
单次指令流水后能处理更多数据，掩盖不必要的流水线时间
GPU
"宽而浅"
指令流水线浅，ALU宽度宽
流水本身比较简单，直接对更多的数据进行并行计算，同一时刻处理更多数据
Slide 4-14: SISD、MIMD与SIMD
SISD、MIMD与SIMD阵列处理机、SIMD向量处理机的时间对比图
Slide 4-16: 向量
科学计算领域的很多问题需要处理大批量操作重复且互相没有关联的计算
向量：一组由一维数组所组成的数据
标量：单个数据
在流水线处理机中，设置向量数据表示及相应的向量指令，称为向量处理机。
不具有向量数据表示和相应的向量指令的流水线处理机，称为标量处理机。
Slide 4-17: 向量处理方式
Slide 4-18: 横向（水平）计算
向量计算是按行的方式从左到右横向地进行。
先计算： d1←a1×(b1＋c1)
再计算： d2←a2×(b2＋c2)
……
最后计算： dN←aN×(bN＋cN)
组成循环程序进行处理。
ki←bi＋ci
di←ki×ai
数据相关：N次 功能切换：2N次
不适合于向量处理机的并行处理。
Slide 4-19: 纵向（垂直）处理方式
向量计算是按列的方式从上到下纵向地进行。
k1←b1＋c1 d1←k1×a1
先计算 …… 再计算 ……
kN←bN＋cN dN←kN×aN
表示成向量指令：
K＝B＋C
D＝K×A
两条向量指令之间：
数据相关：1次 功能切换：1次
适合于向量处理机的并行处理。
Slide 4-20: 纵向处理方式
纵向处理方式N的大小不受限制
每次计算都需要访问到向量中的全部元素
纵向处理方式对处理机结构的要求：存储器－存储器结构
向量指令的源向量和目的向量都存放在存储器中，运算的中间结果需要送回存储器。
Slide 4-21: 纵横 (分组)处理方式
当前计算机体系结构的存储结构往往是层次化的，指令操作数一般都会加载到寄存器中，而寄存器的数量一般不会太多（相比于可以无限增长的向量长度N来说）
分组处理方式：把向量分成若干组，组内按纵向方式处理，依次处理各组。
对于上述的例子，设：
N=S×n＋r
其中N为向量长度，S为组数，n为每组的长度，r为余数。
若余下的r个数也作为一组处理，则共有S+1组。
Slide 4-22: 纵横 (分组)处理方式
运算过程为：
先算第1组：
k1～n←b1～n＋c1～n
d1～n←k1～n×a1～n
再算第2组：
k(n+1)～2n←b(n+1)～2n＋c(n+1)～2n
d(n+1)～2n←k(n+1)～2n×a(n+1)～2n
依次进行下去，直到最后一组：第S+1组。
每组内各用两条向量指令。
数据相关：1次 功能切换：2次
适合于向量处理机的并行处理。
Slide 4-23: 纵横处理方式
纵横处理方式N的大小不受限制，但n值固定
对处理机结构的要求：寄存器－寄存器结构
设置能快速访问的向量寄存器，用于存放源向量、目的向量及中间结果，让运算部件的输入、输出端都与向量寄存器相联，构成寄存器－寄存器型操作的运算流水线。
典型的寄存器－寄存器结构的向量处理机
美国的CRAY-1、我国的YH-1巨型机
Slide 4-24: 向量体系结构Vector Architecture
向量体系结构应当具有很大的顺序寄存器堆 (Register File)，可加载更多向量元素以支持纵向计算
向量体系结构从内存中收集散落的数据，将其放入寄存器堆中，并对寄存器堆中的数据们进行操作，然后将这些结果放回内存（一次传输一组数据，LD/ST流水化）
一条指令能够对一个向量的数据进行操作，也就对向量中诸多独立数据元素进行了操作（纵向计算，功能单元流水化）
Slide 4-25: 功能部件流水化
把流水线技术应用于运算的执行过程，就形成了运算操作流水线，也称为部件级流水线。
例如：浮点加法流水线
把浮点加法的全过程分解为求阶差、对阶、尾数相加、规格化4个子过程。
Slide 4-26: 向量体系结构的一些优势
由于向量的Load与Store是深度流水线化的，大型寄存器堆充当了缓冲的作用，因此其能够掩盖访存延迟并充分利用内存带宽；
乱序的超标量处理器往往具有复杂的设计，且乱序程度越高，其复杂性和功耗也会越高，在此方向发展很容易触及功耗墙；
将顺序的标量处理器扩展为向量处理器则不会带来复杂度和功耗的大幅升高，且开发者也能很容易适应和转换到向量指令。
Slide 4-27: 实例：VMIPS=标量MIPS+逻辑向量扩展
向量寄存器：
64bit*64元素*8个VR，足够多读写口
向量功能单元：
全流水化，每周期1操作
需要一个控制单元检测结构冒险和数据冒险
向量load/store单元：
全流水化，每周期1个字，可操作标量
标量寄存器集合：
32通用，32浮点，可存地址和数据
可以为向量功能单元提供数据，也可以为向量load/store单元提供地址
Slide 4-28: 向量专用特殊寄存器
向量长度寄存器VL
表示向量的长度，值≤64。
VL控制所有向量运算的长度，包括ld/st
作用：将软件层程序中实际向量长度N与硬件层向量寄存器中的元素数目64相适配
向量屏蔽寄存器VM
当向量长度小于64时，或者条件语句控制下对向量某些元素进行单独运算时使用
即使maskcode中有大量的0，使用VM的向量指令速度依然远远快于标量计算模式。
Slide 4-29: VMIPS向量指令格式
VMIPS指令=MIPS指令+Op1类型+Op2类型.精度
mips：add.d
vmips：addvv.d(V1,V2,V3)，addvs.d(V1,V2,F0)
Slide 4-30: 代码：VMIPS vs MIPS
a*X[i]+Y[i]
MIPS代码：
L.D F0,a ；载入标量
DADDIU R4,Rx,#512 ；载入最后地址
LOOP: L.D F2,0(Rx) ；载入X[i]
MUL.D F2,F2,F0 ；a * X[i]
L.D F4,0(Ry) ；载入Y[i]
ADD.D F4,F4,F2 ；载入a * X[i] + Y[i]
S.D F4,0(Ry) ；存入Y[i]
DADDIU Rx,Rx,#8 ；递增X
DADDIU Ry,Ry,#8 ；递增Y
DSUBU R20,R4,Rx ；计算范围
BNEZ R20,LOOP ；检查是否完成
VMIPS代码：
L.D F0,a ；载入a
LV V1,Rx ；载入X
MULVS.D V2,V1,F0 ； a x X
LV V3,Ry ；载入Y
ADDVV.D V4,V2,V3 ； +Y
SV Ry, V4 ；存入Y
减少的仅仅是代码的行数么？
不只是行数，减少了指令条数，大量节约取指令和指令译码所消耗的时间
Slide 4-31: 向量体系结构的相关概念
循环间相关：对一个循环来说，如果各轮迭代之间存在相关性，则称为循环间相关，否则为循环间无关
可向量化：针对一组MIPS指令描述的循环，如果满足循环间无关，则循环称为可向量化的，编译器可为其生成向量指令。
指令编队（convoy）：由一组不包含结构冒险的向量指令组成，一个编队中的所有向量指令在硬件条件允许时可以并行执行。
Slide 4-32: 向量体系结构的性能优化
多车道技术
链接技术
编队技术
分段开采技术
Slide 4-33: 多车道技术
刚刚讨论的是向量计算在串行化中尽可能进行优化的结果
从并行的角度去考虑，增加功能单元（ALU）的数量也能大大提升向量的计算速度（多车道）
这就好像，将一条单车道的窄巷扩宽成四车道的公路，吞吐量自然会提升
Slide 4-34: 链接技术
链接技术：当两条指令出现"写后读"相关时，若它们不存在功能部件冲突和向量寄存器(源或目的) 冲突，就有可能把它们所用的功能部件头尾相接，形成一个链接(长)流水线，进行流水处理。
链接过程：无链接情况下，后面的功能需要等到前一个功能的n个结果都产生才能开始；而链接情况下，后面的功能只需要等到前一个功能的第一个结果产生就可以开始，即向量数据的生产与向量数据的消费进行延迟的重叠。
链接实质：把流水线定向的思想引入到向量执行过程，对两条流水线进行联合控制，没有改变寄存器和运算电路。
以下的讨论中假设各个部件之间传递一个结果需要一拍时间。
Slide 4-35: 冲突例子
Vi冲突：并行工作的各向量指令的源向量或结果向量使用了相同的Vi。
例如：源向量相同
V3←V1＋V2
V5←V4∧V1
功能部件冲突：并行工作的各向量指令要使用同一个功能部件。
例如：都需使用乘法功能部件
V3←V1×V2
V5←V4×V6
Slide 4-36: 链接技术例题
假设功能单元的时间开销为：浮点加减(6 cycle)，浮点乘法(7 cycle)，浮点存储操作(6 cycle)。为了同步要求，将向量元素送往功能部件，以及把结果存入向量寄存器需要一拍时间，从存储器中把数据送入访存功能部件也需要一拍时间。为以下程序画出链接示意图，并分析非链接执行和链接执行两种情况下的执行时间。假设向量长度为N，N≤64。
V3 ← 存储器 // 访存取向量A
V2 ← V0 ＋ V1 // 向量B和向量C进行浮点加
V4 ← V2 × V3 // 浮点乘，结果存入V4
Slide 4-37: 链接技术例题解答
1. 三条指令全部用串行方法执行，则执行时间为：
[（1＋6＋1）＋N－1]＋[（1＋6＋1）＋N－1]
＋[（1＋7＋1）＋N－1] = 3N ＋22 （拍）
2. 前两条指令并行执行，然后再串行执行第3条指令，则执行时间为：
[（1＋6＋1）＋N－1]＋[（1＋7＋1）＋N－1] = 2N ＋15 （拍）
3. 前两条指令并行执行，并与第3条指令链接执行。
从访存开始到把第一个结果元素存入V4所需的拍数
（亦称为链接流水线的建立时间）为：
[（1＋6＋1）] ＋[（1＋7＋1）] = 17 （拍）
生成所有结果的执行时间为：
17 ＋（N－1） = N＋16 （拍） 解毕。
Slide 4-38: 进行向量链接的要求
保证：无向量寄存器使用冲突和无功能部件使用冲突
只有在前一条指令的第一个结果元素送入结果向量寄存器的那一个时钟周期才可以进行链接。
当一条向量指令的两个源操作数分别是两条先行指令的结果寄存器时，要求先行的两条指令产生运算结果的时间必须相等，即要求有关功能部件的通过时间相等。
要进行链接执行的向量指令的向量长度必须相等，否则无法进行链接。
一次链接行为通常仅发生在分组内部，即不对整个N进行链接，而对个分组内的n个向量元素的计算过程进行链接
Slide 4-39: 编队技术
几条能在同一个时钟周期内一起开始执行的向量指令集合称为一个编队；
同一个编队中的向量指令之间
不存在结构冲突；
不存在数据冲突；
存在数据冲突，但是可以链接。
Slide 4-40 ~ 4-41: 编队技术例题
LV V1,Rx ；从存储器中载入向量Rx
MULTSV V2,R0,V1 ；向量V1和标量寄存器R0中的内容相乘
LV V3,Ry ；从存储器中载入向量Ry
ADDV V4,V2,V3 ；向量V2和V3相加，结果保存在向量寄存器V4中
SV Ry, V4 ；将结果存入Ry
假设每种向量功能部件只有一个，下面一组向量指令，在不使用链接技术和使用链接技术的情况下如何编队？
写后读，通过链接可编入同一编队
结构冒险，
需要放入新编队
写后读，通过链接可编入同一编队
结构冒险，
需要放入新编队
编队
Ry=R0*Rx+Ry
使用链接技术分为3个编队
第一编队：LV；MULTSV；
第二编队：LV；ADDV；
第三编队：SV
解：不使用链接技术分为4个编队
第一编队：LV
第二编队：MULTSV； LV
第三编队：ADDV
第四编队：SV
Slide 4-42: 分段开采技术
当向量的长度N大于向量寄存器的长度n时，必须把长向量N分成长度固定为n的段，然后循环分段处理，每一次循环只处理一个向量段。这种技术称为分段开采技术。
由系统硬件和编译软件合作完成控制，对程序员是透明的。
Slide 4-43: 向量体系结构的性能影响因素
操作数向量的长度
向量启动时间
操作之间的数据相关，是否采用链接
操作之间的结构性相关，发射限制，车道数量，是否采用编队
Slide 4-45: GPU
GPU=Graphics Processing Unit，图形处理器
具有极高的计算吞吐率和内存带宽
专用图形图像处理器
适合纹理图像处理
GPGPU=General-purpose computing on GPUs，通用图形处理器
可进行通用计算编程：例如CUDA, OpenCL
适合处理SIMD程序
Slide 4-46: CPU vs GPU
多核 vs 众核：几个核 vs 上千个核
GPU将芯片中更多的晶体管用于计算
Slide 4-47: GPU处理模式
流水线 vs 并发执行
Slide 4-48: SPMD编程模型
Single Program Multiple Data (SPMD)单程序多数据
程序采用多线程的模型，而非向量指令
与SIMD编程模型不同
不同的进程/线程运行同一个程序源代码（SP），但是分别使用不同的输入数据进行各自的计算（MD）
不同进程/线程相互独立，没有执行顺序的要求
常用的并行编程模型多数采用SPMD模式
例如：CUDA，MPI，openMP
Slide 4-49: GPU编程模型
CUDA=Compute Unified Device Architecture, 统一计算设备架构
由NVIDIA提出的通用并行计算平台和编程模型
CUDA软件环境使用C++作为高级编程语言
采用SPMD编程方式
Slide 4-50: 异构编程
Host：
主机(CPU)
运行C++程序
Device：
物理上分离的协处理器(GPU)
运行kernel程序
运行thread
运行CUDA目标代码
Slide 4-51: CUDA程序示例
向量加法
// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
int main()
{
...
// Kernel invocation with N threads
VecAdd<<<1, N>>>(A, B, C);
...
}
Slide 4-52: Host-Device数据传输
Host代码
… Allocate h_A, h_B, h_C …
void vecAdd(float *h_A, float *h_B, float *h_C, int n)
{
int size = n * sizeof(float); float *d_A, *d_B, *d_C;
cudaMalloc((void **) &d_A, size);
cudaMalloc((void **) &d_B, size);
cudaMalloc((void **) &d_C, size);
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
// Kernel invocation code – to be shown later
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
cudaFree(d_A); cudaFree(d_B); cudaFree (d_C);
}
… Free h_A, h_B, h_C …
Slide 4-53: GPU执行方式
Single Instruction Multiple Thread单指令多线程
SIMD与多线程相结合
线程按照固定的方式运行（与传统SPMD不同）
固定个数的线程一起执行（SIMD）
使用标量指令（与传统SIMD不同）
Slide 4-54: GPU中处理单元的执行方式
Single Instruction Multiple Data (SIMD)
Slide 4-55: SPMD, SIMT, 和SIMD
SPMD
GPU的编程模型
SIMT
GPU的执行方式
SIMD
GPU计算单元的处理方式
Slide 4-56: SIMT的优点
编程灵活
任意大小的工作量vs任意硬件SIMD宽度
每个线程可以被单独对待
分支控制
Warp细节对程序员透明
Warp大小
哪些线程由同一个warp执行
Slide 4-57: NVIDIA H100 GPU
16896 FP32 CUDA cores
66.9 TFLOPS (FP64 Tensor core)
33.5 TFLOPS (FP64)
3352 GB/s memory bandwidth
700W TDP
TSMC's 4N customized for NVIDIA
Slide 4-64: GPGPU微体系结构
线程以warp为单位进行调度
同一warp中的不同线程按照SIMT方式执行
大容量寄存器组用于存放所有线程的数据
SP（Streaming Processor）进行数据的并行计算，i.e., ALU
标量指令处理
SIMD数据访问及计算
Slide 4-67: Warp
在Nvidia GPU中，warp为32个标量线程构成的SIMT执行单位
GPU SM中以warp为最小单位进行调度和执行
GPU无法控制单独执行某一个thread
Warp中不同的Thread在执行时保持同步
即，所有Thread永远执行相同PC所对应的指令
Slide 4-68: 通过细粒度多线程并行掩盖延迟
细粒度多线程执行
Slide 4-69: Warp调度
LD R1, 0x00027 PC=1
LD R2, 0x00034 PC=2
ADD R3, R1, R2 PC=3
ST R3, 0x00045 PC=4
Cycle 100: 调度 warp0 PC1
Cycle 101: 调度 warp1 PC2
Cycle 102: 调度 warp2 PC1
…
如果访存时间长，则PC3对应的指令需要等待，GPU将大量warp交替调度执行，可以一定程度上掩盖长时间访存
Slide 4-70: Warp执行
32-thread warp executing ADD A[tid],B[tid] → C[tid]
Slide 4-72: Streaming Processor
CUDA core
SIMD（Single Instruction Multiple Data）执行
大量SP提供极高并行算力
Slide 4-73: Warp指令级的并行执行
可以将不同指令交替执行
假设机器中有8个通道，1个warp中有32个线程
每周期完成24个操作，但是每周期流出1个warp
Slide 4-74: SIMT的访存模式
执行同一指令的不同线程使用线程ID来访问不同的数据
假设N=16, 每个warp中有4个线程 → 共有4个warp
Slide 4-75: Coalescing
将一个warp中不同thread的内存访问合并成更少的访问次数
如果一个warp中的32个thread访问内存中连续的4B大小的位置
合并成一个128B的访存请求(coalescing)
而不是发送32个4B的访存请求
有效降低SM和DRAM之间的访存次数
减少片上网络、片上存储划分、DRAM的工作量
Slide 4-77 ~ 4-79: GPU程序层次模型
一个程序kernel grid由多个block组成
例如，256个block
一个block由多个thread组成
例如，1024个thread
Thread为标量线程，执行完整的程序源代码
一个kernel以block为单位分配到硬件中
多个block可以分配到同一个SM中，只要有足够的硬件资源
在SM内部，来自同一个block中的不同线程（编号相邻）组成一个warp
例如，1 warp = 32 thread
Block和thread对程序员可见，warp对程序员不可见，由硬件进行控制
同一kernel中的所有线程都共享相同的代码
不同线程执行进度可能不同
但是同一warp中的线程一同执行
可以在Block级别进行同步
Slide 4-80: GPU程序层次例子
要处理所有元素的代码在一个grid中
将一个grid分割成可处理的大小（Block）
例如：一个block中有512个线程
一次以SIMD方式处理32个元素
Block与向量处理循环中的向量长度为32的分段开采类似
Slide 4-82: GPU与传统SIMD的区别
传统SIMD采用单线程
向量指令之间锁步执行：一个向量指令完成后后续指令才能开始执行
编程方式为SIMD，软件需要知道向量长度
ISA中有向量/SIMD指令
GPU以SIMD方式执行大量标量线程
可以非锁步执行
每个线程可以被独立看待，由硬件进行线程组合(Warp)
编程方式为SPMD
ISA是标量的
Slide 4-84: Nvidia ISA
"Parallel Thread Execution (PTX)"
NVIDIA编译器的指令集目标时硬件指令集的一种抽象，可实现各代GPU的兼容性
opcode.type d,a,b,c;
使用虚拟寄存器
由软件将其翻译成机器语言
Slide 4-86: 分支处理
GPU分支硬件也使用掩码，类似向量处理机中的向量屏蔽寄存器
此外，还使用硬件栈处理分支
分歧后push
汇聚后pop
Slide 4-87: 例子（分支处理）
if (X[i] != 0)
X[i] = X[i] – Y[i];
else X[i] = Z[i];
ld.global.f64 RD0, [X+R8] ; RD0 = X[i]
setp.neq.s32 P1, RD0, #0 ; P1 is predicate register 1
@!P1, bra ELSE1, *Push ; Push old mask, set new mask bits
; if P1 false, go to ELSE1
ld.global.f64 RD2, [Y+R8] ; RD2 = Y[i]
sub.f64 RD0, RD0, RD2 ; Difference in RD0
st.global.f64 [X+R8], RD0 ; X[i] = RD0
@P1, bra ENDIF1, *Comp ; complement mask bits
; if P1 true, go to ENDIF1
ELSE1: ld.global.f64 RD0, [Z+R8] ; RD0 = Z[i]
st.global.f64 [X+R8], RD0 ; X[i] = RD0
ENDIF1: <next instruction>, *Pop ; pop to restore old mask
*Push， *Comp， *Pop为由PTX汇编程序插入的分支同步标记
Slide 4-88 ~ 4-89: 分支分歧(branch divergence)
同一warp中的不同thread在遇到分支指令时可能会有分歧
例如，if-else
分支分歧：当同一warp中的不同线程在遇到分支指令时执行不同的路径
GPU使用简化的控制逻辑来减少控制部分所占的面积
thread无法单独控制和调度
并行的处理单元同一个时钟周期只能处理相同的操作
Slide 4-90 ~ 4-92: 分支分歧处理
每个warp用一个栈来处理分支分歧
分支指令中可以加入一个汇聚指令域
A: v = foo[tid.x];
B: if (v < 10)
C: v = 0;
else
D: v = 10;
E: w = bar[tid.x]+v;
foo[] = {4,8,12,16};
One stack per warp
每个warp使用一个栈来处理分支分歧的执行
Slide 4-93: 存储层次
Local memory
一个thread
Shared memory
一个block中的所有thread
Global memory
所有thread
Texture memory
只读，所有thread
适合特定地址访问模型
Constant memory
只读，所有thread
Slide 4-94: GPGPU Virtualization虚拟化
企业级做云计算时常常采用虚拟化技术：
一个GPU的SM很多，但是对于单用户只能用很少的SM,这样利用率很低；
硬件上把多个SM进行分割成多个切片；
软件上为每个用户分配一个SM切片；
使得每个用户都能满负荷使用SM;
每个用户感觉自己拥有了一个GPU。