矩阵乘法运算在路径方案求解、线性方程组求解、图像处理等领域一直有着广泛应用，普通的迭代式串行算法的时间复杂度为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-A15 CPU上的串行方案和Mali-T604 GPU上的并行方案的测试结果如面的表1所示，为不失一般性，测试时输入矩阵内容为随机值，每种不同矩阵大小的测试项进行10次测试，将测试值的平均值作为测试结果。

上表仅列出了输入量较大时的测试结果，笔者实际测试时，发现输入数据量较小的时候，并行方案没有串行方案的效率高，因为计算过程大部分都消耗在数据的传输上，由于计算量小，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-T604 GPU的硬件特点，并基于OpenCL设计了二维矩阵乘法的并行方案，在Mali-T604上获得了惊人的加速比，结果表明Mali GPU对于庞大输入量的计算密集型高度可数据并行化通用计算问题有显著的加速能力，且并行优化结果正确可靠。