面向英特尔® 至强融核™ 协处理器的 OpenCL* 设计和编程指南

关于本文档

本文旨在对开发面向英特尔® 至强融核™ 协处理器的高性能 OpenCL 应用所需的设计和编码指南进行介绍。 阅读本文,您可以了解英特尔至强融核协处理器架构和微架构、关键的 OpenCL 结构,以及学习如何通过高效地使用这些架构最大限度地利用英特尔至强融核协处理器硬件。 由于使用硬件的并行性对高性能应用至关重要,因此我们将为您展示如何在英特尔至强融核协处理器上改善 OpenCL 应用的并行性。 这些知识将帮助您对您的应用进行有效的设计和编程,以便在英特尔至强融核协处理器上最大限度地利用 OpenCL。

为何需要本文?

OpenCL 是一个易于使用的编程模型,尽管如此,使用该模型并不意味着一定能获得最高的性能。 传统的 GPU 和英特尔至强融核协处理器有着不同的硬件设计。 这种硬件上的明显差异导致它们只能采用不同的应用优化。 例如,传统的 GPU 要求必须使用快速共享本地内存,因为编程人员需要使用这种内存进行显示编程。 英特尔至强融核协处理器包含了完全一致的高速缓存结构,这与常见的 CPU 高速缓存类似,而该结构可以自动加速内存访问。 另一个例子: 一些传统的 GPU 基于许多小线程的硬件调度构建而成,而英特尔至强融核协处理器则通过设备操作系统对中等大小的线程进行调度。 所有这些差异表明只有对应用进行调整使其符合其将要运行的硬件时这些应用才能发挥最大性能。

我们需要针对不同的设备进行不同的 OpenCL 优化吗?

不一定。 您想过只要在您的代码中添加一个小的 #ifdef 就能在英特尔至强融核协处理器上实现 50% 的加速吗? 为此您愿意复制一个 1000 行的文件吗? 为了仅仅 10% 的加速,您愿意这样做吗? 或者,您可能更愿意无条件地进行优化,宁可让其他设备减速 10% 也要在英特尔至强融核协处理器上将性能提升 50%? 这完全由您来决定。 在某些情况下,您需要在设备性能和 OpenCL 应用的可维护性之间做出权衡。 我们希望开发人员通过使用本文中指南对英特尔至强融核协处理器的潜在性能展开探索,然后根据性能数做出自己的选择。 本文并不会针对所有的问题给出答案,而是为您提供一些有益的工具,帮助您自己回答这些问题。

英特尔至强融核协处理器高级硬件综述

 

图 1. 英特尔® 至强融核™ 协处理器微架构

一个英特尔至强融核协处理器包含许多内核,每个内核具备一个 512 位向量运算单位,可以执行 SIMD 向量指令。 每个内核包含一个一级高速缓存(32 KB 数据 + 32 KB 指令)。 每个内核与一个二级高速缓存相关联(512 KB 组合数据和指令、包含一级 D 高速缓存)。 通过一个高速互联在二级高速缓存和内存子系统间进行数据传输。 每个内核都可以同时执行多达 4 个硬件线程。 这种同时进行的多线程处理有助于隐藏指令和内存延迟。 OpenCL 可以对多数这种细节进行隐藏,因此编程人员不能看到这些细节。

如欲了解更多有关英特尔至强融核协处理器硬件的信息,请访问:

http://software.intel.com/en-us/mic-developer

影响英特尔至强融核协处理器性能的主要方面

多线程并行性

 

受 SKU 影响,英特尔至强融核协处理器硬件包含了许多内核(本文中假定为 60 )。 每个内核可以运行多达四个硬件线程。 在多数情况下,使 240 个线程满负荷运行是确保性能最大化的关键。 硬件线程的准确数量可以通过 clGetDeviceInfor(NUM_COMPUTE_UNITS) 进行查询;接口。

芯内矢量化

英特尔至强融核协处理器中的向量大小是 512 位宽 SIMD。 通常,这种向量代表 8 个双精度浮点数,或者 16 个单精度浮点数。 每个英特尔至强融核协处理器内核每次循环时能发布一个向量计算指令。

PCI Express* (PCIe) 总线接口

英特尔至强融核协处理器位于 PCIe 总线上。在 PCIe 总线上传输数据有着最高的延迟和最低的带宽。 正如您在任何其他 PCIe 设备中做的一样,您应当将这种流量减至最小值。

