【论文解读】Whole-Function Vectorization

[ysh329]

project: https://compilers.cs.uni-saarland.de/projects/wfv/ papar: https://compilers.cs.uni-saarland.de/papers/karrenberg_wfv.pdf slide: https://compilers.cs.uni-saarland.de/projects/wfv/wfv_cgo11_slides.pdf

---

作者：要术甲杰 链接：https://www.zhihu.com/question/65708935/answer/475572657 来源：知乎 著作权归作者所有。商业转载请联系作者获得授权，非商业转载请注明出处。

从vectorization的发掘角度来讲，可以是面向循环的，也可以是面向基本块的，也可以是straight-line的，这给vectorization的技术发展创造了很多研究课题。例如对内层循环的向量化“Compiling for vector-thread architectures”（CGO 2008）和对外层循环的向量化“Outer-Loop vectorization—revisited for short SIMD architectures”（PACT 2008）。像SLP技术，你可以把它用于基本块内的向量挖掘。另外还有针对函数级的向量化，如“Whole-function vectorization”（CGO 2011）。另外还有向部分向量化“A Compiler Approach for Exploiting Partial SIMD Parallelism”（TACO 2016）、混合向量化“ Exploiting mixed SIMD parallelism by reducing data reorganization overhead”（CGO 2016）等针对不同对象的向量化发掘方法。

[ysh329]

关键词：SIMD、Vectorization、Data Parallelism、Code Generation、OpenCL。

全函数向量化

数据并行编程语言是当今并行计算领域的一个重要组成部分。其中包括特定于领域的语言，如图形中的着色语言（ HLSL 、 GLSL 、 RenderMan 等）和“通用"语言，如 CUDA 或 0penCL 。这些语言在 CPU 上的当前实现仅依赖于多线程来实现并行性，而忽略了这些处理器的 SIMD 指令集（如英特尔的 SSE 和即将推出的 AVX 或 Larrabee 指令集）提供的额外内核内并行性。

在本文中，我们讨论了在具有 SIMD 指令集的机器上实现数据并行语言的几个方面。我们的主要贡献是一种独立于语言和平台的代码转换，它对 SSA 形式的控制流图给出的低级中间代码执行全函数关量化。

在两种情况下评估我们的技术：首先，将其纳入实时光线跟踪中使用的特定领域语言的编译器中。其次，在一个独立的0penCL驱动程序中。我们观察到射线示踪剂的平均加速因子为3.9，不同0penCL内核的平均加速系数在0.6和5.2之间。

[ysh329]

1. 介绍

数据并行编程是当今并行计算领域的一个重要概念，在过去的几十年里，已经创建有各种编程语言来支持数据并行编程。这些语言包括领域特定语言比方图形中的着色语言，例如 HLSL、GLSL、RenderMan 等。或者 Nvidia GPU 通用计算语言 CUDA 或 OpenCL 。

这些语言都具有 SIMD 的概念，即一段代码在存放数据的数组并行执行某些操作，需要注意的是这并不意味着这段代码每个实例都执行完全相同的指令，如下面图中的示例，左边的是控制流图，给出的代码并行应用于 4 个不同的输入，这导致了 4 种不同的执行轨迹。如图所示，例如实例 1 执行的路径是a-b-c-e-f，实例 3 执行的路径是a-b-c-e-b-c-e-f。

分散的控制流使得 SIMD 的执行变得复杂：

• 实例 1 在分叉的 b 之后执行块 c 。
• 实例 3 则在 b 之后执行块 d 。

假设块 c 和 d 包含不同的指令，则它们不能以 SIMD 方式执行。通常的解决方案是 SIMD 程序执行两个分支并且补偿不必要的影响：

• 实例 1 忽略块 d 的计算。
• 实例 2 忽略块 c 的影响。

以此类推，这也被称为预测执行（Predicated Execution）（属于modulo sheduling技术，用于并行化带条件语句的最内层循环。处理器也有称为预测执行-Speculative Execution 的概念，编译器和处理器都可以实现或利用预测执行技术来提高程序的执行效率）。

实际中通过不同的方式来实现这个预测执行的行为。本文讲的是 GPU 的 Predicated Execution，有这么几点：

• 硬件层面隐式实现这种预测执行，每个数据并行程序的实例被映射到多核处理器的标量核上。
• 所有实例以锁步的方式来执行。lock-step，是一种常用在 GPU 上用于提高效率和简化编程模型的一种执行模型，表示并强调多个处理单元如线程或核心同步地执行相同地指令，在并行任务时特别有效，局限性是控制流路径过多的情况。
• 中央控制单元（central control unit）将处理器停在实例处于非活动状态的代码区域上。概念补充，这里的中央控制单元在 GPU 的上下文中，可以理解为调度器或者控制器，负责管理和调度 GPU 上的执行单元（如核心或线程块）的执行。总得来说，在数据并行程序的执行过程中，中央控制单元可以根据执行单元实例，是否活跃来决定是否暂停其执行，从而优化 GPU 的性能和资源利用率。

