Mali－T604GPU的二维浮点矩阵运算并行优化方法※

龚若皓，杨斌

（西南交通大学 信息科学与技术学院，成都610031）

龚若皓（硕士研究生）、杨斌（教授），主要研究方向为嵌入式系统及异构并行运算。

引 言

GPGPU 技术早年主要用于在超级计算机平台上进行高性能计算，而近年该技术逐渐被引入嵌入式领域。但在过去的移动GPU 平台上没有专门针对通用计算的软件框架和编程接口，软件设计者难以对数据的同步和计算的并行功能进行控制，所以移动GPU 在通用计算领域一直难以应用。本文基于Exynos5250SoC 平台详述了Mali GPU 的硬件特性以及将其应用于通用计算的编程方法，最后将二维浮点矩阵乘法并行化作为优化实例，验证Mali GPU 的并行能力，为计划使用嵌入式GPU 的GPGPU技术进行优化工作的研究人员和应用开发者提供技术参考。

1 MaliT－604GPU的硬件结构和编程特性

Mali是由ARM 研发设计的移动显示芯片组（GPUs）系列，能够在移动端提供强大的图像渲染能力，同时对通用计算进行了良好的软／硬件支持。

1.1 Mali－T604GPU的硬件结构

Mali－T604是Mali系列中首款使用统一渲染架构Midgard的移动GPU，Mali－T604GPU 包含4个着色器核心，采用AMBA 4ACE－LITE 总线接口，该总线以Cache Coherent Interconnect技术为特色，在多个处理器之间提供Cache一致性，通过ARM 的一致性和互连技术，计算任务在异构系统中进行共享处理时，可以轻松跨越CPU、GPU 和其他可用计算资源，更高效地访问数据。图1展示了Mali－T604 GPU 的 基 本 框 架。Cortex－A15CPU 核心以及Mali GPU 核心物理上共享了片外的RAM 存储器，并保持了L2Cache的一致性。Exynos5250处理器框图略——编者注。

Mali－T604GPU 在硬件层面优化了对任务管理和事件依赖的处理，并将这部分功能完全集成在其硬件的任务管理单元之中，可将计算任务从CPU 卸载到GPU，并在活动的着色器核心之间实现无缝负载平衡。

1.2 Mali GPU的并行化线程结构特征

Mali GPU 进行通用计算的技术核心是以多核多线程的思想将密集的计算任务进行拆解，将大量的计算线程分配于众多计算核心中，GPU 可以同时处理成百上千的线程，大量晶体管用于ALU［1］。GPU 适合做高密度数据的并行运算，只有在运算的并行粒度足够大的时候才能发挥出强大的并行运算能力［6］。图2 展示了CPU 和Mali GPU 之间工作调配的过程。

图1 Mali－T604基本硬件框图

图2 Cortex－A15CPU和Mali GPU之间的工作调配

Mali GPU 中每个计算线程会占用着色器核心的一部分资源（存储器和ALU 等），每个线程占用资源的多少影响了同时并行处理的活动线程的数量。对Mali GPU，每一个线程都有自己的程序计数器，这意味着Mali GPU 和桌面GPU 平台不同，程序分支的发散不是影响效率的重要问题［2］。每个Mali－T604GPU 的着色器核心最多可以同时容纳256个线程［2］，Mali GPU 在进行通用计算时需要大量的线程进行切换才能保证得到计算效率上的收益，对于Mali－T604 而言，这个最少的总工作项数量是4 096［2］。如果分配于单个着色器核心上的线程数目不足128，很可能带来并行效率的下降［2］，这时需要将工作拆分为不同的步骤，简化每个步骤的线程复杂度，让单个着色器核心并行容纳的线程数量足够多，以保证并行度。

2 Mali GPU的并行化计算模型构建

Mali－T600系列的GPU 对OpenCL 1.1Full Profile标准进行了良好的支持［1］，OpenCL是真正意义上的跨平台异构并行框架，能够真正挖掘出Mali GPU 的并行计算特性。

2.1 Mali GPU在OpenCL框架下的并行任务抽象及线程规划

OpenCL是一个由编程语言规范、应用程序接口、库函数和运行时系统组成的跨平台异构并行计算框架，Mali－T604GPU 在OpenCL下的抽象层次如图3所示。

