视频云转码-GPU矩阵乘-ARM CPU 分析

视频云转码-GPU矩阵乘-ARM CPU 分析

参考文献链接

https://mp.weixin.qq.com/s/oroRvkeafKqXxR-ad0WOLA

https://zhuanlan.zhihu.com/p/158161592

https://zhuanlan.zhihu.com/p/353975826

AI时代视频云转码的移动端化：让模型设计/推理库/硬件成为一体

一套轻量级AI算法模型，需要结合移动端硬件特性，差异化优化前馈推理库，让算法模型、推理库、硬件成为一体，使得视频云转码移动端化成为可能。本文将分享ZEGO在产品架构、移动端视频转码、移动端智能视频处理、四位一体网络模型设计上的思考。

 

 编者按：AI技术的落地是渐渐地从服务器端、云端落地，逐步到移动端及边缘设备上。这些年随着AI技术的进步，轻量级算法模型开始在移动端实时跑起来，并且移动端算法也在不断进行迭代和完善，而对于实时直播场景，也有越来越多的AI算法落地。

ZEGO即构科技提出了一套极轻量级AI算法模型，结合移动端硬件特性，差异化优化前馈推理库，让算法模型、推理库、硬件成为一体，使得视频云转码移动端化成为可能。本文将分享产品架构、移动端视频转码、移动端智能视频处理、四位一体网络模型设计以及具体实施Demo。

疫情三年我们的模型会变得越来越小、越来越快、效果越来越好，这就是“更快、更好、更低、更广”。

快 —— 速度越来越快，好 —— 效果越来越好，低 —— 耗费很低的码率，广 —— 尽可能覆盖更多的机型。曾经在云端转码遇到大批量的任务，本文讨论在移动端如何处理，分享内容更侧重在移动端如何把任务快速跑起来，把服务器端的事情放在手机端做。

 

 

 本次分享按以下方向逐一展开：一是 ZEGO 产品概貌；二是要在移动端去做的理由，这是帮助解决客户提出的需求问题；三是在移动端可以做到哪种程度，能不能面对用户提出的苛刻需求，解决客户问题；四是如何解决行业内难题，这是任何一个做音视频厂商、做视频前后处理都孜孜不倦追求的问题；五是简单实验结果分享。

ZEGO 于2015年6月成立，是一家全球云通讯服务商。企业/开发者通过接入ZEGO提供的服务即可低门槛地获得实时音视频通讯能力，尤其在弱网环境下仍然可以实现高质音画、稳定可靠。专注自研音视频引擎，在音频前处理、网络自适应和跨平台兼容性等方面，达到国际一流水平，同时充分利用基础云服务商的能力，构建了MSDN海量有序自学习数据网络，服务覆盖全球，涵盖上百个音视频互动业务场景，单日时长突破30亿分钟。

1、产品架构

 

 

 产品应用方面，会议端的产品小艺帮主要针对考试类的应用；直播类有游戏直播，KTV唱歌等；方案有社交娱乐、泛娱乐、以及这几年很火的在线教育。以及考试类应用，不限于在线教育的方案。我们同时会提供基础技术服务，比如视频转码、画质增强、美颜特效、Avatar 元宇宙底层技术、背景分割等。 

1.1视频处理

 

 本次分享的重点是视频处理。上图是在IPhone XR上不同分辨率下的耗时。

在 iPhoneXR 这样的中端机型上，超分辨率 640*480 两倍率做 AI 超分到 1280*960，21ms 可以搞定。同时也测试了安卓机型，大概可以覆盖 2500 款，1000 元以上机型基本能够覆盖。如果机型好一点，超分到 1080p 没有问题，但国家政策鼓励出海去东南亚、印度、非洲。这些国家市场非常大，但最大的现状为使用者大多为千元机以下，这就面临一个很大的难题，千元机难做超分。
去编码速度比超分速度快很多，超分天然的数据量很大，不能把图片缩小，将图片原始分辨率不做改动的情况下放大。去编码效应的优势在原始分辨率上做，处理速度比超分快很多，基本可覆盖千元机。而在印度、东南亚、非洲这些带宽很低的地方，这时做去编码效应比超分，更具实际价值。
最难的点是视频插帧，两帧插一帧，基础数据是在 IPhoneXR 上 960*540，32ms，这一部分依然面临难覆盖千元机，也是我们一直去追求的点，后续可以演示插帧在移动端上手机具体效果。
随后也会与大家分享如何实现低照度增强，大家也可以跟着这个思路复现低照度增强，速度可以覆盖到 2ms，很低端的机型例如 2015 年的机型小米2S，2ms 可以实现。道理很简单，就是 3D或 2D 查找表，非常快，效果在最后一部分演示。 