[ysh329]

本文研究数据并行语言在需要_显式矢量化_的处理器上的实现。这样的处理器特点是具有矢量寄存器文件和专用 SIMD 指令（如 Intel 的 SSE 和 IBM AltiVec 指令集）。

• 矢量寄存器文件，主要用于支持并行数据处理的处理器上，可以在单个指令周期内并行处理多个数据元素，通常支持 SIMD 或类似的执行模型，允许同时对多个数据点执行操作，如一个 256 位的矢量寄存器可以存储 8 个 32 位的浮点数。
• 矢量指令集，用于向量处理器或阵列处理器作为一种中央处理单元的指令集，其指令被设计用于在大型一维数组上做高效执行。与标量处理器形成对比，后者指令只操作单个数据项。

要在存在发散的控制流情况下使用 SIMD 指令，通常程序员或者编译器必须使用数据流代替控制流。

图 《Conversion of Control Dependence to Data Dependence》

对于没有循环的代码，这通常称为 if 转换，来自《Conversion of Control Dependence to Data Dependence》这篇文章，这篇文章探讨了如何将程序中的控制依赖（Control Dependence）转换为数据依赖（Data Dependence），通过消除 goto 语句并引入逻辑变量来控制程序中语句的执行，从而将控制依赖转换为数据依赖。

图 程序员如何使用 Intel SSE 指令集的函数来实现对同一个代码的矢量化变体

例如，上面图中的函数f，以及其没有控制流f'的版本，那么函数通过 select(mask, t, s) 指令根据掩码值选择 t 或者 s，其中，t 和 s 分别表示当条件为真或假时选择的值或者表达式。select 函数是一个假设的或者概念函数，用于根据布尔向量 mask 值来选择 s 或 t 。

后文将会讲到几种将控制流转换为数据流以执行矢量化的方法。这些方法往往源于并行计算社区，这种并行化是在编译过程的早期执行的（这里的“早期”是相较于其他编译优化和变化的步骤而言，编译过程通常包含：预处理、词法分析、语法分析、语义分析、中间表示生成、优化、代码生成等步骤，“早期”执行意味着在这些步骤中的优化阶段之前或之中进行），通常已经在源代码层级上，但这有两个缺点：

1. 过早地将控制流转换为数据流。编译器中使用控制流信息的所有分析和优化都将变得毫无用处，这些分析比方条件常数传播或循环不变代码运动，都需要重写，因为需要适配矢量代码。但是你早早将控制流转为数据流，分析不到了。
2. 现代编译器的基础设施会使用虚拟指令集作为代码交换格式，如 LLVM 的 bit-code 或者 Nvidia 的 PTX，目的是将前端与代码生成器解耦。这些表示通常使用指令的控制流图来表示代码。而从前端（如源码级别）将控制流转换为数据流，会破坏生成代码的可移植性，原因有二：首先，要在不需要控制流->数据流转换的架构上使用这个转换流程，这个转换需要撤销；其次，即使目标架构需要控制流到数据流的转换，前端也需要考虑诸如向量宽度、掩码实现细节等架构参数。因此，代码不再具有可移植性。

图 本文的目标

[ysh329]

贡献

我们认为，矢量化应该在编译过程的后期进行（注：矢量化属于性能优化的一种且与架构支持相关，编译的早期阶段包含控制流信息，并不适于过早做向量化，而且过早也会带来冗余计算，不利于移植性），我们提出了一种代码转换，对当今常见的中间表示（如 LLVM bit-code ）中表示的代码执行控制流到数据流的转换，实现是通过将矢量化视为静态单赋值（SSA）形式的控制流图的程序转换，总结为如下四点贡献：

• 针对 SIMD 指令处理器，提出了 SSA 形式控制流图的全函数矢量化转换。SSA 对处理器的矢量化特别有用，因为在 SSA 形式中，φ 函数（phi functions）非常有用，这些 φ 函数指示了在向量化过程中需要插入选择（“blend”，原文这里用“blend”应该是想表达np.where(mask, true_value, false_value)的“混合”不同分支的意思）代码的位置，以便在不同的执行路径中选择正确的数据。
  • 这里补充 φ 函数作用的知识点：①合并变量值，在程序存在多条执行路径的情况下如 if-else 结构，不同路径对同一个变量会进行不同赋值，在 if-else 后的 φ 函数将会合并这些不同路径下的值，方式是通过创建一个新的变量，其值取决于不同路径上的赋值；②维护 SSA 形式：SSA 要求每个变量在程序中只被赋值一次， φ 函数帮助对存在复杂控制流的场景仍然遵守这一形式；③简化数据流分析， φ 函数提供了一种清晰的方式来表示和分析变量在不同路径上的只，简化编译器的数据流分析和优化过程。
