导读：本文将论述OpenCL的一些深刻的常识，同时联合挪动端支流GPU厂商之一的高通Adreno芯片设计，论述挪动端OpenCL编程的优化的一些通用的伎俩。_全文5201字，预计浏览工夫14分钟。_

一、前言

在《挪动端异构运算技术-GPU OpenCL编程（根底篇）》中，介绍了挪动端GPU计算现状以及OpenCL编程的根底概念。本文将进一步论述OpenCL的一些深刻的常识，同时联合挪动端支流GPU厂商之一的高通Adreno芯片设计，论述挪动端OpenCL编程的优化的一些通用的伎俩。
备注：高通GPU系列之外，华为麒麟、联发科天机等芯片采纳的是ARM公司设计的Mali系列GPU，因为篇幅限度本文不做独自论述。

二、根底概念

丨OpenCL

OpenCL是由Khronos组织开发和保护的凋谢的、收费的规范，服务于异构零碎中的跨平台并行编程。这种设计能够帮忙开发者在利用古代异构零碎来施展出弱小的并行计算能力，同时肯定水平上也是跨平台的。

丨OpenCL高通

高通是最早齐全反对挪动端OpenCL运算能力的芯片厂商之一，并且在国内及国内市场上占据了肯定的市场份额。

三、OpenCL构造

一个形象的OpenCL利用，通常蕴含以下几个局部：

  1. CPU Host：作为整个OpenCL利用的治理、调度者，管制整个OpenCL的执行流程。
  2. OpenCL Devices：具体的OpenCL硬件设施，比方GPU、DSP、FPGA等等。
  3. OpenCL Kernels：承当执行工作的OpenCL内核代码，会被OpenCL Host进行编译，并在对应的硬件进行执行。

丨挪动端设施OpenCL

目前的教训来看，Android营垒中的挪动端设施，OpenCL通常应用GPU作为硬件加速端。高通的倡议是在挪动端抉择GPU作为OpenCL的减速设施（注：其实少数状况并无抉择可能，有且只能取得到一个GPU的device）。

四、OpenCL兼容性

丨程序可移植性

OpenCL提供了不错的程序兼容性，一套OpenCL的代码，在不同的设施上都能够失常运行。当然，少部分基于硬件拓展能力，取决于以后硬件的反对状况。

丨性能可移植性

与程序兼容性不同，OpenCL的性能可移植性通常是比拟差的。作为一种高级的计算规范，OpenCL硬件局部的实现是依赖厂商的，每个厂商都会有各自的长处和毛病。因而，针对不同的硬件平台，如高通Adreno 或者Arm Mali，同样的代码的性能体现是不同的。即便是雷同的产商，随着硬件的迭代，相应的驱动也会有对应的微调，以充分利用新一代硬件的全副能力。针对不同的设施或者硬件针对性优化是十分必要的。当然这个是一个ROI（投入回报）问题。

丨向后兼容性

OpenCL的设计尽可能的保障向后兼容性。如果要应用曾经过期的能力的话，只须要引入特定的头文件即可。值得注意的是：OpenCL的拓展是不齐全向后兼容的，这些拓展通常由硬件厂商联合硬件个性来提供，因而利用时须要思考到不同硬件之间的拓展兼容性。

五、高通 Adreno OpenCL架构

图中为高通Adreno GPU OpenCL（Adreno A5x GPUS）下层架构，OpenCL在执行过程中波及到几个要害的硬件模块。

丨Shader (or streaming) processor (SP) (着色器、流处理器)

  • Adreno GPU的外围模块，蕴含泛滥硬件模块，如算数逻辑单元、加载存储单元、控制流单元、寄存器文件等。
  • 运行图形着色器（如顶点着色器、片元着色器、计算着色器等），运行计算负载，如OpenCL内核等。
  • 每个SP对应一个或多个OpenCL的运算单元。
  • Adreno GPU可能蕴含一个或者多个SP，取决于芯片的品位，上图中展现的是单个SP的状况。
  • SP加载和读取Buffer类型或者带有\_\_read\_write标记的Image类型数据对象时，能够利用L2缓存。
  • SP加载只读的Image类型的数据对象时，能够利用L1缓存或者纹理处理器。