1.2其他AI技术产品

 

 技术人员要经常与销售打交道，面对一线用户，技术人员就能了解是否要解决此问题，知道需求在哪儿，做的技术有无价值。当然我们也会做其他的应用。一是抠图，其中有隐私反向打码，例如小米摄像头监控功能，老人在家监控但不放心，就可以把老人反向打码，将老人虚化而背景不虚化，或只虚化老人身体的一部分，反向抠图，关键看你怎么去做去用。二是小目标检测，曾经有用户给一个条件：小分辨率在 640*360，只有 10 像素的手机，考场监控场景下，看拿手机作弊的小动作能否被检测出。这样的难题是实际中碰到过的很多基础应用。

2、移动端视频转码

 

 第二部分是移动端视频转码。 

2.1 移动端转码

 

 关于移动端视频转码，也有一个来自于真实用户提出的需求：为什么要将视频上传至云服务，云服务下载后又要转一遍，用户自己拍的视频或是转的视频不愿意上传至服务器，那么移动端是否能将视频修复呢？

视频修复需要做很多任务，首先要从云转码角度分析是什么场景下的视频，游戏、体育或是日常生活，质量如何、块的效应如何、编码状态如何。一个视频块很多，但也许是录屏的，初步分析后考虑是否要做超分、做编码块效应、做降噪，那这一整套能不能在手机端做到？用户提出此问题后，如果可以解决用户就愿意付费，但用户不想放在云端，ToB 公司不乐意放在云端为单个服务器付费，在用户的手机将事情落地对于 ToB 公司来说是乐意付费的。所以我们很多的技术实现都是源自于具体生活。到底能不能实现，实现到什么程度，后面会分享具体一些数据，看能否支撑能否做到。 2.2 移动端 AI 技术性能一览
  

在上图的超分性能数据中可以看到，我们可以覆盖 2500 款机型，能给到厂商实测。最难的是移动端实时视频插帧，需要解决人物不扭曲、大运动不虚、字幕不花等场景，并且机型覆盖程度仍在努力中。刚才提到的低照度增强，提升通道通透性指的是视频进入编码之后，会有朦朦一层，没有深度参差，将其处理，使之纵深变强，立体感更强，而不是像有一层雾，当然也可以用去雾通道处理此问题。在降噪方面，我们本身也在做各种优化。抠图和小目标检测在前文中也有提到，我们主要做安卓端。

3、移动端智能视频处理

 

 第三部分是移动端智能视频处理。 

3.1 全链路自研可控

 

  上图中是公司大型团队端到端做的事情。我们所负责的是与视频相关的业务，包括视频前后处理部分、麦克风采集、涉及端到端的网络，以及部分编解码。 

3.2 移动端视频处理
  

 

 上图是我们部门在移动端视频处理做的事情。整个数据流是这样走的，核心重点还是在视频后处理或视频前处理上，包括前端推理库的优化，后续会讲到前端推理库的一部分优化和网络模型如何设计得更加小型化。 

3.3 移动端超分辨率

 

 再看一下移动端超分辨率具体技术细节。 

  • 超分辨率-性能PK

 

 超分是我们去年上半年做的事情，与行业内进行 PK，在上图数据上做测试，越向左模型越小越好，越往上效果越好，左下角都是一些小模型的效果。上图中 PAN 在2020年模型参数量最小，效果排第二。RFDN 在 2020 年效果第一，与之作初步对比大模型的位置如图，但实际覆盖大概 2500 款机型使用不同的模型。 

  • 超分辨率网络模型-设计

 

 上图是超分很简单的网络模型设计，期初花了很多心思精力阅读大量文章加上自己的奇思妙想、胡思乱想，我们做了一些通道处的处理，包括通道注意力机制，但最终发现流程越多，模型加速可能性越低。后来回到初衷，原始设计出来的模型，最后还是这样的，这样做加速会非常容易。如果模型效果好但没法加速也不行，结构比较复杂。

 

 

 这一部分我们也发了一篇文章，影响因子5.9（https://doi.org/10.1016/j.neucom.2022.07.050）。模型怎么训练，其中就有一些弯弯绕绕，涉及到后续会统一提到的技术点。 

