一種面向OpenCL架構(gòu)的矩陣-向量乘并行算法與實現(xiàn)

肖 漢,周清雷,姚鵬姿

1(鄭州師范學(xué)院 信息科學(xué)與技術(shù)學(xué)院,鄭州 450044)2(鄭州大學(xué) 信息工程學(xué)院,鄭州 450001)

1 引 言

科學(xué)技術(shù)的發(fā)展日新月異,當(dāng)今許多應(yīng)用領(lǐng)域如電路仿真模擬、數(shù)學(xué)建模和網(wǎng)絡(luò)分析等都需用到矩陣-向量乘法算法[1].隨著現(xiàn)代工程問題復(fù)雜度的提升,計算過程運算尺度大、運算數(shù)據(jù)量劇增,造成矩陣-向量乘法運算過程的效率低,導(dǎo)致無法滿足應(yīng)用系統(tǒng)實時處理的要求,限制了其發(fā)展.因此,人們對矩陣-向量乘法的運算速度的要求越來越高[2,3].

運用昂貴的服務(wù)器、工作站等設(shè)備,多核計算平臺和計算機集群系統(tǒng)能夠在一定程度上提供了較強的并行處理能力,但是在數(shù)據(jù)量較大時容易出現(xiàn)傳輸瓶頸,制約了性能的進一步提高.同時,將算法部署在工作站和計算機集群系統(tǒng)上,將造成硬件成本上升,軟件研發(fā)過程復(fù)雜;而采用多核處理器進行性能提升,則存在器件集成度較低、功耗較高、價格昂貴等不足[4,5].而通過圖形處理器(Graphic Processing Unit,GPU)的并行計算,將使計算機的運算效率得到很大的提升,GPU以其高速的浮點運算能力用于通用計算越來越引起人們的重視,GPU加速的矩陣-向量乘并行算法的優(yōu)化與實現(xiàn)將是一個研究的熱點[6].

開放式計算語言(Open Computing Language,OpenCL)不僅僅是一種編程語言,更是一個能幫助開發(fā)者充分而合理的利用整個計算機系統(tǒng)的所有計算資源,并挖掘和發(fā)揮出其中計算能力的跨平臺異構(gòu)框架,即完整的在異構(gòu)平臺上并行編程的開放框架標(biāo)準(zhǔn),它包括編程語言、API、函數(shù)庫以及運行時系統(tǒng)來支持軟件在整個平臺上的開發(fā).

2 相關(guān)研究工作

近年來,關(guān)于矩陣-向量乘法算法的并行化處理和優(yōu)化已有許多工作,代表性的工作主要有:Zhang Jilin等[7]提出了采用緩存遺忘的擴展四叉樹存儲結(jié)構(gòu)的稀疏矩陣-向量乘法,與CSR格式相比,達到了1.5倍加速.鄔貴明等[8]提出了稀疏矩陣二維均勻分塊方法和相應(yīng)的表示方法嵌套分塊CSR,將有效開發(fā)定制結(jié)構(gòu)的并行性.Gao Zhen等[9]提出了一種基于FPGA的整數(shù)并行矩陣-向量乘法的容錯設(shè)計方法,顯著降低了算法開銷和容錯成本.蘇錦柱等[10]提出了一種基于FPGA的二元域大型稀疏矩陣向量乘SpMV的環(huán)網(wǎng)硬件系統(tǒng)架構(gòu),相對于PR模型獲得了2.65倍的加速比.薛永江等[11]提出了一種基于FPGA的浮點數(shù)格式,二叉樹數(shù)據(jù)流的矩陣向量乘法器,與軟件執(zhí)行相比,雙精度運算速度能達到其3倍.Malovichko M等[12]使用OpenMP加速矩陣向量乘法,提出了一種用積分方程方法進行頻域聲學(xué)建模的并行算法.Hassan Somaia A等[13]通過使用AVX指令集、內(nèi)存訪問優(yōu)化和OpenMP并行化,提出了一種單精度矩陣-向量乘法的優(yōu)化算法,性能提升了18.2%.吳長茂等[14]通過在CPU和MIC以及網(wǎng)絡(luò)間實現(xiàn)了異步計算和異步數(shù)據(jù)傳輸,設(shè)計了SpMV并行算法,使得行星流體動力學(xué)數(shù)值模擬在CPU-MIC異構(gòu)眾核環(huán)境下獲得了6.93倍的加速比.張愛民等[15]針對Intel Xeon Phi的體系結(jié)構(gòu)特點,提出了一種通用的分塊壓縮存儲表示的SpMV并行算法,性能比MKL的CSR算法平均快2.05倍.李政等[16]根據(jù)油藏模擬線性解法器的算法要求,設(shè)計了基于一種塊混合結(jié)構(gòu)稀疏存儲格式及其對應(yīng)的GPU線程組并行策略的SpMV并行算法,相對串行BCSR格式的SpMV的加速比可達19倍.Liang Yun等[17]提出了基于OpenCL硬件遺忘實現(xiàn)和基于至強Phi本地編程語言的硬件有意識實現(xiàn)的稀疏矩陣向量乘法,并分別獲得了2.2倍和3.1倍加速比.

