C-for-Metal: High Performance SIMD Programming on Intel GPUs
C-for-Metal: High Performance SIMD Programming on Intel GPUs
C-for-Metal: 面向Intel GPU的高性能SIMD编程
作者/机构: Guei-Yuan Lueh, Kaiyu Chen, Gang Chen, Joel Fuentes, Wei-Yu Chen, Fangwen Fu, Hong Jiang, Hongzheng Li, and Daniel Rhee, Intel Corporation, Santa Clara, CA, USA
A1 主要贡献
本文针对在Intel GPU上，主流的“单指令多线程”（SIMT）编程模型（如CUDA和OpenCL）与底层“单指令多数据”（SIMD）硬件架构之间的性能鸿沟问题，提出了一种名为 C-for-Metal (CM) 的显式SIMD编程框架。
核心问题:
主流的SIMT模型（如OpenCL）虽然简化了并行编程，但在Intel GPU（代号Gen）这种原生SIMD架构上存在严重的性能瓶颈。开发者编写的标量代码由编译器和硬件隐式并行化，导致无法充分利用硬件的关键能力，具体表现为：
1. 寄存器文件控制困难：OpenCL语言缺乏直接控制寄存器文件的机制，而有效利用寄存器以减少内存访问对Gen架构至关重要。编译器的复杂优化使得源码级别的寄存器压力估算极不准确。
2. 跨通道数据共享受限：SIMT模型抽象阻止了同一硬件线程内不同SIMD通道间的数据共享，即使它们共享同一个寄存器文件。这导致了冗余的计算和内存操作。虽然OpenCL 2.0引入了subgroup来缓解此问题，但仍不够高效。
3. 向量长度控制缺失：Gen ISA支持可变SIMD大小的指令，这对于优化寄存器压力和处理分支分化非常重要。而OpenCL编译器对整个内核进行隐式矢量化，通常使用固定的SIMD分派大小，缺乏指令级的灵活性。
研究目标与创新点:
为了解决上述性能差距，本文介绍了C-for-Metal (CM) 开发框架，这是一个专为在Intel GPU上实现“接近硬件”性能而设计的显式SIMD编程模型。其核心创新点包括：
1. 提供显式SIMD编程接口：CM语言是C/C++的扩展，通过引入特殊的vector和matrix类型，提供了直观的接口来表达数据并行性。这些变量默认分配在寄存器中，使开发者能在源码层面更精确地控制寄存器使用。
2. 直接映射硬件特性：CM内核描述的是整个硬件线程的算法，而非单个工作项。其内置操作，特别是select操作符，能够直接映射到Gen ISA强大的区域（regioning）寻址能力，实现零开销的跨通道数据访问和重组。
3. 精细的SIMD大小控制：程序员可以通过改变select操作返回的元素数量，显式地控制每条指令的SIMD大小，从而根据寄存器需求和分支分化情况进行动态优化。
4. 编译器与运行时支持：CM编译器（CMC）基于LLVM，能将高级的向量和矩阵操作高效地生成Gen ISA SIMD指令。CM内核与Intel GPU的OpenCL运行时和oneAPI Level Zero完全兼容，可直接启动。
通过在不同领域的代表性应用上的实验评估，本文证明了CM相比于使用Intel特定扩展优化过的最佳OpenCL实现，仍能取得最高达2.7倍的性能提升，并且在开发效率上（考虑性能调优时间）比OpenCL高出2-3倍。
A3 新编程模型的动机
SIMT模型在Gen架构上面临的挑战。本文描述了以OpenCL为代表的SIMT模型在Intel GPU上遇到的三个主要挑战，以此正式阐明对CM的需求。
1. 寄存器文件控制。有效利用寄存器文件以减少不必要的内存流量是Intel GPU上最重要的优化策略【索引27，Register allocation for Intel processor graphics，2018，CGO】。在OpenCL中，精确管理寄存器压力非常困难，因为语言将寄存器分配的决策完全交给了编译器。一个OpenCL内核在编译成Gen汇编代码的过程中会经过数百个编译器转换和优化过程；其中大部分都会对寄存器压力产生显著影响，但其行为对程序员来说既不透明，也通常无法控制。例如，分化分析【索引28，Divergence analysis and optimizations，2011，PACT】是SIMT GPU编译器的关键分析环节，其结果可用于通过为变量分配标量寄存器来减少寄存器使用（如果能证明所有通道持有相同的值）。但在复杂的数据和控制依赖关系存在时，分析结果往往过于保守，且没有提供任何机制让程序员辅助分析。相比之下，CM变量默认在寄存器中分配，并且向量和矩阵可以在硬件限制内拥有任意大小。因此，CM开发者可以直接将他们的统一变量分配在一个寄存器中，也可以将多个变量合并成一个大的矩阵以进行显式的生命周期管理。
2. 跨通道数据共享。SIMT执行模型一个众所周知的局限性是硬件线程中的工作项之间缺乏数据共享。尽管一个线程中的SIMD通道共享寄存器文件，但SIMT抽象阻止了一个通道访问另一个通道的寄存器数据，这不可避免地导致了冗余的计算和内存操作。CUDA和OpenCL都引入了显式的SIMD原语来促进跨通道通信，提供的功能包括shuffle、reduction和barrier操作【索引29，Using CUDA Warp-Level Primitives，2018，NVIDIA Developer Blog】、【索引30，The OpenCL Extension Specification，2018，Khronos】。这些扩展帮助弥合了SIMT模型与底层SIMD硬件之间的差距，但它们并不代表实际的硬件能力。相比之下，CM的select操作直接映射到硬件的区域寻址（regioning），并且可以直接用于计算指令中，从而消除了不必要的shuffle移动操作。
3. 向量长度控制。每条Gen ISA指令都有其自己的执行大小，而指令级的SIMD大小可以成为一项重要的优化技术。改变向量大小的一个直接用途是控制寄存器压力。大多数应用程序会经历高寄存器需求和低寄存器需求的阶段，一个内核应该混合使用其SIMD大小，以避免在高压力区域发生寄存器溢出，同时在向量内存的gather/scatter操作中实现最大带宽。类似地，分支分化会显著降低程序的效率【索引31，Taming control divergence in gpus through control flow linearization，2014，CC】、【索引32，Reducing branch divergence in gpu programs，2011，GPGPU】；在没有硬件机制的情况下，非活动通道在控制流重新汇合之前不会执行。通过在分化区域内以较低的SIMD大小运行，内核可以减少浪费的工作量。由于CM的显式SIMD模型，程序员可以通过向量和矩阵select的大小轻松控制每条指令的SIMD大小。然而，SIMT模型不提供此类功能，因为OpenCL GPU编译器对内核执行隐式矢量化。一个OpenCL内核可以指定其分派大小，但所有非统一指令将默认使用该大小。
线性滤波器示例对比。本文使用一个简单的3x3盒式模糊滤波器（即线性滤波器）来对比CM和OpenCL的编程模型。首先展示一个直接的OpenCL实现，并指出其在Intel GPU上的低效之处。
// Algorithm 1 Linear filter in OpenCL with SIMT model
1: kernel LINEAR(image2d src, image2d dst, int width, int height)
2: int x = get_global_id(0);
3: int y = get_global_id(1);
4: float4 pixel1 = 0.0f;
5: float4 pixel = 0.0f;
6: int tempx, tempy;
#pragma unroll
7: for i = -1; i <= 1; i++ do
#pragma unroll
8: for j = -1; j <= 1; j++ do
9: tempx = min(width-1, max(0, x+j));
10: tempy = min(height-1, max(0, y+i));
11: pixel1 = read_imagef(src, sampler, (int2)(tempx, tempy));
12: pixel.z += pixel1.z;
13: pixel.y += pixel1.y;
14: pixel.x += pixel1.x;
15: end for
16: end for
17: uint4 p = convert_uint4(pixel * 0.1111f);
18: write_imageui(dst, (int2)(x, y), p);
19: end kernel
OpenCL实现分析。在算法1中，每个工作项通过计算其在输入图像中邻居像素的平均值来计算一个像素的结果，该像素的位置由工作项的x和y全局ID指示。Intel的OpenCL编译器将此内核矢量化为SIMD16指令，其中每个通道对应于输入和输出图像中的一个像素。输入和输出图像都是3通道RGB格式，硬件图像读取单元将每个通道的8位整数转换为结构数组（SoA）格式的归一化浮点值。图像写入操作则执行相反的格式转换。生成的汇编代码包括9次图像gather加载（第11行）、27次浮点加法（第12-14行）和一次图像scatter写入（第18行）。这个简单的实现在每个硬件线程中存在严重的冗余加载，因为在一次迭代中，每个工作项都在读取其相邻通道在之前迭代中已经加载过的像素值。一种更高效的方法是让一个线程中的工作项协作加载图像的一个2D块的原始格式数据（即像素被加载到寄存器中而不进行格式转换），然后将每个通道转换为浮点值以进行后续计算。这种特殊的2D块读/写功能由Intel的cl_intel_media_block_io扩展提供。
SIMT模型的局限性。然而，这种方法的有效性仍然受到SIMT模型的限制，因为内置函数的返回数据必须均匀地分布在一个子组（subgroup）中的工作项之间。因此，需要一个子组shuffle操作来读取邻近通道的像素，并将它们从结构数组（AoS）转换为数组结构（SoA）布局。OpenCL编译器通常无法优化掉这些昂贵的移动操作，因为它必须维持SoA格式的计算值以满足SIMT模型。作为最后的手段，可以通过在主机代码中转置输入图像来避免shuffle移动，但这会增加CPU开销，并且现实世界的应用程序不一定能控制其输入布局。
A2 方法细节
IV. CM编程语言
CM语言基础。CM编程语言使用Clang实现，支持标准C++的一个子集，并有一些限制（更多细节见CM语言规范【索引6，C-for-Metal Compiler，2019，GitHub】的2.6节）。在Clang的基础类型系统中添加了两个容器类型：vector和matrix。这两个新的基础类型构成了CM显式SIMD编程模型的基础。在这两种类型之上，我们添加了与Gen指令集非常相似的操作和内置函数。这些新类型和函数共同构成了在Gen上进行“接近硬件”编程的抽象接口。
A. Vector和Matrix类型
类型定义。这些类型使用类似于C++模板类的语法定义。参数是数据元素的类型和向量/矩阵的大小。元素类型必须是CM支持的基本类型之一，大小必须是正整数和编译时常量。
vector<short, 8> v; // 一个包含8个short的向量
matrix<int, 4, 8> m; // 一个4x8的整数矩阵
引用类型。此外，CM提供了两种引用组件数据类型：vector_ref和matrix_ref。它们定义了对基本向量或矩阵对象的引用。引用变量不会分配额外的内存空间。例如，矩阵m的第二行可以定义为一个引用变量：vector_ref<int, 8> vref(m.row(2));
硬件映射与访问。向量或矩阵变量映射到Gen硬件通用寄存器文件（GRF）中一系列连续的元素。向量或矩阵变量不能被取地址；间接访问通过引用类型来执行。引用变量通常由对基础变量的操作构造而成，这些操作为基础对象提供了不同的视图。读取一个引用变量会直接映射到Gen的基于区域（region）的寻址方案，该方案在两个寄存器内提供零开销的数据打包、解包和重排。
核心操作。对于向量、矩阵及其对应的引用变量，CM支持成员函数和操作，包括构造函数和赋值；算术、移位、逻辑和比较；以及行、列和元素访问。CM向量和矩阵类型独有的主要操作包括：
* select: 支持一系列select函数，用于引用向量/矩阵元素的子集。每个select操作返回对基础对象元素的引用，并且可以作为左值表达式使用。select操作的形式为（其中v是向量，m是矩阵）：
v.select<size,stride>(i)
m.select<vsize,vstride,hsize,hstride>(i,j)
- iselect: CM允许用户对另一个向量进行索引访问。间接
select总是右值表达式。例如，考虑一个包含16个浮点数的基础变量v，以及一个包含4个元素{0, 1, 2, 2}的向量idx。那么表达式v.iselect(idx)可以用来创建一个新向量，其元素为{v[0], v[1], v[2], v[2]}。此函数暴露了Gen的寄存器间接寻址能力。 - merge: 提供了两种形式的
merge操作以支持条件更新：v.merge(x, mask)和v.merge(x, y, mask)。前者在对应掩码位为真时，将x的元素复制到v。后者在对应掩码位为真时，将x的元素复制到v；否则，将y的元素复制到v。第一种merge映射到Gen的谓词mov指令，而第二种merge映射到sel指令。 - format: 此操作允许重新解释矩阵/向量变量的元素类型并改变其形状。例如，对于一个包含8个浮点数的向量v，表达式
v.format<char, 4, 8>()的类型为matrix_ref<char, 4, 8>，意味着v被重新解释为一个4行8列的char类型矩阵。 - replicate: 此操作提供通用的区域操作，用于从向量或矩阵中收集元素。表达式
v.replicate<K, VS, W, HS>(i)从输入向量v的位置i开始，收集K个块，每个块有W个元素。VS和HS是垂直和水平步长。例如，对图1中的向量v执行v.replicate<2, 4, 4, 0>(2)将收集元素{v[2], v[2], v[2], v[2], v[6], v[6], v[6], v[6]}。
混合操作与类型转换。CM也支持不同形状的向量和矩阵对象的混合操作，只要每个操作数具有相同数量的元素即可。操作数形状的一致性在编译时使用向量/矩阵类的模板特化规则进行检查。CM编译器根据源操作数数据类型，遵循标准的C++类型提升规则（使用模板特化机制）来确定目标操作数的元素类型。与标准C++一样，用户可能需要添加显式类型转换来改变默认的类型提升和转换规则。一个简单的隐式和显式转换示例如下：
vector<float, 8> f;
vector<int, 8> i;
f = i; // 隐式转换
f = vector<short, 8>(i); // 显式转换
全局变量与寄存器绑定。CM允许将向量和矩阵声明为文件作用域变量，这些变量被视为线程私有变量。它们可用于促进同一线程中主函数及其调用函数之间的数据共享。可选地，CM支持两种全局变量用法。第一种，用 _GENX_VOLATILE_ 限定符表示，通知编译器对这些变量执行保守优化，以降低寄存器压力并提高代码质量。第二种，用 _GENX_VOLATILE_BINDING_(Offset) 限定符表示，指示该全局变量应映射到从指定字节偏移量开始的GRF块。这种寄存器绑定功能使程序员能够实现细粒度的寄存器分配控制，并有效应对其他挑战，如性能关键型应用中的bank冲突。
B. Memory Intrinsics
内存访问函数。CM提供了一组内存访问函数，这些函数与底层的Gen硬件操作相似。默认情况下，使用基于缓冲区索引的寻址模式。一个内核包含多个SurfaceIndex参数，每个参数代表一个底层内存对象的句柄。一个read或write内建函数接受一个surface index，并访问由偏移量指定的元素。应用程序主机代码负责通过运行时API调用将每个内核参数绑定到一个内存对象。最有用的内建函数包括：
* 2D-block read/write: 对于由其SurfaceIndex标识的图像，块读取操作将给定x/y位置的一个像素块加载到一个矩阵中。2D块写入操作将一个矩阵存储到图像中给定x/y位置的一个像素块中。以下是2D块读取的内建函数定义。
template<typename T, int N, int M>
void read(SurfaceIndex index, CmBufferAttrib attr, int X, int Y, matrix_ref<T, N, M> output)
- Oword-block read/write: 对于线性寻址的缓冲区，块读取操作在给定偏移处读取一个连续的oword序列（每个oword为16字节）到一个向量中。块写入操作将一个向量写入到缓冲区给定偏移处的一个连续oword序列中。以下是Oword块读取的内建函数定义。
template<typename T, int N>
void read(SurfaceIndex idx, CmBufferAttrib attr, int offset, vector_ref<T, N> output)
- Scattered read/write: 还支持各种粒度的向量gather和scatter。要读取/写入的每个元素的零基偏移量（相对于一个全局偏移量）在一个向量中指定。对于scattered read和write函数，地址、源数据和返回数据的类型必须是大小相同的向量。以下是scattered read的内建函数定义。
template <typename T, int N>
void read(SurfaceIndex index, uint globalOffset, vector<uint, N> elementOffset, vector_ref<T, N> ret)
- Atomics: CM支持Gen上所有的原生原子操作，包括and、add、max、inc、compxchg等。与scattered read/write类似，原子函数的参数也必须是向量类型。以下是atomic inc的内建函数定义。
template<CmAtomicOp Op, typename T, int N>
void write_atomic(vector<ushort, N> mask, SurfaceIndex index, vector<uint, N> element_offset)
平坦寻址模型。除了SurfaceIndex，CM还支持平坦寻址模型，其中内核参数是一个指针，可以直接用于内存访问。这允许主机和内核代码共享数据结构并并发访问它们。
C. Boolean Reductions
布尔归约函数。为了方便对掩码向量进行布尔归约，CM提供了两个预定义的布尔函数：
ushort vector<ushort, size>::any(void)
ushort vector<ushort, size>::all(void)
any()在掩码中有任何非零值时返回1；否则返回0。all()在掩码中所有值都非零时返回1；否则返回0。请注意，这些函数也适用于矩阵类型。任一函数的结果都可以用作标量值，并用于标准的C++控制流结构中。归约函数被高效地映射到Gen的比较指令。
D. SIMD Control Flow
默认标量控制流。在CM中，默认的控制流语句就是C++的标量控制流语句——条件语句（if-else/switch）、循环语句（for/while/do-while）、跳转语句（break/continue/goto/return）或函数调用。对于这些语句，条件必须是标量，并且所有SIMD通道都统一分支。
每通道SIMD控制流。除此之外，CM还提供了每通道SIMD控制流机制，利用Gen的simd-goto和simd-join指令，支持SIMD执行下的分化控制流【索引33，IGC: the open source Intel Graphics Compiler，2019，CGO】。此功能为对长指令序列进行谓词化提供了一种替代方案，因为非活动通道在SIMD控制流区域内不执行。
SIMD控制流宏。CM中的SIMD控制流通过预定义的C++宏来表示。例如，一个分化的if由SIMD_IF_BEGIN和SIMD_IF_END宏表示，使用方式如下：
vector<uint, 16> v(0);
vector<ushort, 8> cond = ...
SIMD_IF_BEGIN(cond > 0){ // ...
v.select<8, 2>(0) = 1;
}SIMD_ELSE{
v.select<8, 2>(1) = 1;
}SIMD_IF_END;
比较cond > 0产生一个向量掩码，该掩码决定了哪个通道是活动的。then语句和else语句都可能为其活动的通道执行。如果所有通道都处于非活动状态，则SIMD控制流块将被跳过。注意，SIMD控制流内的SIMD操作的大小必须与掩码大小相同或者是标量。
E. Linear Filter in CM
CM实现线性滤波器。现在我们描述如何在CM中实现线性滤波器（算法2）。CM内核中的每个线程读取一个8x32字节的矩阵，并输出一个对应于6x8像素的6x24字节矩阵。虽然我们只需要8x30字节来表示8x10的输入像素，但为每行添加两字节的填充可以在寄存器文件中为计算提供一个良好的布局。select操作的作用如下：在输入像素被加载到8x32字节的矩阵m中后，在每个步骤中，我们通过一个select操作提取一个6x24字节的子矩阵，将所有元素转换为浮点数，然后将它们加到累加和中，累加和是一个6x24的浮点矩阵。图2展示了算法2中执行的第一个6x24字节子矩阵select操作。
内存操作优化。2D块读/写函数用于执行第5行和第18行的加载和存储操作。如第III节所述，对于此滤波器，专门的2D块消息比原生OpenCL实现（算法1）中的图像gather/scatter操作效率高得多，因为它消除了冗余的内存流量。
V. CM编译器
编译器架构。与Intel图形编译器（IGC）【索引33，IGC: the open source Intel Graphics Compiler，2019，CGO】类似，CM编译器由三层组成：
* 前端：clang前端编译器【索引34，LLVM and Clang: Next generation compiler technology，2008，BSD conference】将CM源代码转换为LLVM中间表示（IR）【索引3，LLVM: A compilation framework for lifelong program analysis & transformation，2004，CGO】。
* 中端：中端执行通用和CM特定的优化与转换，然后将LLVM IR转换为虚拟ISA（vISA）汇编语言。vISA与Gen ISA非常接近，但作为编译目标更方便，因为它有无限的虚拟寄存器并隐藏了各种硬件特定的限制。
* 后端（Finalizer）：vISA finalizer【索引27，Register allocation for Intel processor graphics，2018，CGO】是Intel GPU的代码生成器。它以vISA汇编作为输入，执行局部优化、寄存器分配和调度，以生成目标Intel GPU的最终指令。
编译流程。CM自定义优化的总体流程如图3所示（在中端模块内）。输入对应于由LLVM通用优化生成的LLVM IR。降低（lowering）过程逐渐将高级的CM语言结构转换为更接近目标Gen ISA的代码序列。之后，在每个IR级别上执行若干优化以提高代码质量。
LLVM IR扩展。Gen ISA具有独特的特性，如可变的执行大小、混合数据类型、灵活的寄存器区域寻址和修饰符支持【索引33，IGC: the open source Intel Graphics Compiler，2019，CGO】。向量和矩阵数据类型及其区域选择操作需要被仔细建模，以便能够直接映射到这些独特特性，而无需额外的移动指令。由于LLVM基于静态单赋值（SSA）形式，其中每个值只被定义一次，我们用以下两个内建函数扩展其IR，以在SSA形式下建模对向量/矩阵变量的部分读/写，使其能受益于常见的LLVM优化。
* Read region (rdregion)：从一个向量中提取选定的元素，以创建一个新的较小向量。
* Write region (wrregion)：将元素插入到选定位置，并为旧向量返回一个新值。
IR内建函数示例。以下是一个简化的例子来说明该设计。原始向量a被定义为一个8 x i32的值%a0。rdregion内建函数根据给定参数（垂直步长=0，宽度=4，水平步长=2，起始字节偏移=4）从%a0中提取4 x i32个元素。wrregion内建函数根据其他给定参数（垂直步长=0，宽度=4，水平步长=2，起始字节偏移=0）将%b的元素插入到a的旧值（%a0）中。SSA属性得以维持，因为wrregion内建函数返回一个不同的%a1来表示向量a的新值。
vector<int, 8> a(init_v);
vector<int, 4> b;
b = a.select<4, 2>(1);
a.select<4, 2>(0) = b;
%a0 = <8xi32> ...
%b = call <4xi32> @llvm.genx.rdregioni...(<8xi32> %a0, i32 0, i32 4, i32 2, i16 4);
%a1 = call <8xi32> @llvm.genx.wrregioni...(<8xi32> %a0, <4xi32> %b, i32 0, i32 4, i32 2, i16 0);
打包（Baling）与合法化（Legalization）。由于其表达能力，一条vISA指令可能在LLVM IR中由多条指令表示。打包（Baling）是确定哪一组LLVM指令可以组合（baled）在一起并高效映射到vISA的过程。一个bale有一个根指令以及源和目标操作数上的可选修饰符和区域指令。打包分析过程构建一个映射来标记哪些IR指令被选中以及它们在最终bale中扮演的角色。一个bale的根是该bale中所有指令在程序顺序中的最后一条指令，也是其值在bale外部被使用的唯一指令。由于打包过程可能会决定将一个有多个用途的指令打包为非根指令，因此会克隆该指令以确保其在bale内部只有一个用途。vISA被设计为接近Gen ISA，并继承了类似的限制（例如，一个操作数的大小不能超过两个GRF）。在初始打包分析之后，合法化（legalization）过程可能会将一个bale拆分成多条指令以符合vISA的限制。通常，拆分必须小心进行，以利用目标平台允许的最大SIMD宽度。此处执行的其他转换示例包括：由于冲突的合法化要求而取消打包指令、为内存访问操作对齐操作数，以及将字节类型操作提升为等效的短整型操作以绕过硬件限制。
向量优化。向量优化过程基于rdregion和wrregion执行专为向量和矩阵定制的优化。以下是几个例子：
* 常量折叠：我们扩展了LLVM的常量折叠，使其能够通过rdregion和wrregion折叠和传播向量常量。
* 将C数组提升为LLVM向量：虽然不推荐，但用户可以在CM中使用C数组而不是CM向量。CM编译器可以用rdregion和wrregion替换C数组的加载和存储。
* 区域折叠：这可以看作是特定于rdregion和wrregion的指令合并转换。
* 死向量移除：这是对向量值进行的一种更通用的死代码消除形式。通过跟踪每个向量元素的使用情况来确定整个向量是否已死。
* 向量分解：给定一个大向量，如果编译器能证明它可以被分成多个段，且这些段上的rdregion和wrregion是不相交的，那么这个大向量就可以转换成多个小向量，这增加了寄存器分配器的灵活性。
代码生成示例。再次考虑算法2中提出的线性CM实现。图4说明了如何通过一个select操作（算法2中的第7行）完成6x24子矩阵的char到float转换。
这个select操作被编译成如下所示的9条SIMD16指令：
1) mov (16|M0) r11.0<1>:f r4.3<8;8,1>:ub
2) mov (16|M0) r13.0<1>:f r4.19<16;8,1>:ub
3) mov (16|M0) r15.0<1>:f r5.11<8;8,1>:ub
4) mov (16|M0) r17.0<1>:f r6.3<8;8,1>:ub
5) mov (16|M0) r19.0<1>:f r6.19<16;8,1>:ub
6) mov (16|M0) r21.0<1>:f r7.11<8;8,1>:ub
7) mov (16|M0) r23.0<1>:f r8.3<8;8,1>:ub
8) mov (16|M0) r25.0<1>:f r8.19<16;8,1>:ub
9) mov (16|M0) r27.0<1>:f r9.11<8;8,1>:ub
在Gen ISA中，源操作数的区域是一个行主序的二维数组，格式为<V;W,H>，其中W（宽度）是一行中元素的数量，H（水平步长）是一行中两个元素之间的步长，V（垂直步长）是两行之间的步长。这个例子展示了在Gen上进行CM编程的强大之处；程序员使用高级的矩阵操作来表达他们的算法，而编译器则利用基于区域的寻址方案将它们生成为多条SIMD指令，以高效地访问寄存器数据。
A4 实验环境
- 硬件配置：实验在一台搭载Intel IceLake (ICL) 处理器的系统上进行。该系统包含一个4核Intel Core i7 CPU，16GB系统内存，以及一个集成的Gen11 GPU。该GPU拥有64个执行单元（EU）。
- 软件配置：
- CM内核使用CM编译器进行编译。
- OpenCL内核来自Intel OpenCL SDK【索引35，Intel SDK for OpenCL Applications，2019，Intel】（除了Histogram和K-means是由内部专家开发）。
- 所有内核都使用
-O2优化级别进行编译，以作为基线。
- 数据集与应用：
- 使用了来自不同领域的代表性应用进行基准测试，包括：线性滤波器、双调排序、直方图、K-均值聚类、稀疏矩阵向量乘法（SpMV）、矩阵转置、通用矩阵乘法（SGEMM和DGEMM）、前缀和。
- 为每个应用使用了典型的输入参数进行基准测试。具体的输入数据集在每个应用的结果部分有描述，例如不同尺寸的图像、随机数据、真实世界图像（如City, Earth）以及来自公开数据集的稀疏矩阵（Protein, Nd24k, Webbase）。
A5 实验结果
实验通过测量总执行时间来比较CM和OpenCL的性能。CM在所有测试应用中均表现出优于最佳优化过的OpenCL版本的性能，速度提升范围从10%到2.7倍不等。
各应用性能分析 (参考图5):
1. 双调排序 (Bitonic Sort): CM实现利用了大的寄存器空间来暂存256个数据元素，在寄存器内部完成多个排序阶段，从而避免了频繁的全局内存访问和同步。相比OpenCL，CM性能提升了1.6倍至2.3倍，输入规模越大，优势越明显。
2. 直方图 (Histogram): CM内核将每个线程的局部直方图高效地存储在寄存器中，而OpenCL则使用共享本地内存（SLM）。CM避免了OpenCL中更新全局直方图所需的原子操作争用。因此CM性能最高可达OpenCL的2.7倍，尤其是在真实图像（如Earth）这种具有大量同质背景、容易引发原子操作冲突的场景下。
3. K-均值聚类 (K-means): CM内核通过在寄存器文件中高效共享质心和其他辅助数据结构，避免了使用SLM和线程同步。CM编译器还能有效重叠分散内存读取操作以隐藏延迟。在三个不同数据集上，CM比OpenCL快30%到50%。
4. 稀疏矩阵向量乘法 (SpMV): CM实现利用了动态改变指令SIMD大小的能力，根据数据块的大小选择最佳执行宽度，从而提高了内存和计算效率。此外，使用布尔归约来检测并跳过全零行的计算。对于非零元素密度和方差不同的矩阵，CM性能比OpenCL高出10%到160%不等（Webbase矩阵提升最大）。
5. 矩阵转置 (Matrix Transpose): CM实现完全绕过了SLM和同步开销，直接在寄存器上通过select和merge操作的组合来完成转置。相比基于SLM的OpenCL实现，CM取得了最高2.2倍的加速。
6. 通用矩阵乘法 (GEMM): 尽管OpenCL和CM都采用了类似的寄存器分块策略，但CM凭借更高效的寄存器文件管理，每个线程能处理更多数据。因此，CM在DGEMM中性能高出8.5%，在SGEMM中高出约10%。
7. 前缀和 (Prefix Sum): CM实现将并行的归约和部分求和操作完全在寄存器中进行，避免了SLM和多次全局/局部内存数据移动及同步。相比OpenCL，CM取得了1.6倍的加速。
生产效率 (Productivity):
为了评估编程模型对开发效率的影响，本文对比了在一个未来GPU平台上开发几个深度学习核心内核的投入。结果表明，尽管OpenCL可以快速开发出功能性原型，但后续性能调优过程非常耗时，开发者需要与编译器和编程模型“斗争”以生成期望的汇编代码。而CM提供了对关键硬件资源的显式控制，代码性能可预测性高，初版性能就已接近甚至超过目标，大大减少了调优时间。如表I所示，对于这些深度学习内核，CM的开发效率平均是OpenCL的2-3倍，同时还取得了更好的性能。
A6 结论
本文介绍了C-for-Metal (CM)，一种为Intel GPU设计的高级但又“接近硬件”的编程语言。通过阐述其主要特性，如代表寄存器的向量/矩阵变量、映射到寄存器区域寻址的select操作、高效内存访问的块读/写以及支持混合SIMT/SIMD模型的分化控制流结构，展示了CM如何暴露底层硬件能力。实验评估表明，CM与OpenCL之间的性能差距显著，CM的性能优势范围从20%到超过100%。
本文的目的并非攻击SIMT编程模型，而是指出SIMT模型的便利性带来了性能成本，即使是专家级编程也很难克服。一个原生设计用于充分利用硬件能力的编程模型填补了重要的空白，这种“金属级”的表达能力对于性能关键型应用尤为重要。
CM被定位为Intel GPU的底层编程工具，并且已有其他语言的前端开始使用CM作为其后端。例如，DPC++-ESIMD将一些CM语言特性集成到DPC++中，而ISPC也生成CM向量内建函数并依赖CM的优化和代码生成。随着向量和矩阵数据类型在神经网络编程中日益重要，预计类似于CM的rdregion和wrregion的IR扩展也可能被添加到LLVM中以支持其他目标机器。
A7 附录
A. 摘要
制品内容。本文的制品（artifact）包含了CM编译器（CMC）的实现，以及实验评估部分使用的应用程序和基准测试。提供了编译和执行基准测试所需的脚本，从而允许在任何配备Intel Gen9（Skylake）GPU或更高版本GPU的系统上复现我们的结果。
B. 制品元信息
- 程序: CM编译器（C++实现）；CM应用程序；OpenCL应用程序（包含所有源码和二进制文件）。
- 编译: 通过提供的脚本使用gcc/g++进行编译。
- 数据集: 应用程序使用的输入数据集要么作为独立文件包含，要么在运行时生成。对于前者，它们位于每个应用程序的目录中。
- 运行时环境: Linux Ubuntu 18.04或更高版本，CM运行时和OpenCL运行时。
- 硬件: Intel Gen9 GPU或更高版本。
- 输出: 为每个用CM和OpenCL评估的应用程序生成文本格式的性能结果文件。
- 公开可用性: CM编译器以及所有的CM和OpenCL示例都是公开的，除了生产力部分（第6.1节）中列出的那些。
- 代码许可证: Intel(R) CM编译器和示例在MIT许可证下分发。
C. 描述
1. 如何获取: CM编译器可在Github上找到：https://github.com/intel/cm-compiler。CM和OpenCL示例，以及构建和运行所有基准测试的脚本可在https://github.com/jfuentes/C-for-Metal_CGO2021上找到。CM编译器和基准测试的二进制文件也包含在制品库中。
2. 硬件依赖: 推荐在Intel Gen11 GPU (Icelake)上运行基准测试，但任何高于Gen9 (Skylake)的Intel GPU应能得到相似结果。注意，由于硬件配置差异，可能需要在不同的Gen平台上进行特定于应用的进一步调优以达到峰值性能。
3. 软件依赖: 此制品使用Ubuntu 18.04准备。类似的Linux发行版也应兼容。制品库包含CM编译器构建及其依赖项，用于编译所有基准测试。要从源码构建CM和IGC编译器，关于依赖项和构建方法的具体细节可在它们的仓库中找到。运行基准测试需要CM运行时和OpenCL运行时，这些都可以在相应的仓库中找到。
D. 安装
基础依赖安装。首先，安装此制品的基本依赖：g++, git, make, cmake和jansson。
CM编译器、运行时和基准测试。下载制品库，其中包含了CM编译器的构建和所有基准测试。如果需要从源码构建CM编译器，请访问其仓库获取更多细节。注意需要通过git lfs pull下载大文件。然后，需要构建并安装包含CM运行时的media driver。这包括安装libva及其依赖。
OpenCL编译器(IGC)和运行时。从compute-runtime仓库的release页面下载并安装IGC和NEO运行时包。然后，安装OpenCL头文件和OpenCL C++头文件。最后，安装OpenCL ICD Loader。
E. 实验流程
构建基准测试。安装完上述软件包后，在制品库的benchmarks目录下，运行build_CM_all.sh和build_OCL_all.sh脚本来构建所有CM和OpenCL基准测试。CM是离线编译，会询问目标GPU平台（如SKL, ICL等）。
运行基准测试。之后，运行run_CM_all.sh和run_OCL_all.sh脚本来执行所有基准测试。
F. 评估和预期结果
结果报告。基准测试完成后，性能结果会报告到标准输出，并保存在results目录下的文本文件中。对每个基准测试，会报告内核执行时间和总执行时间。性能结果以毫秒为单位，并按输入数据组织。
方法细节中的引用汇总
-
引用段落: A3. 新编程模型的动机 - 第2段 (1. 寄存器文件控制)
- 索引: 【27，Register allocation for Intel processor graphics，2018，CGO】
- 原文描述: "Effective use of the register file to reduce unnecessary memory traffic is perhaps the most important optimization strategy for Intel GPUs [27]." (有效利用寄存器文件以减少不必要的内存流量可能是Intel GPU最重要的优化策略[27]。)
- 索引: 【28，Divergence analysis and optimizations，2011，PACT】
- 原文描述: "For example, divergence analysis [28] is a critical analysis for SIMT GPU compilers..." (例如，分化分析[28]是SIMT GPU编译器的关键分析...)
-
引用段落: A3. 新编程模型的动机 - 第3段 (2. 跨通道数据共享)
- 索引: 【29，Using CUDA Warp-Level Primitives，2018，NVIDIA Developer Blog】
- 原文描述: "...functionalities provided include shuffle, reduction, and barrier operations [29], [30]." (...提供的功能包括shuffle、reduction和barrier操作[29], [30]。)
- 索引: 【30，The OpenCL Extension Specification，2018，Khronos】
- 原文描述: "...functionalities provided include shuffle, reduction, and barrier operations [29], [30]." (...提供的功能包括shuffle、reduction和barrier操作[29], [30]。)
-
引用段落: A3. 新编程模型的动机 - 第4段 (3. 向量长度控制)
- 索引: 【31，Taming control divergence in gpus through control flow linearization，2014，CC】
- 原文描述: "Similarly, branch divergence can significantly reduce a program’s efficiency [31], [32];" (类似地，分支分化会显著降低程序的效率[31], [32];)
- 索引: 【32，Reducing branch divergence in gpu programs，2011，GPGPU】
- 原文描述: "Similarly, branch divergence can significantly reduce a program’s efficiency [31], [32];" (类似地，分支分化会显著降低程序的效率[31], [32];)
-
引用段落: A2. 方法细节 - IV. CM编程语言 - 第1段 (CM语言基础)
- 索引: 【6，C-for-Metal Compiler，2019，GitHub】
- 原文描述: "...with some restrictions (more details in section 2.6 of the CM language specification [6])." (...并有一些限制（更多细节见CM语言规范[6]的2.6节）。)
-
引用段落: A2. 方法细节 - IV. CM编程语言 - D. SIMD Control Flow - 第2段 (每通道SIMD控制流)
- 索引: 【33，IGC: the open source Intel Graphics Compiler，2019，CGO】
- 原文描述: "...utilizing the Gen simd-goto and simd-join instructions that support divergent control-flow under SIMD execution [33]." (...利用Gen的simd-goto和simd-join指令，支持SIMD执行下的分化控制流[33]。)
-
引用段落: A2. 方法细节 - V. CM编译器 - 第1段 (编译器架构)
- 索引: 【33，IGC: the open source Intel Graphics Compiler，2019，CGO】
- 原文描述: "Like Intel Graphics Compiler (IGC) [33], the CM Compiler consists of three layers:" (与Intel图形编译器（IGC）[33]类似，CM编译器由三层组成：)
- 索引: 【34，LLVM and Clang: Next generation compiler technology，2008，BSD conference】
- 原文描述: "The clang front-end compiler [34] converts CM source code into LLVM intermediate representation (IR) [3]." (clang前端编译器[34]将CM源代码转换为LLVM中间表示（IR）[3]。)
- 索引: 【3，LLVM: A compilation framework for lifelong program analysis & transformation，2004，CGO】
- 原文描述: "...converts CM source code into LLVM intermediate representation (IR) [3]." (...将CM源代码转换为LLVM中间表示（IR）[3]。)
- 索引: 【27，Register allocation for Intel processor graphics，2018，CGO】
- 原文描述: "The vISA finalizer [27] is a code generator for Intel GPU." (vISA finalizer[27]是Intel GPU的代码生成器。)
-
引用段落: A2. 方法细节 - V. CM编译器 - 第3段 (LLVM IR扩展)
- 索引: 【33，IGC: the open source Intel Graphics Compiler，2019，CGO】
- 原文描述: "Gen ISA has distinct features such as varying execution size, mixed data types, flexible register regioning, and modifier support [33]." (Gen ISA具有独特的特性，如可变的执行大小、混合数据类型、灵活的寄存器区域寻址和修饰符支持[33]。)
💬 评论讨论
欢迎在这里分享您的想法和见解！