3.4 移动端低照度

  • 低照度设计

 

 

 

 低照度方面，第一版是用 2D 查找表方式，发现其维度不够，后改用 3D 查找表，只要进来数据，有三个维度分别是图像输入的原始像素点、当前像素点周围邻域的均值、根据自己定义例如当前像素点Y通道的亮度与周边像素点亮度差异值，在 0-1 之间。将其变成 0-255，此表最大为 255*255*255，最终结果是几毫秒。 

  • 网络模型-设计指导方向

 

 回到视频超分辨率的问题，如何将视频超分辨率模型变小，可以做知识蒸馏；重参数，这在最近几年很火，美团的 YOLOv6 就用了重参数、台湾的YOLOv7 同样使用重参数，重参数同样变成了在 CV 基础业务领域的骨干网络，重参数的意思是训练中模型非常大，推理模型非常小，但训练模型和推理模型是等价的，最终效果一样，推理速度更快，依据这个可以看更多的文章。

还有就是做量化 PQT 和 QAT，训练前和训练后量化会带来不同效果，后续会提到一个模型的设计是否可以转换成量化，是否适合，精度是否下降，模型本身与推理库是否紧密耦合，模型在硬件上借助推理库能否达到最佳，这是个考验人的事情。对所有底层计算机视觉视频增强任务大部分可以用查找表来做，例如锐化，二维三维都可以实现，对照度、对比度增强，甚至做得好可以做 8bit 到 10bit 转换，速度性能非常快。这里仅仅是给了大致的方向。

4、四位一体网络模型设计

 

 第四部分比较关键，四位一体网络模型设计。

 

 

  四位一体网络模型设计，虽然是四个，但上图中只有三个，大家可以猜一下还有哪一个。

第一维度是模型设计；第二维度是模型推理引擎，例如 TNN、NCNN、MNN 等很多，基于别人的推理引擎去改和优化，推理引擎与网络模型是否能很好结合，模型本身在不同的推理引擎上速度差别大，模型在 iPhone 上是跑在 CPU 还是 GPU 还是 NPU上，iPhone 的 Core ML 是个黑盒子，如果 OPS 不适合在 NPU上跑，数据就有可能一会儿在 GPU，一会儿在 NPU，模型效率低；第三维度是 Peak，意思是需要去度量模型，操作算子在指定硬件上，例如指定的硬件如高通、联发科、华为芯片上，它的性能到底如何，一个模型在同一时代的机型上跑出的数据效率不同，最重要的是要将 roofline 模型图画出，反向指导我们设计模型。
因为目标就是性能和精度，我们认为第四个维度是数据，数据占了功效的 60%。举个例子，超分可以刷榜刷得很厉害，但到实际场景就效果不好，一定是要根据业务指定场景强相关。我们用自定义数据集，标注好的 3 万张图片，将所有 YOLO 系列的模型参数全部设置在差不多的数量级，初始化一样，训练结果发现 YOLOv5 精度较高。但在 YOLOv7 文章中表示从零开始直接训练比多次训练效果好，这表明可能是在特定场景效果好，但在另外场景下是不一定的。
  

 

 这里更细致地讲一下在推理引擎端，模型设计尽可能契合硬件特点优化（下面这段文字，我大改了一下，原公众号的内容根本没法看，写的是个什么东西）。

  • 模型设计要结合roofline model，将roofline model在单算子上的特点反馈到模型中优化模型设计；
  • 算子本身可以合并也可以拆分，将复杂构重新拆分组合成合适的算子；
  • 再就是优化指令流水线得到最佳的组合。CLPeak 和 RAMMER 做了硬件 CPU、GPU 在移动端 NPU 的性能度量，还有旷视 MegPeak ，都在做说明做类似的MircoBenchMark是值得的。

 

 模型方面，我们认为数据是最重要的，要保证百分之百准确，比如做目标检测的标注，首先运用大模型预标注，人工筛检，再训练，再预标注，反复几次。标注的时候只标注一类，比如生成部署标注十类，实际上做的时候一类一类标注，这样就保证每一类精度非常高，但没人愿意标注，觉得是粗活脏活，但这是把数值刷最高最有效直接快速的方法，丢给外包质量不好，包括算法同事标注质量也差。所以要想出各种办法通过算法本身将数据筛选出来做更多的事情，达到效率提升（看样子是内包给各位开发了）。  

 

 上面分享的效率、速度需要同时上升，主观感受质量不能下降，但模型可以裁剪。

走向ARM CPU 1-bit推理的极致道路

深度学习的端侧学习对轻量化有着极致的追求，从2016到现在，我们看到业界在推理时使用的数据类型位数一降再降：FP32、FP16、INT8……。显而易见，低比特能节省模型占用的存储空间和功耗，在某些精心设计的算法上面也能获得性能优势。在这一点上，1-bit毫无疑问是经典计算机上可以做到的极致。

