Chapter 04 : Data-Level Parallelism Architecture¶

约 5571 个字 31 行代码 26 张图片预计阅读时间 28 分钟

Abstract

本章主要介绍的内容分为三大块，它们都是 SIMD 的变体：

向量架构（Vector Architecture）
多媒体 SIMD 指令集扩展（Multimedia SIMD Instruction Set Extensions）
图形处理器（Graphic Processing Units, GPUs）

Vector Architecture¶

向量架构的大致原理可以概括为：获取分散在内存中的数据元素集，将它们放在一个较大的顺序寄存器堆内，在寄存器堆上对数据进行操作，最后将结果分散放回到内存中。

单条指令作用在多个向量数据上，导致在独立的数据元素上有多个寄存器 - 寄存器运算
作为由编译器控制的缓冲区，这些寄存器堆能够隐藏内存时延，并且能够利用好内存带宽
功率墙（Power Wall）使得架构师看重那些能够以不达到无序超标量处理器那么高的成本为代价，就能获取高性能的架构，而向量架构正符合这样的要求——仅需在一个简单的标量处理器上用到向量指令（Vector Instructions），就能够提升性能

RV64V¶

下图展示了一个向量处理器：

其包含了以下组件：

向量寄存器（Vector Registers）：
- RV64V 有 32 个向量寄存器，每个向量寄存器的大小为 64 位，保存一个向量
- 向量寄存器堆需要提供足够量的端口，以供给向量函数单元；这些端口允许不同向量寄存器的向量运算间的高度重叠
- 至少有 16 个读端口和 8 个写端口，它们通过一堆纵横开关和函数单元的输入和输出相连
- 提升寄存器堆带宽的一个方法是使用多个分区（Multiple Banks）
向量函数单元（Vector Functional Units）：
- 所有单元都是完全流水线化的，并且每个时钟周期都可以开始一条新的指令
- 控制单元用于检测冒险，包括函数单元的结构冒险和寄存器访问的数据冒险
- 如上图所示，这里仅提供 5 个函数单元，但本节我们仅关注浮点数单元
向量加载 / 存储单元（Vector Load/Store Unit）：
- 这些单元同样是完全流水线化的，这样能保证在初始时延后，每个时钟周期的带宽为一个字
- 它们也能用于处理标量的加载和存储
标量寄存器组（Scalar Registers）：
- 它们主要提供数据和计算好的地址
- 在 RV64G 中，有 31 个通用目的寄存器，以及 32 个浮点数寄存器

向量架构中的向量能够容纳不同大小的数据——假如某个向量寄存器有 32 个 64 位的元素，那么它也能容纳 128 个 16 位的元素，或者 256 个 8 位的元素。RV64V 支持的数据大小为：

整数：8, 16, 32, 64 位
浮点数：16, 32, 64 位

RV64V 的向量指令如下：

Example

其中 DAXPY 指的是形如 Y = a * X + Y 的运算，其中 X, Y 为向量，a 为标量，开头的 D 表示双精度浮点数（Double）

两种方法的主要区别在于向量处理器大大减少了指令的带宽，仅需 8 条指令就可以完成，而 RISC-V 需要 258 条指令

编译器产生向量指令序列时，代码会在向量模式中运行很多时间，这样的代码被认为是向量化的（Vectorized）。当迭代之间没有依赖时，循环代码就可以被向量化，这称为循环携带依赖（Loop-Carried Loop）

RV64G 和 RV64V 的另一个区别是流水线互锁（Interlock）发生的频率。由于存在数据依赖，RV64G 指令会有很多停顿；但在向量处理器中，每个向量指令仅需为向量的第一个元素停顿，之后的元素就会很丝滑地流过流水线，因此这样的元素依赖运算被称为链（Chaining）

Example

同上例，分别给出 RV64V 的单精度浮点数和整数版本的指令序列

单精度浮点数整数

Vector Execution Time¶

向量运算序列的执行时间取决于以下因素：

操作数向量的长度
运算间的结构冒险
数据冒险

如果我们能够知道向量长度和初始化速率（Initiation Rate）（即向量单元接收运算符并产生结果的速率）的话，我们就能计算单条向量指令的计算时间

所有现代的向量计算机都具备带有多个并行流水线（称为通道，Lane）的向量函数单元，能够在每个时钟周期中产生两个及以上的结果，但也存在一些没有完全流水线化的函数单元