這些研究表明,并行計算的應(yīng)用能夠極大地減少處理時間,系統(tǒng)性能得到了有效提高.但是這種高計算效率多獲益于一些比較傳統(tǒng)的并行處理技術(shù)的運用,卻沒有與目前主流的并行計算技術(shù)所產(chǎn)生的效率優(yōu)劣勢的對比,以至于并行化方案并沒有充分發(fā)揮出系統(tǒng)的效率.此外,部分研究雖然采用了當(dāng)今比較流行的硬件架構(gòu)和計算模型,但是多針對單一的并行計算技術(shù)進行算法的改進,多算法跨平臺系統(tǒng)的性能評估較少.本文將根據(jù)矩陣-向量乘算法的特點和OpenCL架構(gòu)特征,實現(xiàn)了基于GPU加速的矩陣-向量乘并行算法和算法在不同計算平臺上的性能移植.

3 OpenCL并行計算引擎

OpenCL并行計算引擎的執(zhí)行模型分為兩個部分,一部分是通過主機端CPU系統(tǒng)執(zhí)行,另一部分是通過設(shè)備端在若干個OpenCL設(shè)備上執(zhí)行的內(nèi)核.OpenCL通過使用主機端系統(tǒng)對多個支持OpenCL的計算設(shè)備進行定義上下文并管理內(nèi)核的調(diào)度,通過主機端內(nèi)存和設(shè)備端存儲器,共同完成整個系統(tǒng)的運行[18].

出于對系統(tǒng)的可移植性的考慮,OpenCL定義了一個抽象的存儲器模型,它能夠?qū)⒃摯鎯ζ髂Ｐ陀成涞綄嶋H的設(shè)備存儲器系統(tǒng)中.OpenCL將內(nèi)核程序中用到的存儲器分為了四類不同的類型:本地存儲器(Local Memory)、常量存儲器(Constant Memory)、全局存儲器(Global Memory)和私有存儲器(Private Memory),如圖1所示[19,20].

圖1 OpenCL存儲器模型Fig.1 OpenCL memory model

① 私有存儲器:只從屬于當(dāng)前的工作節(jié)點,私有存儲器中的數(shù)據(jù)通常映射到寄存器.一個工作項內(nèi)部的私有緩沖區(qū)對于其他工作項是完全不可見的.

② 本地存儲器:從屬于一個工作組的存儲器,同一個工作組中所有的工作項都可以共享使用本地存儲器.可用來分配一些由該工作組中所有工作項共享的數(shù)據(jù),它通常位于片上,與全局存儲器相比較,帶寬更高,訪問延遲更短,但容量非常有限.

③ 常量存儲器:帶有緩存的全局存儲器,是全局存儲器的一部分,工作空間內(nèi)所有的工作項可以只讀常量存儲器中的任意元素,而不是特意針對某種只讀數(shù)據(jù)類型.

④ 全局存儲器:工作空間內(nèi)所有的工作項都可以讀/寫全局存儲器中的任意位置元素.

4 在OpenCL中的并行化實現(xiàn)和優(yōu)化

4.1 并行算法描述

矩陣-向量乘法將一個n×n階的方陣A=[aij]乘以n×1的列向量B=[b1,b2,…,bn]T可得一個n×1的列向量C=[c1,c2,…,cn]T:

根據(jù)OpenCL架構(gòu)的特點,在行帶狀劃分矩陣-向量乘法中,設(shè)每個工作組內(nèi)工作項數(shù)量為S,對方陣A按行劃分為p塊,每塊含有連續(xù)的S行行向量,p=「n/S?,這些行塊依次記為A0,A1,…,Ap-1,分別存放在索引號為0,1,…,p-1的工作組中,同時所有工作組將接收來自系統(tǒng)廣播的列向量B.存儲于局部數(shù)組a中的行塊Ai和列向量B在各工作組中并行做乘積運算,矩陣-向量乘并行算法具體描述如下:

