OpenCL 通用编程与优化（7）

---

6 关于Adrenogpu的6个顶级内核优化技巧

本章为Adrenogpu提供了一些顶级的OpenCL优化技巧，更多的细节和其他信息将在下面的章节中描述。本章中的所有建议都应该具有最高的优先级，开发人员在进行内核优化时应该在尝试其他建议之前尝试它们。

6.1工作组绩效优化

内核的工作组大小和形状对性能影响很大，而调整工作组大小是一种简单而有效的性能优化方法。本节介绍了关于工作组的基本信息，包括如何获得给定内核的工作组大小，为什么需要调整工作组大小，以及关于最佳工作组大小调整的标准实践。

6.1.1获取最大工作组大小

开发人员应该在运行clBuild程序后，通过使用以下API函数查询设备上内核的最大工作组大小：

size_t maxWork-groupsize;
clGetKernelWorkgroupInfo(myKernel, 
myDevice, 
CL_KERNEL_WORK_GROUP_SIZE, 
sizeof(size_t), 
&maxWork-groupsize, 
NULL);

克隆队列范围内核使用的实际工作组大小不能超过最大工作组大小。如果应用程序未指定该工作组的大小，则Adreno OpenCL软件将选择一个默认的和工作工作组的大小。

6.1.2 必需的和首选的工作组规模

内核可能需要或希望某的工作组大小来正常或有效地工作。OpenCL提供了一些属性，允许内核从编译器中请求或请求特定的工作组大小：

• 使用reqd_work_group_size属性。
  reqd_work_group_size（X、Y、Z）属性显式地要求一个特定的工作组大小。如果编译器不能满足指定的工作组大小，则会返回一个编译错误。例如，要要求一个16x16的工作组大小：

__kernel __attribute__(( reqd_work_group_size(16, 16, 1) )) 
 void myKernel( __global float4 *in, __global float4 *out) 
 { . . . }

• 使用work_group_size_hint属性。
  OpenCL软件尝试使用暗示的工作组大小，但不能保证实际的工作组大小与提示相匹配。例如，要提示工作组大小为64x4：

__kernel __attribute__(( work_group_size_hint (64, 4, 1) )) 
 void myKernel( __global float4 *in, __global float4 *out) 
 { . . . }

在大多数情况下，由于工作组大小的限制，编译器不能保证它生成最优的机器代码。此外，如果编译器不能使用片上的寄存器来满足所需的工作组大小，那么它可能不得不将寄存器泄漏到系统内存中。因此，不鼓励开发人员使用这两个属性，除非内核需要工作组大小才能正确地工作。
编写一个依赖于固定工作组大小或布局的内核不适合用于跨平台兼容性和可移植性的目的。

6.1.3影响最大工作组大小的因素

如果没有指定工作组大小属性，则内核的最大工作组大小取决于许多因素：

• 内核的寄存器占用空间（所需的寄存器数）。通常，内核越复杂，寄存器占用就越大，最大工作组规模就越小。可能增加寄存器占用的因素如下：
  • 为每个工作项打包更多的工作负载。
  • 控制流。高精度数学函数（例如，不使用本机数学函数或快速数学编译标志）。
  • 本地内存，如果这导致临时分配额外的寄存器来存储加载/存储指令的源和目标。
  • 专用内存，例如，为每个工作项定义的数组。
  • 循环正在展开。
  • 内联函数。
• 通用寄存器（GPR）文件的大小。
  • Adreno低层可能有更小的寄存器文件大小。
• 内核中的障碍。
  • 如果内核不使用障碍，其最大工作组大小可以设置为Adreno A4x、A5x、A6x和A7x系列中的设备最大值，无论寄存器占用面积如何。
  • 详见第6.1.4节。
  • 详见第6.1.4节。

无屏障的6.1.4内核（蒸汽模式）

传统上，工作组中的所有工作项目都必须同时驻留在GPU上。对于寄存器占用空间较大的内核，这可能会将其工作组大小限制在设备最大值以下。

从Adreno A4x系列开始，没有障碍的内核可以拥有Adreno支持的最大工作组大小，通常是1024，尽管它们很复杂。由于波之间没有同步，当这些类型的内核的旧波完成时，新波可以开始执行。

在这种情况下，拥有最大的工作组大小并不意味着它们具有良好的并行性。一个没有障碍的内核可能非常复杂，以至于只有有限数量的波在SP内部并行运行，从而导致较差的性能。开发人员应该继续优化和最小化寄存器占用，而不管从clGetKernelWorkgroupInfo 中获得的最大工作组大小。

6.1.5工作组的大小和形状的调整

本节介绍了选择最佳工作组大小和形状的一般指导原则。

6.1.5.1避免使用默认的工作组大小

