BLOG

由于多指令多数据 （MIMD） 架构需要为每个数据操作获取一条指令，因此单指令多数据 （SIMD） 可能更节能，单条指令可以启动多个数据操作。这两个答案使 SIMD 对个人移动设备和服务器都具有吸引力。

SIMD 的三种变体：矢量体系结构（vector architecture）、多媒体 SIMD 指令集扩展（SIMD）和图形处理单元 （GPU）

• vector architectures比其他两种变体早 30 多年，扩展了许多数据操作的流水线执行。与其他 SIMD 变体相比，矢量架构更易于理解和编译，但直到最近，它们还被认为对于微处理器来说过于昂贵。该费用的一部分是晶体管，另一部分是足够的动态随机存取存储器（DRAM）带宽的成本，因为人们普遍依赖缓存来满足传统微处理器的内存性能需求。
• SIMD表示基本上同时进行并行数据操作，现在在大多数支持多媒体应用的指令集架构中都可以找到。对于 x86 架构，SIMD 指令扩展始于 1996 年的 MMX（多媒体扩展），随后在接下来的十年中出现了几个 SSE（流式 SIMD 扩展）版本，并且一直持续到今天 AVX（高级矢量扩展）。要从 x86 计算机获得最高的计算速率，通常需要使用这些 SIMD 指令，尤其是对于浮点程序。
• 第三种变体来自图形加速器社区，它提供了比当今传统多核计算机更高的潜在性能。尽管 GPU 与矢量架构共享功能，但它们也有自己独特的特征，部分原因是它们发展的生态系统。除了 GPU 及其图形内存外，此环境还具有系统处理器和系统内存。事实上，为了认识到这些区别，GPU 社区将这种类型的架构称为异构架构。

矢量架构抓取散布在内存中的数据元素集，将它们放入大型顺序寄存器文件中，对这些寄存器文件中的数据进行操作，然后将结果分散回内存。单个指令对数据矢量进行操作，这会导致对独立数据元素执行数十个寄存器-寄存器操作。

这些大型寄存器文件充当编译器控制的缓冲区，既可以隐藏内存延迟，又可以利用内存带宽。由于矢量加载和存储是深度流水线化的，因此程序仅在每次矢量加载或存储时支付一次较长的内存延迟，而不是每次元素支付一次，从而将延迟摊销到（例如）32 个元素上。事实上，矢量程序努力使内存保持繁忙。

功耗墙导致架构师重视那些能够在不增加高度乱序超标量处理器的能耗和设计复杂性成本的情况下提供良好性能的架构。矢量指令是这种趋势的自然匹配，因为架构师可以使用它们来提高简单有序标量处理器的性能，而不会大幅增加能耗需求和设计复杂性。实际上，开发人员可以更有效地将许多在复杂乱序设计上运行良好的程序表示为矢量指令形式的数据级并行性。

RV64V 指令集架构的主要组件如下：

• 向量寄存器——每个向量寄存器存储一个向量，RV64V 有 32 个这样的寄存器，每个宽度为 64 位。向量寄存器文件需要提供足够的端口来为所有向量功能单元供电。这些端口将允许向量操作在多个向量寄存器之间实现高度重叠。读端口和写端口总数至少为 16 个读端口和 8 个写端口，它们通过一对交叉开关连接到功能单元的输入或输出。
• 向量功能单元——在我们的实现中，每个单元都是完全流水线的，并且它可以在每个时钟周期开始新的操作。需要一个控制单元来检测风险，包括功能单元的结构风险和寄存器访问的数据风险。
• 向量加载/存储单元—向量内存单元用于将向量加载或存储到内存中。在我们的假设 RV64V 实现中，向量加载和存储是完全流水线的，因此可以在初始延迟后，以每个时钟周期一个字的带宽在向量寄存器和内存之间移动字。这个单元通常也会处理标量加载和存储。
• 一组标量寄存器—标量寄存器同样可以向向量功能单元提供数据输入，以及计算地址以传递给向量加载/存储单元。这些是 RV64G 的正常 31 个通用寄存器和 32 个浮点寄存器。向量功能单元的一个输入会在标量值从标量寄存器文件中读出时锁存这些值。

假设输入操作数都是向量寄存器，但这些指令还有一些版本，其中一个操作数可以是标量寄存器 (xi 或 fi)。当两者都是向量时，RV64V 使用后缀 .vv；当第二个操作数是标量时，使用 .vs；当第一个是标量寄存器时，使用 .sv。因此，这三个都是有效的 RV64V 指令：vsub.vv、vsub.vs 和 vsub.sv。（加法和其他交换运算只有前两个版本，因为 vadd.sv 和 vadd.sv 是冗余的。）因为操作数决定指令的版本，所以我们通常让汇编器提供适当的后缀。向量功能单元在指令发出时获取标量值的副本。

RV64V 的创新之处在于将数据类型和数据大小与每个矢量寄存器相关联，而不是采用指令提供该信息的常规方法。因此，在执行矢量指令之前，程序会配置正在使用的矢量寄存器以指定其数据类型和宽度。

动态寄存器类型化的一个原因是，对于支持这种多样性的传统向量体系结构，需要许多指令。

• 动态寄存器类型还允许程序禁用未使用的向量寄存器。因此，已启用的向量寄存器被分配所有向量内存作为长向量。例如，假设我们有 1024 字节的向量内存，如果启用了 4 个向量寄存器并且它们是 64 位浮点数类型，则处理器会给每个向量寄存器 256 字节或 256/8¼32 个元素。此值称为最大向量长度 (mvl)，由处理器设置，软件无法更改。
• 动态寄存器类型的另一个好处是，当不使用向量寄存器时，程序可以将其配置为禁用，因此无需在上下文切换时保存和恢复它们。
• 动态寄存器类型的第三个好处是，根据寄存器的配置，不同大小的操作数之间的转换可以是隐式的，而不是作为其他显式转换指令。

example

Y = a * X + Y

X 和 Y 是向量，最初驻留在内存中，a 是标量。

RV64G code:
fld f0,a           # Load scalar a
addi x28,x5,#256   # Last address to load
Loop: fld f1,0(x5) # Load X[i]
fmul.d f1,f1,f0    # a ! X[i]
fld f2,0(x6)       # Load Y[i]
fadd.d f2,f2,f1    # a ! X[i] + Y[i]
fsd f2,0(x6)       # Store into Y[i]
addi x5,x5,#8      # Increment index to X
addi x6,x6,#8      # Increment index to Y
bne x28,x5,Loop    # Check if done

RV64V code:
vsetdcfg 4*FP64    # Enable 4 DP FP vregs
fld f0,a           # Load scalar a
vld v0,x5          # Load vector X
vmul v1,v0,f0      # Vector-scalar mult
vld v2,x6          # Load vector Y
vadd v3,v1,v2      # Vector-vector add
vst v3,x6          # Store the sum
vdisable           # Disable vector regs

在直接的RV64G代码中，每个fadd.d都必须等待一个fmul.d，每个fsd都必须等待fadd.d。在矢量处理器上，每个矢量指令只会为每个矢量的第一个元素停顿，然后后续元素将顺利流过流水线。因此，流水线停顿仅在每个矢量指令中需要一次，而不是每个矢量元素一次。矢量架构师将元素相关操作的转发称为"chaining"，因为相关操作被"链接"在一起。在这个例子中，在RV64G上的流水线停顿频率将约高32倍于RV64V。软件流水线，循环展开（附录H）或乱序执行可以减少RV64G上的流水线停顿；然而，指令带宽的大差异不能显著减少。

向量运算的执行时间主要取决于三个因素：（1）操作数向量的长度，（2）操作之间的结构性风险，以及（3）数据依赖性。给定向量长度和启动率，即向量单元消耗新操作数并产生新结果的速率，我们可以计算单个向量指令的执行时间。

所有现代向量计算机都配备了具有多个并行流水线（或通道）的向量功能单元，每个时钟周期可以产生两个或更多结果，但它们也可能有一些未完全流水线化的功能单元。

为了简化向量执行和向量性能的讨论，我们使用"纵队"（convoy）这一概念，它指的是可能一起执行的向量指令的集合。纵队中的指令不能包含任何结构性的数据冒险；如果存在这样的冒险，这些指令就需要在不同的纵队中顺序执行。

为了简化对向量执行和向量性能的讨论，我们使用了护航队（convoy）的概念，护航队是一组可能一起执行的向量指令。护航队中的指令不能包含任何结构性危害；如果存在此类危害，则需要对指令进行序列化并在不同的护航队中启动。因此，前面示例中的 vld 和以下 vmul 可以位于同一护航队中。正如我们很快就会看到的，您可以通过计算护航数量来估计一段代码的性能。为了保持这种分析的简单性，我们假设一组指令必须先完成执行，然后才能开始执行任何其他指令（标量或向量）。

由于链式操作允许向量操作在其向量源操作数的各个元素可用后立即开始，因此链式操作允许它们位于同一个队列中：链中第一个功能单元的结果被“转发”到第二个功能单元。链式的早期实现就像标量流水线中的转发一样，但这限制了链中源指令和目标指令的时序。最近的实现使用灵活的链式，它允许一个矢量指令与任何其他活动矢量指令链式连接，假设我们不会产生结构性危害。所有现代矢量体系结构都支持灵活的链式。

为了将队列转换为执行时间，我们需要一个度量来估计队列的长度。它被称为时钟（chime），它只是执行一个队列所花费的时间单位。因此，由 m 个队列组成的向量序列在 m 个时钟中执行；对于长度为 n 的向量，对于我们简单的 RV64V 实现，这大约是 m*n 个时钟周期。

时钟近似忽略了一些特定于处理器的开销，其中许多开销取决于向量长度。因此，以时钟测量时间对于长向量来说比短向量来说是一个更好的近似。我们将使用时钟测量，而不是每结果的时钟周期，以明确表示我们正在忽略某些开销。

如果我们知道向量序列中的队列数，我们就知道以时钟为单位的执行时间。在测量时钟时忽略的一个开销来源是对在单个时钟周期内启动多个向量指令的任何限制。如果只能在一个时钟周期内启动一个向量指令（大多数向量处理器中的现实情况），则时钟计数将低估队列的实际执行时间。因为向量的长度通常远大于队列中的指令数，所以我们将简单地假设队列在一个时钟中执行。

矢量指令集的一个关键优势在于，它允许软件仅使用一条简短的指令将大量并行工作传递给硬件。一条矢量指令可以包含数十个独立操作，但编码的位数与常规标量指令相同。矢量指令的并行语义允许实现使用深度流水线功能单元执行这些基本操作，就像我们迄今为止研究的 RV64V 实现一样；一系列并行功能单元；或并行和流水线功能单元的组合。

RV64V 指令集具有以下特性：所有向量算术指令只允许一个向量寄存器的元素 N 与其他向量寄存器的元素 N 参与运算。这极大地简化了高度并行向量单元的设计，该单元可以构建为多个并行通道。与交通公路一样，我们可以通过增加通道数来提高向量单元的峰值吞吐量。图 4.5 显示了四通道向量单元的结构。因此，从一个通道增加到四个通道，将钟声的时钟数从 32 个减少到 8 个。为了使多通道具有优势，应用程序和体系结构都必须支持长向量；否则，它们将执行得如此之快，以至于您会用尽指令带宽，需要 ILP 技术（参见第 3 章）来提供足够的向量指令。