我们的 RV64V 实现只有一个通道，其单条指令的初始化速率为每时钟周期一个元素，因此单条向量指令的执行时间（以时钟周期数为单位）近似为向量长度
指令组（Convoy）：一组可以一起执行的向量指令
- 要求这些指令不能包含任何结构冒险，如果存在的话需要将它们序列化，放在不同指令组中执行
- 我们还假设指令组内的指令必须在其他指令开始执行前完成执行
- 链（Chaining）的存在允许同一指令组中存在 RAW 依赖冒险，因为向量运算能够在源操作数变得可用时马上开始执行——链上的第一个函数单元的结果会被“前递”到第二个函数单元上。实际上，我们通过允许处理器在相同时间内读写特定向量寄存器来实现链的思想
- 最近的实现中用到灵活链（Flexible Chaining），它允许向量指令和其他活跃的向量指令链接起来（假设没有结构冒险）
时钟间隔（Chime）：执行一个指令组所需的时间
- 因此 \(m\) 个指令组就要执行 \(m\) 个时钟间隔；假如向量长度为 \(n\)，那么共计 \(m\times n\) 个时钟周期数
- 因为这种计算方式忽略了一些处理器特定的开销（很多取决于向量长度），以及在单时钟周期内初始化多条向量指令的局限，因此这种近似的测量法在长向量上更准确
- 此外，该方法忽略的最重要的开销是向量启动时间（Start-Up Time），即在流水线完全充满前的时延，取决于向量函数单元的流水线时延

Example

Optimizations¶

下面给出一些用于提升性能或增加一些在向量架构上能够运行的程序类型的优化方法：

多通道（Multiple Lane）：让向量处理器能够在单个时钟周期内处理向量中的多个元素
向量长度寄存器（Vector-Length Registers）：应对向量长度不等于最大向量长度的情况（大多数情况）
谓词寄存器（Predicate Registers）：高效处理条件语句，使得更多的代码能够被向量化
内存分区（Memory Banks）：为向量处理器提供足够的内存带宽
步幅（Stride）：处理多维矩阵
聚集 - 分散（Gather-Scatter）：处理稀疏矩阵
程序向量架构（Programming Vector Architecture）

Multiple Lanes¶

向量指令集的一个关键优势是能让软件仅通过一条较短的指令（包含多个独立运算，编码长度同传统的标量指令），就能向硬件传递大量的并行工作。而向量指令的并行语义使得通过高度流水线化（Pipelined）或并行（Parallel）的函数单元来执行这些元素运算的实现成为可能。下图展示的是通过多个并行函数单元提升向量加法指令性能的示意图：

RV64V 指令集的一个性质是：所有的向量算术指令仅允许向量寄存器中的 N 个元素和另一个向量寄存器中的 N 个元素进行运算，从而简化了并行向量单元的设计——这些单元被结构化为多并行通道（Lanes），通道的增加能够提升向量单元的吞吐量峰值。下图展示的是 4 通道向量单元：

Vector-Length Registers¶

一般情况下，向量的长度不会和最大向量长度（\(mv_l\)）匹配的上，并且有时我们无法在编译时知道向量的长度。比如对于下面的代码：

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

其中向量长度取决于 \(n\)，但是 \(n\) 一般要等到运行的时候才知道，为了解决这个问题，RV64V 提供的方案是增加一个向量长度寄存器 \(v_l\)，它控制任何向量运算的长度，其值不超过 \(mv_l\)。这个参数意味着向量寄存器的长度在之后增长的时候无需改动指令集

但如果 \(n\) 在编译时未知，且它的值可能会超过 \(mv_l\) 的话，那就需要用到条带挖掘（Strip Mining）技术了，它能够生成向量运算长度不超过 \(mv_l\) 的代码

具体来说，一个循环处理任何数量为 \(mv_l\) 倍数的迭代，另一个循环处理剩余的迭代，并且数量要小于 \(mv_l\)
RISC-V 提供了更好的解决方案：setvl 指令将一个小于 \(mv_l\) 的值和循环变量 \(n\) 写入到 \(v_l\)（或其他寄存器）中
- 如果循环迭代数量超过 \(n\)，那么循环最快能一次计算 \(mv_l\) 个值，因此 setvl 将 \(v_l\) 设置为 \(mv_l\)
- 若 \(n\) 小于 \(mv_l\)，则应在循环的最后这次迭代中仅对末尾的 \(n\) 个元素进行计算，故 setvl 将 \(v_l\) 设为 \(n\)
- setvl 还会写入另一个标量寄存器，以协助后续的循环管理