如果内核调用没有指定工作组大小，则OpenCL软件将使用某种简单的机制找到工作工作组大小。开发人员应该知道，默认的工作组大小不太可能是最佳的。手动尝试不同的工作组大小和形状布局（对于2D/3D），并找到最佳的布局总是一个很好的做法。

6.1.5.2工作组规模大，性能更好？

对于大多数内核来说都是这样的，因为增加工作组大小允许更多的波驻留在SP上，这通常意味着更好的延迟隐藏和提高SP利用率。但是，随着工作组规模的增加，一些内核的性能可能会更差。一个例子是，由于糟糕的数据位置和访问模式，导致缓存抖动增加。纹理访问的局部性问题也很严重，因为纹理缓存通常小于L2缓存。找到最好的工作组规模和形状需要大量的尝试和错误。

6.1.5.3固定与动态工作组大小

为了实现跨设备的性能可移植性，请避免假设一个工作组大小适合所有工作组和硬编码的工作组大小。在一个GPU上工作得最好的工作组大小和布局在另一个GPU上可能并不理想。因此，开发人员应该为内核可以执行的所有设备配置不同的工作组大小，并在运行时为每个设备选择最好的工作组大小。

6.1.5.41vs.2vs.三维（1D/2D/3D）内核

一个内核最多可以支持三个维度。而其维度的选择可能对性能有影响。与每个工作项只有1D索引（例如，全局ID、本地ID等）的1D内核相比，一个二维内核有一组额外的这些内置索引，如果这些索引有助于节省一些计算，那么它可能会提高性能。

根据按工作项划分的数据访问模式，2D内核可能在缓存中具有更好的数据本地化性，从而导致更好的内存访问和性能。而在其他情况下，2D内核可能比1D内核导致更严重的缓存抖动。最好是尝试不同维度的内核的最佳性能。理想情况下，第一个维度上的工作组大小应该是子_group_size的倍数，如果内核有分歧，这一点尤为重要。

6.1.6在工作组上的其他主题

6.1.6.1全局工作规模和填充物

OpenCL1。x要求内核的全局工作大小是其工作组大小的倍数。如果应用程序指定的工作组大小不满足此条件，则clenqueeuend范围内核调用将返回一个错误。在这种情况下，应用程序可以填充全局工作大小，使其成为用户指定的工作组大小的倍数。

OpenCL 2.0解除了此限制，并且全局工作大小不必是工作组大小的倍数，这被称为非统一工作组。
理想情况下，在其第一个维度中的工作组的大小应该是波的大小的倍数（例如，32），以充分利用波的资源。如果不是这样，请考虑填充工作组大小以满足此条件。

6.1.6.2 Brute力搜索

由于工作小组规模选择的复杂性，实验通常是找到最佳规模和形状的最佳方法。
一种选择是使用与实际工作负载（但可能是较小的工作负载）的复杂性相似的预热内核，在应用程序开始时动态搜索最优工作组大小，然后为实际内核使用选定的工作组大小。商业基准测试依赖于这种方法。

6.1.6.3避免跨工作组的工作负载不均匀

应用程序在各个工作组之间的工作负载分布可能不均匀。例如，基于区域的图像处理可能具有比其他区域需要更多的资源来处理的区域。将它们均匀地分配给工作组可能会产生一个平衡问题。如果单个工作组花费的时间太长而无法完成，那么它也会使上下文切换复杂化。

避免这一问题的一种方法是采用两阶段的处理策略。第一阶段可以收集刺激点，并为第二阶段的处理准备数据。工作负载更加确定性，更容易在工作组之间平均分配。

6.1.6.4工作组同步

OpenCL不保证工作组的执行顺序，也没有定义工作组同步的机制。开发人员永远不应该假设在gpu上运行的工作组的顺序。
在实践中，可以使用原子函数或其他方法跨工作组进行有限的同步。例如，应用程序可以分配一个由来自不同工作组的工作项原子更新的全局内存对象。一个工作组可以监视由其他工作组更新的内存对象。这样，就有可能实现有限的工作组同步。

6.1.6.5持久线程

启动一个工作组需要GPU硬件的时间，而如果工作组的数量很大，成本就会影响性能。如果每个工作组的工作负载都很轻，那么这一点就会特别昂贵。因此，开发人员可以减少工作组的数量，并增加每个工作组的工作量，而不是启动大量的工作组。在极端情况下，内核可以在每个SP中只使用一个工作组，并经过多次迭代，以完成许多工作组所完成的相同任务。这种所谓的“持久线程”可以最小化硬件启动的成本，并提高性能。这种方法的一个需要注意的是，上下文切换可能会受到影响，如第3.3节中所讨论的，工作组的工作负载不应该影响其他高优先级的应用程序，例如用户界面（UI）的呈现。

6.2 使用图像对象，而不是使用缓冲区对象