丨Texture Processer (TP) (纹理处理器)

  • 依据内核的调度来进行纹理操作，如纹理的读取、过滤等。
  • TP和L1缓存相结合，缩小从L2缓存中读取数据时的缓存失落几率。

丨Unified L2 Cache (UCHE) (对立L2缓存)

  • 响应SP对于Buffer类型的读取和加载，以及L1对于Image类型的数据的加载操作。

六、如何编写高性能OpenCL代码

丨性能兼容性

前文提及了OpenCL的性能兼容性，因为不同硬件的个性并不相同，因而在一块芯片上的调优后的OpenCL代码在另一块芯片上性能可能并非最优的。须要参考对应硬件的文档来进行特异性的优化工作。对于不同的芯片，针对性的优化是必要的。

丨伎俩总览

OpenCL程序的优化通常能够分为以下三类：

  • 程序、算法级别优化
  • API级别优化
  • OpenCL内核优化

程序算法以及API层级的优化伎俩是较为通用的，此处次要开展OpenCL内核的优化伎俩。

OpenCL的优化问题实质上一个如何利用内核带宽和计算能力的问题。即正当的利用全局内存、本地内存、寄存器、多级缓存等，以及正当的利用逻辑运算单元、纹理单元等等。

丨程序是否实用OpenCL

开发者须要确定程序是否适宜应用OpenCL编写，能够通过以下几个方面来判断：

  • 是否存在较大的数据输出
  • 程序自身是否是计算密集型
  • 程序是否对并行计算亲和
  • 程序中的控制流操作绝对较少

丨将CPU代码革新为GPU代码时性能Tips

明确了上述的几个关键点之后，开发者能够着手将CPU的代码转化为OpenCL的代码，为了达到一个最优的性能，须要关注以下几个方面：

  • 一些状况下，将多个CPU的操作合并到一个OpenCL内核当中能够失去性能收益。这个形式通常实用于缩小GPU和主存之间的内存拷贝。
  • 一些状况下，将一个简单的CPU程序拆分成几个简略的OpenCL内核，能够失去更好的程序并行性，进而达到全局性能最优。
  • 开发者须要思考从新设计整体的数据架构，便于缩小数据传递的开销。

这些状况要结合实际的状况进行考量，通常也是高性能异构编程自身的难点所在。

丨并行化CPU和GPU的工作流

充沛的利用芯片的计算性能，该当正当的布局工作，在GPU执行一些计算工作的同时，CPU也能够同时承当局部工作。通常能够总结为以下几点：

  • 使CPU去执行CPU长于执行的局部，比方分支管制逻辑，以及一些串行的操作。
  • 尽可能防止GPU进入闲置状态，期待CPU下达进一步工作的状况。
  • CPU和GPU之间的数据传递老本极高，为了缩小这部分老本，能够将一些自身适宜CPU进行的工作放到GPU进行。

七、性能剖析

丨性能Profile

能够联合Profile伎俩来分析程序性能。因为OpenCL程序分为宿主的CPU的调度逻辑，以及GPU硬件上的执行逻辑。开发者能够别离从CPU调度流程以及GPU执行两个层面去进行性能的Profile。通常CPU Profile是用来掂量整个流程端到端的性能，GPU Profile用来掂量OpenCL内核性能。

CPU Profile

能够采纳规范的c++编程形式，例如通过 gettimeofday 之类的api去进行CPU流程间的工夫统计。
本文中列出局部示例代码，具体demo可参考OpenCL Profile（https://github.com/xiebaiyuan...\_cook/tree/master/profile）。

#include <time.h>#include <sys/time.h>void main() {    struct timeval start, end;    // get the start time    gettimeofday(&start, NULL);     // execute function of interest    {        . . .        clFinish(commandQ);    }    // get the end time    gettimeofday(&end, NULL);     // Print the total execution time    double elapsed_time = (end.tv_sec - start.tv_sec) * 1000. +       \                (end.tv_usec - start.tv_usec) / 1000.;    printf("cpu all cost %f ms \n", elapsed_time);