OpenCL的并行基于SMT（同时多线程）的思想，由用户指定自定义数目的线程，并根据线程的标识符设计计算线程与数据关联的映射法则，SMT 架构主要用于隐蔽访存 的 延 时［3］。OpenCL 框 架 下，CPU 主 机 端 程 序 由OpenCL的API编写，实现计算平台的初始化、存储器的分配和交互的控制，并决定分配的计算线程的维度和每一维的数量。设备端的内核程序由OpenCL C 语 言 编 写，Mali GPU 会根据内核对象创建主机端请求数量的线程实例，每个线程的运算工作都由图2中一个对应的PE（Procoss Element，处理单元）进行处理，线程的工作逻辑决定了线程标识号和数据的关联关系。多个线程被组织为工作组的形式，每一个工作组固定分配到一个CU（Compute Unit，计算单元）上进 行 处 理［4－5］，同 一 个 工 作 组中的线程会在对应的CU 上由Mali GPU 的任务管理单元进行快速的切换和调度，保证一个CU 上的PE最大限度保持忙碌。

2.2 Mali GPU多核环境下的存储器空间映像方法

如图3所示，Mali GPU 和Cortex A15CPU 所共用的RAM 在逻辑上被OpenCL 框架切割成了4 种不同的类型，Mali－T600系列的GPU 使用统一存储器模型，四种类型的存储器都映射到片外RAM 上，Cortex－A15CPU 和Mali－T604GPU 共享物理RAM，相对桌面GPU 平台而言，在Mali平台上将数据从全局存储器拷贝到局部或者私有存储器并不能使访存性能得到提升，但也不用像桌面GPU 一样进行从主存到显存的数据拷贝［2］。Mali GPU有三种访问RAM 的方式，由传入clCreateBuffer函数中的不同参数决定，其示意图如图4所示。

图3 OpenCL针对Mali－T604的抽象层次

图4 OpenCL框架下Mali GPU对存储器的不同访问方式

Cortex－A15CPU 和Mali－T604 GPU 使 用 不 同 的 虚拟地址空间，在主机端由malloc函数分配的缓存，Mali GPU 无法访问。Mali GPU 可以访问clCreateBuffer函数分配出的缓存，CPU 借助OpenCL 中的map映射操作也可实现对这类缓存的读写。图4（b）需要主机端的缓存进行数据拷贝来初始化，图4（b）和（c）类似，但只在OpenCL的内核函数首次使用该缓存时才进行数据拷贝，在CPU端进行map操作时GPU 还会将数据拷贝回主机端的缓存，对于Mali GPU 而言，多余的数据拷贝操作会降低访存效率。图4（a）是ARM 官方建议的访存方式，CPU 和GPU 共享一块物理缓存，高速实现数据交互。

2.3 Mali GPU的向量处理特性

Mali－T604GPU 内部有128位宽度的向量寄存器［2］，使用OpenCL C 中的内建向量类型可以让数据自动以SIMD的形式在Mali GPU 的ALU 中进行并行计算，Mali GPU 中将数据以16个字节对齐可以使得数据的长度和高速 缓 存 适 配，加 快 数 据 存 取 速 度［2］，Mali－T600 系 列GPU 中加载一个128位的向量和加载一个单字节数据花费的时间是一样的。将数据以128位进行对齐，能够最大限度发挥Mali－T604GPU 的访存和运算效率。

3 基于Mali－T604GPU 的快速浮点矩阵乘法并行化实现

矩阵乘法运算在路径方案求解、线性方程组求解、图像处理等领域一直有着广泛应用，普通的迭代式串行算法的时间复杂度为O（n3），对于大型的矩阵乘法，特别是浮点类型的矩阵乘法，计算量非常惊人，传统的算法基于CPU 进行设计，CPU并不能提供大型的并行度和强大的浮点计算能力，对于大型浮点类型矩阵乘法的处理力不从心。

AB两个矩阵的乘法的结果矩阵中，每个数据均依赖于A 中的一行和B中的一列的点积结果，每个计算结果没有依赖和相关，显然是高度可数据并行的计算，适合使用GPU 做并行处理，使用GPU 上的多个线程可以并行进行矩阵A 和B中不同行和列的点积。