内存子系统

英特尔至强融核协处理器包含三个等级的内存(GDDR、二级高速缓存

 

和一级高速缓存)。 以下表格中包含了重要的高速缓存信息:

  一级(数据 + 指令) 共享二级高速缓存

总大小:

32 KB + 32 KB

512 KB

延迟

15-30 循环

500-1000 循环

访问一级高速缓存包含仅为一个循环的延迟。

 

由于英特尔至强融核协处理器是一种按顺序处理的机器,内存访问的延迟将对软件的性能产生重大的影响。 幸运的是,编程人员可以降低这些延迟。 借助预取等工具,我们就能隐藏内存延迟。 稍后我们将详细讨论这种工具。

数据访问模式

连续访问内存是访问英特尔至强融核协处理器上内存的最快方式。 它提升了高速缓存的效率,降低了 TLB(转换后备缓冲器)的数量,并支持硬件预取程序发挥作用。

将 OpenCL 结构映射至英特尔至强融核协处理器

 

理解在英特尔至强融核协处理器上如何实施关键的 OpenCL 结构将有助于您更好地设计您的应用,以便利用协处理器的硬件。 此外它还能帮助您避免协处理器的性能缺陷。

从概念上讲,在初始化时 OpenCL 驱动程序会创建 240 条软件线程并将它们固定到硬件线程上(针对一个 60 核心的配置而言)。 在完成了一个 clEnqueueNDRange() 调用之后,驱动程序将对 240 条线程上当前的 NDRange 工作组 (WG) 进行调度。 一个工作组就是调度到线程上的最小任务。 因此当调用 clEnqueueNDRange() 的工作组少于 240 个时,协处理器就没有得到充分的利用。

OpenCL 编译器创建一个可以执行一个工作组的优化例程。 这个例程通过多达三个嵌套循环构建而成,如以下虚拟码所示:

 
1 __Kernel ABC(…)
2 For (int i = 0; i < get_local_size(2); i++)
3     For (int j = 0; j < get_local_size(1); j++)
4         For (int k = 0; k < get_local_size(0); k++)
5             Kernel_Body;

请注意最里面的循环用于 NDRange 的 0 维。 这将对您的性能关键代码的访问模式产生直

 

接的影响。 同时它还会影响隐式矢量化的效率。

OpenCL 编译器根据 0 维循环对工作组例程进行隐式向量化,即通过向量大小展开 0 维循环。 执行矢量化后的工作组代码如下所示:

 
1 __Kernel ABC(…)
2 For (int i = 0; i < get_local_size(2); i++)
3     For (int j = 0; j < get_local_size(1); j++)
4         For (int k = 0; k < get_local_size(0); k += VECTOR_SIZE)
5             Vector_Kernel_Body;

无论内核中使用的数据类型如何,英特尔至强融核协处理器的向量大小都为 16。 未来,我们可能还会增加矢量化的尺寸以支持更高的指令级并行性。

显示算法的并行性

尽管 OpenCL 规范提供了各种展示并行性和并发性的方法,其中的一些方法并不能有效地映射到英特尔至强融核协处理器。 我们将为您展示如何将关键的 OpenCL 结构映射到协处理器上,以便帮助您设计出可以利用平行性的应用。

多线程

为了充分利用这 240 条硬件线程,最好使每个 NDRange 拥有超过 1000 个工作组。 每个 NDRange 上拥有 180‒240 个工作组时只能确保基本的线程利用;同时,执行任务时还可能受制于较差的负载平衡和较高的调用开销。

建议: 确保每个 NDRange 拥有至少 1000 个工作组,以便充分利用英特尔至强融核协处理器的硬件线程。 当 NDRange 拥有 100 个工作组甚至更少时,应用的性能将会因为线程极低的利用率受到严重的制约。

单工作组执行的持续时间也会影响线程的效率。 在此也不建议使用轻量级工作组,因为这种工作组可能会受制于较高的开销。

矢量化

