Mali－T604GPU的二維浮點矩陣運算并行優(yōu)化方法※

龔若皓，楊斌

（西南交通大學 信息科學與技術學院，成都610031）

龔若皓（碩士研究生）、楊斌（教授），主要研究方向為嵌入式系統(tǒng)及異構并行運算。

引 言

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

1 MaliT－604GPU的硬件結構和編程特性

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

1.1 Mali－T604GPU的硬件結構

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

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

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

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

圖1 Mali－T604基本硬件框圖

圖2 Cortex－A15CPU和Mali GPU之間的工作調(diào)配

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

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

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

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

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

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

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

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

圖3 OpenCL針對Mali－T604的抽象層次

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

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

2.3 Mali GPU的向量處理特性

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

3 基于Mali－T604GPU 的快速浮點矩陣乘法并行化實現(xiàn)

矩陣乘法運算在路徑方案求解、線性方程組求解、圖像處理等領域一直有著廣泛應用，普通的迭代式串行算法的時間復雜度為O（n3），對于大型的矩陣乘法，特別是浮點類型的矩陣乘法，計算量非常驚人，傳統(tǒng)的算法基于CPU 進行設計，CPU并不能提供大型的并行度和強大的浮點計算能力，對于大型浮點類型矩陣乘法的處理力不從心。

AB兩個矩陣的乘法的結果矩陣中，每個數(shù)據(jù)均依賴于A 中的一行和B中的一列的點積結果，每個計算結果沒有依賴和相關，顯然是高度可數(shù)據(jù)并行的計算，適合使用GPU 做并行處理，使用GPU 上的多個線程可以并行進行矩陣A 和B中不同行和列的點積。

進行實驗時，以N×N 的兩個浮點矩陣A 和B 進行乘法，得出N×N 的浮點結果矩陣matrixResult，利用Mali GPU 進行并行化時，總共分配N×N 個線程，以二維方式進行排布，標識號為（i，j）的線程提取出矩陣matrixA 的第i行和矩陣matrixB 的第j列，利用OpenCL 中長度為128位的float4向量類型快速實現(xiàn)兩個一維向量的點積，再將該點積結果存儲到matrixResult［i］［j］位置。主機端分配線程的代碼段如下：

筆者將clEnqueueNDRangeKernel函數(shù)中工作組大小參數(shù)設置為NULL，由Mali GPU 硬件自動確定最佳的工作組大小。由于內(nèi)核中每次會連續(xù)讀取4個浮點數(shù)值、湊成float4類型的數(shù)據(jù)，所以對于矩陣的寬度不是4的倍數(shù)的情況需要進行特殊處理。首先可在主機端將輸入矩陣A 修改為N 行（N／4＋4）列，將矩陣B修改為（N／4＋4）行N 列，多出的矩陣部分均以0補齊，這樣既不影響計算結果，又不會影響線程的分配方案，實現(xiàn)并行方案的內(nèi)核函數(shù)略——編者注。

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

表1 二維浮點矩陣乘法優(yōu)化效果對比

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

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

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

結 語

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

編者注：本文為期刊縮略版，全文見本刊網(wǎng)站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實戰(zhàn)［M］.北京：人民郵電出版社，2012.

［6］梁霞.基于GPU 的H.264并行解碼器設計［D］.大連：大連理工大學，2010.