进行实验时，以N×N 的两个浮点矩阵A 和B 进行乘法，得出N×N 的浮点结果矩阵matrixResult，利用Mali GPU 进行并行化时，总共分配N×N 个线程，以二维方式进行排布，标识号为（i，j）的线程提取出矩阵matrixA 的第i行和矩阵matrixB 的第j列，利用OpenCL 中长度为128位的float4向量类型快速实现两个一维向量的点积，再将该点积结果存储到matrixResult［i］［j］位置。主机端分配线程的代码段如下：

笔者将clEnqueueNDRangeKernel函数中工作组大小参数设置为NULL，由Mali GPU 硬件自动确定最佳的工作组大小。由于内核中每次会连续读取4个浮点数值、凑成float4类型的数据，所以对于矩阵的宽度不是4的倍数的情况需要进行特殊处理。首先可在主机端将输入矩阵A 修改为N 行（N／4＋4）列，将矩阵B修改为（N／4＋4）行N 列，多出的矩阵部分均以0补齐，这样既不影响计算结果，又不会影响线程的分配方案，实现并行方案的内核函数略——编者注。

采用Arndale Board开发板作为测试平台，软件平台采用Linaro机构为Arndale Board定制的基于Ubuntu的嵌入式Linux操作系统，其内核版本为3.10.37，实验时使用arm－linux－gnueabihf工具链对程序进行编译。不同规模的二维浮点矩阵乘法运算在ARM Cortex－A15CPU 上的串行方案和Mali－T604GPU 上的并行方案的测试结果如表1所列。为不失一般性，测试时输入矩阵内容为随机值，每种不同矩阵大小的测试项进行10次，将测试值的平均值作为测试结果。

表1 二维浮点矩阵乘法优化效果对比

表1仅列出了输入量较大时的测试结果，笔者实际测试时，发现输入数据量较小时，并行方案没有串行方案的效率高，这是因为计算过程大部分都消耗在数据的传输上，由于计算量小，GPU 端的计算瞬间完成，没有办法将Mali GPU 访存的延迟掩盖，所以此时访存速度较快的CPU 端的串行方案反而效率更高。

当计算量逐步增加的时候，Mali GPU 的并行能力逐渐体现出其优势，加速比有显著提升。当计算量大到一定程度的时候，加速比趋于稳定，此时Mali GPU 上有大量的线程切换，不仅隐蔽了访存的延迟，也使得Mali GPU上的计算单元满载，其计算效率已达到硬件能够承受的极限，Mali GPU 可以提供接近40倍的惊人的加速比。

实际测试时，笔者使用top指令观察矩阵进程的CPU占用量，串行方案的CPU 占用量在98%左右，而基于Mali GPU 的并行方案对CPU 几乎没有占用量，说明并行方案不仅可以提升计算效率，而且降低了CPU 的负担，大大提升了系统实时性。实验结果与GPU 异构运算特点吻合。

结 语

本 文 针 对Mali－T604 GPU 论 述 了 基 于OpenCL 的Linux平台上进行通用计算并行优化的方法，论述了Mali－T604GPU 的硬件特点，并基于OpenCL 设计了二维矩阵乘法的并行方案，在Mali－T604 上获得了惊人的加速比。实验结果表明，Mali GPU 对于庞大输入量的计算密集型高度可数据并行化通用计算问题具有显著的加速能力，且并行优化结果正确可靠。

编者注：本文为期刊缩略版，全文见本刊网站www.mesnet.com.cn。

［1］Owens J D，Houston M，Luebke D，et al.GPU computing［J］.Proceedings of the IEEE，2008，96（5）：879－899.

［2］ARM.ARM Mali－T600 Series GPU OpenCL Developer Guide Version 2.0，2013.

［3］ARM.Roberto Mijat.Take GPU Processing Power Beyond Graphics with Mali GPU Computing，2012.

［4］Benedict R Gaster，Lee Howes，David R Kaeli，et al.OpenCL异构计算［M］.2版.北京：清华大学出版社，2013.

［5］Matthew Scarpino.OpenCL实战［M］.北京：人民邮电出版社，2012.

［6］梁霞.基于GPU 的H.264并行解码器设计［D］.大连：大连理工大学，2010.