每个通道包含一个矢量寄存器文件的一部分和来自每个矢量功能单元的一个执行流水线。每个矢量功能单元以每周期一个元素组的速度执行矢量指令，使用多个流水线，每个通道一个。第一个通道为所有矢量寄存器保存第一个元素（元素 0），因此任何矢量指令中的第一个元素的源操作数和目标操作数都位于第一个通道中。此分配允许通道本地的算术流水线在不与其他通道通信的情况下完成操作。避免通道间通信降低了构建高度并行执行单元所需的布线成本和寄存器文件端口，并有助于解释为什么矢量计算机每时钟周期最多可以完成 64 次操作（16 个通道上的 2 个算术单元和 2 个加载/存储单元）。

添加多个通道是一种流行的技术，可提高矢量性能，因为它只需要少量增加控制复杂性，并且不需要更改现有机器代码。它还允许设计人员在不牺牲峰值性能的情况下权衡裸片面积、时钟速率、电压和能量。如果矢量处理器时钟速率减半，通道数增加一倍将保持相同的峰值性能。

矢量寄存器处理器的自然矢量长度由最大矢量长度 （mvl） 决定。这个长度（在上面的示例中为 32 个）不太可能以匹配程序中的实际向量长度。此外，在实际程序中，特定向量运算的长度在编译时通常是未知的。事实上，一段代码可能需要不同的向量长度。例如，请考虑以下代码：

for (i=0; i<n; i=i+1)
	Y[i] = a*X[i] + Y[i];

所有向量运算的大小都取决于 n，它甚至可能要到运行时才能知道。n 的值也可能是包含前面循环的过程的参数，因此在执行过程中可能会更改。

解决这些问题的办法是添加一个向量长度寄存器 (vl)。vl 控制任何向量操作的长度，包括向量加载或存储。但是，vl 中的值不能大于最大向量长度 (mvl)。只要实际长度小于或等于最大向量长度 (mvl)，这就能解决我们的问题。此参数意味着向量寄存器的长度可以在以后的计算机世代中增长，而无需更改指令集。正如我们将在下一节中看到的那样，多媒体 SIMD 扩展没有 mvl 的等效项，因此它们每次增加向量长度时都会扩展指令集。

如果在编译时不知道 n 的值，因此可能大于 mvl，该怎么办？为了解决向量长度大于最大长度的第二个问题，传统上使用一种称为带状挖掘的技术。带状挖掘是生成代码，以便每个向量操作以小于或等于 mvl 的大小完成。一个循环处理任何数量的迭代，这些迭代是 mvl 的倍数，另一个循环处理任何剩余的迭代，并且必须小于 mvl。RISC-V 针对带状挖掘有一个比单独循环更好的解决方案。指令 setvl 将 mvl 和循环变量 n 中较小的一个写入 vl（以及另一个寄存器）。如果循环的迭代次数大于 n，那么循环计算最快的速度是每次 mvl 个值，因此 setvl 将 vl 设置为 mvl。如果 n 小于 mvl，则它应该仅计算此循环的最后 n 个元素，因此 setvl 将 vl 设置为 n。setvl 还写入另一个标量寄存器，以帮助进行后续的循环记账。

以下是适用于任何 n 值的向量 DAXPY 的 RV64V 代码。

vsetdcfg 2 DP FP    # Enable 2 64b Fl.Pt. registers
fld f0,a            # Load scalar a
loop: setvl t0,a0   # vl = t0 = min(mvl,n)
vld v0,x5           # Load vector X
slli t1,t0,3        # t1 = vl * 8 (in bytes)
add x5,x5,t1        # Increment pointer to X by vl*8
vmul v0,v0,f0       # Vector-scalar mult
vld v1,x6           # Load vector Y
vadd v1,v0,v1       # Vector-vector add
sub a0,a0,t0        # n -= vl (t0)
vst v1,x6           # Store the sum into Y
add x6,x6,t1        # Increment pointer to Y by vl*8
bnez a0,loop        # Repeat if n != 0
vdisable            # Disable vector regs

根据阿姆达尔定律，我们知道具有低到中等矢量化水平的程序的加速将非常有限。循环中条件语句（IF 语句）的存在和稀疏矩阵的使用是矢量化级别较低两个主要原因。

由于 IF 语句在循环中引入了控制依赖关系，因此无法使用我们迄今为止讨论过的技术以矢量模式实现。同样，我们无法使用我们迄今为止见过的任何功能有效地实现稀疏矩阵。我们在此研究处理条件执行的策略，稍后讨论稀疏矩阵。

for (i=0; i<64; i=i+1)
	if (X[i] != 0)
		X[i] = X[i] - Y[i];

由于主体条件执行，此循环通常无法矢量化；但是，如果可以针对 X[i]≠0的迭代运行内部循环，则可以矢量化减法。

此功能的通用扩展是矢量掩码控制。在 RV64V 中，谓词寄存器保存掩码，并实质上提供矢量指令中每个元素操作的条件执行。这些寄存器使用布尔矢量来控制矢量指令的执行，就像条件执行指令使用布尔条件来确定是否执行标量指令一样。当谓词寄存器 p0 设置时，所有后续矢量指令仅对谓词寄存器中相应条目为 1 的矢量元素进行操作。目标矢量寄存器中对应于掩码寄存器中 0 的条目不受矢量操作的影响。与矢量寄存器一样，谓词寄存器已配置并且可以禁用。启用谓词寄存器会将其初始化为全 1，这意味着后续矢量指令对所有矢量元素进行操作。现在，我们可以对前一个循环使用以下代码，假设 X 和 Y 的起始地址分别在 x5 和 x6 中：

vsetdcfg 2*FP64     # Enable 2 64b FP vector regs
vsetpcfgi 1         # Enable 1 predicate register
vld v0,x5           # Load vector X into v0
vld v1,x6           # Load vector Y into v1
fmv.d.x f0,x0       # Put (FP) zero into f0
vpne p0,v0,f0       # Set p0(i) to 1 if v0(i)!=f0
vsub v0,v0,v1       # Subtract under vector mask
vst v0,x5           # Store the result in X
vdisable            # Disable vector registers
vpdisable           # Disable predicate registers

编译器编写人员使用 IF 转换术语将 IF 语句转换为使用条件执行的直线代码序列。

但是，使用矢量掩码寄存器确实有开销。对于标量体系结构，当条件不满足时，条件执行的指令仍需要执行时间。尽管如此，消除分支和相关的控制依赖关系可以使条件指令更快，即使它有时会做无用功。同样，即使对于掩码为零的元素，使用矢量掩码执行的矢量指令仍需要相同的时间。同样，尽管掩码中有很多零，但使用矢量掩码控制可能仍然比使用标量模式快得多。

正如我们在第 4.4 节中所见，矢量处理器和 GPU 之间的一个区别在于它们处理条件语句的方式。矢量处理器使谓词寄存器成为体系结构状态的一部分，并依靠编译器显式地操作掩码寄存器。相比之下，GPU 使用硬件来操作对 GPU 软件不可见的内部掩码寄存器，从而获得相同的效果。在这两种情况下，硬件都会花费时间来执行矢量元素，无论相应的掩码位是 0 还是 1，因此在使用掩码时，GFLOPS 速率会下降。

加载/存储矢量单元的行为比算术功能单元的行为复杂得多。加载的启动时间是从内存中获取第一个字到寄存器中的时间。如果可以不暂停地提供矢量的其余部分，则矢量启动速率等于获取或存储新字的速率。与更简单的功能单元不同，启动速率不一定是一个时钟周期，因为内存库暂停可能会降低有效吞吐量。

通常，加载/存储单元的启动惩罚高于算术单元的启动惩罚——在许多处理器上超过 100 个时钟周期。对于 RV64V，我们假设启动时间为 12 个时钟周期，与 Cray1 相同。（最近的矢量计算机使用高速缓存来降低矢量加载和存储的延迟。）

为了保持每个时钟周期获取或存储一个字的启动速率，内存系统必须能够产生或接受这么多数据。跨多个独立内存库扩展访问通常会提供所需速率。正如我们很快会看到的，拥有大量库对于处理访问数据行或列的矢量加载或存储很有用。

大多数矢量处理器使用内存库，它允许几个独立访问，而不是出于三个原因进行简单的内存交错：

1. 许多矢量计算机每时钟周期支持许多加载或存储，并且内存库周期时间通常比处理器周期时间大几倍。为了支持来自多个加载或存储的并发访问，内存系统需要多个库，并且需要能够独立地控制对库的访问。
2. 大多数矢量处理器支持加载或存储非顺序数据字的能力。在这种情况下，需要独立的库寻址，而不是交错。
3. 大多数矢量计算机支持多个处理器共享相同的内存系统，因此每个处理器将生成其自己的独立地址流。

综合起来，这些特性导致了对大量独立内存库的需求，如下例所示

从更高层次的角度来看，矢量加载/存储单元在标量处理器中扮演着与预取单元类似的角色，因为两者都试图通过为处理器提供数据流来提供数据带宽。

向量中相邻元素在内存中的位置可能不是连续的。考虑一下 C 语言中矩阵乘法的简单代码：

for (i = 0; i < 100; i=i+1)
	for (j = 0; j < 100; j=j+1) {
		A[i][j] = 0.0;
		for (k = 0; k < 100; k=k+1)
			A[i][j] = A[i][j] + B[i][k] * D[k][j];
	}

我们可以将 B 的每一行与 D 的每一列的乘法矢量化，并使用 k 作为索引变量对内部循环进行带状挖掘。

为此，我们必须考虑如何寻址 B 中的相邻元素和 D 中的相邻元素。当为数组分配内存时，它会被线性化，并且必须按行优先顺序（如在 C 中）或列优先顺序（如在 Fortran 中）进行布局。这种线性化意味着行中的元素或列中的元素在内存中不是相邻的。例如，前面的 C 代码按行优先顺序分配，因此内部循环中的迭代访问的 D 的元素由行大小乘以 8（每个条目的字节数）隔开，总共为 800 个字节。在第 2 章中，我们看到阻塞可以提高基于缓存的系统中的局部性。对于没有缓存的矢量处理器，我们需要另一种技术来获取内存中不相邻的矢量元素。

将元素聚集到单个向量寄存器中所间隔的距离称为跨距。在此示例中，矩阵 D 的跨距为 100 个双字（800 字节），矩阵 B 的跨距为 1 个双字（8 字节）。对于 Fortran 使用的列优先顺序，跨距将相反。矩阵 D 的跨距为 1，或 1 个双字（8 字节），间隔连续的元素，而矩阵 B 的跨距为 100，或 100 个双字（800 字节）。因此，如果不重新排序循环，编译器无法隐藏 B 和 D 的连续元素之间的长距离。

一旦将向量加载到向量寄存器中，它将表现得好像具有逻辑相邻的元素。因此，向量处理器可以使用仅具有跨距功能的向量加载和向量存储操作来处理大于 1 的跨距，称为非单位跨距。访问非顺序内存位置并将它们重新整形为密集结构的能力是向量体系结构的主要优势之一。

