Chart 5 性能优化概述

---

前言

这一章节提供了对 OpenCL 应用程序优化的高层次概述。更多的细节将在接下来的几章中讨论。
对 OpenCL 应用程序进行优化可能是一项具有挑战性的任务。通常，它需要比初始开发更多的工作。

---

5.1 可移植性

正如第2.4.2节所讨论的，OpenCL的性能严重依赖于硬件平台的底层体系结构，通常在不同体系结构间不具备可移植性。在其他平台上优化的OpenCL应用，特别是在离散GPU上优化的应用，可能在移动GPU上表现不佳。因此，其他OpenCL供应商的最佳实践可能不适用于Adreno GPU。对Adreno GPU的优化工作至关重要，需要仔细阅读整个文档。另外，针对Adreno GPU优化的OpenCL应用可能需要进一步调整和优化，以在其他Adreno GPU上实现最佳性能。

性能、生产力和可移植性通常是开发者需要平衡的三个因素。同时实现这三个目标是困难的，因为它们取决于时间、预算和目标。例如，更好的可移植性需要更通用的代码，因为某些特定的优化技巧可能仅适用于某些GPU。同样，实现更好的生产力，即能够快速开发、优化和部署应用程序，通常需要在性能或可移植性上做出妥协。开发者必须根据他们的优先级和目标做出谨慎的决策。

5.2 优化梗概

对于 OpenCL 应用程序的优化可以从高到低大致分为以下三个层次：

• 应用程序/算法。
• API 函数。
• 内核优化。

一个OpenCL的优化问题本质上是如何最优地利用内存带宽和计算能力的问题，包括：

• 如何最优地使用全局内存、局部内存、寄存器、缓存等。
• 如何最优地利用计算资源，例如ALU（算术逻辑单元）和纹理操作。
• 应用程序级别的优化策略在本章的其余部分讨论，而其他层次的内容将在接下来的章节中详细介绍。

5.3 OpenCL 移植的初始评估

在盲目移植之前，开发者必须评估一个应用程序是否适合使用OpenCL。以下是适合在GPU上使用OpenCL加速的应用程序的典型特征：

• 大规模的输入数据集。
  • 对于小规模输入数据集，CPU与GPU之间的通信开销可能会掩盖OpenCL的性能优势。
• 计算密集型。
  • GPU具有许多计算单元（ALUs），它们的峰值计算能力通常远高于CPU。为充分利用GPU，应用程序应具有相当高的计算复杂性。
• 并行计算友好。
  • 工作负载可以被分割成小的独立单元，每个单元的处理不会影响其他单元。
  • GPU需要并行化任务以隐藏延迟。
• 控制流分支有限。
  • 通常情况下，GPU不能像CPU一样高效地处理分歧控制流。如果使用情景需要大量条件检查和分支操作，可能更适合使用CPU。

5.4 将CPU代码移植到OpenCL GPU

通常，开发者可能已经有了一个基于CPU的用于OpenCL移植的参考程序。假设该程序由许多小的功能模块组成。虽然将每个模块逐个转换为OpenCL内核似乎很方便，但性能可能不会达到最佳。考虑以下因素是至关重要的：

• 在某些情况下，将多个CPU功能模块合并为一个OpenCL内核可能会导致更好的性能，尤其是在减少GPU与内存之间的数据流量时。
• 在某些情况下，将复杂的CPU功能模块拆分为多个简单的OpenCL内核可能会产生更好的（内核）并行化和更好的整体性能。
• 开发者可能需要修改数据结构以调整数据流，减少整体流量。

5.5 并行化GPU和CPU工作负载

为了充分利用SOC的计算能力，应用程序可能会在GPU执行内核的同时将一些任务委派给CPU。在设计这种拓扑结构和分配工作负载时，需要考虑以下几点：

• 让CPU执行最适合CPU的部分，比如分歧控制流和顺序操作。
• 避免GPU处于空闲状态等待CPU完成，反之亦然。
• CPU和GPU之间的数据共享可能会很昂贵，特别是如果它们需要大量的同步或数据传输。相反，尝试将轻量级的CPU任务转移到GPU，即使它可能不适合GPU，以消除开销。

5.6 瓶颈分析

瓶颈是开发者应该花费最多时间优化的最慢的阶段。无论其他阶段效率多高，应用程序的性能都受到其瓶颈的限制。识别和分析瓶颈是至关重要的，而且并不总是直截了当的。本节简要讨论了如何识别和解决瓶颈。

5.6.1 识别瓶颈

通常，一个内核要么是受内存限制的，要么是受计算限制的（也称为ALU受限）。一个简单的技巧是操纵内核代码并在设备上运行内核，方法如下：

• 如果增加更多的计算不改变性能，那可能不是计算受限的。
• 如果过多的数据加载不改变性能，那可能不是内存受限的。

正如第4.3节所讨论的，Snapdragon性能分析工具可用于识别瓶颈。

5.6.2 解决瓶颈

