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

龔若皓，楊斌

（西南交通大學(xué) 信息科學(xué)與技術(shù)學(xué)院，成都610031）

龔若皓（碩士研究生）、楊斌（教授），主要研究方向?yàn)榍度胧较到y(tǒng)及異構(gòu)并行運(yùn)算。

引 言

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

1 MaliT－604GPU的硬件結(jié)構(gòu)和編程特性

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

1.1 Mali－T604GPU的硬件結(jié)構(gòu)

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

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

1.2 Mali GPU的并行化線程結(jié)構(gòu)特征

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

圖1 Mali－T604基本硬件框圖

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

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

2 Mali GPU的并行化計(jì)算模型構(gòu)建

Mali－T600系列的GPU 對(duì)OpenCL 1.1Full Profile標(biāo)準(zhǔn)進(jìn)行了良好的支持［1］，OpenCL是真正意義上的跨平臺(tái)異構(gòu)并行框架，能夠真正挖掘出Mali GPU 的并行計(jì)算特性。

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

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

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

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

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

圖3 OpenCL針對(duì)Mali－T604的抽象層次

圖4 OpenCL框架下Mali GPU對(duì)存儲(chǔ)器的不同訪問方式

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

2.3 Mali GPU的向量處理特性

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

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

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

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

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

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

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

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

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

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

實(shí)際測(cè)試時(shí)，筆者使用top指令觀察矩陣進(jìn)程的CPU占用量，串行方案的CPU 占用量在98%左右，而基于Mali GPU 的并行方案對(duì)CPU 幾乎沒有占用量，說(shuō)明并行方案不僅可以提升計(jì)算效率，而且降低了CPU 的負(fù)擔(dān)，大大提升了系統(tǒng)實(shí)時(shí)性。實(shí)驗(yàn)結(jié)果與GPU 異構(gòu)運(yùn)算特點(diǎn)吻合。

結(jié) 語(yǔ)

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

編者注：本文為期刊縮略版，全文見本刊網(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異構(gòu)計(jì)算［M］.2版.北京：清華大學(xué)出版社，2013.

［5］Matthew Scarpino.OpenCL實(shí)戰(zhàn)［M］.北京：人民郵電出版社，2012.

［6］梁霞.基于GPU 的H.264并行解碼器設(shè)計(jì)［D］.大連：大連理工大學(xué)，2010.