• 提出一个针对任意控制流生成预测代码（predicated code，我理解这里应该是说向量布尔 mask）的 Pass，即使这个硬件不支持预测执行（predicated execution，注：我理解这里应该是说 bool mask），这个 pass 也会谨慎地进行混合操作。
• 提出一种数据流分析（analysis），可以确定在不同并行实例中变量值的某些约束限制。这个分析有助于遵从当前硬件 SIMD 体系结构的内存访问的指令对齐和连续性的限制，如确保数据在内存中的正确对齐，以便 SIMD 指令可以高效地访问。此外，如果某个计算在每个实例中都产生相同的值，那么这个计算就不需要向量化，可以保持为标量（scalar）形式，从而避免不必要的向量化开销。
• 在 LLVM 编译器中实现，并在两个场景做了性能评估。第一个场景是在 RenderMan 着色语言的编译器中添加了该算法， RenderMan 是一种在电影行业广泛用于视觉效果的计算机图形 DSL（领域特定语言），作者使用实时光线追踪器进行了实验。第二个场景是实现了一个 OpenCL 程序，该程序使用所介绍的向量化算法，并展示了在不同 OpenCL 应用程序中的加速效果，速度提升在 0.6 到 5.2 倍之间，而实时光线追踪器的平均加速为3.9倍。

本文结构

• 第 2 节内容：讨论与本文相关的当前 SIMD （单指令多数据）架构的设计特点。这可能包括 SIMD 指令集的特性、寄存器结构、数据对齐要求、向量宽度等，这些特性对于实现和优化向量化处理至关重要。
• 第 3 节内容：简要概述了他们考虑的数据并行程序。这可能涉及程序的类型、结构、特性或它们在特定领域（如科学计算、图形渲染等）的应用。
• 第 4 节内容：展示了本文的核心贡献，即面向 SSA （静态单赋值）形式程序的全函数向量化技术，可以将整个程序函数转换成适合 SIMD 执行的形式，提高程序的执行效率。
• 第 5 节内容：讨论了相关工作，这通常包括对先前研究的回顾，作者将比较自己的工作与现有技术或解决方案的异同，展示自己工作的独特之处或改进。
• 第 6 节内容：介绍作者的实验评估。这一部分通过实验结果来验证所提出方法的有效性，可能包括性能测试、基准测试、与现有技术的比较等。

[ysh329]

2. SIMD 指令集

见 https://github.com/ysh329/OpenCL-101/issues/65

[ysh329]

3. 数据并行程序

见 https://github.com/ysh329/OpenCL-101/issues/65

[ysh329]

贡献

我们认为，矢量化应该在编译过程的后期进行（注：矢量化属于优化的一种且与架构支持相关，编译的早期阶段包含控制流信息，并不适于过早做向量化，而且过早也会带来冗余计算，不利于移植性），我们提出了一种代码转换，对当今常见的中间表示（如 LLVM bit-code ）中表示的代码执行控制流到数据流的转换，实现是通过将矢量化视为静态单赋值（SSA）形式的控制流图的程序转换，总之贡献如下：

• 针对 SIMD 指令处理器，提出了 SSA 形式控制流图的全函数矢量化转换。SSA 对处理器的矢量化特别有用，因为在SSA形式中，φ 函数（phi functions）非常有用，这些 φ 函数指示了在向量化过程中需要插入选择（“blend”，原文这里用“blend”应该是想表达np.where(mask, true_value, false_value)的“混合”不同分支的意思）代码的位置，以便在不同的执行路径中选择正确的数据。这里补充 φ 函数作用的知识点：①合并变量值，在程序存在多条执行路径的情况下如 if-else 结构，不同路径对同一个变量会进行不同赋值，在 if-else 后的 φ 函数将会合并这些不同路径下的值，方式是通过创建一个新的变量，其值取决于不同路径上的赋值；②维护 SSA 形式：SSA 要求每个变量在程序中只被赋值一次， φ 函数帮助对存在复杂控制流的场景仍然遵守这一形式；③简化数据流分析， φ 函数提供了一种清晰的方式来表示和分析变量在不同路径上的只，简化编译器的数据流分析和优化过程。
• 提出一个针对任意控制流生成预测代码（predicated code，我理解这里应该是说向量布尔 mask）的 Pass，即使这个硬件不支持预测执行（predicated execution，注：我理解这里应该是说 bool mask），这个 pass 也会谨慎地进行混合操作。
• 提出一种数据流分析（analysis），可以确定在不同并行实例中变量值的某些约束限制。这个分析有助于遵从当前硬件 SIMD 体系结构的内存访问的指令对齐和连续性的限制，如确保数据在内存中的正确对齐，以便 SIMD 指令可以高效地访问。此外，如果某个计算在每个实例中都产生相同的值，那么这个计算就不需要向量化，可以保持为标量（scalar）形式，从而避免不必要的向量化开销。
• 在 LLVM 编译器中实现，并在两个场景做了性能评估。第一个场景是在 RenderMan 着色语言的编译器中添加了该算法， RenderMan 是一种在电影行业广泛用于视觉效果的计算机图形 DSL（领域特定语言），作者使用实时光线追踪器进行了实验。第二个场景是实现了一个 OpenCL 程序，该程序使用所介绍的向量化算法，并展示了在不同 OpenCL 应用程序中的加速效果，速度提升在 0.6 到 5.2 倍之间，而实时光线追踪器的平均加速为3.9倍。