英特尔至强融核协处理器上的包含一个隐式矢量化模块。 OpenCL 编译器可对 0 维工作项上的隐式工作组循环自动进行矢量化操作(见上面的示例)。 无论在内核中使用的数据类型如何,当前的矢量化宽度都是 16。 今后我们还计划对 32 个元素进行矢量化操作。 由于 OpenCL 工作项一定是独立的工作项,因此 OpenCL 矢量化器不需要可行性分析就能实施矢量化操作。 然而,经过矢量化处理的内核只有当 0 维的 local size 大于或等于 16 时才能使用。否则,OpenCL 运行时将对每个工作项运行标量内核。 如果 0 维的工作组尺寸不能被 16 整除时,那么工作组的末端需要被标量代码执行。 这对大型的工作组(例如 0 维有 1024 项)来说不算是一个问题,但对 0 维上大小为 31 的工作组来说有着显著的影响。

建议 1: 不要通过手动的方式对内核进行矢量化,这是因为 OpenCL 编译器会对您的代码进行标量化,为隐式矢量化做好准备。

建议 2: 避免使用一个不能被 32 整除的工作组大小(可以使用被 16 整除的工作组)。

工作项 ID 非均匀控制流

在本节中我们将在隐式向量化的环境中讨论均匀控制流与非均匀控制流之间的差异。 理解这一点非常重要,因为均匀控制流可能会对性能产生一些微小的负面影响。 然而非均匀控制流将在 NDRange 的最里面的一维造成极高的性能开销。 因此矢量化循环(0 维)的均匀性至关重要。

如果能够确保在一个工作组内的所有工作项执行一个分支的同一端,那么该分支就是均匀的。

均匀分支示例:

 
1 //isSimple is a kernel argument
2 Int LID = get_local_id(0);
3 If (isSimple == 0)
4     Res = buff[LID];

非均匀分支示例:

 
1 Int LID = get_local_id(0);
2 If (LID == 0)
3     Res = -1;

另一个均匀分支示例:

 
1 Int LID = get_local_id(1);
2 //Uniform as the IF is based on dimension one, while vectorization on dimension on.
3 If (LID == 0)                      
4     Res = -1;

即使执行了矢量化,编译器也需要通过预测对任何被非均匀控制流控制的代码执行线性化处理(即均匀化)。 预测的成本主要集中在分支两端的执行方面。 而隐藏执行又会导致额外的成本支出。

假设以下代码就是原始的内核代码:

 
1 Int gid = get_global_id(0);
2 If(gid % 32 == 0)
3     Res = HandleEdgeCase();
4 Else
5     Res = HandleCommonCase();
6 End

执行矢量化(和预测)之后,代码如下所示:

 
1 int16 gid = get16_global_id(0);
2 uint mask;
3 Mask = compare16int((gid % broadcast16(32)), 0)
4 res_if = HandleEdgeCase();
5 res_else = HandleCommonCase();
6 Res = (res_if & mask) | (res_else & not(mask));

请注意已经针对所有的工作项执行了 IF 和 ELSE。

建议: 避免分支,尤其是 0 维上的那些非均匀分支。

数据对齐

由于各种原因,向量尺寸经过对齐处理的内存访问相比未对齐的内存访问会更快一些。 在英特尔至强融核协处理器中,OpenCL 缓冲器可以确保在一个向量尺寸经过对齐处理的地址启动。 然而,这只能确保第一个工作组能在一个对齐的地址启动。 为了确保所有的工作组可以在一个经过正确对齐的位置启动,工作组的大小 (local size) 要求被 16 整除,或者要求被 32 整除(此时您将会获得潜在的产品性能提升)。 将 local size 设为 NULL,调用 EnqueueNDRange,让 OpenCL 驱动程序为您选择最佳的工作组大小。 该驱动程序需要确保足够智能,以便选择出一个可以满足对齐要求的工作组大小。 然而,编程人员需要确保全局尺寸能够被 VECTOR_SIZE 整除,同时商数足够大以允许运行时可被高效地拆分成多个工作组。 “足够大”指的是在小内核的情况下应达到 1,000,000,而在包含 1000 个迭代循环的一个大内核的情况下应达到 1000。 此外,NDRange 偏移也会破坏这种对齐。

建议 1: 不要使用 NDrange 偏移。 如果您必须要使用偏移,就需要使其成为 32 的倍数,或者至少是 16 的倍数。

建议 2: 使用是 32 的倍数的 local size,或者至少是 16 的倍数。

设计您的算法,使其可以利用英特尔至强融核协处理器内存子系统的各项优势