举一个实际的例子，Birealnet18（二值化的Resnet18）用FP32类型存储大概需要45MB，而通过Bolt转换后得到的1-bit+FP16存储只需要占用2.7MB，exciting！

低位数与性能提升的可能

那么，权重位数走向极致，是不是就能带来计算速度的提升呢？答案取决于你使用的硬件。

现阶段深度学习的核心计算还是矩阵乘法，这里面涉及到的就是乘法和加法。大多数硬件的ALU可以处理的最小数据类型就是INT8，要是按照严格的矩阵乘法定义，低于8位的数据类型都得先展开到8位来操作（INT4和INT2哭晕在厕所）。但是1-bit可以来个“降维打击”：权重只有0和1，干脆省掉乘法，直接全都做加减法，多么HAPPY。看到这里，你已经领会到Binary Weight Net的精神了，在经典CPU上已经可以靠着加法相比于乘法的优势行走江湖了。

还没有换台的观众们，估计眉头一锁，觉得此事没有那么简单。是的，深度学习早已经皈依了黄老板的皮衣神教，移动端CPU的绝对霸主ARM也推出了SIMD单元。这些个SIMD单元，乘法和加法用的就是同一个单元，甚至可以用一个指令完成乘加操作。

在现有的SIMD上，1-bit要想占有一席之地，必须再次发动降维。当我们把 feature map 也二值化了后，乘法的两端只会出现0和1，而最终的乘积也只会是0和1，直接位操作就行了，总有一款适合你，而且还不用改动现有的硬件——这就是BNN的精髓。

试想一下，中低端SoC是不支持FP16和INT8的专门优化的，能找到的估计就是跑FP32的CPU。部署FP32的模型，不仅存储感人，推理时间也会感人。但你要是训出一个BNN模型，存储一下子就下去了，而且还能轻松在原有芯片上打败FP32。有些读者可能在担心精度问题，但BNN已经取得长足的进步了，不搞识别，搞个实时检测还是很有价值的。

如何用Bolt基于ARM NEON进行BNN推理

接下来，本文将会介绍Bolt是怎样基于ARM NEON进行BNN推理的。

怎样基于ARM NEON进行BNN推理的。在这个事情上，Bolt没有抢到先发，我们被大缺弦老师的BNN框架（daBNN）截胡了。但我们也要感谢大老师选择了开源，特别是把他训练的birealnet18模型也放到了repo上，做BNN框架才有了真正意义上的benchmark。

在我们调试该网络的过程中，他也知无不言，衷心感谢（https://github.com/JDAI-CV/dabnn）。

在一款搭载A55小核心的开发板上，daBNN实测birealnet18（带stem优化版本）单张图片是243ms，而Bolt的单张推理时间是78ms——对于这个结果我们当时也是很震惊的。

关于Bolt相比daBNN的性能提升原因，大家容易想到的第一个因素是在模型的浮点部分，Bolt用的是FP16啊，性能优势是这样来的吗？但是，我们最初开始开发的时候，直接把整个网络按FP16来跑，也需要160ms。其实，在A55小核心上，由于关键指令无法双发射，daBNN是在客场作战的——Bolt几乎所有汇编kernel都有针对A55进行指令发射的优化（注重能耗的开发者记得pick我们）。那么，性能对比在A76上又如何呢？请容许我卖个关子。

接下来先进入技术解密，看完你就更能理解Bolt BNN推理为什么有这么大的优势了。

技术解密：ARM CPU做BNN推理的基本思路

正如大老师在他的论文中讲的，ARM上面做BNN推理最重要的是以下两个指令：

  • AND：逻辑门位操作，要是使用XNORnet那就换一条逻辑门指令；
  • CNT：在向量寄存器上每8个比特统计1的数目，得到0到8之间的数字，按照UINT8存回到那8个比特上

这两条指令应该怎么理解呢？在64位的芯片上，向量寄存器的宽度是128bit。一条AND指令就可以完成128个乘法，而一条CNT指令就完成了16组8个数字的加法。换言之，这两条指令就相当于16组8x8的向量内积。

ARM CPU做BNN推理：daBNN的做法及Bolt的改进

那么，接下来就要说到Bolt和daBNN的重要区别了。大老师的128bit装的都是对应同一个输出通道的操作数，所以16组8x8内积得到的16个UINT8数字，还得通过ADDV指令加到一起，然后才加到对应的输出位置上，并行度突然就砍掉了。大老师也发现了ADDV指令是很吃亏的，于是他尽可能先把很多组16个UINT8数字通过向量ADD加起来——毕竟UINT8最大能存255嘛，只用8很浪费——然后再用ADDV。