本文结构

• 第 2 节内容：讨论与本文相关的当前 SIMD （单指令多数据）架构的设计特点。这可能包括 SIMD 指令集的特性、寄存器结构、数据对齐要求、向量宽度等，这些特性对于实现和优化向量化处理至关重要。
• 第 3 节内容：简要概述了他们考虑的数据并行程序。这可能涉及程序的类型、结构、特性或它们在特定领域（如科学计算、图形渲染等）的应用。
• 第 4 节内容：展示了本文的核心贡献，即面向 SSA （静态单赋值）形式程序的全函数向量化技术，可以将整个程序函数转换成适合 SIMD 执行的形式，提高程序的执行效率。
• 第 5 节内容：讨论了相关工作，这通常包括对先前研究的回顾，作者将比较自己的工作与现有技术或解决方案的异同，展示自己工作的独特之处或改进。
• 第 6 节内容：介绍作者的实验评估。这一部分通过实验结果来验证所提出方法的有效性，可能包括性能测试、基准测试、与现有技术的比较等。

[ysh329]

4. Whole-Function Vectorization

Whole-function vectorization 算法可简写为 vectorizer，包含六个阶段：

1. 变换准备
2. 向量化分析
3. 掩码 mask 生成
4. 选择 select 生成
5. 控制流图 CFG 线性化
6. 指令向量化

4.1 变换准备

在进行向量化之前，会执行一些准备性的转换。最值得注意的是，循环会被简化，以确保每个循环恰好有一个进入边和一个回边。这保证了存在一个唯一的循环头部（循环开始的地方），一个唯一的循环前头部（进入循环的区块）以及一个唯一的循环锁存器（一个边从这里回到头部的区块）。第 4.6 节描述了算法如何在具有不可约控制流的图上工作。

这里的“不可约控制流”指的是在程序的控制流图中，存在无法通过任何序列的转换来简化的循环或条件结构。在处理这类控制流时，编译器需要采用特定的算法来确保向量化的转换能够正确进行。

[ysh329]

4.2 向量化分析

大多数指令可以通过简单地与它们的向量对应物交换来向量化（例如，加法运算被转换成向量加法）。然而，当前的 SIMD （单指令多数据）架构和我们的语言（见第 3.1 节）具有接受地址向量的地址算术（gep，Get Element Pointer，获取元素指针，是一种在编译器中用于计算数组或结构体元素地址的操作）和内存访问（加载/存储）指令。一种简单的方法是将包含偏移的向量分开，复制标量代码 W 次，并将结果重新插入到向量中。

然而，我们可以做得更好。在向量化之前，我们执行一个简单的前向数据流分析来推断向量化变量形状的不变量。例如，如果我们能够证明基地址加上 gep 使用的偏移量在每个实例中都包含相同的值，我们就不需要复制 gep （Get Element Pointer）以及可能跟随的内存访问指令。在这种情况下，我们可以将值作为标量加载，并在需要时将其广播到向量。

更常见的例子是具有连续值的向量元素。假设有一个基址，并且知道对于连续的实例 id，偏移量的值是连续的。将偏移量放在一个向量中得到{n, n + 1, …, n + W − 1}。在 gep （Get Element Pointer）中使用这样的偏移向量会得到一个连续地址的向量。因此，可以直接使用连续加载 vector 元素的 vector load 和 vector store 指令。

我们的分析跟踪每个变量的以下信息：

• 变量是否为连续实例包含连续值（用缩写"c"表示）。consecutive 表示变量的值对于连续的实例来说是连续的。例如，如果有一个数组，它的索引0, 1, 2, 3的值是连续的，那么这些值可以被一次性加载到一个 SIMD 寄存器中。
• 以及它是否对齐到 SIMD 宽度（用缩写"ca"表示）。consecutive aligned 表示变量的值不仅连续，而且对齐到 SIMD 寄存器的边界。这意味着可以使用 SIMD 指令集中的对齐内存访问指令，这些指令通常比未对齐的访问更快。

后者允许使用更快的、对齐的内存指令。前者虽然避免了分割，但需要进行未对齐的内存访问。