英特尔至强融核协处理器由于是一种按顺序处理的机器,因此对内存延迟非常敏感。 针对应用层面而进行了内存优化有助于实现 2X-4X 的性能加速。

内部工作组数据再使用

内存优化的第一步就是对您的应用进行设计,以便最大限度地增加来自高速缓存的数据再使用的数量。 然而,只有某些算法才需要重新使用数据。 例如,增加两个矩阵时将导致无法重新使用任何数据。 但是,对两个矩阵执行矩阵乘法操作(GEMM)就可以重新使用许多数据。 因此,是否可以利用优化完全由您决定。 更多细节,请参见 面向 OpenCL 应用 XE 的英特尔 SDK – 优化指南

为了利用数据再使用,您需要考虑工作组隐式循环(本文前面部分已对此进行过描述)。 编程人员对这些循环的控制是通过 local size 定义进行的。 编程人员可以在内核中添加额外的循环(显式)。

交叉工作组数据再使用

交叉组数据再使用意味着更大的挑战。 当前,英特尔至强融核协处理器上的 OpenCL 不允许对工作组调度实行完全的控制。 因此,交叉工作组数据再使用几乎是不可能实现的。 我们将在今后的开发中考虑如何解决这种挑战。

数据访问模式

通过连续的数据访问通常可以实现最佳的内存系统性能。 就实现连续的内存访问而言,首先最关键的一步就是了解工作组隐式循环的结构。 最里面的隐式循环即为 0 维之上的循环。 如果您的内核没有使用额外的(显式)循环,那么您需要考虑使大多数的内存访问与 0 维隐式循环保持一致。 例如:

以下代码可以在内存中连续访问 2D 缓冲器(推荐):

 
1 __kernel ABC(…){
2 int ID1 = get_global_id(1);
3 int ID0 = get_global_id(0);
4 res[ID1][ID0] = param1 * buffer[ID1][ID0];
5 }

以下代码没有在内存中连续访问 2D 缓冲器(不推荐):

 
1 __kernel ABC(…){
2 int ID1 = get_global_id(1);
3 int ID0 = get_global_id(0);
4 res[ID0][ID1] = param1 * buffer[ID0][ID1];
5 }

第二个示例代码扫描 2D 缓冲器“column major”。 在使用矢量化的情况下,这会导致双重错误,即: 1) 输入向量数据需要沿着该纵列从 16 个连续行中收集。 通过将指令分散到 16 个不同的行中存储结果。 两种运算方法执行起来都很缓慢。 2) 内存访问不是连续的逐级迭代。 这两个错误在 TLB 上施加了更大的压力同时阻碍了预取。

简单的 1 维示例(推荐):

连续访问:

 
1 Int id = get_global_id(0);
2 A[id]= B[id];

不连续访问(不推荐):

 
1 Int id = get_global_id(0);
2 A[id*4] = B[id*4]

建议: 使用 ID(0) 在行内连续索引内存。 针对显式 2D 缓冲器: 缓冲器[ID1][ID0]。 针对进入 1D 缓冲器的 2D 索引: 缓冲器[STRIDE * ID1 + ID0]

如果您的内核包含一个显式循环,您应记住隐式矢量化仍然是基于 ID(0) 隐式循环进行的。 因此通过 OpenCL 的 ID 访问缓冲器应当遵循以上的建议(缓冲器[ID1][ID0])。 这样就能保证向量访问的连续性和效率。 通过内循环索引 (idx)访问缓冲器在内循环 (缓冲器[ID1][idx]) 内部应当是连续的,而对矢量化的循环应当是均匀的! 然而,应当避免将 ID0 与 idx 混合。 例如,缓冲器[ID0][idx] 进入矢量化循环会产生聚集或分散情况。

数据布局

SOA (阵列结构) 数据布局可以实现简单且高效的向量加载和存储。 但是这会导致空间局部性更低、TLB 上的压力更高,以及同时使用的页面数量更多。

而在使用 AOS(结构阵列)数据布局时,生成的矢量化内核需要通过聚集和分散指令加载和存储数据,这与简单的向量加载和存储相比效率更低。 然而,就随机访问模式而言,AOS 布局由于具备更好的空间局部性,因此可以实现比 SOA 布局更高的效率。 请注意 SOA 数据布局的随机访问也会创建聚集和分散指令。