算法.行帶狀劃分的矩陣向量乘法并行算法

輸入:矩陣An×n,向量Bn×1.

輸出:向量Cn×1.

Begin

forall work-grouppar-do

fori=0top-1do

c[i]=0.0

forj=0ton-1do

c[i]=c[i]+a[i,j]*b[j]

endfor

endfor

endfor

End

4.2 并行算法執(zhí)行模式

基于OpenCL的矩陣-向量乘GPU并行化處理執(zhí)行模式如圖2所示.

矩陣向量乘并行算法實現(xiàn)的主要工作過程如下:

1)進行平臺創(chuàng)建和初始化OpenCL設(shè)備.

2)創(chuàng)建上下文.通過在主機端創(chuàng)建這種抽象容器,實現(xiàn)在主機與設(shè)備之間進行相應(yīng)的協(xié)調(diào)、交互,并對設(shè)備上的可用存儲器進行管理、安排.

圖2 矩陣-向量乘并行算法執(zhí)行模式Fig.2 Matrix-vector multiplication parallel algorithm execution mode

3)創(chuàng)建命令隊列.將提交的命令與設(shè)備進行通信,命令隊列是主機端向相應(yīng)的設(shè)備端傳遞消息和發(fā)送請求的一種行為機制.

4)創(chuàng)建設(shè)備緩存.創(chuàng)建三個buffer用來讀入矩陣A、列向量B和輸出列向量C.

5)將主機端數(shù)據(jù)寫入設(shè)備存儲器.將矩陣A和列向量B分別寫入全局存儲器和常量存儲器中,用來完成數(shù)據(jù)的傳遞.

6)編譯程序?qū)ο?通過clCreateProgramWithSource()將kernel轉(zhuǎn)化為一個cl_program對象,再使用clBuildProgram()編譯kernel.

7)創(chuàng)建kernel對象.將主機端傳來的矩陣A和列向量B對象進行乘積操作.

8)將輸出讀回主機端并釋放OpenCL資源.將輸出結(jié)果列向量C讀到主機端內(nèi)存.

4.3 矩陣-向量乘并行算法設(shè)計

1)兩級并行計算

根據(jù)矩陣-向量乘并行算法執(zhí)行模式,需要將矩陣A進行分塊,將大規(guī)模的矩陣劃分為小規(guī)模的行塊Ai進行處理.每個工作組按照獲得的坐標(biāo)get_group_id(0)位置將相應(yīng)的行塊Ai調(diào)度到計算單元中.每個工作組中的每個工作項按照獲得的坐標(biāo)get_local_id(0)位置將相應(yīng)的矩陣行和常量存儲器中的列向量B做乘―加運算.在矩陣-向量乘并行算法中,工作組負(fù)責(zé)行塊Ai的數(shù)據(jù)處理,每個工作項負(fù)責(zé)將相應(yīng)坐標(biāo)位置上的行向量和列向量B調(diào)度到相應(yīng)處理單元中做乘―加運算,從常量存儲器中取出列向量B并將列向量C存儲到全局存儲器中.同一個工作組中的工作項可以通過work_group_barrier()進行同步處理.

矩陣-向量乘法算法的并行計算按照在工作組進行粗粒度并行計算和在工作項進行細(xì)粒度并行計算的兩個層次進行組織.粗粒度并行是存在于同一N維工作空間NDRange內(nèi)的不同工作組之間,不同工作組之間不需要進行數(shù)據(jù)的通信和交換,使得并行算法系統(tǒng)可擴展性很好:由于一個計算單元上可以任意運行每個工作組中的子任務(wù),因此,基于OpenCL的矩陣-向量乘并行計算系統(tǒng)在核心數(shù)量不同的GPU上都能正常運行.細(xì)粒度并行是存在于同一工作組內(nèi)的工作項之間,同一工作組內(nèi)的工作項之間需要進行數(shù)據(jù)的交換和通信.

2)內(nèi)核函數(shù)配置