GPU Profile

OpenCL提供了对GPU Kernel Profile的API，别离获取OpenCL工作的各个环节的工夫节点，便于开发者进行性能优化。

// opencl init codes ...// cl gpu time profilecl_event timing_event;cl_ulong t_queued, t_submit, t_start, t_end;// add event when clEnqueueNDRangeKernelint status = clEnqueueNDRangeKernel(runtime.queue, runtime.kernel, 1, nullptr, &ARRAY_SIZE,nullptr, 0, nullptr, &timing_event);check_status(status, "clEnqueueNDRangeKernel failed");clWaitForEvents(1, &timing_event);clGetEventProfilingInfo(timing_event, CL_PROFILING_COMMAND_QUEUED,sizeof(cl_ulong), &t_queued, nullptr);clGetEventProfilingInfo(timing_event, CL_PROFILING_COMMAND_SUBMIT,sizeof(cl_ulong), &t_submit, nullptr);clGetEventProfilingInfo(timing_event, CL_PROFILING_COMMAND_START,sizeof(cl_ulong), &t_start, nullptr);clGetEventProfilingInfo(timing_event, CL_PROFILING_COMMAND_END,sizeof(cl_ulong), &t_end, nullptr);printf("t_queued at %llu  \n""t_start at %llu  \n""t_submit at %llu  \n""t_end at %llu  \n""kernel execute cost %f ns \n""", t_queued, t_start, t_submit, t_end, (t_end - t_start) * 1e-0);

通过上述的api能够失去OpenCL Kernel从进去队列，提交、开始、完结的各个工夫点，并且能够计算出Kernel运算时长：

t_queued at 683318895157  t_start at 683318906619  t_submit at 683318897475  t_end at 683318907168  kernel execute cost 549.000000 ns
丨性能瓶颈

辨认和定位整个程序的性能瓶颈是十分重要的，没有找到性能的瓶颈，即便其余的环节性能失去优化，也无奈使得整个利用性能失去晋升。

瓶颈定位

对于OpenCL内核，瓶颈通常是内存瓶颈与计算瓶颈二者之一。
这里提供两个简略的形式，略微批改代码即可验证：

  • 退出额定的计算逻辑，如何没有影响性能，那该当不是计算瓶颈。
  • 反之，退出更多的数据加载逻辑，如何没有影响性能，那该当不是数据瓶颈。

解决性能瓶颈

胜利的定位到性能瓶颈之后，有一系列的伎俩能够去针对性的解决：

  • 如果是计算瓶颈，能够尝试一些升高计算复杂度的形式、缩小计算数的形式，或者应用 OpenCL提供的 fase relax math 或者 native math 等。在精度不高的时候能够应用fp16代替fp32进行计算。
  • 如果是内存瓶颈，能够尝试去优化内存的拜访策略，如应用向量化的内存加载和存储，利用本地内存或者纹理内存等。在可能的状况下应用更短的数据类型，能够无效的升高内存带宽。

八、总结

本文中以高通Adreno GPU举例，更加深刻的论述了OpenCL的设计思维，同时讲述了OpenCL高性能编程时一些通用的方法论。因为篇幅无限更多细节的内容没有充沛开展，对这个方向趣味感兴趣的小伙伴能够持续关注「百度Geek说」公众号。

九、参考文献

[1] OpenCL-Guide

https://github.com/KhronosGro...\_programming\_model.md

[2]OpenCL-Examples

https://github.com/rsnemmen/O...

[3]Mali-GPU

https://zh.wikipedia.org/wiki...\_%28GPU%29

[4]Adreno-GPU

https://zh.wikipedia.org/wiki...

举荐浏览：

大规模C++编译性能优化零碎OMAX介绍

百度智能小程序巡检调度计划演进之路

挪动端异构运算技术-GPU OpenCL 编程（根底篇）

云原生赋能开发测试