Experiences Building an MLIR-based SYCL Compiler

Ettore Tiotto (Intel Corporation), Vı́ctor Perez (Codeplay Software), Whitney Tsang (Intel Corporation), Lukas Sommer (Codeplay Software), Julian Oppermann (Codeplay Software), Victor Lomuller (Codeplay Software), Mehdi Goli (Codeplay Software), James Brodman (Intel Corporation)

A1 主要贡献

本文旨在解决当前SYCL编译器面临的两大挑战：一是由于过早地将高级语言结构和语义降低到低级中间表示（如LLVM IR），导致高级信息丢失，影响了优化效果；二是主机代码和设备代码的编译过程相互隔离，无法进行联合分析和优化。为了应对这些挑战，本文介绍了基于MLIR编译器框架设计和实现一个SYCL编译器的实践经验。

核心问题与研究目标
* 核心问题：现有的SYCL编译器（通常基于LLVM）在编译早期阶段就将SYCL的高级语义（如工作项并行执行、设备内存访问模型）转换为低级IR，这使得编译器难以利用这些高级信息进行有效的优化。此外，主机和设备代码的独立编译阻碍了跨边界的优化。
* 研究目标：构建一个基于MLIR的SYCL编译器（SYCL-MLIR项目），利用MLIR的方言（dialect）机制来捕获SYCL编程模型的高级语义，并利用其操作嵌套能力来同时表示和分析主机与设备代码，从而实现更强大的设备代码优化。

创新点与主要贡献
本文的主要贡献如下：
* 扩展MLIR框架：设计并实现了一个MLIR方言（SYCL dialect），用于捕获SYCL编程模型的关键语义。
* MLIR编译流程架构：构建了一个基于MLIR的编译流程，该流程支持对主机和设备代码进行联合分析，从而能根据内核的调用上下文进行更优的设备代码优化。
* SYCL代码分析与优化：设计并实现了针对SYCL主机和设备代码的分析方法，以及一系列设备代码优化技术。
* 性能评估：通过一系列SYCL基准应用程序评估了基于MLIR的编译流程，并与现有的基于LLVM的SYCL编译器进行了比较。
* 实践经验分享：报告了在为一个基于C++的并行编程模型构建MLIR编译器过程中的实践经验和遇到的主要挑战。

A3 背景知识/关键Observation/设计原则

II. 背景知识

A. SYCL

SYCL编程模型概述。SYCL 【索引编号2，SYCL 2020 specification (revision 7)+2023+Khronos+https://registry.khronos.org/ SYCL/specs/sycl-2020/pdf/sycl-2020.pdf】 是一个由Khronos制定的开放标准，它定义了一个基于C++的并行编程模型，用于异构系统，为用户提供了访问加速器资源的丰富API。

与其它异构编程模型的区别。尽管SYCL与其他异构并行编程模型（如OpenCL、OpenMP或CUDA）共享一些底层概念和抽象，但它们之间存在一些有趣的差异。与这些编程模型不同，SYCL不依赖于对C++语言的扩展。

SYCL内核的提交与执行。SYCL内核使用仿函数（functor，即定义了调用运算符的普通C++类或lambda表达式）来表示。内核通过一个命令组函数（command-group）提交到队列中执行。命令组允许用户表达内核运行所需的依赖关系，这些依赖关系可能包括等待另一个内核完成或请求内存传输。命令组最终会提交内核，并指定其运行的索引空间。

SYCL执行模型。内核以一个索引空间（称为ND-range）提交执行，该空间被细分为多个工作组（work-group）。对于该空间中的每个点，一个实例（称为工作项，work-item）与其他工作项并行执行内核。工作项也被捆绑到工作组中，从而允许同一工作组内的工作项之间进行一些资源共享、通信和通过屏障（barrier）进行同步。每个SYCL内核仿函数都会接受一个id、item或nd_item类型的参数，该参数封装了工作项在ND-range中的位置。

SYCL内存模型。SYCL还定义了一个具有不同可见性级别的内存层次结构：全局（global）、局部（local）和私有（private）内存。全局内存由所有工作项共享，而局部内存由一个工作组内的所有工作项共享，从而实现工作组范围内的协作，并且通常比全局内存访问速度更快。最后，私有内存仅对给定的工作项可见。

SYCL内存管理机制。SYCL提供了两种在主机和设备之间管理内存的方式：统一共享内存（Unified Shared Memory, USM）和缓冲区-访问器（buffer and accessor）模型。使用USM时，用户直接操作指针。内存使用专用的类malloc和free函数进行分配和释放，用户负责手动在设备与主机之间传输数据。而缓冲区-访问器模型允许SYCL运行时不仅处理跨设备的内存管理（内存创建、删除、与设备之间的数据传输）以确保内存一致性，还负责创建内核之间的依赖关系。

缓冲区与访问器模型详解。SYCL缓冲区（buffer）是一个多维容器，它拥有内存并跟踪其在主机和设备上的副本位置，但它不直接提供对内存的访问。为了访问数据，用户必须在命令组内创建一个访问器（accessor）对象。该对象的创建将在SYCL调度器中创建一个依赖关系，以确保数据在内核执行前是可用的。在内核内部，用户可以像使用常规的C++ vector或mdspan容器一样使用访问器对象。

访问器的复杂性。访问器是一个重量级对象，封装了几个动态信息：数据指针、数据的完整范围，以及子范围和偏移量。子范围和偏移量仅在范围访问器（ranged accessor）的情况下有用，它允许用户只将缓冲区的一部分传递给内核。然而，范围访问器和非范围访问器之间的区别是通过调用不同的构造函数来实现的，并且从未反映在C++类型上。其他信息也通过模板参数静态嵌入，包括它是只读、只写还是读写。

SYCL的高级抽象优势。虽然底层的执行和内存模型与OpenCL或CUDA相似，但缓冲区-访问器模型是SYCL如何利用现代C++的强大功能来自动化异构编程中一些繁琐任务的一个很好的例子。如果应用程序使用缓冲区和访问器抽象，SYCL运行时可以完全自动化内核之间的依赖关系跟踪以及主机与一个或多个设备之间必要的数据移动，而在像OpenCL或CUDA这样更低级的编程模型中，开发人员需要手动处理这些细节。这个例子以及SYCL中其他更高级的抽象，如强大的乱序队列，使得典型的SYCL代码更简洁，从而提高了可维护性和程序员的生产力。