OpenCL支持缓冲区和映像对象（以及来自OpenCL2.0中的管道对象）。一维缓冲区对象对于许多开发人员来说是一个自然的选择，因为它们具有简单性和灵活性，例如支持指针、字节可寻址访问等。图像对象，以预定义的图像格式存储一、二或三维数据，adrgpu首选Adreno而不是缓冲对象：

• Adrenogpu有一个强大的纹理引擎和专用的一级缓存，可以有效地在图像对象中加载数据。
• 使用图像允许硬件自动处理边界外的读取。
• Adrenogpu支持多对图像格式和数据类型组合。
• Adrenogpu支持双线性或三线性插值操作。

在许多用例中，开发人员在用图像对象替换缓冲区对象时，可以期望有良好的性能提高，但代价是稍微复杂一些的主机代码和失去灵活性。更详细的描述见第7.1.5.3节

6.3矢量负载/存储和合并负载/存储

Adrenogpu支持高达128位的每次加载/存储事务的全局/本地内存和映像的读/写。为了最大限度地提高内存加载效率，每个工作项都应该使用向量化加载/存储函数，例如包含4个32位数据的vload4/v存储4，以及CL_RGBA和浮动/int32、f、i、h)、f、i、int32/f、ui、h)/半数据类型。这对与内存绑定的用例很有帮助。

Adrenogpu支持硬件合并负载/存储。例如，假设每个工作项加载在内存中具有连续的地址的16位数据。Adrenogpu可以组合一定数量的相邻工作项的请求，以最小化加载请求的数量。然而，与128位矢量化的负载/存储相比，这种合并的效果较差。更详细的描述见第7.2.2节

6.4恒定内存

Adrenogpu支持快速的芯片上恒定内存，使用它可以大大减少内核执行时间。在大多数情况下，编译器可以自动使用常量内存来存储一些变量，例如常量数组。但是，在某些情况下，开发人员需要提供更多的信息，以便编译器可以决定是否可以使用固定的内存。例如，对于下面的内核，