Example

对于任意 \(n\)，DAXPY 向量方式的 RV64V 代码如下：

Predicate Registers¶

如果程序中有包含 IF 语句的循环的话，就不能以（上述介绍过的）向量模式运行程序，因为 IF 语句引入了控制依赖。考虑以下代码：

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

该循环就不能被向量化；但如果内层循环能够在迭代为 X[i] != 0 时被运行的话，那么这个减法操作就能被向量化了

我们称这种扩展能力为向量掩码控制（Vector-Mask Control），而编译器设计者则将其称为 IF- 转换（IF-Conversion）。在 RV64V 中，谓词寄存器保存掩码，并且为每个在向量指令中的元素运算提供条件执行。这些寄存器使用布尔向量来控制向量指令的执行。当谓词寄存器 p0 被设置时，所有向量指令仅操作对应项在谓词寄存器里的值为 1 的向量元素，对应值为 0 的元素就不会发生改变。类似向量寄存器，谓词寄存器也能被启用和禁用，在启用时里面的值均被初始化为 1，意味着之后的向量指令会对所有元素进行操作

Memory Banks¶

向量加载 / 存储单元的行为相比算术函数单元会复杂得多——它们的初始化速率不一定是 1 个时钟周期，因为内存分区的停顿会减少有效的吞吐量；并且启动损失会比一般的函数单元高很多。为了维持理论上的初始化速率，内存系统必须能够产出或接收大量数据，具体做法是将访问分散在多个独立的内存分区中。

很多向量计算机支持单个时钟周期内的多次加载和存储，且内存分区周期时间通常几倍于处理器周期时间。为支持多个同步访问，内存系统需要多个分区，并能够独立控制分区地址
多数向量处理器支持不按顺序的加载或存储字数据的能力，因而需要独立分区寻址的支持，简单的内存交错无法做到这一点
多数向量计算机支持多处理器共享相同的内存系统，因此每个处理器将会产生自己单独的地址流

Stride¶

向量中相邻元素在内存中的位置不一定是按顺序的，考虑以下矩阵乘法的代码：

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

我们能够向量化 B 的每行以及 D 的每列之间的乘法运算，并将 k 作为索引变量，对内层循环进行条带挖掘。然而，为数组分配内存时，数组会被线性化，这意味着同一行或同一列的元素在内存中不一定是相邻的。以上述 C 代码为例，由于 C 语言是行主序（Row-Major Order）的，所以在迭代中被访问 D 的元素之间间隔了 800 个字节（每行 100 个元素 * 8 字节）。因此，我们需要一种能够获取不在内存中相邻的向量元素的技术

步幅（Stride）：将待收集到单个向量寄存器中的元素之间的间隔距离
- 上述例子中，矩阵 D 的步幅为 100 个双精度字，B 的步幅为 1 个双精度字；如果是列主序的话，两者正好互换
当向量被加载到向量寄存器时，这个向量看起来就像有逻辑上相邻的元素。因此向量处理器通过具备步幅能力的向量加载和存储运算，能够处理非单元步幅（Nonunit Stride）（步幅 > 1），而这种能够访问非顺序内存位置并将其重塑为紧密结构的能力正是向量架构的一大优点
由于步幅可能和向量长度一样在编译时是未知的，因此我们可以将向量步幅放在一个通用目的寄存器内，然后 RV64V 指令 vlds （Load Vector With Stride）将向量放在向量寄存器内；对于存储，也有对应的 vsts（Store Vector with Stride）指令
当多个访问在单个分区中发生竞争时，内存分区冲突就发生了，因此要停顿一个访问。当满足以下关系时，我们认为分区冲突发生：

\[ \frac{\text{Number of banks}}{\text{Least common multiple(Stride, Number of banks)}}<\text{Bank busy time} \]

Gather-Scatter¶

在稀疏矩阵中，向量元素通常以紧凑的形式被存储着，然后需要间接访问。假设有一个简化的稀疏结构，对应的代码如下：

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

这段代码实现了在数组 A 和 C 上的稀疏向量求和，用到了索引向量 K 和 M 来指定 A 和 C 中的非零元素，其中 A 和 C 必须有相同数量（n）的非零元素，K 和 M 也得大小相同。

