王 平 全吉成 王宏偉
(空軍航空大學航空航天情報系 長春 130022)
?
GPU圖像并行處理設(shè)計模型研究*
王 平 全吉成 王宏偉
(空軍航空大學航空航天情報系 長春 130022)
GPU通用計算的發(fā)展為圖像處理提供了新的技術(shù)手段,目前一些圖像處理算法已經(jīng)移植到GPU中且得到了較好的加速效果,但是并沒有形成通用的并行設(shè)計模型,推廣性差。為此,論文在研究GPU結(jié)構(gòu)特點以及已有的GPU圖像處理并行算法的基礎(chǔ)上,提出了單GPU和多GPU圖像并行處理設(shè)計模型,并通過圖像平滑并行算法在GPU中的執(zhí)行驗證了該模型的有效性。試驗結(jié)果表明,該模型在GPU圖像處理中通用性較好,為更多算法的并行化處理提供了理論借鑒。
圖像算法; 并行化; GPU; 模型; 通用性
Class Number TN911
隨著圖像獲取手段的發(fā)展,獲取圖像的效率逐漸加快以及圖像分辨率急劇增長,現(xiàn)有的圖像處理工具已經(jīng)無法滿足海量圖像數(shù)據(jù)的快速處理,造成影像數(shù)據(jù)的大量堆積。因此如何快速、高效地進行海量圖像數(shù)據(jù)處理是圖像處理領(lǐng)域亟待解決的問題。
圖形處理器(Graphic Process Unit,GPU)在通用計算方面取得了快速發(fā)展,計算能力已經(jīng)遠遠超過了CPU,成為實現(xiàn)高性能計算的重要部件[1]。近年來,不少專家學者將圖像處理算法移植到GPU中[2~5],并取得了很高的算法加速比,加快了圖像算法并行化處理的發(fā)展。這些映射到GPU中算法只是一些特定的算法,并沒有形成通用的GPU圖像處理并行設(shè)計模型,難以推廣到其他的圖像處理算法中。
本文根據(jù)GPU體系結(jié)構(gòu)的特點以及對已有的GPU圖像處理并行算法的分析,提出了通用的GPU圖像并行處理設(shè)計模型,并通過圖像平滑算法進行GPU并行處理來驗證模型的有效性,進而可以將模型擴展到更多圖像處理算法中。
GPU圖像處理并行設(shè)計模型主要描述圖像串行算法進行GPU并行處理的一般性的解決方案。 這樣可以讓我們復(fù)用已經(jīng)建立的設(shè)計模型,為自己的算法并行化問題找到有效的解決方法。在大量圖像處理算法并行化分析的基礎(chǔ)上,本文提出的針對單個GPU的通用的GPU圖像處理并行設(shè)計模型包括以下四個步驟:串行算法的并行化分析,并行粒度劃分,映射到GPU中,程序的優(yōu)化,如圖1所示。如果含有多個GPU,并行設(shè)計模型應(yīng)該包括五個步驟:串行算法的并行化分析,任務(wù)劃分,并行粒度劃分,映射到GPU中,程序的優(yōu)化,如圖2所示。
圖1 GPU圖像并行處理設(shè)計模型
圖2 GPUs圖像并行處理設(shè)計模型
1) 串行算法的并行化分析:串行算法的并行化分析就是找尋串行算法中計算量最大、處理時間最耗時、數(shù)據(jù)相關(guān)性小的部分,然后對其進行并行化分析。這主要以下三個方法[6]:
(1)根據(jù)算法中已有的并行性將算法改變成并行算法;
(2)從算法本身要處理的問題特征著手把傳統(tǒng)算法并行化;
(3)根據(jù)已有經(jīng)典并行算法的設(shè)計思路,改造傳統(tǒng)圖像處理算法。
2) 任務(wù)劃分:在并行算法中,將整個任務(wù)分成若干個子任務(wù),每個GPU對應(yīng)一個子任務(wù)。在多核CPU上啟動多個線程,每個線程控制一個設(shè)備,實現(xiàn)設(shè)備間的并行計算。
3) 并行粒度劃分:GPU中執(zhí)行的核函數(shù)Kernel可以看做一個線程網(wǎng)格,一個線程網(wǎng)格包含若干個線程塊,每個線程塊包含若干個線程,如圖3所示。在GPU具體執(zhí)行中,線程塊和塊中線程被調(diào)度分別實現(xiàn)粗粒度和細粒度并行[7]。并行粒度劃分就是將圖像并行算法映射到線程塊和塊中的線程中,從而實現(xiàn)GPU內(nèi)粗、細粒度并行執(zhí)行。
圖3 GPU內(nèi)部線程組織方式
4) 映射到GPU中:將并行算法映射到GPU眾核的架構(gòu)上,通過多線程并行執(zhí)行來實現(xiàn)算法的并行運行。具體步驟如下:
(1)分配顯存空間,將影像數(shù)據(jù)從主存?zhèn)鬏數(shù)斤@存中;
(2)設(shè)置執(zhí)行配置即設(shè)置線程塊大小和數(shù)量,啟動核函數(shù)Kernel;
(3)將結(jié)果影像數(shù)據(jù)從顯存?zhèn)鬏數(shù)街鞔妗?/p>
5) 程序優(yōu)化:在以上幾個步驟完成后,程序能夠正確運行,并且結(jié)果驗證準確,這是程序優(yōu)化的前提,也是程序優(yōu)化的起點。程序優(yōu)化包括存儲器優(yōu)化、指令優(yōu)化和編碼優(yōu)化。存儲器優(yōu)化是指優(yōu)化存儲器利用率和最大化全局存儲器吞吐量[8]。指令優(yōu)化是指使用GPU提供的存儲器控制、線程控制等各種優(yōu)化指令,能顯著提升程序的性能。編碼優(yōu)化指程序編寫時,盡可能少用if分支語句;使用#unroll指令展開循環(huán)語句;在保證精度的情況下,盡可能的使用專用的算術(shù)指令等[9~10]??偠灾绦騼?yōu)化的核心思想是改進計算訪存比以及各存儲層次上的延遲隱藏[11]。
3.1 算法并行化分析
圖像平滑算法是針對單個像素進行計算,各個像素之間的操作沒有相關(guān)性,而且計算量大,并行程度高,適合移植到GPU中進行并行處理。
3.2 并行粒度劃分
平滑后的圖像與GPU線程網(wǎng)格對應(yīng),將圖像進行分塊,使每個圖像塊與一個線程塊對應(yīng),每個圖像分塊中像素與線程塊中的線程相對應(yīng),保證線程塊中線程對每一個像素進行處理。這樣就實現(xiàn)了線程塊之間的在SM中粗粒度并行執(zhí)行,以及線程塊內(nèi)的線程在CUDA核心中細粒度并行執(zhí)行。
3.3 映射到GPU中
根據(jù)圖像大小分配顯存大小,將圖像數(shù)據(jù)從主存?zhèn)鬏數(shù)斤@存。然后設(shè)置線程塊大小和數(shù)量,啟動核函數(shù),最后將結(jié)果影像數(shù)據(jù)從顯存中傳輸?shù)街鞔嬷泻铣山Y(jié)果影像。部分代碼如下所示:
cudaMalloc((void**)&ImageDevice , ImageSize); //在GPU上分配內(nèi)存空間
cudaMemcpy(ImagedataDevice,ImagedataHost,ImageSize,cudaMemcpyHostToDevice);//將圖像數(shù)據(jù)由內(nèi)存復(fù)制到顯存中
dim3 dimBlock(m,n)
dim3 dimGrid((width+dimBlock.x-1) /dimBlock.x) , (height+dimBlock.y-1) /dimBlock.y) //劃分線程塊大小和數(shù)量
Kernel<<
CopyDataFormDevice(ImagedataOut,ImagedataDevice); //將GPU上運算結(jié)果傳輸?shù)街鳈C內(nèi)存
FreeDeviceAndHost();//釋放主機端和GPU端的內(nèi)存
3.4 程序優(yōu)化
經(jīng)過將并行算法映射到GPU中能夠正確運行且結(jié)果驗證準確,才可以在此基礎(chǔ)對程序做進一步的優(yōu)化,來提高程序的性能。平滑算法中每一個像素的計算都使用相同的平滑算子,而且共享存儲器讀寫速度非???,可以被同一線程塊的所有線程訪問。因此為降低對全局存儲器DRAM內(nèi)存的存取,將平滑算子傳入共享存儲器中,從而提高數(shù)據(jù)的存取速度。
實驗平臺使用GTX580顯卡,全局存儲器大小為1.5G,共享存儲器為48KB,有16個流多處理器,512個流處理器。CPU為AMD Athlon (tm) II X4 640,內(nèi)存為10G。分別在CPU和單GPU上進行圖像平滑操作。
測試數(shù)據(jù)為256×256、512×512、1024×1024、2048×2048、4096×4096、8192×8192的灰度圖。將不同分辨率的圖像分別在CPU、原始GPU和程序優(yōu)化后GPU進行平滑處理,對處理時間進行計時,取實驗10次的結(jié)果得到平均執(zhí)行時間。實驗結(jié)果如表1所示。
表1 算法在不同環(huán)境下的運行時間
從表1可以看出,在GPU環(huán)境下不同分辨率的圖像進行平滑操作的時間明顯縮短,經(jīng)過程序優(yōu)化之后,執(zhí)行時間又得到了進一步減小。不同分辨率圖像在GPU和優(yōu)化后GPU上運行時間與在CPU上運行時間的加速比曲線如圖4所示。
圖4 GPU優(yōu)化前后并行算法加速比
從圖4可以看出隨著圖像分辨率的增大,優(yōu)化前GPU和優(yōu)化后GPU的并行算法加速比曲線都呈上升趨勢,當圖像分辨率為8192×8192時,優(yōu)化前GPU和優(yōu)化后GPU的并行算法加速比都分別達到最大的34.3倍和37.3倍。
圖5 雙GPU并行處理流程圖
對圖像平滑算法進行雙GPU試驗時,需要按GPUs圖像并行處理設(shè)計模型的步驟2對任務(wù)進行劃分,將圖像數(shù)據(jù)拆分為兩部分,分別在兩個GPU上并行處理。在雙GPU環(huán)境下,CPU需創(chuàng)建兩個線程分別控制不同的GPU設(shè)備進行圖像處理,然后在CPU中合成最終結(jié)果,處理流程如圖5所示。
對雙GPU圖像平滑算法進行測試時,使用單GPU實驗圖像數(shù)據(jù),結(jié)果如表2所示。
表2 單GPU和雙GPU的處理時間比較
從表2可以看出,當圖像分辨率較低時,雙GPU處理速度稍低于單GPU處理速度,隨著分辨率增加到一定程度,雙GPU處理速度優(yōu)勢才越來越明顯。
從單GPU和雙GPU兩個實驗結(jié)果中表明,并行設(shè)計后的圖像平滑算法在GPU中取得了很好的加速效果,從而驗證了GPU圖像處理并行設(shè)計模型的有效性。
本文在對大量圖像并行算法研究的基礎(chǔ)上,提出了GPU圖像處理并行設(shè)計模型,通過并行平滑算法在GPU中的執(zhí)行從而驗證了模型的有效性。該模型簡單易懂,通用性強,為更多圖像處理算法并行化并移植到GPU中提供了理論基礎(chǔ)。
[1] 袁濤,馬艷,劉定生.GPU在遙感圖像處理中的應(yīng)用綜述[J].遙感信息,2012,27(6):110-117.
[2] 方留楊,王密,李德仁.CPU和GPU協(xié)同處理的光學衛(wèi)星遙感影像正射校正方法[J].測繪學報,2013,42(5):668-675.
[3] 李朋龍,鄧非,何江,等.GPU-CPU協(xié)同航空影像快速正射糾正方法[J].測繪地理信息,2016,41(2):44-47.
[4] 王化喆,魏先勇.基于GPU的遙感圖像前期處理算法研究與應(yīng)用[J].現(xiàn)代電子技術(shù),2016,39(3):47-50.
[5] D GRANATA, U AMATO, B ALFANO. MRI denoising by nonlocal means on multi-GPU[J]. Journal of Real-Time Image Processing,2016:1-11.
[6] 吳頎峰.基于GPU的數(shù)字圖像處理算法研究[D].贛州:江西理工大學,2011.
[7] 肖漢,張祖勛.基于GPGPU的并行影像匹配算法[J].測繪學報,2010,39(1):46-51.
[8] 肖漢.基于CPU_GPU的影像匹配高效能異構(gòu)并行計算研究[D].武漢:武漢大學,2011.
[9] 張舒,褚艷利.GPU高性能運算之CUDA[M].北京:中國水利水電出版社,2009.
[10] NVIDIA Corporation. CUDA Programming Guide 1.0[EB/OL]. http://www.nvidia.com,2009(8):11.
[11] 馬安國,成玉,唐遇星,等.GPU異構(gòu)系統(tǒng)中的存儲層次和負載均衡策略研究[J].國防科學技術(shù)大學學報,2009,31(5):38-43.
Parallel Design Model of GPU Image Processing
WANG Ping QUAN Jicheng WANG Hongwei
(Department of Aeronautic and Astromautic Intelligence, Aviation University of Air Force, Changchun 130022)
The development of GPU general computing provided a new technique for image processing, and some image processing algorithms had been transplanted to GPU and got a better acceleration effect. But it could not form a general parallel design model. Thus, based on the study of the characteristics of GPU structure and the existing parallel algorithms of GPU image processing, this paper proposed a parallel image processing design model of single GPU and multi GPU. And this paper took the algorithm design of parallel image smooth as an example to testify the effectiveness of this model. The experimental results showed that the proposed model had good generality, and it could provide a theoretical reference for the parallel processing of more algorithms in the GPU image processing.
image algorithms, parallel, GPU, model, generality
2016年5月6日,
2016年6月25日
王平,男,碩士研究生,研究方向:數(shù)字圖像處理和GPU并行計算。全吉成,男,博士,教授,博士生導師,研究方向:二維可視化技術(shù)。王宏偉,男,碩士,講師,研究方向:數(shù)字圖像處理和地理信息系統(tǒng)。
TN911
10.3969/j.issn.1672-9730.2016.11.021