但是，大家有没有想过可以完全省掉ADDV呢？AND+CNT输出的16个数字，要是对应的是16个输出通道，那我们就只需要直接ADD起来了。这其实是经典的gemm tiling思想，就不在此赘述了。

可能有课代表发现了，直接ADD起来，顶天也就只能加到255啊，不可能一路ADD下去。是的，BNN卷积的结果大多数情况至少也要用INT16才能存储，我们的确也是需要累加到INT16向量的。跟上面一样的道理，为了节省INT16累加的指令，我们需要在ADD的时候尽量用满UINT8的空间。理论上，8x8内积最多产生数字8，而UINT8最大能存255，也就是最多约可以累加32次，追求极致的我们需要想这个32怎么来。

卷积的时候的每一组加法次数，等于ic*fh*fw（input_channel*filter_h*filter_w）。在Bolt中，我们要求BNN卷积层的输入通道数必须是32的倍数，每次8x8内积会用掉ic里的8，那么ic/8我们只能保证是4的倍数，距离32还差一个8。这个8要怎么来呢？欢迎大家想一想。

daBNN和Bolt的性能对比

 

 

 图 1-bit推理性能时间（ms）对比

由于stem版模型的onnx比较晚才放出来，做这张图的时候就没加上去。Bolt+stem在A55上是70.9ms，而在A76上是17.0ms。

Bolt里面的每一个kernel，每一个精度，甚至每一个功能，都经历了如上所示的严密推敲和不断优化，欢迎大家多多体验，这也是我们开源的初衷。大家如果想到了更好的做法，那就太好了，贡献出来，我们一起走向极致。

相关成果我们已经开源到Github：https://github.com/huawei-noah/bolt，方便大家复现和使用，欢迎大家试用反馈体验效果，在社区积极讨论。

高性能，精确，轻量级，易用和安全是我们不懈追求的目标，未来我们会继续利用高性能计算优化技术和编译技术，加速深度学习，在计算机视觉和自然语言处理领域发力，将更多的研究结果带到社区。

移动端GPU矩阵乘优化

移动端GPU目前主要有3家供应商， Qualcomm的Adreno系列，Arm的mali系列和Imagination的PowerVR GPU。主流开发语言包括OpenCL、OpenGL以及Vulkan，本文不对各个语言的应用进行讨论，仅以OpenCL为例。不同设备的体系结构差异很大，即使同一供应商的设备，也存在多个系列，因此优化策略也有不同。本文仅介绍纹理内存在Adreno和Mali设备上所带来的性能提升。

在GPU上的内存一般分为两种，一种是普通内存，OpenCL中叫做buffer内存，一种是纹理内存(Texture内存)， OpenCL中叫做Image内存。纹理内存和普通的buffer内存是通过不同的硬件单元来加载和写入的。除此之外，移动端设备上，不同的GPU架构下，对纹理内存和buffer内存的访问都存在差异；例如高通设备上，纹理内存的读可以使用L1 Cache，Mali设备上虽然无此差异，但是Mali确在最近几代GPU架构的迭代中不断的增强纹理内存的访存能能力。

本文主要从以下几个方面展开：

  1. 测试环境介绍
  2. 基础优化版本
  3. Adreno设备的Texture方案
  4. Mali设备的Texture方案
  5. Mali(ValHall)的FMA方案
  6. 其他优化方案简介

测试环境及指标介绍

本文测试设备使用Qualcomm 865芯片和MTK的天玑1000芯片，对应GPU为Adreno 650 及Mali的G77 MP9，峰值数据是实际测试乘加计算的吞吐，非理论峰值.

GPU

FP16峰值(FP32峰值)

Adreno 650

1407(764)

Mali G77(MP9)

881(445)

测试数据

矩阵维度为：A的维度为M x K, B的维度为K x N, C维度为M x N， 其中(M=N=K=1024); 测试数据采用float16 随机数进行测试。

指标计算

评价指标采用GFLOPS , 计算方式为(M * N * K) * 2 / 1024 / 1024 / 1024 / computeTime(s);

使用OpenCL的event机制对计算kernel计时，计时之前会循环调用10次该kernel进行warm up；随后对该kernel循环调用20次，取平均值作为执行时间。

基础优化版本

直接实现版本