可移植性。更高级别的抽象也可以促进不同架构之间的可移植性。与只能在特定供应商设备上运行的专有模型（如CUDA）相比，SYCL的开放性使其实现能够针对来自不同供应商的硬件，从而使SYCL在许多不同的平台上可用。

与其他模型的详细比较。虽然本文不打算更详细地比较SYCL与其他编程模型，但可以在【索引编号3，Performance study of gpu applications using sycl and cuda on tesla v100 gpu+2021+IEEE High Performance Extreme Computing Conference (HPEC)】–【索引编号8，Experiences migrating cuda to sycl: A molecular docking case study+2023+Proceedings of the 2023 International Workshop on OpenCL】中找到这类比较。

B. MLIR

MLIR框架的动机。MLIR框架【索引编号1，MLIR: Scaling compiler infrastructure for domain specific computation+2021+IEEE/ACM International Symposium on Code Generation and Optimization (CGO)】的诞生源于两个洞见：一是过早地将高级编程模型降低到低级表示会严重损害编译器执行强大优化的能力；二是许多现有的IR表示（例如LLVM IR）难以扩展。许多优化遍（pass）如果能在一个保留程序结构（例如循环）并能编码领域知识的更高级别表示上工作，将会受益匪浅。

MLIR与本工作的关系。这一洞见也是本工作的一个重要动机，我们寻求通过一个允许直接编码SYCL语义的中间表示来改进SYCL代码的优化。

MLIR之前的多层次IR思想。在MLIR出现之前，在编译器基础设施中拥有多个中间表示的好处已在文献【索引编号9，Xil and yil: The intermediate languages of tobey+1995+Papers from the 1995 ACM SIGPLAN Workshop on Intermediate Representations】–【索引编号11，Generic and gimple: A new tree represen-tation for entire functions+2003+https://api.semanticscholar.org/ CorpusID:58211542】中得到广泛认可。像Swift或Rust这样的编程语言后来也引入了更高级别的表示，目标相同，最终再降低到LLVM IR。MLIR现在旨在为这类表示提供一个平台和通用基础设施，以实现不同前端之间更好的可重用性。

MLIR的核心概念：方言（Dialects）。使MLIR能够覆盖如此广泛的抽象级别和领域的核心概念是方言（dialects）。一个方言封装了与特定领域表示相关的属性（attributes）、类型（types）和操作（operations）。上游MLIR项目收集了许多在不同抽象层次上有用的可组合方言，例如用于算术操作的arith方言或用于结构化控制流（如循环）的scf方言。除了上游方言，MLIR的用户可以通过定义专门的方言来自由扩展框架，以适应他们的问题领域。为此，本工作定义了一个捕获SYCL语义的方言，更多细节在第四节中给出。

MLIR的渐进式降低。基于MLIR的编译流程通常会使用多个方言的组合来表示应用程序在编译过程的每个阶段。当一个抽象级别上的优化完成后，表示通常会被降低到另一组方言，在那里可以进行更多的优化。总的来说，这通过方言转换和模式重写产生了一个渐进式降低（gradual lowering）的过程，本工作也在第四节中描述并实践了这一方面。

MLIR的通用基础设施。除了方言抽象和渐进式降低，MLIR还为创建分析和转换遍（passes）提供了通用基础设施。这个通用基础设施也支撑了本工作中的分析和转换，具体在第五至第七节中描述。

III. 目标与方法

SYCL-MLIR项目的目标。如第一节简要讨论的，本工作提出的SYCL-MLIR项目旨在为SYCL异构编程模型构建一个基于MLIR的编译器。使用MLIR编译器框架来构建SYCL编译器主要有两个动机。

动机一：主机与设备代码的联合表示。首先，MLIR固有的操作嵌套能力允许编译流程在同一个MLIR模块中并排表示SYCL的主机和设备代码，从而能够在其主机调用上下文中更好地分析设备内核。编译流程的细节以及如何实现这种主机和设备代码的并排分析将在第四节和第七节中描述。

动机二：保留高级语义。其次，MLIR的方言概念使得编译流程能够初步保留SYCL编程模型的高级语义。通过这种方式，语言的精确语义可以被分析和转换过程所利用，并且只有在受益于SYCL语义的优化完成后才进行降低。为此，SYCL-MLIR项目定义了一个SYCL方言，将SYCL编程模型的关键实体建模为MLIR的属性、类型和操作。

设备代码的SYCL方言建模。在设备代码中，SYCL方言主要建模了两个概念。一个是当前工作项（见第II-A节）在整个执行网格及其工作组内的位置。为此，SYCL的id, item, nd_item, range, nd_range和group类被建模为SYCL方言中的类型，而获取和操作这些类实例的关键函数则被表示为方言的操作。

访问器（Accessor）的建模。SYCL编程模型在设备上对编译器转换具有高度相关性的另一个重要部分是通过SYCL访问器（见第II-A节）访问设备内存。为了使优化能够推理和转换内核的内存访问行为，SYCL accessor类被建模为SYCL方言中的另一个类型，诸如通过访问器访问内存元素等操作被建模为方言中的MLIR操作。通常，内存访问基于工作项在执行网格中的位置，因此在SYCL方言中表示id及相关类在这里也扮演了另一个重要角色。

主机代码的SYCL方言建模。在主机端，该方言旨在表示内核的调用上下文，包括传递给这些内核的参数和ND-range。为此，代表SYCL命令组和内核函数对象（见第II-A节）构造的操作被添加到方言中。许多为表示设备端SYCL实体而引入的类型，例如用于id或range类的类型，可以在这里重用，并且还为像buffer这样的类添加了额外的类型。

访问器来源的建模与分析。除了内核启动及其参数，设备代码中使用的访问器的来源可以提供有价值的信息。因此，建模访问器及其底层缓冲区（见第II-A节）构造的操作也被添加到方言中。分析用于构造访问器和缓冲区的参数可以为了解它们的行为提供有用的见解，例如两个访问器之间潜在的重叠或别名问题。

主机代码提升过程。主机代码的提升（raising）过程以及上述操作如何在其中使用，将在第七节中更详细地描述。