根據(jù)互不重疊的矩陣行塊數(shù)量,整個矩陣可以劃分成若干個行塊,每個行塊在GPU中可作為一個基本處理單位,分割出來的大量行塊的計算可由工作空間和工作組處理.從數(shù)據(jù)層面上看,GeForce GTX 980中擁有16個計算處理單元,一個計算處理單元可同時處理32個行塊.所以,在GPU中可并行處理512個行塊.矩陣-向量乘并行算法將n作為工作空間在x方向上的長度,行塊的大小S作為每個工作組內(nèi)工作項值,核函數(shù)的網(wǎng)格配置如下:

size_tlocalWorkSize[]=(S)

size_tglobalWorkSize[]=(n)

clEnqueueNDRangeKernel

3)工作項映射模型

為了標(biāo)記行塊中每個工作項在工作空間中的位置,利用OpenCL內(nèi)建函數(shù)建立了工作項的ID索引workItemID,如下所示:

workItemID=get_local_id(0)+get_group_id(0)*get_local_size(0)

4.4 優(yōu)化策略

OpenCL存儲模型包含了可以訪問的多種快速存儲區(qū)域,如私有存儲器、本地存儲器和常量存儲器等,合理地對快速存儲器的運用,將會使算法獲得更高的性能.由于在矩陣-向量乘并行算法的處理過程中,列向量B是不會改變的,而且每個工作項的處理都要用到該列向量.為了提高運行效率,本算法選擇把列向量B的內(nèi)容傳遞到常量存儲器,并且常量存儲器帶有緩存機制,這樣能進一步提高每個工作項訪問列向量B的效率,加速了算法的運算速度.

5 實驗測試及設(shè)計

5.1 測試平臺和實驗數(shù)據(jù)

1)硬件環(huán)境

研究平臺:圖形處理器采用GeForce GTX 980,擁有2048顆計算核心,其流處理器頻率為1216MHz,核心頻率為1126MHz,GPU顯存位寬為256位,GPU顯存帶寬為224GB/s,顯存為4GB GDDR5.

對比平臺1:GPU為AMD Radeon HD5850,共有1440個流處理單元,核心頻率為725MHz;GDDR5顯存,頻率為1050MHz,容量為1GB,位寬為256bits.

對比平臺2:CPU為Intel i5 8400 2.8GHz(六核心),系統(tǒng)存儲器為8.0GB.

2)軟件環(huán)境

操作系統(tǒng)采用Windows7 64bits;集成研發(fā)工具為Visual Studio 2013;環(huán)境支持為CUDA Toolkit 7.5,內(nèi)部支持OpenCL 1.2.

實驗采用矩陣大小分別為5×5、10×10、50×50、100×100、500×500、1000×1000、2000×2000共7組實驗數(shù)據(jù),并將矩陣乘法算法分別運行于基于CPU的OpenMP系統(tǒng)、基于NVIDIA GPU的CUDA系統(tǒng)、基于AMD GPU的OpenCL系統(tǒng)和基于NVIDIA GPU的OpenCL系統(tǒng),并記錄處理時間,如表1所示.

為了驗證各種架構(gòu)下并行算法的效率,可以直觀地用加速比作為加速效果的衡量標(biāo)準(zhǔn),其定義如下:

表1 不同計算平臺下矩陣-向量乘算法執(zhí)行時間(s)

Table 1 Matrix-vector multiplication algorithm execution time under different computing platforms

矩陣大小串行處理時間并行處理時間OpenMPCUDAOpenCL(AMD)OpenCL(NVIDIA)5×50.1580.1270.0280.0230.02410×100.2240.1660.0280.0270.02650×500.8240.2880.0960.0800.077100×1002.9320.6740.3140.2340.223500×5009.4722.0240.8750.6370.6331000×100022.6694.3511.8521.2841.2432000×200032.1315.8742.0281.5951.540

加速比是指CPU串行算法執(zhí)行時間與并行算法執(zhí)行時間的比值.

表2 不同計算平臺下矩陣-向量乘并行算法性能對比
Table 2 Performance comparison of parallel algorithm of matrix-vector multiplication on different computing platforms

序號矩陣大小加速比 OpenMPCUDA OpenCL(AMD)OpenCL(NVIDIA)相對加速比1相對加速比2150×501.245.646.876.585.311.172100×1001.358.008.308.626.391.083200×2002.868.5810.3010.703.741.254300×3004.359.3312.5313.153.021.415500×5004.6810.8214.8714.963.201.3861000×10005.2112.2417.6518.233.501.4972000×20005.4715.8420.1420.863.811.32

相對加速比1是指基于CPU的OpenMP并行算法執(zhí)行時間與基于NVIDIA GPU的OpenCL并行算法執(zhí)行時間的比值.