首先按照矩阵乘法的计算公式，实现最简单版本作为base，如下如图所示，A矩阵的第一行乘以B矩阵的第一列得到C矩阵对应行列的一个元素。

 

 

 代码实现如下：

  1. // global_work_size = {N, M}
  2. #pragma OPENCL EXTENSION cl_khr_fp16 : enable
  3. __kernel void gemm_opt(__global half* A, __global half* B, __global half* C, int M, int N, int K)
  4. {
  5.     int idx = get_global_id(0); // 0--(N-1)
  6.     int idy = get_global_id(1); // 0--(M-1)
  7.     if(idx > N || idy > M) return ;
  8.  
  9.     int a_index = idy * K;
  10.     int b_index = idx;
  11.  
  12.     half cval = 0;
  13.     for(int i = 0; i < K; i++)
  14.     {
  15.         cval += A[a_index + i] * B[b_index + i * N];
  16.     }
  17.  
  18.     int c_index = idy * N + idx;
  19.     C[c_index] = cval;
  20. }

该实现版本性能如下：

 

 

 该版本可以看出，对于矩阵A的访问步长为 K * sizeof(float), 显然不满足GPU访存合并的原则。其次，计算过程中存在大量的数据重复加载，例如A矩阵的第一行数据，会在计算第一行每一列数据的时候被反复加载。

合并访存优化

首先，可以将矩阵A进行转置以达到访存合并，其次单线程可以计算更多的输出点，以减少数据的重复加载，向量化加载也可以更好的提高带宽利用率；

优化方案如下图所示：

 

 

 转置后使用A的一列与B的一列乘累加，得到C的一个点；代码实现如下：

  1. // global_work_size[] = {(N + 3)/4, (M + 3) / 4}
  2. // 单线程计算16个点；
  3. // 读者可以在不同架构的设备上尝试其他方案，虽然单线程计算点越多重复加载数据越小，但是也可能导致寄存器溢出，性能反而下降严重
  4. #pragma OPENCL EXTENSION cl_khr_fp16 : enable
  5. __kernel void gemm_opt(__global half* A, __global half* B, __global half* C, int M, int N, int K)
  6. {
  7.     int idx = get_global_id(0) << 2;
  8.     int idy = get_global_id(1) << 2;
  9.  
  10.     if(idx > N || idy > M) return;
  11.  
  12.     half4 cval[4] = {(half4)(0), (half4)(0), (half4)(0), (half4)(0)};
  13.     for(int i = 0; i < K; i++)
  14.     {
  15.         half4 a = vload4(0, A + idy + i * M);
  16.         half4 b = vload4(0, B + idx + i * N);
  17.  
  18.         cval[0] += a.s0 * b;
  19.         cval[1] += a.s1 * b;
  20.         cval[2] += a.s2 * b;
  21.         cval[3] += a.s3 * b;
  22.     }
  23.     vstore4(cval[0], 0, C + idy * N + idx);
  24.     vstore4(cval[1], 0, C + (idy + 1) * N + idx);
  25.     vstore4(cval[2], 0, C + (idy + 2) * N + idx);
  26.     vstore4(cval[3], 0, C + (idy + 3) * N + idx);
  27. }

该版本实现性能如下：

 

 

 从数据可以看出，目前版本相对于直接实现版本提升了17倍，Adreno的实现提升8倍。其主要原因在于缺少L1 Cache的加持Adreno设备的buffer吞吐远低于Mali设备的吞吐。接下来通过使用Texture内存对两种设备做进一步的优化。

Adreno设备的Texture方案

下图是Qualcomm文档中关于纹理内存的描述，

 

 

 从图中可以看出，shader在加载数据的时候，texture内存和buffer内存是通过不同的通道进行的，texture内存的加载可以使用到单独的Texture Processor/L1 Cache,而buffer内存的加载只能使用L2 Cache，因此合理的使用Texture 内存存储数据可以进一步提升上一版本性能。

Texture和buffer内存一般是通过不同的硬件单元进行加载的，所以，在使用纹理内存的时候，是选择A/B其一存储在Texture 内存，另外一个存储到Buffer内存呢？还是选择两块内存都使用Texture呢？

这里给出结论，Qualcomm上使用双Texture内存，Mali部分机型上使用两种不同的内存类型来存储数据，部分机型使用双Texture内存。感兴趣的读者可以在不同机型上测试不同的case。吐槽一下，Qualcomm的文档更新太慢，以上信息来源于5xx GPU的文档。