A2 方法细节

IV. 编译流程

SYCL标准编译流程：SMCP。SYCL标准（见II-A节）描述了一种单源多编译器遍（Single-Source Multiple Compiler Passes, SMCP）的编译流程。这包括使用不同的编译器多次处理同一份源代码：一次用于主机，一次用于每个目标设备。在Intel DPC++编译器【索引编号12，DPC++ open-source SYCL Compiler+2023+Intel Corporation+https://github.com/intel/llvm】中，生成的设备镜像被链接在一起并包装在主机对象中，以便稍后由SYCL运行时读取（如图1中的虚线路径所示） 。

SMCP方法的局限性。然而，SMCP技术在设计上存在一个限制：由于主机和设备编译器之间共享的信息有限，无法在编译时高效地执行跨主机-设备边界的优化。为了解决这个问题，使用一个同时处理主机和设备代码的自定义SYCL编译器，即通过单源单编译器遍（Single-Source Single Compiler Pass, SSCP）的方法，为在主机代码中使用设备代码分析进行转换（反之亦然）开辟了可能性。

SYCL-MLIR的编译流程方法。为了将C++ AST转换为MLIR，我们使用了自己的Polygeist【索引编号13，Polygeist: Raising C to polyhedral MLIR+2021+30th International Conference on Parallel Architectures and Compilation Techniques】分支作为设备编译器。虽然它在编译设备端代码时效率很高，但在处理主机端所需的C++构造时遇到了困难。例如，虚函数和异常没有得到很好的支持。为了在不需要一个健壮的C++ MLIR前端（见第九节）的情况下获得与SSCP技术类似的结果，我们采取了一种替代方法，即从LLVM IR中获取MLIR主机代码，然后利用它来辅助设备编译。接着，我们利用MLIR的嵌套IR结构，获得了主机和设备代码的联合表示（如图1中的虚线路径所示）。通过这种方式，我们为一个能够利用主机代码分析进行设备代码优化的SYCL编译器奠定了基础。

后续分析与转换。在此之后，会执行一系列使用SYCL领域特定知识的自定义分析和转换遍，包括一个主机提升（host raising）遍，它能检测并提升SYCL主机代码中的相关模式，如第七节所述。与开发一个完整的C++ MLIR前端相比，这种替代方法以很小的成本实现了主机-设备优化。然而，其缺点是，SYCL运行时代码的更改可能导致提升模式匹配失败，迫使这个遍必须与运行时的变化保持同步更新。

V. 编译器分析

分析与工具的扩展。在开发SYCL-MLIR编译器的过程中，我们发现需要用一些在优化编译器中普遍有用的静态分析和工具来扩展MLIR框架。本节简要概述了我们的实现，重点介绍了如何利用几个SYCL方言操作所编码的语义信息。

A. 别名分析

基于MLIR框架的扩展。MLIR提供了一个别名分析框架，可以方便地进行扩展，以考虑不同方言提供的领域知识。为了利用SYCL方言的语义信息，我们创建了一个特定于SYCL的别名分析，扩展了MLIR现有的分析。

利用SYCL方言语义。SYCL方言操作的语义可以被编码到分析中，使得编译器能够在许多情况下证明由SYCL操作产生的值不会发生别名，从而利用SYCL方言使别名分析更加精确。

B. 到达定值分析

分析的定义与实现。经典的到达定值分析是一种数据流分析，旨在提供在给定程序点可能修改了内存中某个值的一组操作。我们的实现利用了MLIR提供的数据流分析框架和我们专门的别名分析。我们考虑一个值的两种到达定值：
* 修饰符 (MODS)：对值本身的定值，或对一个已知明确（must）别名于它的值的定值。
* 潜在修饰符 (PMODS)：对一个已知可能（may）别名于该值的定值的定义。

示例说明。例如，考虑列表1中的MLIR代码片段。第4行的store操作直接更新%ptr1指向的内存位置，而第6行的store操作可能会更新相同的内存位置，因为%ptr2可能与%ptr1别名。因此，在第8行，%ptr1的到达定值是{MODS: a, PMODS: b}。

利用MLIR的内存效应接口。为了定义不同操作的内存效应，MLIR框架为操作提供了一个通用接口，因此像到达定值分析这样的分析可以推理来自不同方言的操作的效应。通过为相关的SYCL方言操作实现这个接口，到达定值分析可以被定制，以考虑每个SYCL方言操作的精确内存语义。

C. 一致性分析

分化分支的定义。在GPU编程中，分化分支（divergent branch）是指分支条件对于工作组（见II-A节）中的所有工作项不会产生相同结果的分支，导致工作项的子集分支到不同的程序点。

一致性跟踪。通过跟踪变量的一致性（uniformity），可以识别出分化分支。如果一个工作组中的所有工作项都为某个变量赋予相同的值，则该变量被称为一致的（uniform），否则为非一致的（non-uniform）。一个非一致值的简单例子是产生工作项全局ID的操作结果。在分化路径下赋给同一内存位置的值会产生数据分化，如果用在条件表达式中，可能会导致分化控制流。

分化分支示例。在列表2的例子中，%gid_x值（第9行）是一个非一致性的来源，因为它是评估sycl.nd_item.get_global_id操作的结果，该操作产生SYCL工作项的全局ID（见II-A节）。因此，分支条件%cond（第12行）和%cond1都是非一致的，前者因为它使用了非一致值，后者因为它从%alloca指向的内存中加载，而该内存在分化分支中被赋予了不同的值（第14和16行）。

一致性分析的实现。为了推理分化控制流，我们实现了一致性分析（Uniformity Analysis），这是一个基于MLIR数据流框架的过程式（inter-procedural）数据流分析。

分析过程。形式参数最初被赋予未知的一致性，但SYCL内核入口点的参数除外，它们根据定义是一致的。然后，分析通过访问函数中的操作来传播值的一致性。一个自定义的特质（trait）会告知分析哪些SYCL操作是已知的非一致性来源。

其他操作的一致性判定。其他操作产生的值的一致性规则如下：
* 非一致 (non-uniform)：如果任何操作数是非一致的。
* 未知 (unknown)：如果任何操作数具有未知的一致性。
* 一致 (uniform)：如果所有操作数都是一致的，并且该操作没有内存效应。