此外，如果一个变量对于每个实例都包含相同的值（用缩写"s"表示，same 表示所有实例的变量值都相同），并且这个值是 SIMD 宽度的倍数（用缩写"sa"表示，same aligned 表示变量的值不仅相同，而且这个值是SIMD寄存器宽度的整数倍，这允许在向量化时更有效地使用内存带宽），这也很重要。这样的变量可以保持标量形式，并在需要时广播到向量中。

图4 lattice L 结构的向量化分析

此外，向包含连续值的向量中添加相同的值会保持连续性（对齐同样适用）。如果无法确定向量的形状，分析值将为"⊤"（⊤: 表示无法确定向量的形状，也就是说，编译器无法从当前的分析中得知变量值的排列方式是否满足 SIMD 处理的要求。）。图 4 显示了相应的格（lattice）结构（“lattice”结构，通常是指在数据流分析中使用的逻辑结构，用于表示变量可能的状态集合，以及这些状态之间的转换关系。这里可能表示变量的不同向量化状态以及它们之间的关系）。

在编译器的上下文中，"lattice L" 通常指的是一个“格（lattice）”数据结构，它用于在数据流分析中表示变量的状态或者属性的集合。在数学和计算机科学中，一个“格”是一个偏序集合，其中每个元素都有明确的顺序关系，并且对于任意两个元素都存在最大下界（meet）和最小上界（join）。 这些状态描述了变量在向量化过程中可能具有的不同属性，例如是否连续（consecutive）、是否对齐（aligned）、是否每个实例的值都相同（same）等。编译器使用这个“格”来跟踪和推断变量的这些属性，并决定如何最有效地进行向量化。例如，"lattice L" 中的每个元素可能代表一个变量的属性集合，而“格”中的顺序关系可能表示这些属性的兼容性或者限制。编译器通过这个“格”来确定哪些变量可以被向量化，以及如何生成最优化的向量代码。通过"lattice L"，这有助于理解变量属性的组织方式以及它们如何影响编译器的向量化决策。

图 5 展示了分析过程中我们所使用如下的转换函数：

这个转换函数a(w)存在着分析信息，这些信息映射变量到 lattice 元素们。其中，标注a | v -> l表示的含义是：

其中需要说明的是：

• 函数 tid 生成的值是连续且对齐的。这是因为~f完成了f函数的W个连续实例的工作。函数 tid 函数是线程标识符（thread identifier）函数，用于生成每个线程的唯一标识。W 可能表示工作组的大小或 SIMD 寄存器的宽度。
• 第一个实例的 ID 总是能被W整除。赋值常量会创建 s （scalar）类型的值；一个能被 W 整除的整数常量会创建一个 sa （aligned scalar，对齐的标量）类型的值。
• 由于每个实例传给f的参数都是相同的，所以在指针参数的情况下，arg（参数获取函数用于传给函数的参数）返回 s 或 sa 类型的值。这意味着数据并行程序的调用者必须将传入的指针数据对齐到 SIMD 寄存器边界上。如果所有实例的地址都相同，那么加载操作就是 s 类型的（s 或 sa）。加法和乘法操作符相应地传播这些信息。
• φ-操作符用于在控制流的合并点合并多个值。所有其他操作符在可能的情况下仅保留相同值信息，以避免过早广播。
• 作者手动检查得出[[·]]^#是单调的。因此，常用的固定点算法会收敛，固定点算法是一种求解数值问题的迭代算法，这里指的是求解程序中某个属性或变量的稳定值。每个程序点的初始值是 ⊥ 。

图5 转换函数

[ysh329]

4.3 掩码生成

正如已经提到的，控制流可能会发散，因为某个条件对于一些标量实例可能是真的，而对于其他实例可能是假的。因此，所有代码都被执行了。控制的显式转移通过控制流边上的掩码变量（简称：掩码 mask ，也常被称为谓词 predicate）来建模。如果在位置 i 上，从 B 到 B' 的 CFG 边上的掩码被设置为 true，那么代码的第 i 个实例就从 B 转移到了 B'。因此，掩码表示的是，在相应的控制流边上，向量中的哪些元素包含有效数据。

控制流边上的掩码隐式地定义了区块的入口掩码：一个区块的入口掩码要么是所有进入边的掩码的析取（即逻辑或），或者——在循环头部的情况下——是一个使用来自循环前驱和循环 latch（循环末尾）的值的 φ-函数。离开一个区块的控制流边上的掩码由区块入口掩码和潜在的条件给出。

名词解释：

• 入口掩码：entry mask 指的是进入某个程序区块如循环或条件语句时，用于指示哪些线程或数据元素应该被包含在内。
• 析取：disjunction，表示只要有一个条件为真，结果就为真。即逻辑或。
• 循环头部：loop header，是循环开始的区块，通常包含循环的入口条件。
• 循环前驱：pre-header，前驱区块，指的是在循环或条件语句之前的区块，它为循环或条件语句提供入口。
• 循环 latch：循环末尾，是循环的最后一个区块，通常包含循环的继续或退出条件。
• φ-函数：φ function ，用于在控制流合并点合并来自不同路径的值。

