ysh329 / OpenCL-101

Learn OpenCL step by step.
123 stars 31 forks source link

SIMD 指令集与数据并行程序 #65

Open ysh329 opened 2 weeks ago

ysh329 commented 2 weeks ago

本文内容来自《Whole-Function Vectorization》的 Introduction 章节的 SIMD 指令集与数据并行程序小节,为作者的主要工作做知识铺垫,本文不涉及作者的主要工作,仅做 SIMD 概念的基本理解与学习。本文目录:

  1. Amdahl's Law 和 Intel MMX
  2. SIMD、数据并行、向量处理器的关系 2.1 向量处理器的优缺点和限制
  3. SIMD 3.1 内存访问 3.2 指针类型于操作 3.3 掩码类型 3.4 副作用
  4. 数据并行程序 4.1 程序表示 由于原文内容较为薄弱,学习过程中,补充了一些内容,不少来自 ece740 《Computer Architecture: SIMD/Vector/GPU》这一节的幻灯片(见文末参考)。

1.Amdahl's Law 和 Intel MMX

理论是推动技术前进的基石,根据阿姆达尔定律,该定律是由计算机工程师吉恩·阿姆达尔提出的,用来描述在多处理器系统中,程序性能提升的理论上限。下图左是该定律公式: image

图 Amdahl 定律:程序中不能并行化的部分将决定整个程序性能提升的最大可能值

需要说明的是:

这个定律指出,程序的加速比受到其串行部分的限制。换句话说,程序中不能并行化的部分将决定整个程序性能提升的最大可能值。 实际需求推动技术向前,随着多媒体和通信软件的发展,对性能要求变高,1996 年, Intel 公司发表了一篇名为《MMX Technology Extension to the Intel Architecture》的文章,想要解决现有的 Intel 架构扩展以支持更高效的多媒体处理,同时保持与现有操作系统和应用程序的兼容性。通过设计一种基于 SIMD 名为 MMX 的技术,达到对现有应用程序的增强。

image

图 MMX 指令集相关指令和描述

该方法也可以在 C 语言中访问 MMX 指令。让加速多媒体和通信以及其他数值密集型应用程序性能成为可能。

2.SIMD、数据并行、向量处理器的关系

在讲 SIMD(单指令多数据)、数据并行之前,有必要讲一下向量处理器,这三个概念是密切相关的,它们都涉及在计算过程中同时处理多个数据点的能力。简单来说:

要说三者的关系,可以这么说:

总结来说,数据并行是一种概念,SIMD是一种实现数据并行的执行模型,而向量处理器则是具备执行SIMD指令能力的硬件。这三者共同为现代计算提供了强大的并行处理能力。

image

图 SIMD 处理(左),向量寄存器(右)

2.1 向量处理器的优缺点与限制

先提一下硬件的 Vector Processor。这里先给出 Vector Processor 优缺点的对比。

image

图 向量处理器 Pros & Cons

优势有:

劣势也很明显:

image 图 向量处理器本身的限制

此外,也存在一些限制:

ysh329 commented 2 weeks ago

3.单指令多数据指令集(SIMD Instruction Sets)

在本节中,我们将讨论 SIMD 指令集设计对数据并行语言实现的影响。

image 图 标量 vs. SIMD 操作

通常, SIMD 单元具有一定位宽的 SIMD 寄存器(例如 SSE 为128 bits 长度),如下图所示,左边是向量寄存器的描述。

image image 图 向量寄存器(左),向量功能单元(右)

对于 Intel AVX,其每个寄存器是 256 bit,可以保存如 8 个 32 bit 的 float32 数或者 int32 数组成的向量。

image

图 Intel® Advanced Vector Extensions Registers 增加了16个寄存器(YMM0-YMM15),每个256位宽,混叠到 16 个 SIMD (XMM0-XMM15)寄存器上。AVX 新指令在 YMM 寄存器上操作。AVX 扩展了某些现有指令在 YMM 寄存器上的操作,定义了一种新的编码方式,在单个指令中多达三个源和一个目的

不同的指令集提供不同的数据类型。例如,SSE 全面支持保存 4 个浮点 float32 数、4个整型 int32 数或 2 个双精度 float64 数。有时,还支持其它数据类型(较小的 int 型,如 8 个 16bit 的 short / ushort,或更小的如 16 个 8 bit 的 char / uchar)。