第三种选择是 AOSOA — 小阵列结构阵列。 适合英特尔至强融核协处理器的小阵列的大小应为 32。 这样就能对多达 32 元素的向量进行矢量化处理。

 
1 struct Point32 {    float x[32], y[32], z[32]; };
2 __kernel void ABC(__global Point32* ptrData)

AOSOA 支持通过简单的向量加载进行高效的矢量化,同时既不会过度加载 TLB,也不会在许多页面上分散访问。 代码可读性是 AOSOA 存在的问题。 大多数人还不了解 AOSOA 的术语。

数据预取

由于英特尔至强融核协处理器是一个按顺序处理的机器,通过数据预取这种基本的方法有助于将数据移至距离内核更近的位置。 借助并行性,加载和存储可被连续执行。 例如,任何两个加载指令完全可被连续执行。 而预取指令是个例外。 它需要和其他指令(包括其他预取指令)并行执行。 因此,没有按时完成的预取指令仍然可以提升性能,这是因为这种内存请求与其他指令并行执行。 一个高速缓存缺失意味着一个线程的停止,同时也意味着需要几个循环才能重新发布指令。 英特尔至强融核协处理器包含一个通向二级高速缓存的简单的自动硬件预取程序。 该硬件预取程序需要花费一些时间才能发挥作用,同时该程序需要在每 4 KB 虚拟页范围上重新启动。

为了方便今后执行迭代时访问数据,OpenCL 编译器为一级和二级高速缓存插入了自动的软件预取,并(通过分析)计算出何时这种插入操作能发挥作用。 beta 版的发布为自动软件预取提供了部分支持。

编程人员可以通过内建的预取特性将手动预取功能插入到 OpenCL 内核之中。当前,手动预取确实已经插入到该位置和编程人员请求的地址中,但是这仅限于二级预取。 未来,OpenCL 编译器可能会针对内建预取添加二级和一级预取。它也可能会改善编程人员指明的位置和步幅。 手动预取应当在数据将被实际使用之前提前至少 500 个循环插入。 通常来说,只有主输入和输出缓冲器需要被预取。

本地内存和 Barriers

传统的 GPU 通常包含需要手动管理的共享本地内存 (SLM),而英特尔至强融核协处理器则包含一个两级高速缓存系统(自动),这与大多数现代的 CPU 相似。 因此,使用 OpenCL SLM 并不会为英特尔至强融核协处理器带来任何好处。 而且,协处理器中的本地内存被分配在普通的 GDDR 内存上,同时像任何其他内存一样需要得到高速缓存系统的支持。 因此,在拷贝和管理冗余数据时就产生了额外的开销。

建议: 避免在英特尔至强融核协处理器上使用共享本地内存。

英特尔至强融核协处理器对 Barriers 不提供特殊的硬件支持。 因此,Barriers 需要通过协处理器上的 OpenCL 模仿。 我们建议应当避免使用 Barriers。 此外,将内核拆分成两个单独的内核将比一个 barrier 更慢,因此我们也不建议采取这种方法。

自 beta 版本发布以后,组合使用 barrier 和不能被 16 整除的工作组大小将导致执行一个标量内核。 请避免使用这种组合。 当前,我们还无法判断在 OpenCL 编译器内对此进行优化是否合理。

总结

为了设计出适合英特尔至强融核协处理器的 OpenCL 应用,您应该特别注意以下几方面:

  1. 使每个 NDRange 包含足够多的工作组 — 建议至少为 1000。
  2. 避免使用轻量级的工作组。 请使用准许的最大 local size(当前为 1024)。 保证工作组大小是 32 的倍数。
  3. 避免使用 ID(0) 非独立控制流。 这允许高效的隐式矢量化。
  4. 最好使用连续数据访问。
  5. 数据布局设置: 针对分散的随机访问使用 AOS;其他则使用纯 SOA 或 AOSOA(32)。
  6. 通过工作组内部的高速缓存利用数据再使用。
  7. 如果自动预取未能发挥作用,在使用之前提前 500‒1000 个循环使用内建预取功能将全局数据移至高速缓存。
  8. 不要使用本地内存。 避免使用 barriers。

延伸阅读

面向 OpenCL 应用 XE 的英特尔 SDK – 优化指南 包含更多细节。

通过 英特尔® VTune™ 放大器 XE 为英特尔至强融核协处理器优化性能。

有关编译器优化的更完整信息,请参阅优化通知