高速缓存本质上处理单位步幅数据；增加块大小有助于降低具有单位步幅的大型科学数据集的未命中率，但对于以非单位步幅访问的数据，增加块大小甚至会产生负面影响。虽然阻塞技术可以解决其中一些问题（参见第 2 章），但有效访问非连续数据的能力对于某些问题上的矢量处理器而言仍然是一个优势，正如我们在第 4.7 节中所看到的。

在 RV64V 上，其中可寻址单位是字节，我们示例的步幅将是 800。必须动态计算该值，因为矩阵的大小在编译时可能未知，或者（就像矢量长度一样）可能因同一语句的不同执行而改变。矢量步幅（如矢量起始地址）可以放入通用寄存器中。然后，RV64V 指令 VLDS（带步幅加载矢量）将矢量提取到矢量，然后 RV64V 指令 VLDS（带步长的加载向量）将向量提取到向量寄存器中。同样，在存储非单位步长向量时，使用指令 VSTS（带步长的存储向量）。

支持大于 1 的步幅会使内存系统复杂化。一旦我们引入非单位步幅，就有可能频繁地请求来自同一存储库的访问。当多个访问争用一个存储库时，就会发生存储库冲突，从而导致一个访问停顿。如果出现以下情况，将发生存储库冲突并因此发生停滞

如前所述，稀疏矩阵很常见，因此有必要有允许具有稀疏矩阵的程序以矢量模式执行的技术。在稀疏矩阵中，矢量的元素通常以某种压缩形式存储，然后间接访问。假设一个简化的稀疏结构，我们可能会看到类似这样的代码：

for (i = 0; i < n; i=i+1)
	A[K[i]] = A[K[i]] + C[M[i]];

此代码使用索引向量 K 和 M 来指定 A 和 C 的非零元素，对数组 A 和 C 执行稀疏向量求和。（A 和 C 必须具有相同数量的非零元素——n 个——因此 K 和 M 的大小相同。）支持稀疏矩阵的主要机制是使用索引向量进行收集-分散操作。此类操作的目标是支持在稀疏矩阵的压缩表示（即不包括零）和常规表示（即包括零）之间移动。收集操作采用索引向量并获取向量的元素，这些元素位于通过将基本地址添加到索引向量中给出的偏移量处。结果是向量寄存器中的密集向量。在密集形式，稀疏向量可以通过散射存储以展开形式存储，使用相同的索引向量。此类操作的硬件支持称为收集-散射，它出现在几乎所有现代矢量处理器上。RV64V 指令是 vldi（加载矢量索引或收集）和 vsti（存储矢量索引或散射）。例如，如果 x5、x6、x7 和 x28 包含前一序列中向量的起始地址，我们可以使用矢量指令对内部循环进行编码，例如：

vsetdcfg 4*FP64     # 4 64b FP vector registers
vld v0, x7          # Load K[]
vldx v1, x5, v0)    # Load A[K[]]
vld v2, x28         # Load M[]
vldi v3, x6, v2)    # Load C[M[]]
vadd v1, v1, v3     # Add them
vstx v1, x5, v0)    # Store A[K[]]
vdisable            # Disable vector registers

这种技术允许具有稀疏矩阵的代码以向量模式运行。简单的矢量化编译器无法自动矢量化前面的源代码，因为编译器不知道 K 的元素是不同的值，因此不存在依赖关系。相反，程序员指令会告诉编译器以向量模式运行循环是安全的。

尽管索引加载和存储（收集和分散）可以流水线化，但它们通常比非索引加载或存储运行得慢得多，因为从指令开始时不知道内存库。寄存器文件还必须提供向量单元的通道之间的通信，以支持收集和分散。

收集或分散的每个元素都有一个单独的地址，因此无法以组为单位处理，并且在整个内存系统中可能存在许多冲突。因此，即使在基于缓存的系统上，每个单独的访问都会产生明显的延迟。然而，正如第 4.7 节所示，内存系统可以通过针对这种情况进行设计并使用更多硬件资源来提供更好的性能，而不是当架构师对这种不可预测的访问采取放任自流的态度时。

正如我们在第 4.4 节中所见，所有加载都是收集，所有存储都是 GPU 中的散射，因为没有单独的指令将地址限制为顺序。为了将潜在的慢速收集和散射转变为更有效的内存单元步幅访问，GPU 硬件必须在执行期间识别顺序地址，而 GPU 程序员必须确保收集或散射中的所有地址都位于相邻位置。

矢量体系结构的一个优点是，编译器可以在编译时告诉程序员一段代码是否会矢量化，通常会给出提示

通常会给出有关为何不矢量化代码的提示。这种简单的执行模型允许其他领域的专家通过修改其代码或在可以假设操作之间独立时向编译器提供提示来学习如何提高性能，例如用于收集-分散数据传输。正是编译器和程序员之间的这种对话，双方互相提示如何提高性能，简化了矢量计算机的编程。

如今，影响程序在矢量模式下运行成功与否的主要因素是程序本身的结构：循环是否具有真正的依赖关系（参见第 4.5 节），或者是否可以对其进行重构以消除此类依赖关系？此因素受所选算法以及在某种程度上受算法编码方式的影响。

为了了解科学程序中可实现的矢量化水平，我们来看看针对 Perfect Club 基准观察到的矢量化水平。图 4.7 显示了在 Cray Y-MP 上运行的两个代码版本中以矢量模式执行的操作的百分比。第一个版本是通过对原始代码进行编译器优化而获得的，而第二个版本则使用了 Cray Research 程序员团队提供的广泛提示。对应用程序在矢量处理器上的性能进行的多项研究表明，编译器矢量化的水平差异很大。

对于编译器本身无法很好地矢量化的代码，提示丰富的版本在矢量化级别方面显示出显著的提升，现在所有代码的矢量化率都高于 50%。矢量化率中位数从大约 70% 提高到大约 90%。

与向量指令一样，SIMD 指令对数据向量指定相同操作。与具有大型寄存器文件（例如 RISC-V RV64V 向量寄存器）的向量机不同，后者可以在 32 个向量寄存器中的每一个中容纳例如三十二个 64 位元素，SIMD 指令倾向于指定更少的操作数，因此使用更小的寄存器文件。

与旨在成为矢量化编译器目标的优雅指令集的矢量体系结构相比，SIMD 扩展具有三个主要遗漏：没有矢量长度寄存器、没有跨步或收集/分散数据传输指令，也没有掩码寄存器。

1. 多媒体 SIMD 扩展修复了操作码中的数据操作数数量，这导致在 x86 体系结构的 MMX、SSE 和 AVX 扩展中添加了数百条指令。矢量体系结构具有矢量长度寄存器，该寄存器指定当前操作的操作数数量。这些可变长度矢量寄存器可以轻松容纳比体系结构支持的最大大小更短的矢量的程序。此外，矢量体系结构与向量长度寄存器结合使用，避免使用许多操作码。
2. 直到最近，多媒体 SIMD 还没有提供矢量体系结构的更复杂的寻址模式，即跨步访问和收集-分散访问。这些特性增加了矢量编译器可以成功矢量化的程序数量。
3. 尽管这种情况正在发生变化，但多媒体 SIMD 通常没有提供掩码寄存器来支持像矢量体系结构中那样的元素条件执行。

此类遗漏使编译器更难生成 SIMD 代码，并增加了用 SIMD 汇编语言编程的难度。

鉴于这些弱点，为什么多媒体 SIMD 扩展如此受欢迎？

首先，它们最初添加到标准算术单元的成本很低，并且易于实现。其次，与矢量架构相比，它们需要的额外处理器状态很少，而这始终是上下文切换时间所关注的问题。第三，您需要大量内存带宽来支持许多计算机不具备的矢量架构。第四，当一条指令可以产生 32 次内存访问并且任何一次访问都可能导致页面错误时，SIMD 不必处理虚拟内存中的问题。最初的 SIMD 扩展使用单独的数据传输，针对内存中对齐的 SIMD 操作数组，因此它们无法跨越页面边界。SIMD 的短固定长度“矢量”的另一个优点是，很容易引入可以帮助处理新媒体标准的指令，例如执行排列的指令或消耗比矢量产生的操作数更少或更多的指令。最后，人们担心矢量架构与缓存配合得如何。最近的矢量架构已经解决了所有这些问题。 然而，最主要的问题是，由于向后二进制兼容性的重要性，一旦一种架构开始走上 SIMD 路径，就很难再走回头路。

鉴于 SIMD 多媒体扩展的临时性质，使用这些指令的最简单方法是通过库或通过汇编语言编写。

最近的扩展变得更加常规，为编译器提供了更合理的目标。通过借用矢量化编译器中的技术，编译器开始自动生成 SIMD 指令。例如，当今的高级编译器可以生成 SIMD 浮点指令，为科学代码提供更高的性能。但是，程序员必须确保将内存中的所有数据对齐到代码运行所在的 SIMD 单元的宽度，以防止编译器为原本可矢量化的代码生成标量指令。

一种直观、可视的方式来比较 SIMD 架构变化的潜在浮点性能是 Roofline 模型（Williams 等人，2009 年）。它生成的图形的水平线和对角线赋予了此简单模型名称，并指明了它的值（参见图 4.11）。它将浮点性能、内存性能和算术强度在一个二维图形中联系在一起。

算术强度是每访问一个字节内存所进行的浮点运算的比率。它可以通过将程序的浮点运算总数除以程序执行期间传输到主内存的数据字节总数来计算。图 4.10 显示了几个示例内核的相对算术强度。

可以通过硬件规格找到峰值浮点性能。本案例研究中的许多内核不适合片上高速缓存，因此峰值内存性能由高速缓存后面的内存系统定义。请注意，我们需要处理器可用的峰值内存带宽，而不仅仅是图 4.27 第 328 页中 DRAM 引脚上的带宽。查找（交付的）峰值内存性能的一种方法是运行 Stream 基准测试。

图 4.11 显示了左侧的 NEC SX-9 向量处理器和右侧的英特尔酷睿 i7 920 多核计算机的屋顶线模型。纵向 Y 轴是可实现的浮点性能，从 2 到 256 GFLOPS/s。横向 X 轴是算术强度，在两个图表中从 1/8 FLOP/DRAM 字节访问变化到 16 FLOP/DRAM 字节访问。请注意，该图表是双对数刻度，并且屋顶线仅对计算机执行一次。

对于给定的内核，我们可以根据其算术强度找到 X 轴上的一个点。如果我们通过该点绘制一条垂直线，则该计算机上内核的性能必须位于该线上的某个位置。我们可以绘制一条水平线，显示计算机的峰值浮点性能。显然，实际浮点性能不能高于水平线，因为那是硬件限制。

我们如何绘制峰值内存性能？因为 X 轴是 FLOP/字节，Y 轴是 FLOP/秒，所以字节/秒只是此图中 45 度角的对角线。因此，我们可以绘制第三条线，给出该计算机的内存系统可以支持的给定算术强度。我们可以用一个公式来表示这些限制，以便在图 4.11 中的图形中绘制这些线：

