Chart 6 Adreno GPUs内核优化建议

半块方便面

已于 2023-12-07 14:31:31 修改

阅读量97

点赞数

分类专栏： OpenCL 文章标签：性能优化 gpu算力

于 2023-12-07 14:30:26 首次发布

本文链接：https://blog.csdn.net/qq_35227495/article/details/134849435

版权

OpenCL 专栏收录该内容

11 篇文章 0 订阅

订阅专栏

前言

这一章节提供了一些针对Adreno GPU的 OpenCL 优化技巧，更多的细节和其他信息将在接下来的章节中描述。本章中的所有建议都应该具有最高的优先级，在进行内核优化时，开发者应该在尝试其他优化方法之前尝试这些方法。

6.1 工作组性能优化

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

6.1.1 获取最大工作组大小

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

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

clEnqueueNDRangeKernel 使用的实际工作组大小不能超过 maxWorkgroupsize。如果应用程序未指定工作组大小，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 影响最大工作组大小的因素

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

内核的寄存器占用（所需寄存器数量）。一般来说，内核越复杂，寄存器占用越大，最大工作组大小就越小。可能导致寄存器占用增加的因素包括：
- 增加每个工作项的工作负载。
- 控制流。
- 高精度数学函数（例如，不使用 native math functions 或者 fast math compilation flag）。
- 局部内存，如果这导致临时分配额外的寄存器来存储加载/存储指令的源和目标。
- 私有内存，例如为每个工作项定义的数组。
- 循环展开。
- 内联函数。
通用寄存器文件（GPR）的大小
- Adreno 低版本设备可能具有较小的寄存器文件大小。
内核中的屏障。
- 如果一个内核没有使用屏障，无论寄存器占用如何，它的最大工作组大小可以在 Adreno A4x、A5x、A6x 和 A7x 系列中设置为 DEVICE MAXIMUM。

6.1.4 没有屏障的内核（流模式， steaming mode）

传统上，工作组中的所有工作项都要求同时驻留在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 工作组大小 fixed vs. dynamic

为了在不同设备上实现性能可移植性，应避免假设一个工作组大小适用于所有情况，以及硬编码的工作组大小。在一个GPU上表现最佳的工作组大小和布局在另一个GPU上可能并不是最优的。因此，开发者应该对内核能够执行的所有设备进行性能分析，为每个设备在运行时选择最佳的工作组大小。

6.1.5.4 1D/2D/3D kernel

一个内核可以支持最多三个维度。内核维度的选择可能对性能产生影响。与每个工作项仅具有1D索引（例如全局ID、本地ID等）的1D内核相比，2D内核具有这些内置索引的额外集合，并且如果这些索引有助于节省一些计算，可能会有性能提升。

根据工作项的数据访问模式，2D内核可能在缓存中具有更好的数据局部性，从而导致更好的内存访问和性能。而在其他情况下，2D内核可能导致比1D内核更糟糕的缓存抖动。尝试使用内核的不同维度以获得最佳性能是一个好的做法。理想情况下，第一个维度上的工作组大小应为 sub_group_size 的倍数，如果内核存在分歧，这一点尤为重要。

6.1.6 工作组的其他事项

6.1.6.1 全局工作大小和填充

OpenCL 1.x 要求内核的 global worksize （参考chart3 3.2.5）必须是其 workgroup size 的倍数。如果应用程序指定的 workgroup size 不符合这个条件，clEnqueueNDRangeKernel 调用将返回错误。在这种情况下，应用程序可以填充全局工作大小，使其成为用户指定的 workgroup size 的倍数。

OpenCL 2.0 解除了这个限制，global worksize 不必是 workgroup size 的倍数，这被称为非统一工作组。

理想情况下，第一个维度上的 workgroup size 应该是 wave 大小的倍数（例如，32），以充分利用 wave 资源。如果不是这种情况，请考虑填充 workgroup size 以满足这个条件。

6.1.6.2 暴力搜索

由于工作组大小选择涉及到的复杂性，实验通常是找到最优大小和形状的最佳方法。

一种选项是使用一个具有与实际工作负载相似复杂性（但可能是较小工作负载）的预热内核，在应用程序开始时动态搜索最佳的工作组大小，然后在实际内核中使用所选的工作组大小。商业基准测试依赖于这种方法。

6.1.6.3 避免工作组之间工作负载不均衡

应用程序可能在工作组之间具有不均匀的工作负载分布。例如，基于区域的图像处理可能存在一些区域需要更多资源来处理。将它们均匀分配给工作组可能会导致平衡问题。如果一个单一的工作组需要太长时间才能完成，还可能使上下文切换变得复杂。