利用MLIR特质机制。通过特质机制，MLIR框架允许利用领域特定知识来告知分析方言操作的效应。该特质可以轻松地添加到其他方言的操作中，从而使一致性分析也能推理其他方言的操作，展示了MLIR框架促进的可重用性。

处理内存效应。具有内存效应的操作会得到进一步分析，使用的是与到达定值分析相同的内存效应接口。这样，分析不仅可以推理SYCL方言，还可以推理所有其他实现了该接口的方言操作。

写内存效应的处理。如果操作具有未知的内存效应，它会被保守地认为具有未知的一致性。否则，每个内存效应都会被分析。对于一个写内存效应，使用到达定值分析（V-B节），未知或非一致的一致性会从（潜在的）修饰符及其支配的分支条件传播过来。

过程间分析。最后，该分析通过使用调用图（call graph）来跨越所有可能的调用点传播实际参数的一致性。如果所有调用点都是已知的（没有外部调用），那么参数的一致性是通过合并实际参数的一致性来计算的。

在SYCL-MLIR中的应用。在SYCL-MLIR编译器中，该分析目前被循环内化优化（VI-C节）用来确定一个循环是否在分化区域内执行。这是一个必要的前提条件，因为该转换需要注入一个组屏障（group barrier），如果在一个分化区域内执行，将会导致死锁。

D. 内存访问分析

分析的重要性。推理内核的内存访问模式是许多转换的关键，包括第六节中介绍的一些优化。为此，我们实现了一个内存访问分析（Memory Access Analysis）来推导SYCL内存在GPU内核中的访问模式。

分析的理论基础与扩展。该分析基于【索引编号14，Exploiting memory access patterns to improve memory performance in data-parallel architectures+2011+IEEE Transactions on Parallel & Distributed Systems】，并对其进行了扩展以考虑SYCL内存访问。对于一个仿射循环，SYCL内存访问模式可以通过使用一个访问矩阵和一个偏移向量来建模。

示例说明。例如，考虑列表3中的MLIR函数。循环中的内存访问（第23行）使用了一个通过sycl.accessor.subscript操作获得的memref，其索引是由第18行sycl.constructor操作构造的三维SYCL id给出的。索引函数是工作项全局ID（%gid_x和%gid_y）和循环归纳变量%i的仿射函数，在分析中，它由一个访问矩阵和一个偏移向量描述如下：

$$\begin{pmatrix} 1 & 0 & 0 \\ 0 & 0 & 2 \\ 0 & 1 & 2 \end{pmatrix} \times \begin{pmatrix} \%gid\_x \\ \%gid\_y \\ \%i \end{pmatrix} + \begin{pmatrix} 1 \\ 0 \\ 2 \end{pmatrix}$$

在SYCL-MLIR中的应用。在SYCL-MLIR编译器中，内存访问分析目前被循环内化优化（VI-C节）用来识别SYCL数组访问，以作为预取到局部内存的候选对象。

VI. 设备优化

优化目标。本节描述了旨在提高SYCL设备代码性能的优化。这些转换为SYCL-Bench套件【索引编号15，SYCL-Bench: A Versatile Cross-Platform Benchmark Suite for Heterogeneous Computing+2020+Euro-Par 2020: 26th International European Conference on Parallel and Distributed Computing】中的多面体基准测试带来了显著的加速，具体细节将在第八节中描述。

A. 循环不变代码外提 (LICM)

基于MLIR的实现。MLIR社区提供了一个工具，可用于将没有内存效应的操作提升出区域（region）。SYCL-MLIR编译器实现了一个LICM转换，该转换也考虑了从内存中读取或存储的操作。该转换使用同样用于到达定值分析的MLIR内存效应接口来确定操作的内存效应，并利用为SYCL专门化的别名分析（V-A节）来确定内存访问是否别名。

操作提升的条件。一个操作可以被提升，如果它的操作数要么已经在循环外定义，要么它们也可以被提升出循环。具有只读内存效应的操作可以被提升，只要编译器能证明循环中没有操作可能写入候选操作正在读取的内存位置。对于一个候选的SYCL读操作（加载一个操作数），该转换会分析所有可能与其别名的值，并确定其中是否有任何一个是在循环外定义的或也可以被提升。表现出写内存效应的操作可以被提升，如果没有后续操作写入同一内存位置，或者没有操作可能从候选操作写入的内存位置读取。

循环版本化。一旦一个循环被分析完毕，并确定了用于提升的候选操作，该转换会通过注入一个版本化条件来保护循环，以保证循环至少执行一次（否则一个被提升的操作可能会导致在原始代码中不存在的副作用）。

处理潜在别名。该转换还能够收集那些除非它们的某些操作数在运行时被证明不别名否则无法提升的候选操作。如果存在这些候选者，它们将通过对转换后的循环进行版本化来处理，版本化条件会检查那些阻止提升的操作数在内存中是否重叠。

B. 归约检测

归约模式的优化。数组归约是科学计算中的常见操作。这个遍（pass）旨在检测这种模式，并将归约累加到一个标量变量中，而不是在每次循环迭代中更新数组元素。例如，考虑列表4中的代码。循环在第2行从内存中加载一个值，并在第5行更新它。如果循环迭代N次，它将执行2N次内存访问。鉴于%ptr是循环不变量，这段代码可以被转换为列表5中所示的代码。

转换后的代码逻辑。转换后的循环通过affine.for操作的iter_args操作数接收数组元素的输入值（第1行）。然后，它将归约结果累加在循环携带的标量变量%red中。最后，在第8行用归约结果更新数组元素。优化后的循环版本不再执行任何涉及%ptr所指向的数组元素的内存访问。减少进出内存的流量是许多类型设备上的一项重要优化，其好处将在第八节中量化。

转换的安全性。别名分析（V-A节）被用来确保转换的安全性，因为在上面的例子中，%ptr和%other_ptr必须不别名，这个转换才是合法的。

range<2> global_size(N, N), wg_size(M, M);
cgh.parallel_for<matrix_multiply>(
nd_range<2>(global_size, wg_size), [=](nd_item<2> item) {
size_t i = item.get_global_id(0),
j = item.get_global_id(1);
for (size_t k = 0; k < N; k++)
C[i][j] += A[i][k] * B[k][j];
});