为内核性能设置了一个上限，具体取决于其算术强度。如果我们将算术强度视为击中屋顶的杆子，那么它要么击中屋顶的平坦部分，这意味着性能受到计算限制，要么击中屋顶的倾斜部分，这意味着性能最终受到内存带宽的限制。在图 4.11 中，右侧的垂直虚线（算术强度为 4）是前者的示例，左侧的垂直虚线（算术强度为 1/4）是后者的示例。给定计算机的 Roofline 模型，您可以重复应用它，因为它不会因内核而异。

我们如何绘制峰值内存性能？由于 X 轴是 FLOP/字节，Y 轴是 FLOP/s，因此字节/秒在此图中只是一条 45 度角的对角线。因此，我们可以绘制第三条线，该线给出该计算机的内存系统在给定条件下可以支持的最大浮点性能

• 为了区分 GPU（设备）函数和系统处理器（主机）函数，CUDA 对前者使用 __device__ 或 __global__，对后者使用 __host__。
• 使用 device 声明的 CUDA 变量分配给 GPU 内存（见下文），所有多线程 SIMD 处理器都可以访问该内存。
• 在 GPU 上运行的函数名称的扩展函数调用语法为name < <<dimGrid, dimBlock>> > (… parameter list…) ,其中 dimGrid 和 dimBlock 分别指定代码的维度（以线程块为单位）和块的维度（以线程为单位）。
• 除了块标识符 (blockIdx) 和块中每个线程的标识符 (threadIdx) 之外，CUDA 还提供了一个用于指定每个块中线程数的关键字 (blockDim)，该关键字来自前面项目符号中的 dimBlock 参数。

GPU 也依赖于单个多线程 SIMD 处理器内的多线程来隐藏内存延迟。

网格是在 GPU 上运行的代码，由一组线程块组成。

举个具体示例，假设我们要将两个向量相乘，每个向量有 8192 个元素：A = B * C。处理整个 8192 个元素乘法的 GPU 代码称为网格（或矢量化循环）。为了将其分解为更易于管理的大小，网格由线程块（或矢量化循环体）组成，每个线程块最多有 512 个元素。请注意，SIMD 指令一次执行 32 个元素。由于向量中有 8192 个元素，因此此示例有 16 个线程块，网格和线程块是 GPU 硬件中实现的编程抽象，可帮助程序员组织其 CUDA 代码。（线程块类似于剥离挖掘的矢量循环，矢量长度为 32。）

线程块被分配给执行该代码的处理器，称之为多线程 SIMD 处理器，由线程块调度程序分配。程序员告诉硬件中实现的线程块调度程序运行多少个线程块。在此示例中，它会向多线程 SIMD 处理器发送 16 个线程块，以计算此循环的所有 8192 个元素。图 4.14 显示了多线程 SIMD 处理器的简化框图。它类似于矢量处理器，但它具有许多并行功能单元，而不是像矢量处理器那样只有几个深度流水线。在图 4.13 中的编程示例中，每个多线程 SIMD 处理器被分配 512 个向量元素来处理。SIMD 处理器是具有独立 PC 的完整处理器，并使用线程进行编程。然后，GPU 硬件包含一组多线程 SIMD 处理器，这些处理器执行线程块网格（矢量化循环体）；也就是说，GPU 是由多线程 SIMD 处理器组成的多处理器。

GPU 可以拥有一个到几十个多线程 SIMD 处理器。

硬件创建、管理、调度和执行的机器对象是 SIMD 指令线程。它是一个传统线程，专门包含 SIMD 指令。这些 SIMD 指令的线程有自己的 PC，并在多线程 SIMD 处理器上运行。SIMD 线程调度程序知道哪些 SIMD 指令的线程已准备好运行，然后将它们发送到调度单元以在多线程 SIMD 处理器上运行。

因此， GPU 具有两个级别的调度器： (1) 分配线程块（矢量化循环的主体）给多线程 SIMD 处理器的线程块调度器和 (2) SIMD 处理器中的 SIMD Thread调度器，它调度 SIMD 命令的线程应该在什么时候运行。

这些线程的 SIMD 指令为 32 位宽，因此此示例中的每个 SIMD 指令线程将计算 32 个计算元素。在此示例中，线程块将包含 16=512/32 个 SIMD 线程。由于线程由 SIMD 指令组成，因此 SIMD 处理器必须具有并行功能单元来执行操作。我们称它们为 SIMD 通道，它们与第 4.2 节中的向量通道非常相似。

对于 Pascal GPU，每个 32 位宽的 SIMD 指令线程都映射到 16 个物理 SIMD 通道，因此 SIMD 指令线程中的每个 SIMD 指令需要 2 个时钟周期才能完成。每个 SIMD 指令线程都以锁定步长执行，并且仅在开始时调度。继续使用 SIMD 处理器作为矢量处理器的类比，可以说它有 16 个通道，矢量长度为 32，时钟周期为 2 个时钟周期。（这种宽而浅的特性是我们使用更准确的术语 SIMD 处理器而不是矢量的原因。）请注意，GPU SIMD 处理器中的通道数可以是线程块中的线程数的任何值，就像矢量处理器中的通道数可以在 1 和最大矢量长度之间变化一样。例如，在 GPU 世代中，每个 SIMD 处理器的通道数在 8 到 32 之间波动。

由于 SIMD 指令的线程定义上是独立的，因此 SIMD 线程调度程序可以选择任何已准备好的 SIMD 指令线程，而无需坚持线程内序列中的下一个 SIMD 指令。SIMD 线程调度程序包括一个记分板（参见第 3 章）来跟踪多达 64 个 SIMD 指令线程，以查看哪个 SIMD 指令已准备就绪。内存指令的延迟因高速缓存和 TLB 中的命中和未命中而变化，因此需要记分板来确定这些指令何时完成。图 4.16 显示了 SIMD 线程调度程序随着时间的推移以不同的顺序选择 SIMD 指令线程。GPU 架构师的假设是，GPU 应用程序具有如此多的 SIMD 指令线程，以至于多线程既可以隐藏对 DRAM 的延迟，又可以提高多线程 SIMD 处理器的利用率。

继续我们的向量乘法示例，每个多线程 SIMD 处理器必须从内存中加载两个向量的 32 个元素到寄存器中，通过读写寄存器执行乘法，并将乘积从寄存器存储回内存。为了容纳这些内存元素，SIMD 处理器拥有令人印象深刻的 32,768-65,536 个 32 位寄存器（图 4.14 中每条通道 1024 个），具体取决于 Pascal GPU 的型号。就像向量处理器一样，这些寄存器在矢量通道或在这种情况下，SIMD 通道上按逻辑划分。

每个 SIMD 线程最多只能使用 256 个寄存器，因此您可以将 SIMD 线程视为最多具有 256 个向量寄存器，每个向量寄存器具有 32 个元素，每个元素的宽度为 32 位。（由于双精度浮点操作数使用两个相邻的 32 位寄存器，另一种观点是每个 SIMD 线程具有 128 个 32 个元素的向量寄存器，每个寄存器的宽度为 64 位。）

寄存器使用和最大线程数之间存在权衡；每个线程的寄存器越少，可能的线程就越多，寄存器越多意味着更少的线程。也就是说，并非所有 SIMD 线程都需要有最大数量的寄存器。Pascal 架构师认为，如果所有线程都有最大数量的寄存器，那么大部分这种宝贵的硅片区域将处于闲置状态。

为了能够执行许多 SIMD 指令线程，在创建 SIMD 指令线程时，会为每个线程动态分配一组 SIMD 处理器上的物理寄存器，并在 SIMD 线程退出时释放这些寄存器。例如，程序员可以拥有一个线程块，该线程块使用每个线程 36 个寄存器，同时拥有 16 个 SIMD 线程，而另一个线程块具有每个线程 20 个寄存器，同时拥有 32 个 SIMD 线程。后续线程块可能会以任何顺序显示，并且必须按需分配寄存器。虽然这种可变性可能导致碎片化并使某些寄存器不可用，但实际上大多数线程块对给定的可矢量化循环（“网格”）使用相同数量的寄存器。硬件必须知道每个线程块的寄存器在大寄存器文件中的位置，并且这是按每个线程块为基础记录的。这种灵活性需要硬件中的路由、仲裁和存储，因为给定线程块的特定寄存器最终可能位于寄存器文件中的任何位置。

请注意，CUDA 线程只是 SIMD 指令线程的垂直切片，对应于一个 SIMD 通道执行的一个元素。请注意，CUDA 线程与 POSIX 线程有很大不同；您无法从 CUDA 线程发出任意系统调用。

NVIDA GPU Instruction Set Architecture

与大多数系统处理器不同，NVIDIA 编译器的指令集目标是硬件指令集的抽象。PTX（并行线程执行）为编译器提供稳定的指令集，并在各代 GPU 之间提供兼容性。硬件指令集对程序员是隐藏的。PTX 指令描述单个 CUDA 线程上的操作，通常与硬件指令一一对应，但一条 PTX 指令可以扩展为多条机器指令，反之亦然。PTX 使用无限数量的一次性写入寄存器，编译器必须运行寄存器分配过程，以将 PTX 寄存器映射到实际设备上可用的固定数量的读写硬件寄存器。优化器随后运行，可以进一步减少寄存器使用。此优化器还会消除死代码、折叠指令，并计算分支可能发散的位置以及发散路径可能收敛的位置。

尽管 x86 微体系结构和 PTX 之间存在一些相似之处，因为两者都转换为内部形式（x86 的微指令），但区别在于此转换在 x86 上是在运行时在硬件中发生的，而在 GPU 上是在软件和加载时发生的。

PTX 指令的格式为

opcode.type d, a, b, c;

c 是源操作数；操作类型是以下之一：源操作数是 32 位或 64 位寄存器或常数值。目标是寄存器，存储指令除外。

图 4.17 显示了基本的 PTX 指令集。所有指令都可以通过 1 位谓词寄存器进行谓词化，该寄存器可以通过设置谓词指令 (setp) 进行设置。控制流指令是函数调用和返回、线程退出、分支以及线程块内的线程的屏障同步 (bar.sync)。在分支指令前放置谓词可为我们提供条件分支。编译器或 PTX 程序员将虚拟寄存器声明为 32 位或 64 位类型值或无类型值。例如，R0、R1 等用于 32 位值，RD0、RD1 等用于 64 位寄存器。回想一下，虚拟寄存器到物理寄存器的分配在加载时使用 PTX 发生。

以下 PTX 指令序列适用于第 292 页上的 DAXPY 循环的一次迭代：

CUDA 编程模型为每个循环迭代分配一个 CUDA 线程，并为每个线程块（blockIdx）提供一个唯一的标识符编号，为块内的每个 CUDA 线程提供一个编号（threadIdx）。因此，它创建了 8192 个 CUDA 线程，并使用唯一编号来寻址数组中的每个元素，因此没有递增或分支代码。前三个 PTX 指令计算 R8 中的唯一元素字节偏移量，该偏移量被添加到数组的基址。以下 PTX 指令加载两个双精度浮点操作数，将它们相乘并相加，然后存储和值。（我们将在下面描述与 CUDA 代码“if (i < n)”对应的 PTX 代码。）

