Snapdragon上的OpenCL介绍 (4)

5 性能优化概述

本章提供OpenCL应用程序优化的高级概述。更详细的讨论可以在接下来的几章中找到。

注意:OpenCL应用程序的优化是具有挑战性的。它通常需要比最初的开发付出更多的努力。

5.1 性能可移植性

正如第2.4.2节所讨论的，OpenCL通常在不同架构之间并不具有良好的性能可移植性。即便OpenCL应用程序已经在其他平台上优化，并不代表着它在其它平台仍能够表现优良，特别是在离散图形处理器上，不太可能在Adreno图形处理器上表现良好。来自其他OpenCL供应商的编程指南和最佳实践可能根本不适用于Adreno gpu。因此，阅读整个文档对于优化Adreno gpu是非常重要的。此外，为一个Adreno GPU优化的OpenCL应用程序可能需要额外的调优或优化，以在其他Adreno GPU上实现最佳性能。

5.2 优化的高级视图

OpenCL应用程序的优化从上到下大致可以分为以下三个层次:

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

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

• 使用全局内存、本地内存、寄存器和缓存等的最佳方式。
• 利用计算资源(如ALU和纹理操作)的最佳方法。

应用程序级优化策略将在本章的其余部分中讨论。其他级别将在下面的章节中介绍。

5.3 对OpenCL移植的初步评估

对于开发人员来说，在盲目地移植应用程序之前，评估应用程序是否适合OpenCL是很重要的。下面是GPU上OpenCL加速候选者的典型特征:

• 大输入数据集

  • CPU和GPU之间的开销可能会远超过OpenCL在小输入数据集上的性能增益。

• 计算密集型

  • GPUs有很多计算单元(ALUs)，其峰值计算能力gflops通常比CPU高很多。为了充分利用GPU，应用程序应该具有相当高的计算复杂度。

• 并行计算友好

  • 工作负载可以划分为独立的小单元，每个单元的处理不影响其他单元。
  • 并行任务需要充分利用GPU的内存延迟隐藏能力，这是使用GPU的一个关键好处。

• 有限的分支控制流

  • GPU并不是像CPU那样有效地处理不同的控制流。如果用例需要很多条件检查和分支操作，CPU可能更合适。

5.4 移植CPU代码到OpanCL GPU

通常，开发人员可能已经有了一个用于OpenCL移植的基于CPU的参考程序。假设程序由许多小的功能模块组成。虽然在逐个映射的基础上将每个模块转换为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 Profiler可以用来识别瓶颈。

5.6.2 解决瓶颈

一旦发现了瓶颈，就可以使用不同的策略来解决它:

• 如果这是一个ALU绑定的问题，找到降低复杂性和计算数量的方法，例如在精度要求不高的地方使用快速放松数学或本机数学，并使用16位浮点格式而不是32位浮点格式。
• 如果这是一个内存绑定问题，尝试改善内存访问，如向量化加载/存储，利用本地内存或纹理缓存(例如，使用只读的图像对象代替缓冲区对象)。使用更短的数据类型在GPU和全局内存之间加载/存储数据有利于节省内存流量。

详细信息将在下面的章节中进行描述。

注意:随着优化的进展，瓶颈可能会改变。如果内存瓶颈得到解决，那么内存限制问题就会变成ALU限制问题，反之亦然。为了获得最佳性能，需要进行多次来回迭代。

5.7 API级性能优化

OpenCL API函数在CPU主机上执行，用于管理资源和控制应用程序的执行。虽然从计算复杂性的角度来看，API函数通常比内核执行轻，但不恰当地使用API函数可能会导致很大的性能损失。这里有几点可以帮助开发人员避免一些常见的陷阱。

5.7.1合理安排API函数调用

开销较大的API函数应该被适当地放置，这样它们就不会阻塞或影响GPU工作负载的启动。一些OpenCL API函数需要很长时间才能执行，应该在执行循环之外调用。例如，下面的函数可能会花费很多时间来执行:

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

为了减少应用程序启动期间的执行时间，使用clCreateProgramWithBinary代替clCreateProgramWithSource。更多细节请参见5.7.3章节

注意:当clCreateProgramWithBinary失败时，不要忘记从源代码进行构建。如果OpenCL软件有不兼容的更新，这种情况可能会发生，但很少发生。

• 避免在NDRange调用之间创建或释放内存对象。clCreate{Image|Buffer}的执行时间与请求的内存数量有关(如果使用了host_ptr)。
• 如果可能，使用Android ION内存分配器。可以使用ION指针创建内存对象，而不是分配额外的内存并复制它。后续7.4章节将讨论如何使用ION memory。
• 尝试在OpenCL中重用内存和上下文对象，以避免创建新的对象。总的来说，主机应该在GPU内核启动期间做一个轻量级的工作，以避免GPU的执行停滞。

5.7.2使用事件驱动的管道

OpenCL排队API函数可以接受一个事件列表，该列表指定了当前API函数开始执行之前必须完成的所有事件。同时，队列API函数还可以发出一个事件ID来标识自己。主机只需将API函数和内核提交给GPU执行，而无需担心它们的依赖性和完整性，如果依赖性在事件列表参数中被正确指定。通过使用这种方法，启动API函数调用的开销显著减少，因为软件可以以自己的最佳方式调度函数，并且主机不必在API函数调用之间进行干预。因此，非常希望使用事件驱动的管道来简化API函数。此外，开发者应该注意:

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

5.7.3 内核加载和构建

在运行时加载和构建内核源代码可能代价昂贵。有些应用程序可能更喜欢实时生成源代码，因为某些参数可能无法预先使用。如果源代码的生成和编译不影响GPU的执行，这可能是好的。但一般来说，不鼓励动态源代码生成。

与实时构建源代码不同，更好的方法是离线构建源代码，并且只使用二进制内核。当应用程序被加载时，二进制内核代码也被加载。这样做将显著减少从磁盘加载代码的开销。

如果应用程序针对不同层次的Adreno设备，则需要不同版本的二进制代码。关于兼容性，需要注意以下几点:

• 二进制代码只能用于编译它的特定GPU。如果一个二进制文件是在拥有Adreno A530 GPU的设备上构建的，那么它就不能在拥有Adreno A540 GPU的设备上使用。
• 向后兼容性可跨编译器版本。如果目标GPU相同，编译器的新版本可能支持较旧的二进制文件。

如果发现不兼容的二进制内核，请使用clCreateProgramWithSource作为回退解决方案。

5.7.4使用顺序命令队列

Adreno OpenCL平台支持无序命令队列。但是，由于实现无序命令队列所需的依赖关系管理，因此会产生更大的开销。Adreno软件将命令输送到有序队列。因此，最好使用有序的命令队列而不是无序的命令队列。