列表 6. 带有循环内化候选的命令组函数片段。

C. 循环内化

利用内存层次结构。如第II-A节所讨论的，SYCL编程模型定义了一个具有不同类型内存的内存层次结构。该层次结构中的不同内存各自具有关于大小和访问延迟的独特性，为优化开辟了可能性。

优化策略。例如，由于局部内存比全局内存小但速度快，它可以被有效地用于预取内存区域的子集，特别是当其访问表现出时间局部性或其访问模式不利于被GPU硬件合并时。

实现方法。通过循环内化（Loop Internalization）遍，利用MLIR的循环分块（tiling）基础设施，使完美嵌套循环中的SYCL内存访问使用局部内存。该遍利用内存访问分析（V-D节）来确定循环内的内存访问模式。

访问模式分析与决策。来自分析的访问模式信息随后被用来确定内存访问是否可以被合并。为此，通过移除对应于循环归纳变量的列来获得描述工作项间访问模式的子矩阵。

示例分析。在第V-D节的例子中，这个子矩阵由前两列给出，因为最右边的列对应于循环归纳变量。

合并与重用判断。如果工作项间的访问矩阵是线性的（Linear）或反向线性的（ReverseLinear）（如【索引编号14，Exploiting memory access patterns to improve memory performance in data-parallel architectures+2011+IEEE Transactions on Parallel & Distributed Systems】中所述），则访问可以被合并。如果工作项内的内存访问矩阵（通过移除对应于线程变量的列获得）不是零矩阵，则存在时间重用。基于这些信息，内存访问被分配为保留在全局内存中或被预取到局部内存中。

转换示例。为了演示该遍执行的转换，考虑列表6中的SYCL代码，它包含一个通过parallel_for启动的SYCL内核内的循环（第6-7行）。对应于访问器A和B的加载操作（第7行）被分析归类为使用局部内存的候选对象，因为它们都表现出时间重用。当存在至少一个要预取的候选访问时，优化会将循环转换为等同于列表7中代码的MLIR表示。请注意，优化后的代码需要两个屏障，因此，在进行转换之前，会使用一致性分析（V-C节）来确保循环不在分化区域内。

优化后的代码逻辑。在优化后的代码中，循环按工作组大小M进行分块（第13行），并为每个全局内存区域分配一个M × M的局部内存瓦片（tile）（第2-3行）。在外层循环中，每个内存区域的一部分被预取到局部内存中（第14-15行）。该优化依赖于工作组中的每个工作项来初始化内层循环使用的局部内存，因此注入了一个组屏障以确保所有线程完成初始化（第16行）。在分块的内层循环中（第17-18行），原始的访问被替换为局部访问器（第18行）。最后，注入了第二个组屏障（第19行），以保证在预取下一个全局内存部分之前，工作组中的所有工作项都完成了内层循环。

性能增益。通过这种转换获得的性能增益将在第八节中讨论。

VII. 主机提升与主机-设备优化

A. 主机提升

提升的动机。SYCL-MLIR编译流程背后的一个基本原理是执行主机代码分析以辅助设备代码编译。然而，从LLVM IR获得的MLIR代码对于分析来说级别太低，因为最终两者模块之间存在一一对应的关系。

实现方法。为了获得更高级别、适合分析的表示，我们定义了一个MLIR转换遍，它匹配DPC++运行时代码中存在的模式，并用SYCL方言中的操作替换它们。为了执行特定于SYCL的主机-设备优化，需要提升的两个主要模式是SYCL对象的构造和内核调度。

提升过程示例。作为提升过程的一个例子，我们将使用列表8中的SYCL程序，该程序在编译、转换为MLIR并提升后，被转换为列表9中的代码。

提升后的结果。正如我们所见，sycl.host.*操作根据前述的转换捕获了原始程序中所有相关的语义。

后续分析。在这些捕获SYCL领域特定语义的更高级别操作被引入代码后，可以运行静态主机代码分析来推断将导致设备代码优化的相关属性。我们的分析利用了第V-B节中描述的到达定值分析，以推断诸如访问器别名或参数常量性等属性，包括隐式参数如访问器范围和偏移量以及内核的ND-range。

B. 主机-设备优化

优化类型。我们联合分析主机和设备代码的方法支持多种优化。虽然在撰写本文时，实现的重点是主机-设备常量传播，但本节末尾讨论的进一步优化可以在未来实现。

利用SYCL方言的优化。利用SYCL方言的领域特定语义，除了传统的常量传播外，还可以实现两种额外的优化：
* 常量ND-range传播：SYCL内核可能在其主体中使用ND-range信息（见第II-A节），例如，为了遍历容器。这些查询通常实现为对平台特定内置函数的调用，这些函数在SYCL-MLIR中被编码为SYCL方言操作。对于常量ND-range信息，相应的getter操作被替换为常量范围或ID。
* 访问器成员传播：DPC++访问器作为四个内核参数传递（参见列表9和第II-A节的上下文）。利用这种参数扁平化，并使用主机代码静态分析，我们不仅可以传播常量访问器成员，还可以推断出两个范围何时相同，从而即使这些范围不是常量，也可以用另一个参数范围的用途替换其中一个。

优化的双重效果。值得注意的是，由于SYCL的异构性，常量传播将导致主机和设备代码的双重优化。在设备端，常量传播可能导致代码优化，如表达式或控制流简化。此外，DPC++的SYCL流水线包括一个SYCL死参数消除（Dead Argument Elimination）遍，它将内核参数标记为未使用。利用这些信息，SYCL运行时将不会将这些参数传递给内核，从而使主机端的内核启动更高效。

性能提升。由主机-设备常量传播带来的性能提升将在第八节中讨论。

未来工作：别名分析的改进。在未来的工作中，可以利用主机和设备代码的联合分析来进行更多的分析和优化。其中一个例子是改进SYCL访问器的别名分析。第五节和第六节中的许多设备分析和优化都依赖于别名分析，因此为SYCL构造提供更好的别名分析结果将支持更强大的分析和优化。