避免这个问题的一种方法是采用两阶段处理策略。第一阶段可能收集感兴趣的点并为第二阶段处理准备数据。工作负载更加确定，更容易均匀分配到工作组中。

6.1.6.4 工作组同步

OpenCL不保证工作组的执行顺序，并且不定义工作组同步的机制。开发者不应该假设在GPU上运行的工作组的顺序。

在实践中，可以使用原子函数或其他方法对工作组之间进行有限的同步。例如，应用程序可以分配一个全局内存对象，由来自不同工作组的工作项进行原子更新。一个工作组可以监视其他工作组更新的内存对象。通过这种方式，可以实现有限的工作组同步。

6.1.6.5 (持久化线程)Persistent thread

启动一个工作组对GPU硬件来说需要时间，如果工作组的数量很大，这个成本会影响性能。如果每个工作组的工作负载较轻，这种情况尤其昂贵。因此，开发者可以减少工作组的数量，增加每个工作组的工作负载，而不是启动大量的工作组。在极端情况下，一个内核可以在每个SP上使用一个工作组，通过多次迭代完成许多工作组执行的相同任务。这种所谓的“持久线程”可以最小化硬件中工作组启动的成本，提高性能。这种方法的一个注意事项是可能会影响上下文切换，如第3.3节所讨论的那样。工作组的工作负载不应影响其他高优先级应用程序，例如用户界面（UI）的渲染。

6.2 用图像代替缓存

OpenCL支持缓冲区对象、图像对象（以及来自OpenCL 2.0的管道对象）。一维缓冲区对象是许多开发者的自然选择，因为它们简单灵活，支持指针、字节寻址访问等。

对于Adreno GPU，由于各种原因，图像对象比缓冲区对象更受欢迎，图像对象存储具有预定义图像格式的一维、二维或三维数据，包括以下优点：

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

在许多使用情况下，开发者可以期望通过用图像对象替代缓冲区对象来获得显著的性能提升，尽管这会牺牲一些主机代码的简单性和灵活性。（详看 7.1.5.3）

6.3 矢量化 load/store 和协同 load/store

Adreno GPU支持每次 读/写 全局/本地内存和图像的 128位的 load/store 事务。为了最大化内存加载效率，每个工作项理想情况下应该使用矢量化的加载/存储函数，例如使用包含四个32位数据的 vload4/vstore4，以及使用带有 CL_RGBA 和float / int32 / uint32 / half数据类型的read_image(f,i,ui,h) / write_image(f,i,ui,h)。这对于内存受限的使用情况非常有帮助。

Adreno GPU支持硬件协同加载/存储。例如，假设每个工作项加载的是内存中连续的地址上的16位数据。Adreno GPU可以合并一定数量的相邻工作项的请求，以最小化加载请求的数量。然而，与128位的矢量化加载/存储相比，这种协同加载/存储的效果较差。（详看 7.2.2）

6.4 Constant memory

Adreno GPU支持快速的片上常量内存，并使用它可以显著减少内核执行时间。在大多数情况下，编译器可以自动使用常量内存来存储一些变量，比如常量数组。然而，在一些情况下，开发者需要提供更多信息，以便编译器能够决定是否可以使用常量内存。例如，对于以下的内核：

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

缓冲区 foo，作为一个全局内存对象，可以在编译器确定其大小不超过可用常量内存的情况下，通过 max_constant_size 属性被提升到快速的片上常量内存中。

这可能导致进一步的性能提升，如果对 foo 中的元素进行的ALU（算术逻辑单元）操作是统一的，即子组或工作组中的工作项在计算时使用 foo 中的相同组件。这是因为常量内存中的内容可以在快速的ALU计算中立即传播到ALUs中。而所有其他内存（global, local, and private）必须通过冗长的加载/存储路径将数据移入寄存器，然后才能用于ALU计算。
（详看 7.1.3）

6.5 Local memory

局部内存是Adreno GPU上的芯片内物理内存。使用局部内存不一定会导致性能提升。以下是开发者应注意的一些事项：

尽量使用128位的矢量化加载/存储。
局部内存应存储最常用的数据。
- 如果只使用一次或很少使用，局部内存可能会降低性能。
推荐使用 Subgroup 函数进行 reduction 和 shuffle 操作。Subgroup 函数允许工作项在不使用局部内存的情况下共享和交换数据。
- Subgroup 函数可能不需要经过冗长的加载/存储路径。
大量使用局部内存可能限制并发工作组执行的数量，从而影响隐藏延迟。
（详看 7.1.2）