请注意，与矢量体系结构不同，GPU 没有用于顺序数据传输、跨步数据传输和收集-分散数据传输的单独指令。

所有数据传输都是聚集-分散的！为了重新获得顺序（单位步幅）数据传输的效率，GPU 包括特殊的地址合并硬件，以识别 SIMD 指令线程内的 SIMD 通道何时集体发出顺序地址。然后，该运行时硬件通知内存接口单元请求 32 个顺序字的块传输。为了获得这一重要的性能改进，GPU 程序员必须确保相邻的 CUDA 线程同时访问附近的地址，以便将它们合并到一个或几个内存或缓存块中，我们的示例就是这样做的。

GPU 中的条件分支

与单元步幅数据传输的情况一样，矢量体系结构和 GPU 处理 IF 语句的方式之间存在很大的相似性，前者主要在软件中实现该机制，硬件支持有限，而后者则利用了更多的硬件。正如我们将看到的，除了显式谓词寄存器之外，GPU 分支硬件还使用内部掩码、分支同步堆栈和指令标记来管理分支何时分歧为多个执行路径以及路径何时收敛。

在 PTX 汇编器级别，一个 CUDA 线程的控制流由 PTX 指令 branch、call、return 和 exit 描述，加上程序员使用每线程通道 1 位谓词寄存器指定的每个指令的单独的每线程通道谓词。PTX 汇编器分析 PTX 分支图并将其优化为最快的 GPU 硬件指令序列。每个都可以对分支做出自己的决定，而无需锁定步调。

在 GPU 硬件指令级别，控制流包括分支、跳转、跳转索引、调用、调用索引、返回、退出以及管理分支同步堆栈的特殊指令。GPU 硬件为每个 SIMD 线程提供自己的堆栈；堆栈条目包含标识符令牌、目标指令地址和目标线程活动掩码。有 GPU 特殊指令可为 SIMD 线程推送堆栈条目，以及可弹出堆栈条目或将堆栈解压缩到指定条目并使用目标线程活动掩码分支到目标指令地址的特殊指令和指令标记。GPU 硬件指令还具有单独的每通道谓词（启用/禁用），由每个通道的 1 位谓词寄存器指定。

PTX 汇编器通常会将使用 PTX 分支指令编码的简单外层 IF-THEN-ELSE 语句优化为仅预测的 GPU 指令，而没有任何 GPU 分支指令。更复杂的控制流通常会导致预测和 GPU 分支指令与特殊指令和标记混合，这些指令和标记使用分支同步堆栈在某些通道分支到目标地址时推送堆栈条目，而其他通道则会直接通过。NVIDIA 表示，当这种情况发生时，分支会发散。当 SIMD 通道执行同步标记或收敛时也会使用此混合，这会弹出堆栈条目并使用堆栈条目线程活动掩码分支到堆栈条目地址。

PTX 汇编器识别循环分支并生成 GPU 分支指令，这些指令分支到循环顶部，以及特殊堆栈指令来处理从循环中中断的各个通道，并在所有通道完成循环后使 SIMD 通道聚合。GPU 索引跳转和索引调用指令将条目压入堆栈，以便当所有通道完成 switch 语句或函数调用时，SIMD 线程聚合。

GPU 设置谓词指令（图 4.17 中的 setp）评估 IF 语句的条件部分。然后，PTX 分支指令取决于该谓词。如果 PTX 汇编器生成没有 GPU 分支指令的谓词指令，它将使用每通道谓词寄存器为每个指令启用或禁用每个 SIMD 通道。IF 语句 THEN 部分内的线程中的 SIMD 指令将操作广播到所有 SIMD 通道。谓词设置为 1 的那些通道执行操作并存储结果，而其他 SIMD 通道不执行操作或存储结果。对于 ELSE 语句，指令使用谓词的补语（相对于 THEN 语句），因此现在处于空闲状态的 SIMD 通道执行操作并存储结果，而它们以前处于活动状态的兄弟通道则不执行操作。在 ELSE 语句的末尾，指令是未设置谓词的，因此原始计算可以继续进行。因此，对于长度相等的路径，IF-THEN-ELSE 的效率为 50% 或更低。

IF 语句可以嵌套，因此可以使用堆栈，而 PTX 汇编器通常会为复杂控制流生成混合的谓词指令和 GPU 分支以及特殊同步指令。请注意，深度嵌套可能意味着大多数 SIMD 通道在执行嵌套条件语句期间处于空闲状态。因此，具有等长路径的双重嵌套 IF 语句以 25% 的效率运行，三重嵌套以 12.5% 的效率运行，依此类推。类似的情况是矢量处理器仅在少数掩码位为 1 时运行。

降低一个详细级别，PTX 汇编器在适当的条件分支指令上设置“分支同步”标记，该指令将当前活动掩码推送到每个 SIMD 线程内的堆栈中。如果条件分支发散（某些通道采用分支，而某些通道则贯穿），它将推送堆栈条目并根据条件设置当前内部活动掩码。分支同步标记在 ELSE 部分之前弹出发散的分支条目并翻转掩码位。在 IF 语句的末尾，PTX 汇编器添加另一个分支同步标记，该标记将先前的活动掩码从堆栈弹出到当前活动掩码中。

如果所有掩码位都设置为 1，那么 THEN 末尾的分支指令将跳过 ELSE 部分中的指令。如果所有掩码位都为 0，则 THEN 部分也有类似的优化，因为条件分支会跳过 THEN 指令。并行 IF 语句和 PTX 分支通常使用一致的分支条件（所有通道都同意遵循相同的路径），以便 SIMD 线程不会分歧到不同的单独通道控制流。PTX 汇编器优化此类分支以跳过 SIMD 线程的任何通道都不执行的指令块。此优化在条件错误检查中很有用，例如，必须进行测试但很少进行测试的情况。

假设 R8 已经具有缩放的线程 ID，则此 IF 语句可以编译为以下 PTX 指令，其中 *Push、*Comp、*Pop 表示 PTX 汇编器插入的分支同步标记，用于推送旧掩码、对当前掩码取反以及弹出以恢复旧掩码：

同样，通常 IF-THEN-ELSE 语句中的所有指令都由 SIMD 处理器执行。只是 THEN 指令仅对部分 SIMD 通道启用，而 ELSE 指令对部分通道启用。如前所述，在各个通道对谓词分支达成一致的惊人常见情况下（例如，对所有通道都相同的参数值进行分支，以便所有活动掩码位均为 0 或均为 1），该分支将跳过 THEN 指令或 ELSE 指令。

这种灵活性使得一个元素似乎有自己的程序计数器；然而，在最慢的情况下，每 2 个时钟周期只能存储一个 SIMD 通道，其余通道处于空闲状态。向量体系结构的类似最慢情况是仅使用一个掩码位设置为 1 来操作。这种灵活性可能导致天真的 GPU 程序员产生较差的性能，但它在程序开发的早期阶段可能会有所帮助。但是，请记住，时钟周期中 SIMD 通道的唯一选择是执行 PTX 指令中指定的操作或处于空闲状态；两个 SIMD 通道不能同时执行不同的指令。

这种灵活性也有助于解释给 SIMD 指令中的每个元素命名的 CUDA 线程，因为它给出了独立操作的假象。一个天真的程序员可能会认为这种线程抽象意味着 GPU 可以更优雅地处理条件分支。一些线程朝一个方向执行，其余的线程朝另一个方向执行，只要您不着急，它似乎是真的。每个 CUDA 线程要么执行与线程块中每个其他线程相同的指令，要么处于空闲状态。这种同步使得处理带有条件分支的循环变得更加容易，因为掩码功能可以关闭 SIMD 通道，并且它会自动检测循环的结束。

由此产生的性能有时会掩盖这种简单的抽象。编写在高度独立的 MIMD 模式下操作 SIMD 通道的程序就像编写在物理内存较小的计算机上使用大量虚拟地址空间的程序。两者都是正确的，但它们运行速度可能非常慢，以至于程序员对结果不满意。

条件执行是 GPU 在运行时硬件中执行矢量架构在编译时执行的操作的一个案例。矢量编译器执行双重 IF 转换，生成四个不同的掩码。执行基本上与 GPU 相同，但针对矢量执行了一些更多的开销指令。矢量架构的优势在于与标量处理器集成，从而允许它们在 0 个案例主导计算时避免时间。尽管这取决于标量处理器与矢量处理器的速度，但使用标量的交叉点可能是当掩码位少于 20% 为 1 时。GPU 在运行时可用的一个优化（但矢量架构在编译时不可用）是在掩码位全部为 0 或全部为 1 时跳过 THEN 或 ELSE 部分。

因此，GPU 执行条件语句的效率取决于分支发散的频率。例如，一个特征值的计算具有深度条件嵌套，但对代码的测量表明，大约 82% 的时钟周期问题在 32 个掩码位中将 29 到 32 个设置为 1，因此 GPU 执行此代码的效率高于预期。

请注意，相同的机制处理向量循环的带状挖掘——当元素数量与硬件不完全匹配时。本节开头的示例表明，IF 语句检查此 SIMD 通道元素编号（存储在前一个示例中的 R8 中）是否小于限制（i < n），并适当地设置掩码。

NVIDIA GPU 内存结构

图 4.18 显示了 NVIDIA GPU 的内存结构。多线程 SIMD 处理器中的每个 SIMD 通道都会获得一段私有的片外 DRAM，我们称之为专用内存。它用于堆栈帧、溢出寄存器以及不适合寄存器的专用变量。SIMD 通道不共享专用内存。GPU 将此专用内存缓存在 L1 和 L2 缓存中，以帮助寄存器溢出并加快函数调用速度。

我们称每个多线程 SIMD 处理器本地的片上存储器为本地存储器。它是一个小型暂存存储器，具有低延迟（几十个时钟）和高带宽（128 字节/时钟），程序员可以在其中存储需要重复使用的数据，无论是由同一个线程还是由同一个线程或同一个线程块中的另一个线程。本地内存大小有限，通常为 48 KiB。它不携带在同一处理器上执行的线程块之间的状态。它由多线程 SIMD 处理器内的 SIMD 通道共享，但此内存不会在多线程 SIMD 处理器之间共享。多线程 SIMD 处理器在创建线程块时会动态地将本地内存的一部分分配给线程块，并在线程块的所有线程退出时释放内存。本地内存的那一部分是该线程块私有的。

最后，我们将整个 GPU 和所有线程块共享的片外 DRAM 称为 GPU 内存。我们的向量乘法示例仅使用了 GPU 内存。

称为主机的系统处理器可以读写 GPU 内存。本地内存对主机不可用，因为它是每个多线程 SIMD 处理器的私有内存。私有内存对主机也不可用。

GPU 内存由所有网格（矢量化循环）共享，局部内存由线程块（矢量化循环的主体）内的所有 SIMD 指令线程共享，而专用内存对单个 CUDA 线程是专用的。Pascal 允许抢占网格，这要求所有局部内存和专用内存都能够保存到全局内存中并从中恢复。为了完整起见，GPU 还可以通过 PCIe 总线访问 CPU 内存。当最终结果的地址位于主机内存中时，通常会使用此路径。此选项消除了从 GPU 内存到主机内存的最终复制。