__kernel void myFastKernel( __constant float *foo 
__attribute__( (max_constant_size(1024)))
 { . . . }

作为一个全局内存对象，如果编译器能够确定通过max_constant_size属性确定它的大小不超过可用的常数内存，则可以将缓冲区foo提升到快速片上常数内存。

如果具有foo中元素的ALU操作是统一的，即子组或工作组中的工作项使用foo中的相同组件进行计算，则可能会进一步导致性能的提高。这是因为固定内存中的内容可以在短时间内广播到ALU，以便快速ALU计算。所有其他内存（全局、本地和私有）必须经过较长的加载/存储路径将数据移动到寄存器中。

更详细的描述见第7.1.3节

6.5本地内存

本地内存是Adrenogpu中的片上物理内存。使用本地内存并不一定会提高性能。以下是开发者应该注意的一些事：

• 尝试做一个128位矢量量化的加载/存储。
• 本地内存应该存储最常用的数据。
  • 如果使用一次或极几次，本地内存可能会影响性能。建议使用
• 子组功能来进行缩减和洗牌操作。子组函数允许工作项共享和交换数据，而不使用本地内存。
  • 子组函数不能经历较长的加载/存储路径。
• 本地内存的广泛使用可能会限制并发工作组执行的数量，从而影响延迟隐藏。

更详细的描述见第7.1.2节

7.内存性能优化

内存优化是最关键和最有效的OpenCL性能技术。许多应用程序都是受内存绑定的，而不是计算绑定的。因此，掌握内存优化对于OpenCL优化至关重要。

7.1在Adrenogpu中的OpenCL内存

OpenCL定义了四种类型的内存（全局内存、本地内存、常数内存和私有内存），理解它们之间的差异对性能优化至关重要。图7- 1说明了这四种内存类型的概念性布局。

图7-1 OpenCL概念记忆层次结构

OpenCL标准只在概念上定义了这些内存类型，而如何实现它们是特定于供应商的。物理位置可能与概念位置不同。例如，私有存储器对象可能位于远离GPU的芯片外系统存储器中。

表7-1列出了四种内存类型的定义及其在gpu中Adreno中的延迟和物理位置。Adrenogpu上的本地内存和恒定内存都是芯片上的，其延迟比片外系统内存要短得多。

通常，内核应该对需要频繁访问的数据使用本地和固定内存，以利用短延迟属性。更多的细节请参见下面的部分。

表7-1在Adrenogpu中的OpenCL内存模型

内存	说明	相对延迟	位置
Local	由一个工作组中的所有工作项所共享	中等	芯片上，SP内部
Constant	工作组中所有工作项的常量	片上分配较低，否则分配较高	如果它能适应的话。否则，在系统RAM
Private	对工作项的私有信息	根据编译器分配内存的位置	在SP中作为寄存器或本地内存，或在系统RAM中（编译器确定）
Global	所有工作组中的所有工作项都可访问	高	系统RAM

7.1.1内存内容的生命周期

一个典型的问题是如何将内存对象的内容从一个内核传递到另一个内核。例如，如何与以下内核共享内核的本地内存的内容。以下是开发人员应该遵循的概念：

• 本地内存是每个工作组的，一旦工作组执行完成，其内容的生命周期就会结束。因此，不可能共享一个工作组的本地内存内容或从一个内核共享到另一个内核。
• 固定内存内容在工作组中的所有工作项中都是一致的。一旦内核执行完成，内容很可能会被在GPU上运行的其他任务所覆盖，比如图形工作负载。
• 单个工作项拥有私有内存，一旦工作项执行完成，就不能共享。但是，
• 全局内存是由主机和GPU所创建和可访问的缓冲区和图像对象支持的。因此，如果没有释放对象，则可以通过不同的内核访问它。

因此，全局内存对象是将数据从一个内核传递到另一个内核的正确方法。对于其他内存，开发人员不应该假设GPU上的一个内核的内容可以通过以下内核访问。

7.1.2本地内存

Adrenogpu支持快速的片上本地内存，而本地内存的大小因系列/层到系列/层而变化。在使用本地内存之前，最好使用以下API查询每个工作组的设备可用多少本地内存：

clGetDeviceInfo(deviceID, CL_DEVICE_LOCAL_MEM_SIZE, ... )

以下是使用本地内存的指导原则。

• 使用本地内存在内核中重复存储两个阶段之间的数据或中间结果。
  • 理想的情况是工作项多次访问同一内容。
    –例如，考虑基于窗口的使用对象匹配的运动估计。假设每个工作项在一个16x16像素的搜索窗口内处理一个8x8像素的小区域，导致相邻工作项之间的数据重叠。在这种情况下，本地内存非常适合存储像素，以减少冗余读取。
• 用于跨工作项的数据同步的障碍可能会很昂贵。
  • 如果工作项之间存在数据交换，例如，工作项A将数据写入本地内存，而工作项B从其中读取，则由于OpenCL的松弛内存一致性模型，需要进行屏障操作。
  • 障碍通常导致同步延迟，导致alu暂停，导致利用率降低。
  • 在某些情况下，将数据缓存到本地内存中会导致同步延迟，从而消除了使用本地内存的好处。在这种情况下，直接使用全局内存来避免障碍可能是一个更好的选择。
• 使用矢量化的本地内存加载/存储器。
  • 建议使用最多为32位对齐的128位（如vload4_float）的矢量化负载。
  • 请参见第7.2.2节中关于向量化数据加载/存储的更多细节。
• 允许每个工作项参与本地内存数据加载，而不是使用一个工作项来完成整个加载
  • 避免只有一个工作项来加载/存储工作组的整个本地内存。
• 避免使用名为async_work_group_copy的函数。编译器通常很难生成加载本地内存的最佳代码，而开发人员更适合编写手动将数据加载到本地内存中的代码。

7.1.3常数内存

Adrenogpu支持片上恒定内存，如果使用得当，可以在四种内存中提供优越的性能。常量内存通常在以下情况下使用：

• 用常量限定符定义的标量和向量变量。
• 如果在程序范围内定义（例如，编译器可以确定其大小），则具有常量限定符的数组，适合于常量内存。具有标量或向量数据类型的
• 内核参数。例如，以下示例中的系数将存储在固定内存中：

__kernel void myFastKernel(__global float* bar, float8 coeffs)
{ //coeffs will be loaded to constant RAM }

• 标量变量和向量变量和具有__常量的数组将被分配到系统内存中。

以下是对持续记忆的基本建议。如果一个内核具有以下两个特征：

• 一个小数组作为核参数，例如，一个5x5高斯滤波器的系数。
• 数组在子组或工作组中统一读取数组的元素

如果数组可以通过一个名为max_constant_size (N)的属性加载到常量内存中，那么它的性能就可以显著提高。该属性是用来指定此数组所需的最大字节数。在下面的例子中，为foo变量在常量内存中分配了1024个字节：

__kernel void myFastKernel( __constant float *foo 
__attribute__( (max_constant_size(1024))) { . . . }

必须指定max_constant_size属性。如果没有此属性，阵列将存储在芯片外的系统内存中，因为编译器不知道缓冲区的大小，并且不能将其提升到芯片上的常量内存中。此特性仅支持16位和32位数组，即不支持8位数组。此外，如果缓冲区太大，无法容纳恒定内存，它将存储在芯片外系统内存中。

对于被工作项动态索引和分散访问的数组，常数内存可能不是最佳的。例如，如果一个工作项获取索引0，而下一个工作项获取索引20，则固定内存效率低下。在这种情况下，使用图像对象是一个更好的选择。