图6 边和块的掩码

如果一个区块通过一个无条件分支退出（无条件分支，即 Goto 这种无论如何都会执行的），其单一退出边的掩码等于入口掩码。如果退出分支是条件性的，那么区块的“真边”（true edge）的退出掩码是其入口掩码和分支条件的合取（逻辑与）。“假边”（false edge）的退出掩码是入口掩码和分支条件的否定的合取。图 6 显示了三个基本区块A、B、C，以及它们对应的区块入口掩码（mA、...）和边掩码（mA→B、...）。

循环掩码（Loop Masks）。循环掩码。每个循环都必须维护一个掩码，该掩码对于循环中仍然活跃的所有实例都为真。只有当所有实例的这个掩码都为假时，循环才能退出。因此，在循环头部生成了一个特殊的φ-函数——循环掩码φ（图7中的mB）。它的第一个传入值是来自前驱头部的进入边的掩码，第二个值是循环回边的掩码。

具有多个退出点的循环需要为每个退出点增加额外的掩码。这些循环退出掩码存储了哪些实例通过相应边离开了循环，并替换了该退出边的先前掩码。它们由两个指令维护：一个掩码更新操作（mexit）和循环头部的一个φ-函数（mphi）。更新操作是相应边的掩码和φ-函数的析取（逻辑或）。

φ-函数有一个来自前驱头部的传入值，还有一个来自循环末尾的值。来自循环末尾的值是更新操作的结果。如果循环是顶层循环，来自前驱头部的值是一个空掩码（所有元素都设置为假）。如果循环是嵌套的，它使用父循环的相应循环掩码φ-函数。

[ysh329]

具有多个退出点的循环需要为每个退出点增设额外的掩码。这些循环退出掩码记录了哪些实例通过相应的边离开了循环，并取代了该退出边上原来的掩码。它们由两个指令维护：一个掩码更新操作（mexit）和循环头部的一个φ-函数（mphi）。更新操作是相应边的掩码和φ-函数的逻辑或（disjunction）。 φ-函数有一个来自前驱头部的传入值和一个来自循环末尾（latch）的值。来自循环末尾的值是更新操作的结果。如果循环是顶层循环，来自前驱头部的值是一个全部设置为假的空掩码。如果循环是嵌套的，它将使用父循环的相应循环掩码φ-函数。

概念解释：

• 循环退出掩码（Loop-Exit Masks）：对于有多个退出点的循环，每个退出点都需要一个掩码来追踪哪些执行线程或实例通过该退出点离开了循环。
• 掩码更新操作（Mask Update Operation, mexit）：这是一个逻辑或操作，它将当前的退出边掩码与φ-函数的结果结合起来，以更新掩码的状态。
• φ-函数（Phi-Function, mphi）：在循环头部的φ-函数用于合并来自循环不同部分的掩码值。
• 前驱头部（Pre-Header）：循环开始之前的区块，为循环提供入口点。
• 循环末尾（Latch）：循环的最后一个区块，通常是循环的退出点或再次进入循环的点。
• 空掩码（Empty Mask）：所有位都设置为假的掩码，表示没有实例应该继续执行循环。
• 顶层循环（Top-Level Loop）：不是任何其他循环的子循环的循环。
• 嵌套循环（Nested Loop）：一个循环内部包含另一个循环，形成循环的嵌套结构。
• 逻辑或（Disjunction）：逻辑或操作，表示只要有一个操作数为真，结果就为真。

在编译器优化和代码生成中，正确地管理和更新这些掩码对于确保循环的正确退出至关重要。通过使用掩码，编译器可以确保在 SIMD 执行环境中，只有应该继续执行循环的实例才会继续执行，而其他实例则可以安全地退出循环。

掩码生成后，每个循环退出点恰好有一个更新操作位于退出分支之前，并且在每个被离开的循环的头部有一个φ-函数。最后，选择生成（见第4.4节）需要一个组合的循环退出掩码。这个掩码结合了当前迭代中离开循环的所有实例的所有信息。在包含更多嵌套循环的循环中，父循环的当前迭代包括所有嵌套循环的所有迭代。因此，组合的循环退出掩码是来自嵌套循环的所有累积循环退出掩码的析取（逻辑或），以及当前循环的退出条件的析取。