使用纹理内存的优化版本如下：

  1. // global_work_size[] = {(N + 3) / 4, (M + 3) / 4}
  2. #pragma OPENCL EXTENSION cl_khr_fp16 : enable
  3. __constant sampler_t default_sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
  4. __kernel void gemm_opt(__read_only image2d_t A, __read_only image2d_t B, __write_only image2d_t C, int M, int N, int K)
  5. {
  6.     int idx = get_global_id(0);
  7.     int idy = get_global_id(1);
  8.  
  9.     if((idx << 2) > N || (idy << 2) > M) return;
  10.  
  11.     half4 c[4] = {(half4)(0), (half4)(0), (half4)(0), (half4)(0)};
  12.  
  13.     for(int i = 0; i < K; i++)
  14.     {
  15.         half4 a = read_imageh(A, default_sampler, (int2)(idy, i));
  16.         half4 b = read_imageh(B, default_sampler, (int2)(idx, i));
  17.  
  18.         c[0] += a.x * b;
  19.         c[1] += a.y * b;
  20.         c[2] += a.z * b;
  21.         c[3] += a.w * b;
  22.     }
  23.  
  24.     idy = idy << 2;
  25.     write_imageh(C, (int2)(idx, idy), c[0]);
  26.     write_imageh(C, (int2)(idx, idy + 1), c[1]);
  27.     write_imageh(C, (int2)(idx, idy + 2), c[2]);
  28.     write_imageh(C, (int2)(idx, idy + 3), c[3]);
  29. }

该版本性能如下：

 

 

 该版本相对于基础版本有3倍的提升，可以看出Texture内存的使用可以极大的提升访存性能，进而发挥GPU的计算能力。

Mali Valhall 设备优化方案

纹理内存方案

 

 

 上图是Mali 各个架构下的GPU型号。Mali设备都是硬件厂商可配置的，同一GPU型号，可能存在多种配置。本文采用Valhall架构下的G77进行测试，SOC为MTK的天玑1000，设备为G77 MP9.

上文最后一个版本是针对Qualcomm架构给出的双Texture版本，那么在mali架构下是否是相同方案最优呢？Bifrost/ValHall架构相关文档中并未提及Texture内存与Buffer内存使用不同的Cache，因此这两个架构下，可以享受不同加载单元可以并行加载所带来的收益。同时， 从G76开始，Arm针对Texture内存的加载进行了加强，所以在Mali架构下，采用单Texture内存的方案进行优化。其他架构下，感兴趣的读者可以查看相应的文档或者相关测试。

实现方案如下：

  1. // global_work_size[] = {(N + 3) / 4, (M + 3) / 4}
  2. #pragma OPENCL EXTENSION cl_khr_fp16 : enable
  3. __constant sampler_t default_sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
  4. __kernel void gemm_opt(__read_only image2d_t A, __global half* B, __write_only image2d_t C, int M, int N, int K)
  5. {
  6.     int idx = get_global_id(0);
  7.     int idy = get_global_id(1);
  8.  
  9.     if((idx << 2) > N || (idy << 2) > M) return;
  10.  
  11.     half4 c[4] = {(half4)(0), (half4)(0), (half4)(0), (half4)(0)};
  12.  
  13.     int idx_ofs = idx << 2;
  14.     for(int i = 0; i < K; i++)
  15.     {
  16.         half4 a = read_imageh(A, default_sampler, (int2)(idy, i));
  17.         half4 b = vload4(0, B + idx_ofs + i * N);
  18.         c[0] += a.x * b;
  19.         c[1] += a.y * b;
  20.         c[2] += a.z * b;
  21.         c[3] += a.w * b;
  22.     }
  23.  
  24.     idy = idy << 2;
  25.     write_imageh(C, (int2)(idx, idy), c[0]);
  26.     write_imageh(C, (int2)(idx, idy + 1), c[1]);
  27.     write_imageh(C, (int2)(idx, idy + 2), c[2]);
  28.     write_imageh(C, (int2)(idx, idy + 3), c[3]);
  29. }

该版本性能如下：

 

 

 该版本相对于基础版本有10%左右的性能提升。因为Mali设备的Image内存相对于buffer内存吞吐优势并不明显，所以从buffer版本到Texture版本，Adreno的性能提升大于Mali设备的性能提升。

以上版本仅通过调整使用的内存类型提升数据吞吐以提升GEMM的性能。在此基础上，可以进一步通过更优的tile划分，更优的LocalWorkSize的配置来进一步提升GEMM性能。这些优化手段会给当前版本带来更大的性能提升，通过更深入的优化，在当前版本基础上，两款GPU都可以有至少50%的性能提升，之后的文章中会逐步介绍。