一旦瓶颈被识别出来，可以采用不同的策略来解决它。

• 如果这是一个受ALU限制的问题，找到降低计算复杂性和计算量的方法。
  • 使用 fast relaxed math 或 native math（类似于fast-math的编译选项）。
  • 使用16位浮点格式而不是32位浮点格式。
• 如果这是一个受内存限制的问题，尝试改进内存访问，比如矢量化加载/存储，利用局部内存或纹理缓存（例如，使用只读图像对象替代缓冲区对象）。在GPU和全局内存之间使用较短的数据类型加载/存储数据可能有助于节省内存流量。

具体细节将在接下来的章节中详细描述。

注意：随着优化的进行，瓶颈可能会发生变化。在解决了内存瓶颈后，一个原本是受内存限制的问题可能会变成受ALU限制的问题，反之亦然。为了获得最佳性能，可能需要多次迭代。此外，不同的GPU可能存在不同的瓶颈。

5.7 API层面性能优化

负责管理资源和控制应用程序执行的OpenCL API函数大部分在CPU主机上运行。尽管这些函数的要求不如内核执行那么高，但不正确使用API函数可能会导致严重的性能惩罚。以下是一些建议，可以帮助开发者避免一些常见的陷阱。

5.7.1 API函数调用的正确安排

API函数应该被正确地放置，以防止它们阻塞或影响将工作负载启动到GPU。一些OpenCL API函数执行时间较长，应该在执行循环之外调用。例如，以下函数可能需要较长时间来执行：

clCreateProgramWithSource()
clBuildProgram()
clLinkProgram()
clUnloadPlatformCompiler()

• 为了减少应用程序启动期间的执行时间，请使用clCreateProgramWithBinary而不是clCreateProgramWithSource。详细信息请参见第5.7.3节。

当clCreateProgramWithBinary失败时，不要忘记回退到从源代码构建。然而，这可能发生在OpenCL软件有不兼容的更新时。

• 在内核调用之间避免创建或释放内存对象。clCreate{Image|Buffer}的执行时间与请求的内存量相关（如果使用host_ptr）。

• 如果可能的话，使用Android ION内存分配器。clCreate{Buffer|Image2D}可以创建具有ION指针的内存对象，而不是分配额外的内存并进行复制。第7.4节讨论了如何使用ION内存。

• 尽量重用OpenCL中的内存和上下文对象，避免创建新对象。主机在GPU内核启动期间应该进行轻量级的工作，以避免阻塞GPU的执行。

5.7.2 使用基于事件驱动的pipeline

OpenCL的enqueue API函数可以接受一个事件列表，该列表指定在当前API函数开始执行之前必须完成的所有事件。同时，enqueue API函数也可以生成一个事件ID来标识自己。主机只需将API函数和内核提交到GPU进行执行，而无需担心它们的依赖关系和完成情况，只要在事件列表参数中正确指定了依赖关系。使用这种方法显著减少了启动API函数调用的开销。软件可以最佳地安排这些函数，主机不必干涉API函数调用。因此，通过使用基于事件的管道来简化API函数是非常可取的。此外，开发者应注意以下几点：

• 避免阻塞式API调用。阻塞调用需要主机CPU等待GPU完成，然后在下一次clEnqueueNDRangeKernel调用之前阻塞GPU。阻塞API调用主要用于调试。
• 使用回调函数。从OpenCL 1.2开始，许多API函数被增强或修改以接受用户定义的回调函数来处理事件。这种异步调用机制使得主机在处理事件方面更加灵活，从而实现更有效的管道执行。

5.7.3 内核编译和构建

在运行时编译和构建内核源代码可能会很昂贵。一些应用程序可能会动态生成源代码，因为一些参数可能在一开始就不可用。如果创建和编译源代码不影响GPU执行，这可能是可以接受的。但总的来说，不建议使用动态源代码生成。

与即时构建源代码相比，更好的方法是离线创建源代码并仅使用二进制内核。在加载应用程序时，二进制内核代码也会被加载。这样做会显著减少从磁盘加载代码的开销。

5.7.4 二进制内核的向后兼容性

如果应用程序面向不同层次的Adreno设备，则需要不同版本的二进制代码。

• 二进制代码只能用于它编译的特定GPU。例如，为Adreno A730 GPU编译的二进制代码不能用于Adreno A740 GPU。
• 不能保证对于同一设备，从OpenCL软件的一个版本构建的二进制代码能够直接在新版本的OpenCL软件中重复使用。
• 如果二进制内核不兼容，可以使用clCreateProgramWithSource作为备用解决方案。

5.7.5 使用有序命令队列

Adreno OpenCL平台支持无序命令队列。然而，实现无序命令队列所需的依赖管理，存在显著的开销。Adreno软件可以高效地为有序队列流水线处理命令。因此，使用有序命令队列而不是无序命令队列是一个良好的做法。

---

总结

本文废话有点多，需要快速开发的同学，可以不用那么细看本章