概念解释： 掩码生成（Mask Generation）：创建用于追踪哪些执行线程或实例应该继续执行循环的掩码。 更新操作（Update Operation）：在循环的退出分支前执行的操作，用于更新循环退出掩码的状态。 φ-函数（Phi-Function）：在循环头部的φ-函数用于合并来自循环不同部分的掩码值。 选择生成（Select Generation）：在 SIMD 执行中，根据掩码选择性地执行指令的过程。 组合循环退出掩码（Combined Loop-Exit Mask）：一个综合的掩码，它结合了当前迭代中所有离开循环的实例的信息。 嵌套循环（Nested Loops）：循环内包含其他循环，形成嵌套结构。 析取（Disjunction）：逻辑或操作，用于组合多个掩码，如果任何一个掩码为真，则组合掩码为真。 退出条件（Exit Conditions）：决定循环何时退出的条件。 在 SIMD 执行环境中，正确地管理和更新循环退出掩码对于确保循环的正确退出至关重要。组合循环退出掩码允许编译器在 SIMD 执行中，根据当前迭代中所有离开循环的实例的信息，来选择性地执行指令。这种方法确保了即使在复杂的嵌套循环结构中，也能正确地处理循环退出逻辑。

[ysh329]

4.4 Select Generation

[ysh329]

4.5 CFG Linearization

[ysh329]

4.6 Irreducible Control-Flow

[ysh329]

4.7 Instruction Vectorization

[ysh329]

5. Related Work

[ysh329]

5 2nd part

[ysh329]

6. Experimental Evaluation

我们将本文中介绍的算法实现在了 LLVM 编译器框架中，并在两个真实世界的场景中评估了向量化程序的运行时间：实时光线追踪中的着色和自定义 OpenCL 驱动程序。所有实验都是在 2.8 GHz 的 Core 2 Quad 上进行的，配备有 4 GB 的 RAM。向量指令集是 Intel 的 SSE 4.1，每个寄存器有 4 个浮点/整数。机器以 64 位模式运行，因此有 16 个向量寄存器可用。

6.1 光线追踪的向量化着色器

我们将我们的向量化工具集成到了实时光线追踪器 RTfact 的着色语言编译器中，该编译器使用 LLVM 将 RenderMan 着色器编译为 x86 机器代码。着色器非常适合整体函数向量化，计算密集，并且只执行对齐和连续的内存访问。因此，拆分向量几乎不会产生开销。

下表显示了每个着色器在 512×512 像素分辨率下应用于被两个点光源照亮的球体的性能。我们将自动向量化的着色器的渲染性能与顺序执行的标量版本的着色器进行比较。

表：十种不同着色器的光线追踪性能，以每秒帧数（fps）表示。“原生”（Native）指的是渲染器内部高度优化的 SIMD 着色器。

向量化的着色器版本在性能上平均比它们的标量对应版本高出 3.9 倍。对于一些着色器，我们观察到超线性加速高达 5.0。这有两个原因：

1. 首先，光线追踪器在内部也使用向量工作，在顺序着色的情况下需要拆分它们。这种开销在向量化着色中是不存在的。
2. 其次，光线追踪器使用的数据布局在向量化情况下能够实现更好的缓存局部性。整个转换和编译过程（包括 LLVM 的即时编译 JIT）对任何着色器的耗时都不到 100 毫秒。这允许在渲染器运行时进行动态重新编译，这在图形学中很重要。

与直接集成到渲染器中的经过高度优化的 SIMD 着色器相比，我们的性能达到了最佳可能性能的 90% 。

[ysh329]

6.2 OpenCL

在我们评估的第二部分中，我们构建了一个自定义的 OpenCL 驱动程序。我们实现了足够完整的 OpenCL API 的一部分，以运行一系列不同的测试应用程序。该驱动程序在内部使用 LLVM 对 OpenCL 特定的运行时代码进行修改、向量化和代码生成。

Vectorization

在 OpenCL 中启用整个函数的向量化需要进行几个程序转换，我们将在这里简要描述。驱动程序的更详细描述超出了本文的范围。

如果应用程序使用多于一个维度的输入数据，驱动程序必须选择一个 SIMD 维度来进行向量化。根据所选择的维度，必须调整内核以适应现在由 SIMD 宽度 W 除的修改后的数据大小。我们的驱动程序目前总是使用第一维度，但可以应用一种启发式方法，通过分析由 get_global_id() 给出的不同维度的索引来访问内存。

我们的数据流分析（第4.2节，用于分析程序中数据的流向和访问模式，以便进行优化）的结果允许我们优化频繁出现的案例，其中数组的连续和对齐元素使用局部 ID 通过局部大小（可以被 W 整除）的倍数递增来访问。

为了支持 OpenCL 的 barrier() 语句，我们在驱动程序内部实现了一种同步方案，这个方案不是基于线程的，而是采用协程风格的函数拆分。同步问题并不简单，因为内核被允许在循环迭代之间进行同步，这要求一个组中的所有线程必须以同步的步伐（lock-step）执行循环迭代。我们的系统避免了代价高昂的回调到驱动程序或操作系统，甚至能够使所需代码与内核本身进行交叉优化。

Benchmarks

