引言

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

1 MaliT604 GPU的硬件结构和编程特性

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

1.1 MaliT604 GPU的硬件结构

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

图1 MaliT604基本硬件框图

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

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

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

图2 CortexA15 CPU和Mali GPU之间的工作调配

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

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

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

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

OpenCL是一个由编程语言规范、应用程序接口、库函数和运行时系统组成的跨平台异构并行计算框架，MaliT604 GPU在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 A15 CPU所共用的RAM在逻辑上被OpenCL框架切割成了4种不同的类型，MaliT600系列的GPU使用统一存储器模型，四种类型的存储器都映射到片外RAM上，CortexA15 CPU和MaliT604 GPU共享物理RAM，相对桌面GPU平台而言，在Mali平台上将数据从全局存储器拷贝到局部或者私有存储器并不能使访存性能得到提升，但也不用像桌面GPU一样进行从主存到显存的数据拷贝[2]。Mali GPU有三种访问RAM的方式，由传入clCreateBuffer函数中的不同参数决定，其示意图如图4所示。

图3 OpenCL针对MaliT604的抽象层次

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

CortexA15 CPU和MaliT604 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的向量处理特性

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

3 基于MaliT604 GPU的快速浮点矩阵乘法并行化实现

矩阵乘法运算在路径方案求解、线性方程组求解、图像处理等领域一直有着广泛应用，普通的迭代式串行算法的时间复杂度为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]位置。主机端分配线程的代码段如下：

size_tglobal_work_size[2]= {MATRIX_WIDTH, MATRIX_WIDTH};

status = clEnqueueNDRangeKernel(command_queue[0],matrix_mul_kernel, 2, 0, global_work_size,NULL, 0, 0, &event);

笔者将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，实验时使用armlinuxgnueabihf工具链对程序进行编译。不同规模的二维浮点矩阵乘法运算在ARM CortexA15 CPU上的串行方案和MaliT604 GPU上的并行方案的测试结果如表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异构运算特点吻合。

结语

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