OpenCL工作组大小性能优化

6 OpenCL工作组大小性能优化

工作组大小调整是许多内核的一种简单有效的优化方法。 本章介绍有关工作组大小的基本信息，例如如何获取工作组大小，为什么工作组大小很重要，还讨论了有关最佳工作组选择和调整的一些常见做法。

6.1 获得最大的工作组规模

运行clBuildProgram之后，可以使用以下API函数查询设备上内核的最大工作组大小：

size_t maxWorkGroupSize;
clGetKernelWorkGroupInfo(myKernel,
myDevice,
CL_KERNEL_WORK_GROUP_SIZE,
sizeof(size_t),
&maxWorkGroupSize,
NULL );

clEnqueueNDRangeKernel使用的实际工作组大小不能超过maxWorkGroupSize。 如果应用程序未指定工作组大小，则Adreno OpenCL软件可以选择最大工作组大小。

6.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)
{
. . .
}

在大多数情况下，如果施加了工作组大小限制，则编译器不能保证生成最佳的机器代码。 另外，如果编译器不能使用片上寄存器满足所需的工作组大小，则可能不得不将寄存器溢出到系统的随机存取存储器（RAM）中。 因此，不鼓励使用这两个属性，除非以某种指定的工作组大小对内核起作用的方式编写内核。

注意：
为了实现跨平台兼容性，编写依赖于固定工作组大小或布局的内核不是一个好习惯。

6.3 影响最大工作组规模的因素

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

注册内核的占用空间。 通常，内核越复杂，寄存器占用空间越大，最大工作组大小也越小。 可能会增加寄存器占用空间的因素如下：

• 为每个工作项目分配更多工作量
• 控制流程
• precision高精度数学函数（例如，不使用本机数学函数或快速数学编译标志）
• 本地存储器，如果这导致分配额外的寄存器来临时存储加载/存储指令的源和目标。
• 专用存储器，例如为每个工作项定义的数组。
• 展开循环
• 内联功能

通用寄存器（GPR）文件的大小

• Adreno低层可能具有较小的寄存器文件大小。

内核中的障碍

• 如果内核不使用障碍，则可以在Adreno A4x和A5x系列中将最大工作组大小设置为DEVICE MAXIUMUM，而不论寄存器占用空间如何。

6.4 无障碍的内核

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

从Adreno A4x系列开始，无障碍内核可以具有Adreno支持的最大工作组大小，通常为1024，尽管它们很复杂。 由于波形之间不需要同步，因此对于这些类型的内核，新波形可以在旧波形完成后开始执行。

在这种情况下，拥有最大的工作组大小并不意味着它们具有良好的并行性。 没有障碍的内核可能非常复杂，以至于SP中只有几波并行运行，这将导致性能下降。 无论从clGetKernelWorkGroupInfo函数获得的最大工作组大小如何，开发人员都应继续优化并最小化寄存器占用空间。

6.5 工作组大小调整

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

6.5.1 避免使用默认工作组大小

如果内核调用未指定工作组大小，则OpenCL软件将使用一些简单的机制来查找工作组大小。 开发人员应注意，默认工作组大小不太可能是最佳的。 手动尝试不同的工作组大小和形状布局（对于2D / 3D）并找到最佳的工作组始终是一个好习惯。

6.5.2 工作组规模较大，性能更好吗？

对于许多内核而言，这是正确的，因为增加工作组的大小可以使更多的波驻留在SP上，这通常可以转化为更好的延迟隐藏和更高的SP利用率。

但是，对于某些内核，性能可能会随着工作组大小的增加而降低。 一个例子是，较大的工作组由于不良的数据局部性和访问模式而导致高速缓存抖动增加。 对于纹理访问，局部性问题也很严重，因为纹理缓存通常小于统一的L2缓存。 最终，决定最佳工作组大小和形状的是内核中数据访问的性质。

6.5.3 固定与动态工作组大小

为了实现跨设备的性能可移植性，请避免假设一个工作组大小适合所有情况，并避免对工作组大小进行硬编码。 在一个设备上最有效的特定工作组大小和布局在另一个设备上可能不是最佳选择。 因此，对于给定的内核，建议为内核可以执行的所有设备配置不同的工作组大小，并在运行时为每个设备选择最佳的工作组大小。

6.5.4 一，二和三（1D / 2D / 3D）工作组

内核的维数可能会影响性能。 根据工作项的数据访问模式，在某些情况下，二维内核可能在缓存中具有更好的数据局部性，从而导致更好的内存访问和更好的性能。 在其他情况下，与1D内核相比，2D内核可能导致更严重的缓存抖动。 尝试对内核使用不同的尺寸以获得最佳性能将是一个很好的选择。

6.6 有关工作组规模的其他主题

6.6.1 全局工作规模和填充

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

注意：
OpenCL 2.0取消了此限制，并且全局工作大小不必是工作组大小的倍数，这称为非统一工作组。

理想情况下，工作组在其第一维度上的大小应为波形大小的倍数（例如32），以充分利用波形资源。 如果不是这种情况，请考虑填充工作组大小以满足此条件，请记住，OpenCL 1.x可能还需要填充全局工作大小。

6.6.2 暴力搜索

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

一种选择是使用具有与实际工作负载类似的复杂性的预热内核（但可能使用更小的工作负载），以在应用程序启动时动态搜索最佳工作组大小。 然后将所选的工作组大小用于实际内核。 许多商业基准测试都依靠这种方法。

6.6.3 避免跨工作组的工作量不均

某些应用程序的编写方式可能会导致跨工作组的工作负载不均衡。 例如，某些基于区域的图像处理用例可能具有一些比其他区域需要更多处理的区域。 应避免这种情况，因为性能可能无法预测。 此外，如果任何单个工作组花费的时间太长，它会使上下文切换复杂化。

解决此问题的方法是采用两阶段处理策略。 第一阶段可以收集有趣的点并准备数据以供第二阶段处理。 工作负载更具确定性，这可能使在工作组之间平均分配更加容易。

6.6.4 工作组同步

OpenCL不保证工作组的执行顺序，也没有定义执行工作组同步的机制。 不建议使该应用程序依赖于工作组的顺序。

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