image 图 基于 SIMD 的求和计算逻辑示意图

为了便于表示,我们只考虑固定大小相等(如 4 个字节 byte)的数据类型。在整篇文章中,我们将通过 SIMD 宽度 W 来表示适合单个寄存器的该类型元素的数量。

3.1 内存访问操作(Memory Access)

通常,通过从(到)对齐的地址连续读(或写)元素来加载(或存储)vector 值。有些指令集还允许访问未对齐的地址。

image 图 对齐的数据读取相比不对齐的更快 | SIMD for C++ Developer

image

图 对地址不规则内存访问

现在不少的新架构还支持分散(scatter)或收集(gather)操作,可以使用(可能是非连续的)偏移量向量从(或到)基址加载(存储)。

image 图 Gather/Scatter 操作

在没有分散 / 收集指令的架构上,这种非连续访问非常慢。这是因为 vector 元素必须一个接一个地加载到不同的寄存器中,然后增量地移动到目标寄存器中。高效连续访问可以看作是索引向量 {0, 1, …, W - 1} 的一种特殊情况。因此,在我们的矢量化过程中,加载和存储中使用的索引的连续性和对齐是很重要的。

3.2 指针类型与操作(Pointers)

一般来说,将指针存储在 SIMD 寄存器中是不可移植的。在 SSE 的情况下,可以将 4 个指针以32位模式放入 SIMD 寄存器中。如果处理器以 64 位模式执行,则只能将 2 个指针放入 SSE 寄存器中。因此,在 SIMD 寄存器中使用指针的指令并不常见。

image 图 Google/Highway 项目

在检索相关内容的时候,发现有一个名为 google/Highway 的 C++ 库,用于 SIMD。这个项目的目的就是为了解决特定平台对 SIMD 代码的移植性的问题。

3.3 掩码类型(Masking)

向量化将控制流转换为数据流。修改后,可能会执行在使用控制流的程序的标量版本中不会执行的代码。在矢量化程序中,控制条件被明确地表示为掩模向量 mask 。

image

图 控制流到数据流的转换

控制流到数据流转换的 select 函数使用该掩码在以前的控制流连接处混合两个值,如下图中间 f‘ 用到了 select(mask, s, t) 方法:

image

图 带控制流的标量函数 f 、数据流的函数 f' 、使用 SSE 的实现的函数 f_sse

从中间到右边的SSE代码,由 _mm_cmpgt_ps 产生掩码替换控制流,变为数据流的 mask:

image

用于某些程序的执行。通过执行:

image

最终的结果通过 _mm_blendv_ps(mask, s, t) 实现对两个分支结果的选择,变量 r 包含每个实例的正确值。

image

图 Larrabee 架构的向量处理单元(VPU),其特点:①支持 3 操作数指令、②允许对寄存器的输入数据重排序、③支持对内存中读取的数据做复制和转换、④掩码寄存器实现对向量操作的条件执行

一些架构如 Larrabee (这是一个专为视觉计算设计的多核 x86 架构,采用顺序执行的 CPU 核心),有一个专用的掩码寄存器文件来存储这些掩码。此外,每个指令都可以被赋予一个掩码寄存器,该掩码寄存器控制目标寄存器中的哪些向量元素将受该指令的影响。这使得使用选择指令的显式混合变得不必要(注:这句话就是说,有些指令就本身带 mask ,就不需要在对向量单独做 select/blend 指令了)。

3.4 副作用

对于可能产生副作用的操作,事后使用 mask 来屏蔽结果是不够的:必须防止它们在非活动实例(即数据)中执行。这是通过将操作拆分为由 if 语句保护的标量顺序执行来实现的。这段话关于副作用,我是这么理解的:

这两种特殊情况,可能就需要考虑分割 SIMD 向量操作提前做条件语句保护,或者分割向量操作为标量操作,规避副作用。

ysh329 commented 2 weeks ago

4.数据并行程序(Data-Prallel Programs)

数据并行性是一种并行计算的形式,其中相同的操作被同时应用于多个数据项。这种并行性是通过使用 SIMD(单指令多数据)架构实现的,它允许一个指令同时对多个数据元素进行操作。例如,在计算两个向量的点积时,可以同时对多个成对的元素执行乘法和加法操作。