支持稀疏矩阵的基本机制是使用索引向量的聚集 - 分散运算（Gather-Scatter Operations）。该运算的目标是支持稀疏矩阵在压缩表示法和正常表示法之间的移动

聚集（Gather）运算接收一个索引向量（Index Vector），并通过将基地址与索引向量中给出的偏移量相加，来获取对应地址处的元素所组成的向量。结果为在一个向量寄存器中的稠密向量
在完成对稠密向量中元素的操作后，稀疏向量可以通过分散（Scatter）存储，以扩展形式保存这些元素，同时使用相同的索引向量
RV64V 提供的对应指令为 vldi（Load Vector Indexed or Gather）和 vsti（Store Vector Indexed or Scatter）。将上述代码转换为 RISC-V 指令序列如下：
- 简单的向量编译器不会自动向量化上述源代码，因为编译器不清楚 K 的元素是否是唯一值（这样的话就没有依赖存在），所以需要程序员发出指示，告诉编译器以向量模式运行上述循环是安全的
- 尽管索引加载和存储能被流水线化，但它们会比没有索引的版本更慢些，因为内存分区无法在指令开始时得知，并且寄存器堆也必须提供向量单元通道之间的通信，以支持聚集和分散
- 聚集和分散中的每个元素都有独立的地址，所以不能将它们按组处理，并且在内存系统中会有多处冲突。因此即使在有高速缓存的系统中，单独的访问还是会导致显著的时延
Multimedia SIMD¶

SIMD 多媒体扩展起源于一个简单的发现：许多媒体应用程序处理的数据类型宽度，比 32 位处理器原本优化的数据类型更窄。下表总结了典型的多媒体 SIMD 指令：

类似向量指令，SIMD 指令指明相同的对向量数据的操作；不同之处在于 SIMD 倾向于指明更少的操作数，因而使用更小的寄存器堆

此外，SIMD 扩展还忽略了以下三样东西：

向量长度寄存器
步幅或聚集 / 分散数据传输指令
掩码（谓词）寄存器

一种直观的比较 SIMD 架构变体的潜在浮点性能的可视化方法是屋顶线模型（Roofline Model）。它用二维图形表示浮点数性能、内存性能以及算术强度之间的关系

算术强度（Arithmetic Intensity）是指每访问一字节内存所执行的浮点运算次数的比率，可通过计算“程序中浮点运算总次数 / 在程序执行时传输到主存中的数据字节数”得到。下图展示了不同场景下的相对算术强度：

峰值浮点性能可通过硬件规格确定
本案例研究中的许多内核无法适应芯片上的高速缓存，因此峰值内存性能由高速缓存背后的内存系统定义，可通过运行 Stream 基准测试得到
- 注意，我们需要峰值内存带宽对处理器而言也是有效的，而不是仅在 DRAM 的引脚上

下图展示了用屋顶线模型比对两个处理器的性能：

我们可以用以下公式来表述屋顶线模型中的曲线：

\[ \text{Atttainable GFLOPs/s}=\min⁡(\text{Peak Memory BW}\times\text{Arithmetic Intensity,Peak Floating−Point Perf.}) \]

考虑模型中对角线和水平线的汇聚点：

如果它在很右侧的位置上，那么只有少数具备高算术强度的内核才能达到计算机的最大性能
如果它在很左侧的位置上，那么几乎所有内核都能达到最大性能

向量处理器相比其他 SIMD 处理器而言，同时具备较高的内存带宽，以及靠左的汇聚点

GPU¶

Programming¶

英伟达（NVIDIA）开发了一种类 C 的语言和编程环境，通过应对上述挑战（异构计算，Heterogeneous Computing，和多面并行，Multifaceted Parallelism）以提升 GPU 程序员的生产力——这个系统被称为 CUDA（Compute Unified Device Architecture）。CUDA 能够生成用于系统处理器的 C/C++ 代码，以及用于 GPU（CUDA 中的 D）的 C/C++ 方言