除了 AOBench 之外的所有基准测试应用程序都是直接取自 AMD 的 Stream 软件开发工具包。AOBench 是一种最小化的环境光遮蔽光线追踪器。该 SDK 作为我们所有性能评估的基础：它使用我们的 OpenCL 驱动程序进行编译，而不是 AMD 的。我们还进一步使用 AMD 的 clc 工具从 OpenCL 代码生成 LLVM 位码文件。

图 我们的OpenCL驱动器在不同应用程序上的平均内核执行时间（100次迭代）。括号内表示使用的线程数，加速比是指标量模式与未进行多线程的向量化模式之间的比较。标记为“AMD”的列表示AMD专有OpenCL CPU驱动器的性能

上表展示了从计算密集型内核（例如 BlackScholes）到主要由内存访问操作支配的内核（例如 Histogram）等多种应用程序的运行时间性能。这表明向量化能够处理真实世界的应用程序，但它也显示了它的局限性：

• 与着色器不同，OpenCL内核经常以非连续的方式修改数组，这会导致拆分向量和顺序加载及存储时受到较大的惩罚。因此，像FastWalshTransform和Histogram这样的基准测试，其中内存访问占主导地位，甚至在启用向量化时会变慢。

• 另一方面，向量化可以显著提高像NBody或Mandelbrot这样合适应用程序的性能：这些基准测试非常计算密集，并且没有复杂的控制流和分散的内存访问模式。BlackScholes基准测试另外几乎没有分支，它的内存占用很小，几乎不发生溢出。这累积成一个5.2倍的加速因子。

观察到的加速比的广泛范围激发了仅在似乎有益的代码区域应用向量化的启发式方法的开发。我们将这留作未来的工作。

为了能够与AMD的参考驱动程序进行比较，我们通过OpenMP实现了一个简单、未优化的多线程方案。除了知道它是多线程的并且也使用LLVM之外，我们没有关于AMD实现的任何信息。尽管如此，请注意，除了Histogram和AOBench之外，我们的定制驱动程序在所有测试用例中都显著优于他们的驱动程序。Mandelbrot基准测试在2.2版本的SDK中被重写，使用OpenCL的float4 SIMD数据类型在手动展开的循环中。尽管性能提高了十二倍以上（6秒），但代码几乎不再可读，编写起来也更加繁琐（从28行代码增加到223行）。更重要的是，我们的驱动程序仍然以7.5倍的因素优于它，而没有给程序员带来那么大的压力。

[ysh329]

7. 总结

在这篇论文中，我们讨论了在具有 SIMD 指令集的机器上实现数据并行程序的方法。我们提出了一种算法，它可以将任意控制流图的函数向量化，该函数以 SSA（静态单赋值）形式给出。该算法基于控制流到数据流的转换，并为没有硬件支持条件执行（predicated execution）的架构生成高效的混合代码。

我们进一步提出了一种数据流分析（analysis），它确定哪些代码区域在向量化方面有关对齐和连续性的约束。该分析（analysis）产生了两个主要优化：

• 首先，我们可以避免由于当前 SIMD 指令集架构的内存访问指令限制而造成的高开销的向量拆分。因为将单个向量操作拆分成多个较小的向量操作。
  • 可能导致①性能下降，并行性的损失：拆分向量意味着更多的指令完成原本单指令做的事情；②内存访问增加：需要多次的 store/load；③寄存器压力：需要额外寄存器存储中间结果；④控制流复杂性：拆分引入额外控制流逻辑，可能导致分支预测失败等从而影响流水线效率；⑤编译器优化负担：拆分导致编译器更多的工作，编译时间延长。
• 其次，我们的向量化算法可以智能地识别和忽略那些在所有执行实例（Instances of the program）中都不变的代码区域，编译器可以减少不必要的工作，从而提高向量化过程的效率。这使得任何重建步骤都变得不必要。
  • 解释一下这句话：这里的“不变”（invariant）指的是某些代码片段在程序的多次执行过程中，其执行结果始终相同，不会受到程序输入或其他变量值的影响。
  • 程序实例（Instances of the program）：这通常指的是程序的一次执行过程。在某些情况下，程序可能因为输入不同或随机数生成等因素，在多次执行时表现不同。然而，如果某个代码区域在所有这些执行过程中都表现出相同的行为，那么它就是不变的。
  • 重建步骤（Reconstruction steps）不必要：在向量化的过程中，编译器有时需要对代码进行重构或调整，以适应 SIMD 指令集的要求。这可能包括重新排列指令、插入额外的指令来处理特殊情况等。但是，如果某个代码区域是不变的，编译器就不需要对这部分代码进行这些额外的重构工作，因为无论程序如何执行，这部分代码的行为都是已知且不变的。

案例研究展示了系统的适用性：将我们的技术集成到一个实时光线追踪器中，提供了比标量版本平均 3.9 倍的线性加速。对 OpenCL 框架的增强所得到的加速在很大程度上取决于目标内核，变化在 0.6 到 5.2 之间。