相對加速比2是指基于NVIDIA GPU的CUDA并行算法執(zhí)行時間與基于NVIDIA GPU的OpenCL并行算法運行時間的之比.

相應(yīng)計算架構(gòu)下GPU并行算法相比CPU串行算法整體性能提升情況可用加速比反映,可用于對系統(tǒng)性能改進程度的衡量,如表2所示.而相對加速比1則反映出基于OpenCL的NVIDIA GPU并行算法相比基于OpenMP的多核CPU并行算法的性能提升程度,相對加速比2則反映出基于OpenCL的NVIDIA GPU并行算法相比基于CUDA的GPU并行算法的性能提升程度.

5.2 并行算法性能分析

1)矩陣-向量乘并行算法加速比性能分析

圖3顯示了在不同計算平臺下矩陣-向量乘并行算法的加速比對比曲線.從中看出,隨著矩陣大小的增加,基于OpenMP算法的加速比不斷地增加,但是由于CPU核心數(shù)有限,加速比較小且增幅較小.而經(jīng)過CUDA和OpenCL加速的矩陣-向量乘并行算法,由于GPU計算資源充裕并且可以創(chuàng)建充足的工作項來滿足大量數(shù)據(jù)的并行處理,通過時間分割的機制將工作項分配到2048個處理單元中,獲得的加速比較大且增幅較大.從中反映出基于GPU運算對于計算密集型的運算不如CPU敏感,在面對大規(guī)模數(shù)據(jù)的計算密集型運算時,GPU計算時間的增幅很小,體現(xiàn)出GPU強大的計算能力.

由圖3可知,隨著矩陣規(guī)模的增加,GPU的加速效果十分明顯,例如:對于2000×2000大小的矩陣-向量乘法運算,獲得了近21倍加速比.然而,GPU并行處理系統(tǒng)性能的提高并非線性增長,而是隨著矩陣規(guī)模的增加,呈現(xiàn)出緩慢提升的趨勢.這是由于當(dāng)矩陣規(guī)模增大時,PCI-E總線帶寬有限,GPU顯存與主機內(nèi)存間數(shù)據(jù)傳輸通信耗時較多,通信瓶頸現(xiàn)象明顯出現(xiàn).內(nèi)存與顯存之間的傳輸時間極大地抵消了GPU并行計算優(yōu)勢,GPU加速比增幅緩慢,制約了系統(tǒng)整體性能的提升.

圖3 矩陣-向量乘并行算法的加速比對比Fig.3 Comparison of acceleration ratios of matrix-vector multiplication parallel algorithm

2)基于OpenCL的矩陣-向量乘并行算法跨平臺性分析

圖3表明,算法源碼級別的可移植性不僅要求源碼可以在不同的平臺上成功的編譯和運行,而且還要求算法保持相當(dāng)?shù)男阅?實驗表明,基于CUDA的矩陣-向量乘并行算法受到硬件平臺的限制,而基于OpenCL的矩陣-向量乘并行算法在多種硬件平臺上具備最大程度的可移植性和兼容性,其算法性能與OpenMP計算平臺和CUDA計算平臺上的算法性能基本相近并略有提升.

6 結(jié)束語

本文對OpenCL計算平臺的并行架構(gòu)和矩陣-向量乘法算法的并行層次進行了詳細(xì)的討論和分析,提出了矩陣-向量乘法算法的GPU并行化實現(xiàn)方法,并對其并行算法和關(guān)鍵技術(shù):兩級并行計算,工作項映射模型與性能優(yōu)化等問題進行了詳細(xì)論述.實驗結(jié)果表明,與基于CPU的串行算法、基于CPU的OpenMP并行算法和基于GPU的CUDA并行算法的性能相比,優(yōu)化后的基于OpenCL的矩陣-向量乘并行算法分別獲得了20.86倍、6.39倍和1.49倍的加速比,不僅實現(xiàn)了高性能,而且實現(xiàn)了在不同計算平臺間的性能移植.本文采用的優(yōu)化方法,充分利用了GPU強大的眾核并行處理能力和OpenCL架構(gòu)良好的可移植性,數(shù)值計算能夠被有效地加速,為進行跨平臺高速大規(guī)模數(shù)據(jù)運算的提供了一條新的思路.今后將進一步研究如何將GPU和集群的計算能力進行整合,形成CPU+GPU多節(jié)點GPU集群,以滿足更大規(guī)模的矩陣對計算能力的需求.