Mali GPU編程特性及二維浮點矩陣運算并行優(yōu)化詳解

　　基于Mali-T604嵌入式GPU的二維浮點矩陣運算并行優(yōu)化

　　ARM Cortex-A15系列處理器是當前最新的嵌入式ARM SoC，該系列處理器首次集成了Mali-T600系列的移動端GPU，該系列GPU支持OpenGL以及OpenCL等計算框架，可以有效加速通用計算，而目前對其應用方法和實際優(yōu)化效果的研究很少。本文基于以三星的Exynos5250處理器為核心的Arndale Board嵌入式開發(fā)平臺，對集成于處理器上的Mali-T604嵌入式GPU的GPGPU（General-Purpose computation on GPU）技術進行研究并對不同運算規(guī)模的浮點矩陣乘法進行并行加速優(yōu)化，提供實際測試結果。

　　GPGPU技術早年主要在超級計算機平臺進行高性能計算，而近年該技術逐漸被引入嵌入式領域。但在過去的移動GPU平臺上沒有專門針對通用計算的軟件框架和編程接口，軟件設計者難以對于數據的同步和計算的并行進行控制，所以移動GPU在通用計算領域一直難以應用。本文基于Exynos5250 SoC平臺詳述Mali GPU的硬件特性和將其應用于通用計算的編程的方法，最后將二維浮點矩陣乘法并行化作為優(yōu)化實例，驗證Mali GPU的并行能力，為計劃使用嵌入式GPU的GPGPU技術進行優(yōu)化工作的研究人員和應用開發(fā)者提供技術參考和借鑒。

　　1.Mali T604 GPU的硬件結構和編程特性

　　Mali是由ARM研發(fā)設計的移動顯示芯片組（GPUs）系列，不僅能夠在移動端提供強大的圖像渲染能力，同時在近期對通用計算進行了良好的軟硬件支支持。

　　1.1 Mali T604 GPU的組成結構

　　Mali-T604是Mali系列中首款使用統(tǒng)一渲染架構Midgard的移動GPU，Mali-T604 GPU包含4個著色器核心，采用AMBA 4 ACE-LITE總線接口，該總線以Cache Coherent Interconnect技術為特色，在多個處理器之間提供完全Cache一致性，通過ARM的一致性和互連技術，計算任務在異構系統(tǒng)中進行共享處理時，可以輕松跨越CPU、GPU和其他可用計算資源，更高效地訪問數據。圖1展示了Mali-T604 GPU的基本框架。如圖2所示，Cortex-A15 CPU核心以及Mali GPU核心物理上共享了片外的RAM存儲器并保持了L2Cache的一致性。

　　圖1 Mali-T604基本硬件框圖

　　圖2 Exynos5250處理器框圖

　　Mali-T604 GPU在硬件層面優(yōu)化了對任務管理和事件依賴的處理，并將這部分功能完全集成在其硬件的任務管理單元之中，可將計算任務從CPU卸載到GPU，并在活動的著色器核心之間實現無縫負載平衡。

　　1.2 Mali GPU的并行化線程結構特征

　　Mali GPU進行通用計算的技術核心是以多核多線程的思想將密集的計算任務進行拆解，將大量的計算線程分配于眾多計算核心中，GPU可以同時處理成百上千的線程，大量晶體管用于ALU.GPU適合做高密度數據的并行運算，只有在運算的并行粒度足夠大的時候才能發(fā)揮出強大的并行運算能力。圖3展示了CPU和 Mali GPU之間工作調配的過程。

　　圖3 Cortex-A15 CPU和Mali GPU之間的工作調配

