Qualcomm_Mobile_OpenCL.pdf 翻译-5-性能优化的概述

这章提供了一个OpenCL应用程序优化的总体概述。更多的细节将会在接下来的章节中找到。

         注意：OpenCL程序的优化是具有挑战性的。相比初始的程序开发工作，经常需要做更多的工作。

      5.1 性能移植性

         就像在2.4.2节中讨论的那样，在不同的架构之间，OpenCL一般都没有很好的性能移植性。针对某一个平台，特别是针对某个GPU优化的OpenCL应用程序，移植到Adreno GPU上后可能没有相同的性能。编程指南和其他OpenCL厂商的最佳做法，可能对Adreno GPU完全不适用。因此，针对在Adreno GPU上的优化，通读整个文档是非常重要的。此外，针对一种Adreno GPU优化的应用程序可能需要经过部分调整或者优化，才能在其他系列的Adreno GPU上达到最佳性能。

      5.2 优化的总体视角

         优化一个OpenCL的应用程序可以简单的分为一下三个级别，从高到低：

•  应用程序/算法
• API 函数
• kernel 函数

 

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

• 以最优方式使用全局内存，本地内存，寄存器和cache。
• 以最优的方式发挥计算资源的作用，比如ALU和texture操作。

        

         这章接下来的部分将会集中在应用程序级别的优化。其他层的优化将会在接下来的章节中讨论。

      5.3 对使用OpenCL进行初始的评估

         在盲目使用OpenCL之前，开发者需要先判断当前的应用程序是否适合用OpenCL优化。下面是一些适合在GPU上加速的程序的典型特点：

•  大量的输入数据
  • 　　对于少量的输入数据，CPU和GPU之间的开销可能抵消掉了OpenCL优化带来性能提升。
• 计算密集
  • 　　GPU拥有很多计算单元，而且他最高的计算能力，gflops，通常比CPU高出很多。为了充分利用GPU，应用程序需要有许多复杂的计算。
• 适合并行化计算
  • 　　工作任务可以被划分为互相独立的小单元，每一个小单元任务的处理并不会影响其他的单元任务。
  • 　　需要使用并行化任务充分利用GPU的隐藏内存延迟的能力，这是GPU最关键的一个能力。
• 有限的分支控制流
  • 　　GPU并没有像CPU那样，设计地能够处理有效的分支控制。如果使用了大量的条件判断和分支操作，CPU可能会更合适。

5.4 将CPU代码移植到GPU OpenCL

         通常情况下，对于需要转成OpenCL的代码，开发者可能已经有一个基于CPU版本的参考程序。假设这个程序包含了许多小的功能模块。将每个模块分别对应一个OpenCL kernel函数，这样看起来很方便，但是，这种情况下的性能可能不是最优的。需要考虑一下几个事实：

• 在某些情况下，将CPU的几个功能模块合并成一个OpenCL函数可能会有更好的性能，如果合并能减少GPU和内存之间的数据流量。
• 在某些情况下，将一个复杂的CPU功能函数模块分解成几个小的简单的OpenCL kernel，可能会对单个的kernel有更好的并行性和对整体有更好的性能。
• 开发者可能需要调整数据结构来适合新的数据流，这种新的数据流方式可以减少整体的数据量。

        

      5.5 GPU和CPU任务的并行

         为了充分使用SOC的计算性能，当GPU执行一个kernel函数时，应用程序可能会将指定的任务分配到CPU上。当设计这种结构和分配任务时，下面是需要考虑的几点：

• 让CPU运行适合在CPU上运行的部分，比如分支控制和顺序操作。
• 避免出现GPU空闲等待CPU执行完成的情况，或者相反情况。
• CPU和GPU之间的数据共享很耗时。所以，试着将一些轻量的CPU任务分配给GPU，尽管这些任务并不是合适GPU，这样是为了避免数据传输。

      5.6 瓶颈分析

         识别和分析瓶颈是至关重要的，因为这会使注意力集中到需要优化的区域。瓶颈导致拖延而且经常是应用程序中最慢的部分。不管其他的部分是多么有效率，应用程序的整体性能将会被最慢的那个部分限制，比如瓶颈部分。在瓶颈解决之前，关注其他部分是没有意义的。

      5.6.1 识别瓶颈

         通常情况，一个kernel要么是内存限制要么是计算限制（也可以说是ALU限制）。一个简单判别技巧是，按如下方式操作kernel代码并将它运行到设备上：         