由同一个线程或同一个线程块中的另一个线程。本地内存大小有限，通常为 48 KiB。它不携带在同一处理器上执行的线程块之间的状态。它由多线程 SIMD 处理器内的 SIMD 通道共享，但此内存不会在多线程 SIMD 处理器之间共享。多线程 SIMD 处理器在创建线程块时会动态地将本地内存的一部分分配给线程块，并在线程块的所有线程退出时释放内存。本地内存的那一部分是该线程块私有的。

最后，我们将整个 GPU 和所有线程块共享的片外 DRAM 称为 GPU 内存。我们的向量乘法示例仅使用了 GPU 内存。

称为主机的系统处理器可以读写 GPU 内存。本地内存对主机不可用，因为它是每个多线程 SIMD 处理器的私有内存。私有内存对主机也不可用。

GPU 传统上使用较小的流缓存，而不是依靠大缓存来包含应用程序的整个工作集，并且由于其工作集可以是数百兆字节，因此依靠 SIMD 指令线程的广泛多线程来隐藏对 DRAM 的长延迟。鉴于使用多线程来隐藏 DRAM 延迟，系统处理器中用于大型 L2 和 L3 缓存的芯片面积反而用于计算资源和大量寄存器，以保存许多 SIMD 指令线程的状态。相比之下，如上所述，向量加载和存储会使延迟摊销到许多元素上，因为它们只支付一次延迟，然后对其余访问进行流水线处理。

尽管隐藏在许多线程背后的内存延迟是 GPU 和矢量处理器的最初理念，但所有最近的 GPU 和矢量处理器都具有缓存以减少延迟。该论点遵循排队理论中的 Little 定律：延迟越长，在内存访问期间需要运行的线程就越多，这反过来又需要更多的寄存器。因此，添加 GPU 缓存以降低平均延迟，从而掩盖寄存器数量的潜在短缺。

为了提高内存带宽并减少开销，如上所述，PTX 数据传输指令与内存控制器协同工作，将来自同一 SIMD 线程的各个并行线程请求合并成一个内存块请求，当地址位于同一块中时。这些限制适用于 GPU 程序，有点类似于系统处理器程序参与硬件预取的准则（参见第 2 章）。GPU 内存控制器还将保存请求并将它们一起发送到同一打开的页面以提高内存带宽（参见第 4.6 节）。第 2 章详细描述了 DRAM，以便读者了解对相关地址进行分组的潜在好处。

Pascal GPU 架构的创新

Pascal 的多线程 SIMD 处理器比图 4.20 中的简化版本更复杂。为了提高硬件利用率，每个 SIMD 处理器有两个 SIMD 线程调度器，每个调度器都有多个指令分派单元（某些 GPU 有四个线程调度器）。双 SIMD 线程调度器选择两个 SIMD 指令线程，并从每个线程向两组 16 个 SIMD 通道、16 个加载/存储单元或 8 个特殊功能单元发出一个指令。由于有多个执行单元可用，因此每个时钟周期都会调度两个 SIMD 指令线程，从而使 64 个通道处于活动状态。由于线程是独立的，因此无需检查指令流中的数据相关性。此创新类似于可以从两个独立线程发出矢量指令的多线程矢量处理器。图 4.19 显示了发出指令的双调度器，图 4.20 显示了 Pascal GP100 GPU 的多线程 SIMD 处理器的框图。

每一代新的 GPU 通常都会增加一些新功能，以提高性能或让程序员更容易编程。以下是 Pascal 的四项主要创新：

Pascal 的双 SIMD 线程调度器框图。将此设计与图 4.16 中的单 SIMD 线程设计进行比较。

• 快速单精度、双精度和半精度浮点运算——Pascal GP100 芯片在三种尺寸上具有显著的浮点性能，所有这些都是浮点 IEEE 标准的一部分。GPU 的单精度浮点运算峰值速度为 10 TeraFLOP/s。双精度速度大约为其一半，为 5 TeraFLOP/s，而半精度速度大约为其两倍，为 20 TeraFLOP/s（以 2 元素向量表示）。原子内存操作包括所有三种尺寸的浮点加法。Pascal GP100 是第一款具有如此高半精度性能的 GPU。
• 高带宽内存——Pascal GP100 GPU 的下一个创新是使用堆叠的高带宽内存 (HBM2)。此内存具有宽总线，4096 条数据线以 0.7 GHz 运行，提供 732 GB/s 的峰值带宽，是以前 GPU 的两倍多。
• 高速芯片到芯片互连——鉴于 GPU 的协处理器性质，当尝试将多个 GPU 与一个 CPU 一起使用时，PCI 总线可能成为通信瓶颈。Pascal GP100 引入了 NVLink 通信通道，支持每个方向高达 20 GB/s 的数据传输。每个 GP100 具有 4 个 NVLink 通道，每个芯片提供 160 GB/s 的峰值总芯片到芯片带宽。具有 2、4 和 8 个 GPU 的系统可用于多 GPU 应用，其中每个 GPU 都可以对通过 NVLink 连接的任何 GPU 执行加载、存储和原子操作。此外，在某些情况下，NVLink 通道可以与 CPU 通信。例如，IBM Power9 CPU 支持 CPU-GPU 通信。在此芯片中，NVLink 提供了所有 GPU 和连接在一起的 CPU 之间内存的相干视图。它还提供高速缓存到高速缓存通信，而不是内存到内存通信。

  

  Pascal GPU 的多线程 SIMD 处理器的框图。64 个 SIMD 通道（内核）中的每一个都有一个流水线浮点单元、一个流水线整数单元、一些用于将指令和操作数分派给这些单元的逻辑，以及一个用于保存结果的队列。64 个 SIMD 通道与 32 个执行 64 位浮点运算的双精度算术逻辑单元 (DP 单元)、16 个加载存储单元 (LD/ST) 和 16 个计算平方根、倒数、正弦和余弦等函数的特殊功能单元 (SFU) 交互。

• 统一的虚拟内存和分页支持——Pascal GP100 GPU 在统一的虚拟地址空间中添加了页面错误功能。此功能为单个系统中所有 GPU 和 CPU 上相同的每个数据结构提供一个虚拟地址。当线程访问远程地址时，内存页会传输到本地 GPU 以供后续使用。统一内存通过提供按需分页而不是在 CPU 和 GPU 之间显式复制内存来简化编程模型或GPU 之间。它还允许分配比 GPU 上存在的内存多得多的内存来解决内存需求大的问题。与任何虚拟内存系统一样，必须注意避免过多的页面移动。

矢量架构和 GPU 之间的异同

正如我们所见，矢量架构和 GPU 之间确实存在许多相似之处。除了 GPU 的古怪术语外，这些相似之处导致了架构界对 GPU 的新颖性产生混淆。现在您已经了解了矢量计算机和 GPU 的内部结构，您就可以理解它们的相似之处和不同之处。由于这两种架构都旨在执行数据级并行程序，但采取了不同的路径，因此为了更好地理解 DLP 硬件的需要，需要深入比较。图 4.21 先显示矢量术语，然后显示 GPU 中最接近的等效术语。

SIMD 处理器类似于矢量处理器。GPU 中的多个 SIMD 处理器充当独立的 MIMD 内核，就像许多矢量计算机具有多个矢量处理器一样。此视图将 NVIDIA Tesla P100 视为具有 56 个内核的机器，并具有对多线程的硬件支持，其中每个内核具有 64 个通道。最大的区别在于多线程，这是 GPU 的基础，而大多数矢量处理器却没有。

从这两个架构的寄存器来看，我们实现中的 RV64V 寄存器文件保存了整个向量，即连续的元素块。相比之下，GPU 中的单个向量将分布在所有 SIMD 通道的寄存器中。RV64V 处理器有 32 个向量寄存器，每个寄存器可能有 32 个元素，总共 1024 个元素。SIMD 指令的 GPU 线程最多有 256 个寄存器，每个寄存器有 32 个元素，总共 8192 个元素。这些额外的 GPU 寄存器支持多线程。

图 4.22 是左侧向量处理器的执行单元和右侧 GPU 的多线程 SIMD 处理器的框图。出于教学目的，我们假设向量处理器有四个通道，多线程 SIMD 处理器也有四个 SIMD 通道。此图显示了四个 SIMD 通道非常类似于四通道向量单元协同工作，并且 SIMD 处理器非常类似于向量处理器。

实际上，GPU 中的通道更多，因此 GPU “时钟周期“更短。

虽然矢量处理器可能有 2 到 8 个通道和一个矢量长度，例如 32——发出 4 到 16 个时钟周期——多线程 SIMD 处理器可能有 8 或 16 个通道。SIMD 线程的宽度为 32 个元素，因此 GPU 时钟周期只有 2 或 4 个。这就是我们使用“SIMD 处理器”作为更具描述性的术语的原因，因为它比传统矢量处理器设计更接近 SIMD 设计。

图 4.22 左侧有 4 个通道的矢量处理器，右侧有 4 个 SIMD 通道的 GPU 的多线程 SIMD 处理器。（GPU 通常有 16 或 32 个 SIMD 通道。控制处理器为标量向量运算提供标量操作数，为对内存的单位和非单位步幅访问递增寻址，并执行其他记帐类型操作。只有当地址合并单元可以发现本地化寻址时，峰值内存性能才会出现在 GPU 中。同样，当所有内部掩码位设置相同时，会出现峰值计算性能。请注意，SIMD 处理器每个 SIMD 线程有一台 PC，以帮助进行多线程处理。

最接近矢量化循环的 GPU 术语是网格Grid，而 PTX 指令最接近矢量指令，因为 SIMD 线程将 PTX 指令广播到所有 SIMD 通道。

关于这两个架构中的内存访问指令，所有 GPU 加载都是收集指令，所有 GPU 存储都是分散指令。如果 CUDA 线程的数据地址引用同时落在同一缓存/内存块中的附近地址，GPU 的地址合并单元将确保高内存带宽。向量架构的显式单位步长加载和存储指令与 GPU 编程的隐式单位步长不同，这就是为什么编写高效的 GPU 代码需要程序员以 SIMD 操作的方式进行思考，即使 CUDA 编程模型看起来像 MIMD。由于 CUDA 线程可以生成自己的地址，因此在向量架构和 GPU 中都可以找到跨步以及收集分散的寻址向量。

正如我们多次提到的，这两种架构采用截然不同的方法来隐藏内存延迟。矢量架构通过深度流水线访问在矢量的所有元素上摊销它，因此您只需为每个矢量加载或存储支付一次延迟。因此，矢量加载和存储就像存储器和矢量寄存器之间的块传输。相比之下，GPU 使用多线程隐藏内存延迟。（一些研究人员正在研究将多线程添加到矢量架构中，以试图获得两全其美的优势。

关于条件分支指令，两种架构都使用掩码寄存器实现它们。两个条件分支路径都会占用时间和/或空间，即使它们不存储结果也是如此。不同之处在于，矢量编译器在软件中显式管理掩码寄存器，而 GPU 硬件和汇编程序使用分支同步标记和内部堆栈隐式管理它们，以保存、补充和恢复掩码。