　　Mali GPU中每個計算線程會占用著色器核心的一部分資源（存儲器和ALU等），每個線程占用資源的多少影響了同時并行處理的活動線程的數量。對Mali GPU，每一個線程都有自己的程序計數器，這意味著Mali GPU和桌面GPU平臺不同，程序分支的發(fā)散不是一個影響效率的重要的問題。每個Mali-T604 GPU的著色器核心最多可以同時容納256個線程，Mali GPU在進行通用計算時需要大量的線程進行切換才能保證得到計算效率上的收益，對于Mali-T604而言，這個最少的總工作項數量是4096.如果分配于單個著色器核心上的線程數目不足128，很可能帶來并行效率的下降，這時需要拆分工作為不同的步驟，簡化每個步驟的線程復雜度，讓單個著色器核心并行容納的線程數量足夠多以保證并行度。

　　2.Mali GPU的并行化計算模型構建

　　Mali-T600系列的GPU對OpenCL 1.1 Full Profile標準進行了良好的支持，OpenCL是真正意義上的跨平臺異構并行框架，能夠真正挖掘出Mali GPU的并行計算特性。

　　2.1 Mali GPU在OpenCL框架下的并行任務抽象及線程規(guī)劃

　　OpenCL是一個由編程語言規(guī)范，應用程序接口、庫函數和運行時系統(tǒng)組成的跨平臺異構并行計算框架，Mali-T604 GPU在OpenCL下的抽象層次如下面的圖4所示：

　　圖4 OpenCL針對Mali-T604的抽象層次

　　OpenCL的并行基于SMT（同時多線程）的思想，由用戶指定自定義數目的線程，并根據線程的標識符設計計算線程與數據關聯(lián)的映射法則，SMT架構主要用于隱蔽訪存的延時。OpenCL框架下，CPU主機端程序由OpenCL的API編寫，實現計算平臺的初始化，存儲器的分配和交互的控制，并決定分配的計算線程的維度和每一維的數量。設備端的內核程序由OpenCL C語言編寫，Mali GPU會根據內核對象創(chuàng)建主機端請求數量的線程實例，每個線程的運算工作都由圖4中一個對應的PE進行處理，線程的工作邏輯決定了線程標識號和數據的關聯(lián)關系。多個線程被組織為工作組的形式，每一個工作組固定分配到一個CU上進行處理，同一個工作組中的線程會在對應的CU上由Mali GPU的任務管理單元進行快速的切換和調度，保證一個CU上的PE最大限度保持忙碌。

　　2.2 Mali GPU多核環(huán)境下的存儲器空間映像方法

　　如圖4所示，Mali GPU和Cortex A15 CPU所共用的RAM在邏輯上被OpenCL框架切割成了四種不同的類型，Mali-T600系列的GPU使用統(tǒng)一存儲器模型，四種類型的存儲器都映射到片外RAM上，Cortex-A15 CPU和Mali-T604 GPU共享物理RAM，相對桌面GPU平臺而言，在Mali平臺上將數據從全局存儲器拷貝到局部或者私有存儲器并不能使訪存性能得到提升，但相對地也不用像桌面GPU一樣進行從主存到顯存的數據拷貝。Mali GPU有三種訪問RAM的方式，由傳入clCreateBuffer函數中的不同參數決定，其示意圖如下：

　　圖5 OpenCL框架下Mali GPU對存儲器的不同訪問方式

　　Cortex-A15 CPU和Mali-T604 GPU使用不同的虛擬地址空間，在主機端由malloc函數分配的緩存，Mali GPU無法訪問。Mali GPU可以訪問clCreateBuffer函數分配出的緩存，CPU借助OpenCL中的map映射操作也可實現對這類緩存的讀寫，圖5中的方式2需要主機端的緩存進行數據拷貝來初始化，方式3和方式2類似，但只在OpenCL的內核函數首次使用該緩存時才進行數據拷貝，在CPU端進行map操作時 GPU還會將數據拷貝回主機端的緩存，對于Mali GPU而言，多余的數據拷貝操作會降低訪存效率。圖5中的方式1是ARM官方建議的訪存方式，CPU和GPU共享一塊物理緩存，高速實現數據交互。