别名分析改进的示例。为了解释主机和设备代码的联合分析如何改进别名分析结果，再次考虑列表8中的SYCL示例。一个只考虑设备代码的别名分析既看不到三个访问器的构造（第4-6行），也看不到底层缓冲区的构造（为简洁起见省略）。因此，它必须假设访问器及其底层指针可能别名，因为SYCL规范允许在同一个缓冲区上定义两个访问器，或者两个不同的缓冲区是另一个缓冲区的重叠子缓冲区。

联合分析的优势。另一方面，主机和设备的联合分析将能够看到主机代码中访问器及其底层缓冲区的构造，因此在许多情况下能够确定两个访问器是否可能别名，这为未来如何进一步扩展主机和设备联合分析提供了一个例子。

未来工作：全局应用优化。另一类通过主机和设备的联合优化可以实现的优化会影响整个应用程序。这类优化的一个例子是设备内核融合。通过合并多个SYCL设备内核，可以减少与内核启动相关的开销，并且通过昂贵的全局内存加载和存储发生的数据流有可能被内化到融合的内核中。对于SYCL，Perez等人在【索引编号16，User-driven online kernel fusion for sycl+2023+ACM Trans. Archit. Code Optim.】中成功地证明了这一点。在他们的工作中，由于在编译时主机和设备是分开编译的，他们必须在运行时使用JIT编译器进行融合，这带来了额外的开销，只能通过编译缓存部分缓解。通过主机和设备代码的联合分析和优化，这类转换可以在编译时完成，从而减少运行时开销。

A4 实验环境

基准测试集
* SYCL-Bench【索引编号15，SYCL-Bench: A Versatile Cross-Platform Benchmark Suite for Heterogeneous Computing+2020+Euro-Par 2020: 26th International European Conference on Parallel and Distributed Computing】，【索引编号17，SYCL-Bench: A Versatile Single-Source Benchmark Suite for Heterogeneous Computing+2020+Proceedings of the International Workshop on OpenCL】：一个代表性的SYCL应用程序基准套件。本文主要关注其中的polybench和single-kernel类别，因为runtime和micro类别主要测试运行时和GPU设备本身性能，而本工作的编译器不改变运行时组件。
* polybench：包含HPC应用中常见的核心负载，如线性代数。问题大小：大多数基准为1024（如2mm, 3mm, GEMM等），Atax和2D Convolution为4096，Bicg、GESUMMV和MVT为16384。
* single-kernel：包含来自图像处理和分子动力学等领域的真实应用和内核。问题大小：大多数基准为1,048,576（如KMeans, Molecular Dynamics等），Linear Regression为65,536，NBody为1024。
* oneAPI Samples【索引编号20，oneAPI samples+2023+oneAPI Samples contributors+https://github.com/ oneapi-src/oneAPI-samples】：由于SYCL-Bench缺少基于模板（stencil）的工作负载，本文从oneAPI示例库中提取了几个基于模板的工作负载作为补充评估。
* 1d HeatTransfer: 一维热传递模拟，包含基于缓冲区-访问器和USM的两种实现。
* iso2dfd: 二维各向同性介质中的波传播模拟。
* jacobi: 使用雅可比迭代求解线性方程组。

硬件配置
* CPU: Intel Xeon Platinum 8480+
* 内存: 503 GiB RAM
* GPU: Intel Data Center GPU Max 1100，48GB RAM

软件配置
* 操作系统: Ubuntu 22.04.2 LTS (Linux kernel 5.15.0)
* SYCL-MLIR/DPC++环境:
* DPC++ 版本: 3482e2d1
* Intel Level Zero驱动版本: 1.3.26690
* AdaptiveCpp环境:
* AdaptiveCpp 版本: c33e8c42
* 构建依赖: LLVM release 17, Boost version 1.77

实验方法
* 比较对象：SYCL-MLIR、Intel DPC++编译器【索引编号12，DPC++ open-source SYCL Compiler+2023+Intel Corporation+https://github.com/intel/llvm】（作为基线）、AdaptiveCpp【索引编号18 ，One pass to bind them: The first single-pass sycl compiler with unified code representation across backends+2023+ser. IWOCL ’23】。
* 性能度量：报告的运行时间和加速比是舍弃第一次运行（用于预热设备驱动）后，对30次运行取平均值的结果。

A4 实验结果

Single-Kernel 基准测试
* 实验内容：在single-kernel类别的基准测试上，比较SYCL-MLIR和AdaptiveCpp相对于DPC++的性能。
* 实验结果：
* AdaptiveCpp在一些情况下实现了加速，但也存在一些小的性能下降，总体几何平均加速比为1.03倍。
* SYCL-MLIR在这些基准上的性能提升虽然不大但很明显，对许多基准有小幅加速，也有少数小幅性能下降，几何平均加速比为1.02倍。
* 图表引用：结果如图2所示。一些基准测试中AdaptiveCpp的结果验证失败，因此图中缺少对应的条形。

Polybench 基准测试
* 实验内容：在polybench类别的基准测试上，比较SYCL-MLIR和AdaptiveCpp相对于DPC++的性能。
* 实验结果：
* AdaptiveCpp相对于DPC++取得了显著的加速，在SYR2K基准上加速接近3倍，其余情况大多持平（除MVT外），几何平均加速比为1.22倍。
* SYCL-MLIR除了少数轻微的性能回归外，相对于DPC++取得了高达4.32倍的加速，几何平均加速比为1.45倍。
* 分析结论：
* 数组归约优化 (VI-B节)：对多个基准测试（如Correlation和Covariance）有显著贡献。这些基准通常包含多个数组归约的机会。
* 循环内化优化 (VI-C节)：对五个基准测试（2mm, 3mm, GEMM, SYR2K, SYRK）的性能提升贡献显著。编译器跟踪显示，优化能够将多个数组引用预取到局部内存中（例如，GEMM中有2个，SYR2K中有4个）。Gramschmidt基准测试因候选循环位于分化区域而未被优化。
* 主机-设备传播 (VII-B节)：也对性能提升有贡献。例如，在Sobel7基准中，声明为常量数组的Sobel滤波器可以被传播到设备代码中以提高性能。
* 图表引用：结果如图3所示。

整体 SYCL-Bench 性能
* 在整个SYCL-Bench测试集上，SYCL-MLIR相对于DPC++实现了1.18倍的几何平均加速比，并且性能优于AdaptiveCpp（几何平均加速比1.13倍）。