Mali(Valhall)的FMA方案

Mali GPU的valhall架构相对于之前的biforst架构做了大幅调整，ValHall架构开始其渲染和计算使用相同的统一的计算单元进行。下图是关于ValHall架构处理单元的介绍，可以看到，一个FMA单元单个周期可以处理16个FP32的FMA和32个FP16的FMA指令。

 

 

 下面为使用FMA指令的优化版本：

  1. // global_work_size[] = {(N + 3) / 4, (M + 3) / 4}
  2. #pragma OPENCL EXTENSION cl_khr_fp16 : enable
  3. __constant sampler_t default_sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
  4. __kernel void gemm_opt(__read_only image2d_t A, __global half* B, __write_only image2d_t C, int M, int N, int K)
  5. {
  6.     int idx = get_global_id(0);
  7.     int idy = get_global_id(1);
  8.  
  9.     if((idx << 2) > N || (idy << 2) > M) return;
  10.  
  11.     half4 c[4];
  12.     for(int i = 0; i < 4; i++)
  13.     {
  14.         c[i] = (half4)(0);
  15.     }
  16.  
  17.     int idx_ofs = idx << 2;
  18.     for(int i = 0; i < K; i += 2)
  19.     {
  20.         half4 a0 = read_imageh(A, default_sampler, (int2)(idy, (i + 0)));
  21.         half4 a1 = read_imageh(A, default_sampler, (int2)(idy, (i + 1)));
  22.  
  23.         half4 b0 = vload4(0, B + idx_ofs + (i + 0) * N);
  24.         half4 b1 = vload4(0, B + idx_ofs + (i + 1) * N);
  25.  
  26.         c[0] = fma(a0.x, b0, c[0]);
  27.         c[1] = fma(a0.y, b0, c[1]);
  28.         c[2] = fma(a0.z, b0, c[2]);
  29.         c[3] = fma(a0.w, b0, c[3]);
  30.  
  31.         c[0] = fma(a1.x, b1, c[0]);
  32.         c[1] = fma(a1.y, b1, c[1]);
  33.         c[2] = fma(a1.z, b1, c[2]);
  34.         c[3] = fma(a1.w, b1, c[3]);
  35.     }
  36.  
  37.     idy = idy << 2;
  38.     write_imageh(C, (int2)(idx, idy), c[0]);
  39.     write_imageh(C, (int2)(idx, idy + 1), c[1]);
  40.     write_imageh(C, (int2)(idx, idy + 2), c[2]);
  41.     write_imageh(C, (int2)(idx, idy + 3), c[3]);
  42. }

该版本对具体性能如下：

 

 

 使用FMA单元后，性能提升有36%左右。而在Adreno650上使用FMA则非常的慢，这是因为在5xx的文档中高通指出其FMA内置函数是通过软件模拟的，非常慢，目前看即使到650设备为止，该指令依然是软件模拟的。

下图是本文各版本之间的性能性能对比图，可以看出不同实现之间的巨大差异，后期通过更细的优化方法，将得到更大比例的性能提升。

 

 

 

 

 其他优化方案简介

前文一直使用的是单线程计算16个点，这是一种分块方案，但未必是最优的；所以在接下来的优化方案中，可以使用在各个维度上的分块策略，提升数据的复用度和cache命中率；合理的分块可以为矩阵乘法带来大幅度的性能提升。

除了分块策略之外，前文的LocalWorkSize一直是NULL，使用编译器的默认work group方案；在GPU优化中work group的划分，对资源划分以及调度都有很大影响。在adreno和mali的文档中也都有描述，默认的local work size未必是最优的。因此更好的local work group划分也将更好的提升性能。

除此之外，高通设备的local memory等其他资源也都有诸多探索空间，之后会逐步展开。下图是目前使用一些细节优化所达到的较优的优化版本性能。

 

 

 本文主要根据Adreno和Mali硬件上访存策略的差异，对初始版本做了简单优化。当前最优版本无论是Adreno还是Mali上距离峰值性能还有很大差异，所以在后续的介绍中会针对具体配置，在tile划分策略，LocalMemory的使用以及Local Work Size的配置等方面进行更细致的优化，进一步提升当前版本性能。

 

 

 

参考文献链接

https://mp.weixin.qq.com/s/oroRvkeafKqXxR-ad0WOLA

https://zhuanlan.zhihu.com/p/158161592

https://zhuanlan.zhihu.com/p/353975826

 

 

 

posted @ 2022-11-22 05:17  吴建明wujianming  阅读(509)  评论(0)  编辑  收藏  举报