　　2.3 Mali GPU的向量處理特性

　　Mali-T604 GPU內部有128位寬度的向量寄存器，使用OpenCL C中的內建向量類型可以讓數據自動以SIMD的形式在Mali GPU的ALU中進行并行計算，Mali GPU中將數據以16個字節(jié)對齊可以使得數據的長度和高速緩存適配，加快數據存取速度，Mali-T600系列GPU中加載一個128位的向量和加載一個單字節(jié)數據花費的時間是一樣的。將數據以128位進行對齊，能夠最大限度發(fā)揮Mali-T604 GPU的訪存和運算效率。

　　3.基于Mali-T604 GPU的快速浮點矩陣乘法并行化實現

　　矩陣乘法運算在路徑方案求解、線性方程組求解、圖像處理等領域一直有著廣泛應用，普通的迭代式串行算法的時間復雜度為O（n3），對于大型的矩陣乘法，特別是浮點類型的矩陣乘法，計算量非常驚人，傳統(tǒng)的算法基于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硬件自動確定最佳的工作組大小。由于內核中每次會連續(xù)讀取4個浮點數值湊成float4類型的數據，所以對于矩陣的寬度不是4的倍數的情況需要進行特殊處理，可在主機端首先將輸入矩陣A修改為N行N/4+4列，將矩陣B修改為N/4+4行N列，多出的矩陣部分均以0補齊，這樣既不影響計算結果，也不會影響線程的分配方案，實現并行方案的內核函數如下所示：

　　本文采用Arndale Board開發(fā)板作為測試平臺，軟件平臺采用Linaro機構為Arndale Board定制的基于Ubuntu的嵌入式Linux操作系統(tǒng)，其內核版本為3.10.37，實驗時使用arm-linux-gnueabihf工具鏈對程序進行編譯。不同規(guī)模的二維浮點矩陣乘法運算在ARM Cortex-A15 CPU上的串行方案和Mali-T604 GPU上的并行方案的測試結果如面的表1所示，為不失一般性，測試時輸入矩陣內容為隨機值，每種不同矩陣大小的測試項進行10次測試，將測試值的平均值作為測試結果。

　　上表僅列出了輸入量較大時的測試結果，筆者實際測試時，發(fā)現輸入數據量較小的時候，并行方案沒有串行方案的效率高，因為計算過程大部分都消耗在數據的傳輸上，由于計算量小，GPU端的計算瞬間完成，沒有辦法將Mali GPU訪存的延遲掩蓋，所以此時訪存速度較快的CPU端的串行方案反而效率更高。

　　當計算量逐步增加的時候，Mali GPU的并行能力逐漸體現出其優(yōu)勢，加速比有顯著提升，當計算量大到一定程度的時候，加速比趨于穩(wěn)定，因為這時Mali GPU上有大量的線程切換，不僅隱蔽了訪存的延遲，也使得Mali GPU上的計算單元滿載，其計算效率已達到硬件能夠承受的極限，此時Mali GPU可以提接近40倍的供驚人的加速比。

　　實際測試時，筆者使用top指令觀察矩陣進程的CPU占用量，串行方案的CPU占用量在98%左右，而基于Mali GPU的并行方案對CPU幾乎沒有占用量，說明并行方案不僅可以提升計算效率，還降低了CPU的負擔，大大提升了系統(tǒng)實時性。實驗的實際測試結果和GPU 異構運算特點吻合。

　　4.結語

　　本文針對Mali-T604 GPU論述了基于OpenCL的Linux平臺上進行通用計算并行優(yōu)化的方法，論述了Mali-T604 GPU的硬件特點，并基于OpenCL設計了二維矩陣乘法的并行方案，在Mali-T604上獲得了驚人的加速比，結果表明Mali GPU對于龐大輸入量的計算密集型高度可數據并行化通用計算問題有顯著的加速能力，且并行優(yōu)化結果正確可靠。