Stencil-based 工作负载
* 实验内容：评估SYCL-MLIR在oneAPI示例中的四个基于模板的工作负载上的性能。
* 实验结果：
* AdaptiveCpp在iso2dfd上实现了1.5倍的加速，但在其他模板工作负载上执行失败。
* SYCL-MLIR在1D heat transfer上表现出轻微性能下降（缓冲区版本为0.86倍，USM版本为0.87倍）。
* 在iso2dfd（0.99倍）和jacobi（1倍）上，SYCL-MLIR与DPC++性能相当。
* 分析结论：初步调查显示，本文描述的设备优化（VI节）目前没有应用于这些工作负载。未来的工作重点是为基于模板和卷积的工作负载添加新的优化或扩展现有优化。

A7 补充细节

IX. 相关工作

C/C++ 到 MLIR 的编译。目前在LLVM社区中有多种将C/C++代码编译到MLIR的方法正在开发中。Polygeist【索引编号13，Polygeist: Raising C to polyhedral MLIR+2021+30th International Conference on Parallel Architectures and Compilation Techniques】支持将Clang AST的一个子集直接转换为标准MLIR方言和少量自定义操作的混合体。输入程序中的仿射循环和其他结构化控制流构造被保留并暴露给分析和转换。我们扩展了Polygeist来编译SYCL设备代码，但它目前仍然缺乏完全翻译主机代码的能力，因为虚函数和异常尚不支持。Polygeist也试图最小化自定义操作的引入，这导致了一些MLIR验证规则的放宽。虽然这在该工具的使用场景中是可行的，但与MLIR生态系统其余部分的集成可能会更加困难。ClangIR【索引编号21，[RFC] An MLIR based Clang IR (CIR)+2022+post on LLVM Discourse+https://discourse.llvm.org/t/ rfc-an-mlir-based-clang-ir-cir/63319/1】是另一个LLVM孵化器项目，旨在从Clang AST实现基于MLIR的静态分析和代码生成。与Polygeist不同，该项目计划开发一个覆盖C/C++全部表面的方言，然后直接降低到LLVM方言，跳过MLIR内置的更高级别方言，这有助于正确表示源语言的语义。ClangIR仍处于早期开发阶段，但是未来替代我们流程中主机代码提升的一个有希望的候选者。VAST【索引编号22，Finding bugs in C code with Multi-Level IR and VAST+2023+post on Trail of Bits blog+https://blog.trailofbits.com/2023/06/15/ finding-bugs-with-mlir-and-vast】将采用一个“IR塔”来在C/C++代码编译的不同抽象级别上进行程序分析。

异构编程模型的主机-设备优化。之前的工作已经尝试过对异构编程模型进行主机-设备优化。Singer等人【索引编号23，Syclops: A SYCL specific LLVM to MLIR converter+2022+IWOCL’22: International Workshop on OpenCL】提出了SYCLOps，一个从LLVM-IR到MLIR标准方言的SYCL专用转换器。虽然在动机上与我们的工作相似，但他们的工具只提升设备代码，并且没有定义一个方言来捕获SYCL特定的操作。Moses等人【索引编号24，High-performance gpu-to-cpu transpilation and optimization via high-level parallel constructs+2022+vol. abs/2207.00257】扩展了Polygeist来将CUDA应用程序编译到MLIR，同时表示主机和设备代码，并保持并行构造和屏障的完整性。这为跨屏障以及并行区域外的代码移动提供了各种机会。Tian等人【索引编号25，Just-in-time compilation and link-time optimization for openmp target offloading+2022+OpenMP in a Modern World: From Multi-device Support to Meta Programming - 18th International Workshop on OpenMP, IWOMP 2022】在LLVM中实现了OpenMP内核的运行时和链接时特化，例如常量传播。

AdaptiveCpp项目。独立于MLIR，AdaptiveCpp项目正在采用一种不同的SYCL编译方法。最初，该项目（以前称为hipSYCL）【索引编号26，Exploring the ¨ possibility of a hipsycl-based implementation of oneapi+2022+IWOCL’22】遵循一种仅库（library-only）的SYCL实现方法，将SYCL设备内核映射到合适的低级编程模型，如OpenMP或CUDA，并依赖相应的编译器，例如Nvidia的nvc++。

AdaptiveCpp的编译器方法。在较新版本中，如第八节所用版本，AdaptiveCpp在其实现中也增加了一种基于编译器的方法【索引编号18，One pass to bind them: The first single-pass sycl compiler with unified code representation across backends+2023+ser. IWOCL ’23】，但这与其他SYCL编译器不同。在所谓的单源多编译器遍（SMCP【索引编号2，SYCL 2020 specification (revision 7)+2023+Khronos+https://registry.khronos.org/ SYCL/specs/sycl-2020/pdf/sycl-2020.pdf】）编译流程中，例如图1所示的DPC++编译流程，源代码被处理多次，一次在主机编译期间，一次在每次设备编译期间（即每个用户指定的目标架构）。而AdaptiveCpp遵循所谓的单源单编译器遍（SSCP【索引编号2，SYCL 2020 specification (revision 7)+2023+Khronos+https://registry.khronos.org/ SYCL/specs/sycl-2020/pdf/sycl-2020.pdf】）方法，源代码只被处理一次。为此，AdaptiveCpp流程使用LLVM IR作为中间交换格式，并在单次编译遍中将设备代码的LLVM IR嵌入到应用程序二进制文件中。在运行时，设备内核的LLVM IR在内核启动时被检索出来，并进一步编译成适当的设备特定格式，例如Intel设备的SPIR-V或Nvidia CUDA GPU的PTX。

编译时机的差异。从SYCL源代码到LLVM IR以及之后到设备格式的编译是DPC++和AdaptiveCpp的共同点，但不同之处在于第二步发生的时间。对于DPC++，这发生在编译时；对于AdaptiveCpp，这发生在应用程序运行时的内核启动时。

AdaptiveCpp与SYCL-MLIR的共同目标与差异。因此，AdaptiveCpp的编译流程与SYCL-MLIR有一个共同目标：通过仅在运行时执行第二步编译，AdaptiveCpp可以将关于设备调用上下文的信息从主机传播到设备编译器，因为这些信息在编译发生的内核启动时是可用的。SYCL-MLIR也试图传播这些信息，但它是在应用程序编译时通过利用MLIR框架来实现主机和设备代码的联合分析（参见图1和第七节）。

