关于后端:移动端异构运算技术GPU-OpenCL编程进阶篇

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

能够采纳规范的 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);

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

// opencl init codes 
...
// cl gpu time profile
cl_event timing_event;
cl_ulong t_queued, t_submit, t_start, t_end;
// add event when clEnqueueNDRangeKernel
int 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 编程（根底篇）

云原生赋能开发测试

关于后端:移动端异构运算技术GPU-OpenCL编程进阶篇

一、前言

二、根底概念

丨OpenCL

丨OpenCL 高通

三、OpenCL 构造

丨挪动端设施 OpenCL

四、OpenCL 兼容性

丨程序可移植性

丨性能可移植性

丨向后兼容性

五、高通 Adreno OpenCL 架构

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

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

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

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

丨性能兼容性

丨伎俩总览

丨程序是否实用 OpenCL

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

丨并行化 CPU 和 GPU 的工作流

七、性能剖析

丨性能 Profile

CPU Profile

GPU Profile

瓶颈定位

解决性能瓶颈

八、总结

九、参考文献