矢量计算机的控制处理器在矢量指令的执行中起着重要作用。它将操作广播到所有矢量通道，并广播矢量标量操作的标量寄存器值。它还执行 GPU 中显式的隐式计算，例如自动递增单位步幅和非单位步幅加载和存储的内存地址。GPU 中缺少控制处理器。最接近的类比是线程块调度器，它将线程块（矢量循环的主体）分配给多线程 SIMD 处理器。GPU 中的运行时硬件机制既能生成地址，又能发现它们是否相邻，这在许多 DLP 应用程序中很常见，但其能效可能不如使用控制处理器。

矢量计算机中的标量处理器执行矢量程序的标量指令;也就是说，它执行的操作在向量单元中执行的操作太慢。尽管与 GPU 关联的系统处理器与矢量架构中的标量处理器最接近，但单独的地址空间加上通过 PCIe 总线传输意味着将它们一起使用需要数千个时钟周期的开销。在矢量计算机中进行浮点计算时，标量处理器可能比矢量处理器慢，但系统处理器与多线程 SIMD 处理器的比率不同（考虑到开销）。

在 GPU 中必须执行您期望使用矢量计算机中的标量处理器执行的计算。也就是说，与其在系统处理器上计算并传达结果，不如使用谓词寄存器和内置掩码禁用除一个 SIMD 通道之外的所有通道，并使用一个 SIMD 通道执行标量工作。矢量计算机中相对简单的标量处理器可能比 GPU 解决方案更快、更节能。如果系统处理器和 GPU 在未来变得更加紧密地结合在一起，那么看看系统处理器是否可以像标量处理器对矢量和多媒体 SIMD 架构所做的那样发挥相同的作用将会很有趣。

多媒体 SIMD 计算机和 GPU 的异同

从高层次来看，具有多媒体 SIMD 指令扩展的多核计算机确实与 GPU 有相似之处。图 4.23 总结了相似之处和不同之处。

两者都是多处理器，其处理器使用多个 SIMD 通道，但 GPU 具有更多处理器和更多通道。两者都使用硬件多线程来提高处理器利用率，但 GPU 具有对更多线程的硬件支持。两者在单精度和双精度浮点运算的峰值性能之间的性能比率大约为 2:1。两者都使用缓存，但 GPU 使用较小的流式缓存，而多核计算机使用大型多级缓存，试图完全包含整个工作集。两者都使用 64 位地址空间，但 GPU 中的物理主内存要小得多。两者都支持页级内存保护以及按需分页，这使它们能够寻址比其板载内存多得多的内存。

除了处理器、SIMD 通道、硬件线程支持和缓存大小方面的巨大数值差异之外，还有许多架构差异。标量处理器和多媒体 SIMD 指令在传统计算机中紧密集成；它们在 GPU 中通过 I/O 总线分隔，甚至具有单独的主内存。GPU 中的多个 SIMD 处理器使用单个地址空间，并且在某些系统上可以支持对所有内存的相干视图，前提是得到 CPU 供应商（例如 IBM Power9）的支持。与 GPU 不同，多媒体 SIMD 指令在历史上不支持收集-分散内存访问，第 4.7 节显示这是一个重大遗漏。

总结

现在，随着面纱被揭开，我们可以看到 GPU 实际上只是多线程 SIMD 处理器，尽管它们比传统的内核计算机拥有更多的处理器、每个处理器拥有更多的通道以及更多的多线程硬件。例如，Pascal P100 GPU 拥有 56 个 SIMD 处理器，每个处理器有 64 个通道，并且硬件支持 64 个 SIMD 线程。Pascal 通过向两组 SIMD 通道发出来自两个 SIMD 线程的指令来支持指令级并行性。GPU 还具有较少的缓存内存（Pascal 的 L2 缓存为 4 MiB），并且它可以与协作的远端标量处理器或远端 GPU 保持一致。

CUDA 编程模型将所有这些形式的并行性都封装在一个抽象（CUDA 线程）中。因此，CUDA 程序员可以考虑对数千个线程进行编程，尽管它们实际上是在许多 SIMD 处理器的许多通道上执行每个 32 个线程的块。想要获得良好性能的 CUDA 程序员要记住，这些线程按块组织并一次执行 32 个，并且地址需要是相邻地址才能从内存系统获得良好性能。

虽然我们在本节中使用了 CUDA 和 NVIDIA GPU，但请放心，OpenCL 编程语言和其他公司的 GPU 中也存在相同的想法。

现在您更了解 GPU 的工作原理，我们将揭示真正的术语。

图 4.24 和 4.25 将本节的描述性术语和定义与官方 CUDA/NVIDIA 和 AMD 术语和定义相匹配。我们还包括 OpenCL 术语。我们认为 GPU 学习曲线陡峭的部分原因是使用了诸如“流式多处理器”表示 SIMD 处理器，“线程处理器”表示 SIMD 通道以及“共享内存”表示本地内存的术语——尤其是因为本地内存不会在 SIMD 处理器之间共享！我们希望这种两步法能更快地让您了解曲线，即使它有点间接。

程序中的循环是我们之前在本章和第 5 章中讨论过的许多类型的并行的源泉。在本节中，我们将讨论用于发现我们可以在程序中利用的并行程度的编译器技术以及对这些编译器技术的硬件支持。我们将精确地定义何时循环是并行的（或可矢量化的），依赖性如何阻止循环并行，以及消除某些类型依赖性的技术。查找和操纵循环级并行对于利用 DLP 和 TLP 以及我们在附录 H 中检查的更激进的静态 ILP 方法（例如，VLIW）至关重要。

从本章中使用的术语到官方 NVIDIA/CUDA 和 AMD 术语的转换。请注意，我们的描述性术语“本地内存”和“私有内存”使用 OpenCL 术语。NVIDIA 使用 SIMM（单指令多线程）而不是 SIMD 来描述流式多处理器。SIMT 优于 SIMD，因为每个线程的分支和控制流不同于任何 SIMD 计算机

循环级并行通常在源代码级别或接近源代码级别进行研究，而大多数 ILP 分析都是在编译器生成指令后进行的。循环级分析涉及确定循环中操作数之间在该循环的迭代中存在的依赖关系。现在，我们只考虑数据依赖关系，数据依赖关系在某个时刻写入操作数并在稍后读取操作数时出现。名称依赖关系也存在，可以通过第 3 章中讨论的重命名技术消除名称依赖关系。

循环级并行的分析重点在于确定后一次迭代中的数据访问是否依赖于前一次迭代中生成的数据值；这种依赖关系称为循环携带依赖关系。

for (i=999; i>=0; i=i-1)
		x[i] = x[i] + s;

在这个循环中，x[i] 的两次使用是相关的，但这种相关性在单个迭代中，并且不是循环携带的。在不同迭代中 i 的连续使用之间存在循环携带相关性，但这种相关性涉及一个可以轻松识别和消除的归纳变量。我们在第 2 章的第 2.2 节中看到了有关循环展开期间如何消除涉及归纳变量的相关性的示例，我们将在本节后面看到更多示例。

因为查找循环级并行性涉及识别诸如循环、数组引用和归纳变量计算之类的结构，所以编译器可以在源代码级别或接近源代码级别更轻松地执行此分析，而不是在机器代码级别。我们来看一个更复杂的示例。

for (i=0; i<100; i=i+1) {
		A[i+1] = A[i] + C[i]; /* S1 */
		B[i+1] = B[i] + A[i+1]; /* S2 */
}

假设 A、B 和 C 是不同的、不重叠的数组。（实际上，这些数组有时可能是相同的或可能重叠。因为这些数组可能作为参数传递给包含此循环的某个过程，所以确定数组是否重叠或是否相同通常需要对程序进行复杂的跨过程分析。）循环中语句 S1 和 S2 之间的数据依赖关系是什么？

存在两种不同的依赖关系：

1. S1 使用 S1 在较早迭代中计算的值，因为迭代 i 计算 A[i + 1]，而 A[i + 1] 在迭代 i + 1 中被读取。S2 对 B[i] 和 B[i + 1] 也是如此。
2. S2 使用 S1 在同一迭代中计算的值 A[i + 1]。

这两个依赖关系是不同的，并且具有不同的影响。为了了解它们的不同之处，我们假设一次只存在其中一个依赖关系。因为语句 S1 的依赖关系是 S1 的较早迭代，所以此依赖关系是循环携带的。此依赖关系迫使此循环的连续迭代按顺序执行。

第二个相关性（S2 依赖于 S1）在迭代中，并且不是循环携带的。因此，如果这是唯一的相关性，则只要迭代中每对语句保持顺序，循环的多个迭代就会并行执行。我们在第 2.2 节的示例中看到了这种类型的相关性，其中展开可以揭示并行性。这些循环内相关性很常见；例如，使用链接的一系列向量指令恰好表现出这种相关性。

也可以有循环携带的相关性，但不会阻止并行性，如下一个示例所示。

for (i=0; i<100; i=i+1) {
		A[i] = A[i] + B[i]; /* S1 */
		B[i+1] = C[i] + D[i]; /* S2 */
}

S1 和 S2 之间存在哪些依赖关系？此循环是否并行？如果不是，请说明如何使其并行。

语句 S1 使用语句 S2 在前一次迭代中分配的值，因此 S2 和 S1 之间存在循环携带依赖关系。尽管存在这种循环携带依赖关系，但此循环可以实现并行。与之前的循环不同，这种依赖关系不是循环的；没有语句依赖于自身，尽管 S1 依赖于 S2，但 S2 不依赖于 S1。如果循环可以编写而依赖关系中没有循环，则该循环是并行的，因为没有循环意味着依赖关系对语句进行了部分排序。

尽管在前面的循环中没有循环依赖关系，但必须对其进行转换以符合部分排序并揭示并行性。以下两点观察对于此转换至关重要：

1. S1 到 S2 没有依赖关系。如果有，则依赖关系中会出现循环，并且循环将不是并行的。由于不存在这种其他依赖关系，因此互换这两个语句不会影响 S2 的执行。
2. 在循环的第一次迭代中，语句 S2 依赖于在启动循环之前计算的 B[0] 的值。

这两个观察结果使我们能够用以下代码序列替换前面的循环：

A[0] = A[0] + B[0];
for (i=0; i<99; i=i+1) {
		B[i+1] = C[i] + D[i];
		A[i+1] = A[i+1] + B[i+1];
}
B[100] = C[99] + D[99];

两个语句之间的依赖不再是循环携带的，因此可以重叠循环的迭代，只要每个迭代中的语句保持顺序即可。

我们的分析需要从查找所有循环携带的依赖开始。这种依赖信息是不精确的，因为它告诉我们这种依赖可能存在。考虑以下示例：

for (i=0;i<100;i=i+1) {
		A[i] = B[i] + C[i]
		D[i] = A[i] * E[i]
}

在此示例中对 A 的第二次引用不必翻译为加载指令，因为我们知道该值由前一条语句计算并存储。因此，对 A 的第二次引用可以简单地引用计算 A 的寄存器。执行此优化需要知道这两个引用始终指向相同的内存地址，并且没有对相同位置的中间访问。通常，数据相关性分析表明只有一个引用可能依赖于另一个引用；需要更复杂的分析来确定两个引用必须指向完全相同的地址。在前面的示例中，此分析的简单版本就足够了，因为这两个引用位于同一个基本块中。