英伟达将上述各类并行统一称为 CUDA 线程（CUDA Thread），作为表示最底层并行的编程原语，这使得编译器和硬件能够将数以千计的 CUDA 线程放在一起，以利用 GPU 内的各种并行。因此英伟达将 CUDA 编程模型归类为单指令，多线程（SIMT）
而这些被分进一个个块内一起执行的一组线程称为线程块（Thread Block）
在这些线程块上执行运算的处理器被称为多线程 SIMD 处理器（Multithreaded SIMD Processor）
为区分来自 GPU 和系统处理器的函数，CUDA 使用 __device__ 或 __global__ 表示前者，__host__ 表示后者
以 __device__ 声明的 CUDA 变量被分配到 GPU 的内存，这样的内存可被所有多线程 SIMD 处理器访问到
运行在 GPU 上的函数 name 的扩展函数调用语法为：name <<<dimGrid.dimBlock>>>( ...parameter list... )，其中 dimGrid 和 dimBlock 分别指明代码（线程块）和块（线程）上的维度
除了块的标识符（blockIdx）和块内每个线程的标识符（threadIdx）外，CUDA 为每个块的线程个数提供了关键字 blockDim

Example

C 代码CUDA 代码

// Invoke DAXPY
daxpy(n, 2,0, x, y);
// DAXPY in C
void daxpy(int n, double a, double *x, double *y) {
    for (int i = 0; i < n; ++i)
        y[i] = a * x[i] + y[i];
}

// Invoke DAXPY with 256 threads per Thread Block
__host__
int nblocks = (n + 255) / 256;
daxpy<<<nblocks, 256>>>(n, 2,0, x, y);
// DAXPY in CUDA
__global__
void daxpy(int n, double a, double *x, double *y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n)
        y[i] += a * x[i];
}

由 GPU 硬件负责处理并行执行和线程管理，而非由应用程序或操作系统。为简化硬件调度，CUDA 要求线程块能够独立地并以任意顺序执行。不同的线程块之间不得直接通信，尽管它们能够通过全局内存的原子内存操作来协调

Computational Structures¶

网格（Grid）是一种在 GPU 上，由一组线程块构成的代码（对应术语为向量化循环，Vectorized Loop）。网格和线程块是在 GPU 上实现的编程抽象，帮助程序员组织 CUDA 代码。下图展示了网格和线程块之间的关系：
下图展示了一个简化的多线程 SIMD 处理器框图：

可以看到，它和向量处理器很像，但它有多个并行的函数单元，而不是少数高度流水线化的函数单元 - 线程块调度器（Thread Block Scheduler）将线程块分配给多线程 SIMD 处理器上 - 为了在具有不同多线程 SIMD 处理器数量的 GPU 模型之间提供透明的可扩展性，线程块调度器将线程块分配给多线程 SIMD 处理器 - 综上，GPU 本质上是一个由多线程 SIMD 处理器构成的多处理器 - 深入一层细节来看，硬件所创建、管理、调度并执行的核心对象是一条 SIMD 指令线程（Thread of SIMD Instructions）。这些指令线程有自己的 PC，并在一个多线程 SIMD 上运行（所以这些线程是相互独立的） - SIMD 线程调度器（SIMD Thread Scheduler）知道哪个 SIMD 指令线程已经准备好运行，并将这样的线程发送给一个即将在多线程 SIMD 处理器上运行的分派单元，不必按顺序获取线程中的下一条指令 - 调度器包含了一个记分板（Scoreboard），用于追踪至多 64 个 SIMD 指令线程，以观察哪些线程是准备好的。由于高速缓存和 TLB 的命中和失效，内存指令的时延是可变的，因此要求记分板确定这些指令何时完成

NVIDIA GPU ISA¶

不同于其他系统处理器，英伟达编译器的指令集目标是一种对硬件指令集（对程序员而言是不可见的）的抽象——PTX（Parallel Thread Execution，并行线程执行）提供了用于编译器的一个稳定的指令集，且在不同代 GPU 上具备兼容性。

PTX 指令描述了在单个 CUDA 线程上的运算，并且通常和硬件指令进行一对一映射；但一条 PTX 指令可扩展至多个机器指令，反之亦然。
PTX 采用无限数量的写一次寄存器
- 编译器必须运行一个寄存器分配过程，将 PTX 寄存器映射到一个固定数量的，在真实设备中可用的读写硬件寄存器。
- 优化器随后运行，以减少寄存器的使用，从而消除无效代码，将指令合并折叠，并计算分支可能分叉的位置以及分叉路径可能重新汇合的地点。
PTX 指令格式为：opcode.type d, a, b, c;，其中 d 是目标操作数（除了在存储指令外都是寄存器），a, b, c都是源操作数（32 位 / 64 位寄存器或者常数值），而运算类型如下所示：

