實(shí)戰(zhàn)技巧，Mali GPU編程特性及二維浮點(diǎn)矩陣運(yùn)算并行優(yōu)化詳解

　　筆者將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-A15 CPU上的串行方案和Mali-T604 GPU上的并行方案的測(cè)試結(jié)果如面的表1所示，為不失一般性，測(cè)試時(shí)輸入矩陣內(nèi)容為隨機(jī)值，每種不同矩陣大小的測(cè)試項(xiàng)進(jìn)行10次測(cè)試，將測(cè)試值的平均值作為測(cè)試結(jié)果。

　　上表僅列出了輸入量較大時(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)定，因?yàn)檫@時(shí)Mali GPU上有大量的線程切換，不僅隱蔽了訪存的延遲，也使得Mali GPU上的計(jì)算單元滿載，其計(jì)算效率已達(dá)到硬件能夠承受的極限，此時(shí)Mali GPU可以提接近40倍的供驚人的加速比。

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

　　4.結(jié)語

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