循环携带相关性通常采用递归的形式。当变量根据该变量在早期迭代（通常是紧接在前一个迭代）中的值定义时，就会发生递归，如下面的代码片段所示：

for (i=1;i<100;i=i+1) {
		Y[i] = Y[i-1] + Y[i];
}

检测循环有以下两个原因很重要：一些架构（尤其是矢量计算机）对执行循环有特殊支持，并且在 ILP 上下文中，仍然有可能利用相当多的并行性。

查找依赖关系

显然，查找程序中的依赖关系对于确定哪些循环可能包含并行性以及消除名称依赖关系都很重要。依赖性分析的复杂性也由于数组的存在而产生以及 C 或 C++ 等语言中的指针，或 Fortran 中的按引用传递参数传递。因为标量变量引用明确地引用一个名称，所以通常可以很容易地分析别名，因为指针和引用参数导致分析中出现一些复杂性和不确定性。

编译器通常如何检测依赖关系？几乎所有依赖性分析算法都基于数组索引是仿射的假设。最简单的说，如果一维数组索引可以写成 a*i+b的形式，其中 a 和 b 是常量，i 是循环索引变量，那么该索引就是仿射的。如果每个维度的索引都是仿射的，那么多维数组的索引就是仿射的。稀疏数组访问通常具有 x[y[i]] 的形式，是非仿射访问的主要示例之一。

因此，确定循环中对同一数组的两次引用之间是否存在依赖关系等同于确定两个仿射函数是否可以在循环边界的不同索引处具有相同的值。例如，假设我们已将值 a 存储到索引值为c*i+d，其中 i 是从 m 到 n 运行的 for 循环索引变量。如果满足两个条件，则存在依赖关系：

1. 有两个迭代索引 j 和 k，它们都在 for 循环的限制范围内。也就是说，m≤j≤n,
2. 循环将数据存储到由 a 索引的数组元素中j + b，然后从当它由 c 索引时获取同一数组元素

通常，我们无法在编译时确定是否存在依赖关系。例如，a、b、c 和 d 的值可能未知（它们可能是其他数组中的值），因此无法判断是否存在依赖关系。在其他情况下，依赖关系测试可能非常昂贵，但在编译时可以决定；例如，访问可能取决于多个嵌套循环的迭代索引。然而，许多程序主要包含 a、b、c 和 d 都是常量的简单索引。对于这些情况，可以设计出合理的编译时依赖关系测试。

例如，不存在依赖关系的一个简单且充分的测试是最大公约数 (GCD) 测试。它基于以下观察：如果存在循环携带的依赖关系，则 GCD (c, a) 必须整除 (d – b)。（回想一下，如果在执行除法 y/x 时得到一个整数商并且没有余数，则整数 x 整除另一个整数 y。）

使用 GCD 测试来确定以下循环中是否存在依赖关系：

GCD 测试足以保证不存在依赖关系；然而，在某些情况下，GCD 测试成功，但不存在依赖关系。例如，由于 GCD 测试不考虑循环边界，因此可能会出现这种情况。

通常，确定依赖关系是否实际存在是 NP 完全的。

然而，在实践中，许多常见情况都可以以低成本进行精确分析。最近，使用准确性与效率兼备的通用性和成本逐渐增加的精确测试层次结构的方法已被证明是准确且有效的。（如果测试准确地确定了依赖关系是否存在，则该测试是精确的。尽管一般情况是 NP 完全的，但对于便宜得多的受限情况存在精确测试。）除了检测依赖关系的存在之外，编译器还希望对依赖关系的类型进行分类。这种分类允许编译器识别名称依赖关系并通过重命名和复制在编译时消除它们。

在循环之后，变量 X 已重命名为 X1。在循环之后的代码中，编译器可以简单地用 X1 替换名称 X。在这种情况下，重命名不需要实际的复制操作，因为它可以通过替换名称或寄存器分配来完成。然而，在其他情况下，重命名需要复制。

依赖性分析是利用并行性以及第 2 章中介绍的类似于转换的阻塞的关键技术。对于检测循环级别的并行性，依赖性分析是基本工具。为矢量计算机、SIMD 计算机或多处理器有效地编译程序严重依赖于此分析。依赖性分析的主要缺点是它仅适用于有限的情况，即单个循环嵌套内的引用以及使用仿射索引函数。因此，在许多情况下，面向数组的依赖性分析无法告诉我们想要了解的信息；例如，使用指针而不是数组索引来分析访问可能困难得多。（这就是为什么对于许多为并行计算机设计的科学应用程序，Fortran 仍然优于 C 和 C++ 的原因之一。）类似地，分析跨过程调用的引用非常困难。因此，虽然对用顺序语言编写的代码进行分析仍然很重要，但我们还需要诸如 OpenMP 和 CUDA 之类的编写显式并行循环的方法。

消除依赖计算

如前所述，最重要的依赖计算形式之一是递推。点积是一个递推的完美示例：

此循环不是并行的，因为它对变量 sum 具有循环携带依赖性。但是，我们可以将其转换为一组循环，其中一个完全并行，另一个部分并行。第一个循环将执行此循环的完全并行部分。它看起来像这样：

请注意，和已经从标量扩展为矢量量（一种称为标量扩展的变换），并且这种变换使这个新循环完全并行。然而，当我们完成时，我们需要执行归约步骤，该步骤对矢量的元素求和。它看起来像这样：

尽管此循环不是并行的，但它具有一个非常特殊的结构，称为约简。约简在线性代数中很常见，正如我们将在第 6 章中看到的那样，它们也是仓库规模计算机中使用的主要并行基元 MapReduce 的关键部分。通常，任何函数都可以用作规约运算符，常见情况包括 max 和 min 等运算符。

在矢量和 SIMD 架构中，有时会通过特殊硬件处理归约，这使得归约步骤比在标量模式下执行的速度快得多。这些工作原理是实现类似于多处理器环境中可以执行的技术。虽然通用转换适用于任意数量的处理器，但为了简单起见，我们假设有 10 个处理器。在归约和的第一步中，每个处理器执行以下操作（其中 p 为处理器编号，范围从 0 到 9）：

此循环在 10 个处理器中的每个处理器上对 1000 个元素求和，完全并行。然后，一个简单的标量循环可以完成最后 10 个和的求和。在矢量处理器和 SIMD 处理器中使用了类似的方法。

需要注意的是，前面的转换依赖于加法的结合律。尽管具有无限范围和精度的算术是结合的，但计算机算术不是结合的，无论是整数算术（由于范围有限）还是浮点算术（由于范围和精度）。因此，使用这些重构技术有时会导致错误的行为，尽管这种情况很少见。出于这个原因，大多数编译器要求明确启用依赖于结合律的优化。

能量和 DLP：慢而宽与快而窄

数据级并行架构的基本功耗优势源自于第一章中的能量方程。假设有足够的数据级并行性，如果我们将时钟速率减半并将执行资源加倍，则性能是相同的：矢量计算机的通道数加倍，多媒体 SIMD 的寄存器和 ALU 更宽，GPU 的 SIMD 通道更多。如果我们能在降低时钟速率的同时降低电压，我们实际上可以降低计算的能耗和功耗，同时保持相同的峰值性能。因此，GPU 的时钟速率往往低于系统处理器，后者依靠高时钟速率来实现性能（参见第 4.7 节）。

与乱序处理器相比，DLP 处理器可以拥有更简单的控制逻辑，以便在每个时钟周期内启动大量操作；例如，矢量处理器中所有通道的控制都是相同的，并且没有逻辑来决定多条指令问题或推测执行逻辑。它们还获取和解码的指令少得多。矢量架构还可以更轻松地关闭芯片的未使用部分。每条矢量指令都明确描述了指令发出时它在一段时间内所需的全部资源。

存储内存和显存

第 4.2 节指出了矢量架构的大内存带宽对于支持单位步幅、非单位步幅以及收集-分散访问的重要性。

为了实现最高的内存性能，AMD 和 NVIDIA 的高端 GPU 中使用了堆叠式 DRAM。英特尔也在其至强融核产品中使用了堆叠式 DRAM。内存芯片也被称为高带宽内存 (HBM、HBM2)，它们被堆叠并与处理芯片放在同一个封装中。较大的宽度（通常为 1024-4096 个数据线）提供了高带宽，而将内存芯片与处理器芯片放在同一个封装中则降低了延迟和功耗。堆叠式 DRAM 的容量通常为 8-32 GB。

鉴于计算任务和图形加速任务对内存的所有潜在需求，内存系统可能会看到大量不相关的请求。不幸的是，这种多样性会损害内存性能。为了应对这种情况，GPU 的内存控制器维护着针对不同存储库的单独队列，等待有足够的流量来证明打开一行并一次传输所有请求的数据。这种延迟会提高带宽，但会延长延迟，并且控制器必须确保没有处理单元在等待数据时挨饿，否则相邻的处理器可能会闲置。第 4.7 节显示，收集-分散技术和考虑内存库的访问技术可以提供比传统基于缓存的架构大幅提升的性能。

步进访问和 TLB 未命中

步进访问的一个问题是它们如何与矢量架构或 GPU 中虚拟内存的转换后备缓冲区 (TLB) 交互。（GPU 也使用 TLB 进行内存映射。）根据 TLB 的组织方式和内存中正在访问的数组的大小，甚至有可能对数组中的每个元素的访问都获得一个 TLB 未命中！缓存中可能会发生相同类型的冲突，但性能影响可能较小。

鉴于图形应用程序的普及，GPU 现在既可以在移动客户端中找到，也可以在传统服务器和重型台式电脑中找到。图 4.26 列出了 NVIDIA Tegra Parker 片上系统的主要特性，该系统在汽车中很受欢迎，以及用于服务器的 Pascal GPU。GPU 服务器工程师希望能够在电影发布后的五年内进行实时动画。而 GPU 嵌入式工程师则希望在五年内在他们的硬件上完成服务器或游戏机今天所做的事情。

NVIDIA Tegra P1 拥有六个 ARMv8 内核和一个较小的 Pascal GPU（能够达到 750 GFLOPS）和 50 GB/s 的内存带宽。它是用于自动驾驶汽车的 NVIDIA DRIVE PX2 计算平台的关键组件。NVIDIA Tegra X1 是上一代产品，用于多款高端平板电脑，例如 Google Pixel C 和 NVIDIA Shield TV。它具有一个能够达到 512 GFLOPS 的 Maxwell 级 GPU。

本章节中广泛讨论的 NVIDIA Tesla P100 是 Pascal GPU。（Tesla 是 Nvidia 为针对通用计算的产品使用的名称。）时钟频率为 1.4 GHz，并包含 56 个 SIMD 处理器。HBM2 内存路径宽 4096 位，并且在 0.715 GHz 时钟的上升沿和下降沿传输数据，这意味着峰值内存带宽为 732 GB/s。它通过 PCI Express 连接到主机系统处理器和内存16 Gen 3 链路，其峰值双向速率为 32 GB/s。

P100 芯片的所有物理特性都非常大：它包含 153 亿个晶体管，芯片尺寸为 16 纳米 TSMC 工艺中的 645 毫米，典型功率为 300 瓦。