移动端异构运算技术 -GPU OpenCL 编程（进阶篇）

2022 年 6 月 01 日
本文字数：4218 字
阅读完需：约 14 分钟

导读：本文将阐述 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 应用，通常包含以下几个部分：

CPU Host：作为整个 OpenCL 应用的管理、调度者，控制整个 OpenCL 的执行流程。
OpenCL Devices：具体的 OpenCL 硬件设备，比如 GPU、DSP、FPGA 等等。
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/opencl_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 进行计算。
如果是内存瓶颈，可以尝试去优化内存的访问策略，如使用向量化的内存加载和存储，利用本地内存或者纹理内存等。在可能的情况下使用更短的数据类型，可以有效的降低内存带宽。