image 图 数据并行的特点

我们的数据并行程序模型接近 CUDA 和 OpenCL 的模型。

image 图 OpenCL 平台模型展示了 OpenCL架构组成:这些粗粒度计算设备中的每一个都由多个“计算单元”(类似于多核 CPU 的执行单元和算术处理单元组)组成。在最底层,这些处理元素都执行OpenCL“内核”

计算单元的具体定义因硬件供应商的不同而不同。在 AMD 硬件中,每个计算单元包含许多“流核”(有时称为 SIMD 引擎),然后包含单独的处理元素。每个流内核都执行 VLIW 4或 5 宽度的 SIMD 指令。有关 ATI 硬件的概述,请参见下图。

image 图 ATI 计算设备的简化结构

在 NVIDIA 硬件中,他们称计算单元为 stream multiprocessors(SM),在一些文档中称为 CUDA 核心。反正无论是是 AMD 亦或是 NVIDIA,执行最底层 SIMD VLIW 指令的都是具有很复杂的硬件层次结构。

image

图 OpenCL 编程模型,特点是执行模型支持数据并行和任务并行

下面来看一个具体的例子。在我们的设置中,数据并行程序由函数 f 给出如下图。执行数据并行程序 f 执行 N 个 f 的实例,这些实例的时间顺序未指定。

image

图 带控制流的标量函数 f 、数据流的标量函数 f' 、使用 SSE 的数据流向量函数 f_sse

此外,它们中的一些可以并行运行。f 的每个实例(即数据)都以实例 ID 如 tid = get_local_id(0) 作为参数获取实例的编号。对于 f 的每个实例,tid 的值是两两不同的,范围从 0 到整个任务需要处理的总元素个数 N。

一个已经利用了一些并行性的直接实现将实例 ID 范围细分为 T 个大小相等的部分,并将它们分配给 T 个线程(内核)。在一个线程内,N / T 个 f 实例依次运行。在本文中,我们还想利用向量指令来开发核内并行性,即图上的函数 f' 和 f_sse 。

image

图 OpenCL 基于队列和事件实现同步。CPU + GPU 情况下的 Kernel 1 和 Kernel 2,左图:Kernel 2在 Kernel 1 还没得到计算结果的时候就开始执行。右图:Kernel 2 等待对应 Kernel 1 的事件结束才开始

在本文中,除了使用上线程特性,还想利用向量指令来开发核内并行性。因此,我们将把 f 转换为同时执行 W 个实例的函数 f',其中 W 是处理器 SIMD 寄存器的宽度:对整个函数 f 进行矢量化。然后,多线程实现按顺序对每个线程应用 f 函数 ( N / W ) / T 次的代码如下:

image

图 调度应用前面的 f 函数的函数的实现

请注意,不同内核之间的多线程与 SIMD 向量化是正交的。换句话说,多线程执行(由多个执行线程组成的工作组并行执行内核函数)和SIMD向量化(利用SIMD指令同时处理多个数据元素)这两种并行执行方式可以独立使用,它们之间没有直接的依赖关系,可以组合在一起以提高性能。

4.1 程序表示(Program Representation)

注:这里讲的是数据并行的程序表示,或者说写法或者在编译器中间的一种表达形式。这里可能和本文没有直接关系,可以忽略。

我们认为函数 f 以一种类型的低级表示形式给出。函数表示为指令的控制流图。此外,我们要求 f 是 SSA 形式的,即每个变量都有一个静态赋值,变量的每次使用都受其定义的支配。

image

图 从 OpenCL 源码经由 Runtime 从前端到 LLVM IR 到最终的 CPU 、GPU 上的表示

这种程序表示的一个突出的例子是 LLVM bit code 如下:

image

文章作者在研究向量化问题的时候,将类型和指令限制在如下图所示的语言子集中,这也是前文标量程序中所用到的。而其它指令,如跳转或算术和比较操作符都是直接的,可以省略了。

image

图 前文中 f 函数涉及到的类型和指令

这种程序表示很好地反映了当今指令集架构对 SIMD 支持的共识。

ysh329 commented 2 weeks ago

参考