PTX 指令集

所有指令均可通过 1 位谓词寄存器进行条件执行，这些寄存器可通过设置谓词指令（setp）来设定
控制流指令为：
- 函数：call, return
- 线程：exit, branch
- 线程块内线程的屏障同步（Barrier Synchronization）：bar.sync
编译器或 PTX 程序员将虚拟寄存器声明为 32 位或 64 位类型或无类型值，比如 R0 , R1 , ... 是 32 位值，而 RD0 , RD1 , ... 是 64 位值

Example

我们用 PTX 指令来实现 DAXPY：

Conditional Branch¶

在 IF 语句的处理上，GPU 相比向量架构会更依赖于硬件支持——除了谓词寄存器外，还会用到内部掩码、分支同步栈以及指令标记器，来管理分支发散成多条执行通路和通路汇集的时间

在 PTX 汇编器级别中，一个 CUDA 线程的控制流由以下内容描述：PTX 指令的分支、调用、返回和退出，以及每个指令的线程通道断言（由程序员使用线程通道的 1 位谓词寄存器指定）。PTX 汇编器分析 PTX 分支图，并将其优化为最快的 GPU 硬件指令序列。这些指令可以在分支上自己做决定，无需被锁住

在 GPU 硬件指令级别中，控制流包括了分支、跳转、索引跳转、调用、索引调用、返回、退出，以及管理分支同步栈的特殊指令。GPU 为每个 SIMD 线程提供一个栈，一个栈元素包含一个标识令牌（Identifier Token），一个目标指令地址以及一个目标线程活跃掩码（Thread-Active Mask）。还有一些 GPU 特殊指令，用于为 SIMD 线程压入栈元素；还有一些特殊指令及指令标记，能够弹出栈元素或将栈回退至指定元素，并依据目标线程活跃掩码跳转至目标指令地址。GPU 硬件指令还有一个每条通道独立的谓词（启用 / 禁用）功能，通过为每个通道分配 1 位谓词寄存器来实现

GPU Memory¶

下图展示了英伟达 GPU 的内存结构：

在多线程 SIMD 处理器上，每个 SIMD 通道被给予一块不在芯片上的 DRAM 的私有区域，称为私有内存（Private Memory）
- 用于存放栈帧（Stack Frame）、溢出寄存器和无法被寄存器容纳的私有变量
- SIMD 通道之间不会共享私有内存
- GPU 将私有内存缓存在 L1 和 L2 高速缓存中，以辅助寄存器溢出和加速函数调用
而在芯片上的，位于每个多线程 SIMD 处理器中的内存称为局部内存（Local Memory）
- 该内存具有低时延，高带宽的特征，程序员可以拿它来存储需要被同一线程或相同线程块内的不同线程复用的数据
- 它的容量不大，一般只有 48 KB
- 不会保存线程块的状态
- 可被多线程 SIMD 处理器的多个 SIMD 通道共享，但无法在多个多线程 SIMD 处理器间共享
- 当多线程 SIMD 处理器创建线程块时，会为这些块动态分配局部内存，并在所有线程块内的线程退出时释放内存
GPU 内存（GPU Memory）：在芯片外的，被整个 GPU 和所有线程块共享的内存
- 系统处理器（称为主机）能够对 GPU 内存进行读写操作，而局部内存和私有内存对主机而言是不可用的

Compare GPU with Vector¶

GPU 和向量架构中对应的术语如下图所示：

下图为 4 通道的向量处理器（左）和 GPU 上 4 SIMD 通道的多线程 SIMD 处理器（右）：

Compare GPU with Multimedia SIMD¶

下表比对了多媒体 SIMD 扩展和 GPU 之间的异同：

Chapter 04 : Data-Level Parallelism Architecture¶

Vector Architecture¶

RV64V¶

Vector Execution Time¶

Optimizations¶

Multiple Lanes¶

Vector-Length Registers¶

Predicate Registers¶

Memory Banks¶

Stride¶

Gather-Scatter¶

Multimedia SIMD¶

GPU¶

Programming¶

Computational Structures¶

NVIDIA GPU ISA¶

Conditional Branch¶

GPU Memory¶

Compare GPU with Vector¶

Compare GPU with Multimedia SIMD¶

评论