• 如果增加许多计算并不改变性能，那么这不是计算限制。
• 如果加载大量的数据并不改变性能，那么这不是内存限制。

 

         在4.3节中讨论的骁龙profiler也可以用来识别瓶颈。

      5.6.2 解决瓶颈

         一旦一个瓶颈被确认了，可以使用不同的策略来解决它：

• 如果是一个ALU计算瓶颈的问题，找到方法减少计算复杂度和计算次数，比如在精度要求不高的情况下啊，使用更快的数学函数和内嵌的数学函数，或者使用16位浮点数代替32位浮点数。
• 如果是一个内存瓶颈的问题，试着提升内存访问效率，比如并行访问/存储，利用本地内存，或者texture cache（比如，用只读的image对象替代缓冲区对象）。使用更短的数据类型来实现在GPU和全局内存中之间存储/装载，这样能够节省内存流量。

         细节的问题将会在接下来的章节中讨论。

 

         注意：随着优化的进展，瓶颈可能会改变。如果内存限制被解决了，内存限制就会变成ALU限制，或者反之。为了获取最佳的性能，需要进行许多来来回回的迭代。

 

      5.7 API层面的性能优化

         OpenCL的API函数是执行在CPU端的，主要是管理资源和控制程序的运行。尽管，一般来说，在计算复杂度方面API函数相对于kernel的执行是很小的，但是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的内存分配。clCreate{Buffer|Image2D}会使用一个ION指针来创建内存对象，而不是分配新内存然后进行拷贝。章节7.4中讨论了如果使用ION内存。
•  在OpenCL中，尝试重复使用内存和上下文对象，避免创建新的对象。总的来说，host端需要做一些轻量级的工作，在启动GPU kernel的时候，避免阻塞GPU的执行。

      5.7.2 使用事件驱动的流水线方式

         OpenCL中入队的API函数可能会接收一个事件列表的参数，这个参数表示在当前的API函数开始执行之前，列表中的所有的事件必须执行完。同时，这个API函数同样可以产生一个时间ID来识别他们自己。如果事件列表参数正确的表示了依赖关系，那么host端只需简单地将API函数和kernel提交给GPU执行，而不需要操心他们之间的依赖关系和完成情况。通过这种方法，启动一个API函数的调用开销将会显著减少，因为软件能够按照最优方式去调度这些函数并且host端不需要在API函数调用之间进行连接（换句话说，就是不需要调用完一个后，host等待他执行完再调用另一个API函数）。因此，通过使用事件驱动的方式使得API函数的执行像流水线的方式，这种方法是非常推荐的。另外，开发者主要注意：

• 避免阻塞的API调用。一个阻塞的调用会是CPU停下来等待GPU执行完成，进而在下一次的clEnqueueNDRangeKernel的调用之前阻塞了GPU。阻塞API调用通常用在调试过程中。
•  使用回调函数。从OpenCL1.2开始，对许多API函数进行了增强和修改，API函数能够接受自定义的回调函数去处理事件，而且因为host端能够更灵活的处理事件，这种异步的调用机制会使流水线更有效地执行。

 

5.7.3 kernel的装载和编译

         实时的装载和编译kernel源码是非常耗时的。因为一些参数可能无法提前获取，所以一些应用程序宁愿运行过程中编译源码。如果生成和编译源码并不影响GPU执行，那么这是可行的。但是，一般情况下，不建议动态地生成源码。

         取代实时编译源码，一个更好的方式是离线编译源码，然后直接使用二进制kernel。当应用程序装载时，二进制的kernel代码也同样被装载。使用这个将会显著降低从磁盘中装载代码的开销。

         如果应用程序是用在不同系列的骁龙设备上，那么就需要不同的版本的二进制代码。考虑到兼容性问题，需要注意以下几点：

• 针对某一种GPU编译的二进制的代码只能在该GPU上使用。如果一个二进制是在Adreno A530的GPU的设备上编译的，那么这个二进制代码不能被用在Adreno A540的GPU上。
• 在编译器版本之间，向后的兼容性是可以达到的。新版本的编译一般会支持旧版本的二进制，不过目标GPU是要一样的。

 

         如果发现了一个不兼容的二进制kernel，使用clCreateProgramWithSource作为一个备用解决方法。

 

      5.7.4 使用有顺序的命令队列

         Adreno OpenCL平台支持乱序的命令队列。然而，在实施乱序的命令队列时需要进行依赖之间的管理，这样会导致很大的开销。Adreno软件流水命令可以发出一个顺序队列。因此，使用顺序的命令队列是比使用乱序的更好的一种的选择。