两种方法的优缺点。AdaptiveCpp方法的优点是，即使是运行时的值也可以被考虑用于设备编译，这可能比SYCL-MLIR在编译时的主机和设备联合分析提供更多信息。另一方面，运行时的编译步骤给内核启动带来了额外的开销。即使这种开销可以通过缓存编译来减少，该缓存也不会在不同的应用程序运行之间持久化，这对于像DPC++或SYCL-MLIR这样只需要编译一次的方法来说是一个缺点。

A5 结论

工作总结
本文展示了构建一个基于MLIR的SYCL异构编程模型编译器的实践经验。通过利用MLIR的方言框架，本文中的SYCL方言在一个高抽象层次上捕获了SYCL API在主机和设备代码中的关键元素，使得编译流程能够访问到SYCL的语义，例如工作项并行执行和设备内存访问。在此方言的基础上，本文实现了强大的设备优化以及能够跨越SYCL主机和设备代码边界进行推理的分析。

成果
利用这些分析和优化，本文的编译流程在一系列SYCL基准应用程序上，相比于一个先进的、基于LLVM的SYCL编译器，取得了高达4.3倍的加速，并且对于基于循环的工作负载尤其有效。

未来展望
尽管在评估中取得了加速，但仍有改进空间：现有C++前端对MLIR的限制目前迫使编译流程对SYCL主机代码执行从LLVM IR到MLIR的提升（raising），因为像C++异常这样的重要构造尚未完全支持。这目前限制了编译器执行跨主机-设备边界优化的能力，例如将设备代码提升到主机。随着第九节中讨论的C++前端的发展，未来的工作可以实现这类优化，并能推理SYCL应用程序的整体结构以执行更高级的优化，如设备内核融合。

参考文献引用说明

• 【索引编号1，MLIR: Scaling compiler infrastructure for domain specific computation+2021+IEEE/ACM International Symposium on Code Generation and Optimization (CGO)】

  • 引用位置：II.B节，第一段
  • 引用原文描述：介绍了MLIR框架的诞生动机，即过早将高级编程模型降低到低级表示会损害优化能力，且现有IR难以扩展。

• 【索引编号2，SYCL 2020 specification (revision 7)+2023+Khronos+https://registry.khronos.org/ SYCL/specs/sycl-2020/pdf/sycl-2020.pdf】

  • 引用位置：II.A节，第一段；IX.节，第四段
  • 引用原文描述：在II.A节中，作为SYCL开放标准的官方规范被引用。在IX.节中，用于解释SMCP和SSCP编译流程术语的来源。

• 【索引编号3-8】

  • 引用位置：II.A节，第十段
  • 引用原文描述：作为SYCL与其他编程模型（如CUDA, OpenCL, OpenMP）进行详细比较的参考文献。

• 【索引编号9-11】

  • 引用位置：II.B节，第三段
  • 引用原文描述：指出在MLIR之前，在编译器中使用多个中间表示的好处已在文献中得到广泛认可，并列举了相关工作。

• 【索引编号12，DPC++ open-source SYCL Compiler+2023+Intel Corporation+https://github.com/intel/llvm 】

  • 引用位置：IV.节，第一段；VIII.节，第四段
  • 引用原文描述：在IV.节中，作为SMCP编译流程的一个具体例子。在VIII.节中，作为实验评估的基线编译器。

• 【索引编号13，Polygeist: Raising C to polyhedral MLIR+2021+30th International Conference on Parallel Architectures and Compilation Techniques】

  • 引用位置：IV.节，第三段；IX.节，第一段
  • 引用原文描述：在IV.节中，被用作将C++ AST转换为MLIR的设备编译器。在IX.节中，作为相关工作中C/C++到MLIR编译方法之一进行讨论。

• 【索引编号14，Exploiting memory access patterns to improve memory performance in data-parallel architectures+2011+IEEE Transactions on Parallel & Distributed Systems】

  • 引用位置：V.D节，第二段；VI.C节，第六段
  • 引用原文描述：在V.D节中，作为内存访问分析的理论基础。在VI.C节中，用于判断内存访问是否可以合并的Linear或ReverseLinear模式的来源。

• 【索引编号15，SYCL-Bench: A Versatile Cross-Platform Benchmark Suite for Heterogeneous Computing+2020+Euro-Par 2020: 26th International European Conference on Parallel and Distributed Computing】

  • 引用位置：VI.节，第一段；VIII.节，第二段
  • 引用原文描述：在VI.节中，提到本文的设备优化能为该基准套件中的多面体基准带来显著加速。在VIII.节中，作为主要的实验评估基准套件。

• 【索引编号16，User-driven online kernel fusion for sycl+2023+ACM Trans. Archit. Code Optim.】

  • 引用位置：VII.B节，第八段
  • 引用原文描述：作为在SYCL中成功实现设备内核融合的一个例子，并指出其工作是在运行时通过JIT编译器完成的。

• 【索引编号17，SYCL-Bench: A Versatile Single-Source Benchmark Suite for Heterogeneous Computing+2020+Proceedings of the International Workshop on OpenCL】

  • 引用位置：VIII.节，第二段
  • 引用原文描述：与[15]共同作为SYCL-Bench基准套件的引用。

• 【索引编号18，One pass to bind them: The first single-pass sycl compiler with unified code representation across backends+2023+ser. IWOCL ’23】

  • 引用位置：VIII.节，第五段；IX.节，第四段
  • 引用原文描述：在VIII.节中，作为实验评估的另一个比较对象。在IX.节中，作为AdaptiveCpp编译器方法的参考文献。

• 【索引编号20，oneAPI samples+2023+oneAPI Samples contributors+https://github.com/ oneapi-src/oneAPI-samples】

  • 引用位置：VIII.节，第九段
  • 引用原文描述：作为补充评估中基于模板的工作负载的来源。

• 【索引编号21-26】

  • 引用位置：IX.节（相关工作）
  • 引用原文描述：在相关工作部分，分别用于介绍ClangIR、VAST、SYCLOps、Moses等人的CUDA到MLIR工作、Tian等人的OpenMP JIT编译工作以及hipSYCL的相关研究。