基于CUDA的圖像預(yù)處理并行化研究

， ，，

(哈爾濱工業(yè)大學(xué)機(jī)器人技術(shù)與系統(tǒng)國(guó)家重點(diǎn)實(shí)驗(yàn)室，黑龍江 哈爾濱 150001)

基于CUDA的圖像預(yù)處理并行化研究

占正鋒，李戈，張學(xué)賀，尹旭悅

(哈爾濱工業(yè)大學(xué)機(jī)器人技術(shù)與系統(tǒng)國(guó)家重點(diǎn)實(shí)驗(yàn)室，黑龍江 哈爾濱 150001)

為加快圖像預(yù)處理算法的執(zhí)行速度，提出了基于計(jì)算統(tǒng)一設(shè)備架構(gòu)(CUDA)的預(yù)處理算法來實(shí)現(xiàn)高速并行處理。分析了圖像灰度化、高斯濾波以及直方圖均衡化等預(yù)處理方法的原理，并對(duì)它們進(jìn)行并行化分析，從而將CUDA并行計(jì)算技術(shù)引入到圖像預(yù)處理算法。實(shí)驗(yàn)結(jié)果表明，此算法充分利用GPU的并行處理能力，與CPU串行處理方法相比，速度提高明顯，有效提高數(shù)據(jù)處理能力。

圖像預(yù)處理；CUDA；并行處理

0 引言

一般情況下，成像系統(tǒng)獲取的圖像(即原始圖像)由于受到種種條件限制和隨機(jī)干擾，往往不能在視覺系統(tǒng)中直接使用，必須在視覺信息處理的早期階段對(duì)原始圖像進(jìn)行灰度變換、噪聲過濾等圖像預(yù)處理[1]。傳統(tǒng)的預(yù)處理算法，當(dāng)數(shù)據(jù)量較大時(shí)，這種方法的效率不高，難以滿足實(shí)時(shí)性的要求。近年來，基于圖像處理器(GPU)的大規(guī)模通用并行計(jì)算，為圖像處理算法的高效分析計(jì)算提供了可能。NVIDIA公司提出的計(jì)算統(tǒng)一設(shè)備構(gòu)架，是一種將GPU作為數(shù)據(jù)并行計(jì)算設(shè)備的軟硬件體系，其特點(diǎn)是將CPU作為主機(jī)，GPU作為協(xié)處理器[2]。這個(gè)模型中，CPU和GPU協(xié)同工作，CPU負(fù)責(zé)邏輯性強(qiáng)的事務(wù)和串行計(jì)算任務(wù)，而GPU處理高度線程化的并行計(jì)算任務(wù)[3]。

在此，以GPU作為計(jì)算平臺(tái)，對(duì)圖像預(yù)處理算法中的圖像灰度化、高斯濾波以及直方圖均衡化處理進(jìn)行并行化分析，并在CUDA上實(shí)現(xiàn)。

1 圖像的灰度化并行

將彩色圖像轉(zhuǎn)化成為灰度圖像的過程稱為圖像的灰度化處理。彩色圖像中的每個(gè)像素的顏色由R,G,B3個(gè)分量決定，而每個(gè)分量有255種值可取，這樣一個(gè)像素點(diǎn)可以有1 600多萬(255×255×255)種顏色的變化范圍。而灰度圖像是R,G,B3個(gè)分量相同的一種特殊的彩色圖像，其一個(gè)像素點(diǎn)的變化范圍為255種，所以在數(shù)字圖像處理中，一般先將各種格式的圖像轉(zhuǎn)變成灰度圖像，以使后續(xù)圖像的計(jì)算量變得少一些。圖像的灰度化處理，一般而言，常用如下幾種方法：最大值法、平均法和加權(quán)平均法等[4]。這里采用加權(quán)平均法，分別賦予R,G,B不同的權(quán)值，然后取其加權(quán)后的平均值作為最后的灰度值。

并行灰度轉(zhuǎn)換實(shí)現(xiàn)過程中，將每個(gè)像素3個(gè)通道的顏色值存放在線程寄存器內(nèi)，1個(gè)block包含256個(gè)thread，則對(duì)于一幅尺寸為width × height的圖像，一共需要width × height/256個(gè)block，創(chuàng)建索引標(biāo)號(hào)為index=blockIdx.x×256+threadIdx.x。

2 高斯濾波并行化

高斯濾波(高斯平滑)是數(shù)字圖像處理和計(jì)算機(jī)視覺里面最常見的操作。高斯濾波器是根據(jù)高斯函數(shù)的形狀來選擇權(quán)值的線性平滑濾波器[5]。高斯平滑濾波器對(duì)去除服從正態(tài)分布的噪聲有很好的效果。對(duì)圖像來說，常用二維高斯函數(shù)表達(dá)式為：

(1)

σ為高斯濾波器的寬度。

高斯濾波時(shí)，在對(duì)整幅圖像進(jìn)行加權(quán)平均的過程中，每一個(gè)像素點(diǎn)的值，都由其本身和鄰域內(nèi)的其他像素值經(jīng)過加權(quán)平均后得到。高斯濾波的具體操作是：用1個(gè)模板(或稱卷積、掩模)掃描圖像中的每一個(gè)像素，用模板確定的鄰域內(nèi)像素的加權(quán)平均灰度值，去替代模板中心像素點(diǎn)的值。

為了實(shí)現(xiàn)高斯濾波的快速運(yùn)行，根據(jù)高斯函數(shù)具有可分離性的特點(diǎn)，提出了可進(jìn)行分離的高斯卷積的方法，把二維高斯卷積轉(zhuǎn)換成2個(gè)一維高斯卷積，將高斯模板放入GPU帶緩存的常量存儲(chǔ)器中。

由高斯函數(shù)的性質(zhì)可知，高斯函數(shù)具有各向同性的特點(diǎn)，二維高斯函數(shù)可進(jìn)行逐次卷積，一個(gè)沿水平方向，另一個(gè)沿垂直方向。即

(2)

在離散系統(tǒng)中，二維高斯卷積是靠n×n的高斯模板對(duì)整幅圖像進(jìn)行卷積的，每個(gè)圖像的像素都需要進(jìn)行n×n次的乘加操作，對(duì)于大小為m×m的圖像就需要m×m×n×n次乘加操作，同時(shí)還要用n×n個(gè)存儲(chǔ)空間來存儲(chǔ)模板。而將二維高斯函數(shù)進(jìn)行分離以后，把二維的n×n高斯模板分解成1 ×n的一維高斯模板，對(duì)圖像先進(jìn)行行卷積再進(jìn)行列卷積，對(duì)于大小為m×m的圖像一共只需要2 ×m×n×n次操作。可以看出，當(dāng)圖像很大時(shí)，使用本文的方法進(jìn)行高斯卷積將提升運(yùn)算速度。以5×5的二維高斯模板為例，如圖1所示。將其分解后對(duì)應(yīng)的一維高斯模板如圖2所示。圖1所示的5×5二維高斯模板可分解為圖2的1×5一維模板和其轉(zhuǎn)置的乘積，這樣就將二維高斯平滑轉(zhuǎn)換成了2次的一維高斯平滑。

在GPU上對(duì)高斯卷積進(jìn)行并行化的主要流程如圖3所示。

圖1 5×5高斯模板

圖2 1×5一維高斯模板

圖3 GPU高斯卷積流程

首先將圖像從主機(jī)端拷入到設(shè)備端，綁定紋理內(nèi)存。選擇綁定紋理內(nèi)存是由于以下原因：紋理內(nèi)存緩存在芯片上，能夠減少對(duì)內(nèi)存的請(qǐng)求并提供更高效的內(nèi)存帶寬，是專門為那些在內(nèi)存訪問模式中存在大量空間局部性的圖形應(yīng)用程序而設(shè)計(jì)的，能在隨機(jī)存取的環(huán)境中，保持較高的性能；圖像在高斯濾波時(shí)要面臨邊界問題，而圖像處理器程序并不擅長(zhǎng)條件判斷語句的執(zhí)行，紋理內(nèi)存能夠通過屬性設(shè)置，比較高效處理此類問題。

高斯卷積模板在主機(jī)端計(jì)算完畢后，傳入設(shè)備端的常量存儲(chǔ)器。在行列2個(gè)維度上使用2個(gè)kernel函數(shù)進(jìn)行卷積。對(duì)行維度卷積時(shí)，將Grid中的block按照二維地址排序，每個(gè)block負(fù)責(zé)處理1幅圖像的1行的部分?jǐn)?shù)據(jù)，大小可以定為16×16，即每個(gè)block有256個(gè)thread，每個(gè)thread負(fù)責(zé)處理圖像的1個(gè)像素點(diǎn)的值，即每個(gè)線程把當(dāng)前的像素點(diǎn)和其周圍的8個(gè)像素點(diǎn)進(jìn)行加權(quán)計(jì)算。同一個(gè)block中的線程在進(jìn)行計(jì)算時(shí)，有很多數(shù)據(jù)是重復(fù)的，將它們放在block的共享存儲(chǔ)器中，這樣就可以加快運(yùn)算速度。行卷積任務(wù)分配如圖4所示。對(duì)列的卷積與行卷積類似，這樣完成行列卷積后，就可以完成對(duì)整幅圖像的高斯二維卷積。

圖4 行卷積任務(wù)分配

3 直方圖均衡化技術(shù)并行化

利用直方圖統(tǒng)計(jì)的結(jié)果，通過使圖像直方圖均衡的方法稱為直方圖均衡化[6]。直方圖均衡化處理是數(shù)字圖像處理中增強(qiáng)圖像的一種重要的方法。

灰度直方圖的計(jì)算十分簡(jiǎn)單，依據(jù)定義，在離散形式下有以下公式成立：

(3)

nk為圖像中出現(xiàn)Sk級(jí)灰度的像素?cái)?shù)；n為圖像像素總數(shù)；nk/n為頻數(shù)。

計(jì)算累積直方圖各項(xiàng)：

(4)

取整擴(kuò)展：

tk=int[(L-1)tk+0.5]

(5)

映射對(duì)應(yīng)關(guān)系：

k?tk

在CUDA上實(shí)現(xiàn)直方圖的均衡化需要進(jìn)行原子操作，而使用原子操作的方式是調(diào)用函數(shù)atomicAdd(addr，y)生成1個(gè)原子的操作序列。這個(gè)操作序列包括讀取地址addr處的地址，將y增加到這個(gè)值，以及將結(jié)果保存回地址addr。底層硬件將確保當(dāng)執(zhí)行這些操作時(shí)，其他任何線程都不會(huì)讀取或?qū)懭氲刂穉ddr上的值，這樣就能確保得到預(yù)計(jì)的結(jié)果。

將直方圖的計(jì)算分為2個(gè)階段。在第1階段中，每個(gè)并行線程塊將計(jì)算它所處理數(shù)據(jù)的直方圖。由于每個(gè)線程塊在執(zhí)行這個(gè)操作時(shí)都是相互獨(dú)立的，因此，可以在共享內(nèi)存中計(jì)算這些直方圖，這將避免每次將寫入操作從芯片發(fā)送到DRAM。因?yàn)樵诰€程塊中的多個(gè)線程之間會(huì)處理相同值的數(shù)據(jù)元素，所以需要原子操作。在第1階段中，給每個(gè)block分配256個(gè)thread，同時(shí)分配1個(gè)共享內(nèi)存緩存區(qū)并進(jìn)行初始化，用來保存每個(gè)線程塊的臨時(shí)直方圖，即_shared_ unsigned int temp[256]；temp[threadIdx.x]=0。由于隨后的步驟將包括讀取和修改這個(gè)緩沖區(qū)，因此，需要調(diào)用_syncthreads()，來確保每個(gè)線程的寫入操作在線程繼續(xù)前進(jìn)之前完成。

在第2階段，要求將每個(gè)線程塊的臨時(shí)直方圖合并到全局緩沖區(qū)histo中。假設(shè)將輸入數(shù)據(jù)分為兩半，這樣就有2個(gè)線程查看不同部分的數(shù)據(jù)，并計(jì)算得到2個(gè)獨(dú)立的直方圖。如果線程A在輸入數(shù)據(jù)中發(fā)現(xiàn)字節(jié)0xFC出現(xiàn)了20次，線程B發(fā)現(xiàn)字節(jié)0xFC出現(xiàn)了5次，那么字節(jié)0xFC在輸入數(shù)據(jù)中共出現(xiàn)了25次。同樣，最終直方圖的每個(gè)元素，只是線程A直方圖中相應(yīng)元素和線程B直方圖中相應(yīng)元素的加和。這個(gè)邏輯可以擴(kuò)展到任意數(shù)量的線程，因此，將每個(gè)線程塊的直方圖合并為單個(gè)最終的直方圖，就是將線程塊直方圖的每個(gè)元素都相加到最終直方圖相應(yīng)位置的元素上。這個(gè)操作需要自動(dòng)完成，即atomicAdd(&(histo[threadIdx.x])，temp[threadIdx.x])。由于使用了256個(gè)線程，并且直方圖中包含了256個(gè)元素，因此，每個(gè)線程將自動(dòng)把它計(jì)算得到的元素只增加到最終直方圖的元素上。本算法并不保證線程塊將按照何種順序?qū)⒏髯缘闹迪嗉拥阶罱K的直方圖中，但由于整數(shù)加法是可以交換的，無論哪種順序都會(huì)得到相同的結(jié)果。

4 實(shí)驗(yàn)分析

實(shí)驗(yàn)平臺(tái)采用Windows 7操作系統(tǒng)，CPU為Intel Core i7-3770，主頻3.40 GHz，系統(tǒng)內(nèi)存為8 GB；GPU為GeForce GTX 570。

4.1 圖像灰度化

分別使用圖形尺寸為256×256，512×512，1 024×1 024的圖像，在GPU上測(cè)試其運(yùn)行速度，并且與預(yù)處理算法在CPU上運(yùn)行的速度進(jìn)行比較。實(shí)驗(yàn)效果如圖5和表1所示。

圖5 灰度化效果

表1 灰度化在GPU和CPU上運(yùn)行時(shí)間的比較

圖像大小CPU/msGPU/ms加速比(CPU/GPU)256×2560．0930．0283．26512×5120．1870．0355．311024×10240．4360．0449．87

4.2 高斯濾波

對(duì)平滑尺度為1，模板尺寸為5×5的高斯濾波算法在GPU上進(jìn)行實(shí)驗(yàn)，分別使用圖形尺寸為256 ×256，512×512，1 024×1 024的加椒鹽噪聲的圖像，在GPU上測(cè)試其運(yùn)行速度，并且與該算法在CPU上運(yùn)行的速度進(jìn)行比較。實(shí)驗(yàn)效果如圖6和表2所示。

圖6 高斯濾波效果

表2 高斯濾波在GPU和CPU上運(yùn)行時(shí)間的比較

圖像大小CPU/msGPU/ms加速比(CPU/GPU)256×2560．7560．04915．21512×5122．3590．08228．631024×10246．7850．18536．65

4.3 直方圖均衡化處理

分別使用圖形尺寸為256×256，512×512，1 024×1 024的圖像，在GPU上測(cè)試其運(yùn)行速度，并且與該算法在CPU上運(yùn)行的速度進(jìn)行比較。實(shí)驗(yàn)效果如表3和圖7所示。

表3 均衡化在GPU和CPU上的運(yùn)行時(shí)間的比較

圖像大小CPU/msGPU/ms加速比(CPU/GPU)256×2560．1240．0158．23512×5120．3870．02515．721024×10240．8360．03324．96

圖7 直方圖均衡化效果

4.4 實(shí)驗(yàn)結(jié)果分析

由表1～表3可以看出，圖像預(yù)處理算法在GPU上運(yùn)行的速度明顯高于在CPU上的速度，并且從圖像灰度化、直方圖均衡化以及高斯濾波這3種算法的對(duì)比來看，算法越復(fù)雜、并行程度越高和圖像的分辨率越高，這種加速效果會(huì)越明顯。

5 結(jié)束語

對(duì)傳統(tǒng)的圖像預(yù)處理算法中的圖像灰度化、高斯濾波以及直方圖均衡化處理進(jìn)行了并行化分析，并且使其在CUDA上實(shí)現(xiàn)。實(shí)驗(yàn)結(jié)果表明，預(yù)處理算法在GPU上運(yùn)行與其在CPU上運(yùn)行相比，明顯提高了其運(yùn)算速度。

[1] 顏 瑞.基于 CUDA 的立體匹配及去隔行算法[D].杭州：浙江大學(xué)，2010.

[2] 鄒治海.GPU 架構(gòu)分析與功耗模型研究[D].上海：上海交通大學(xué)，2011.

[3] 張 舒，褚艷利.GPU高性能運(yùn)算之CUDA[M].北京：中國(guó)水利水電出版社，2009.

[4] Sanders J， Kandrot E.CUDA by example：an introduction to general-purpose GPU programming[M].Addison-Wesley Educational Publishers Inc，2010.

[5] 郭陸峰.雙目視覺系統(tǒng)關(guān)鍵技術(shù)研究[D].泉州：華僑大學(xué)，2007.

[6] Oon-Ee Ng，Ganapathy V.A novel modular framework for stereo vision[C]∥IEEE/ASME International Conference on Advanced Intelligent Mechatronics,2009:857-862.

Parallel Study of Image Preprocessing Based on CUDA

ZHANZhengfeng，LIGe，ZHANGXuehe，YINXuyue

(State Key Laboratory of Robotics and System，Harbin Institute of Technology，Harbin 150001，China)

To speed up the execution speed of image preprocessing algorithm,the paper proposes preprocessing algorithm base on Compute Unified Device Architecture (CUDA) to achieve high-speed parallel processing.It analyzes the principle of image gray,Gaussian filtering and histogram equalization preprocessing methods,and does parallel analysis,and then puts CUDA parallel computing technology into the image preprocessing algorithm. Experimental results show that the image preprocessing algorithm based on CUDA takes full advantage of parallel processing capabilities of GPU,compared with CPU serial processing method,significantly faster,improve data processing capability.

image preprocessing；CUDA；parallel processing

2014-03-10

機(jī)器人技術(shù)與系統(tǒng)國(guó)家重點(diǎn)實(shí)驗(yàn)室(哈爾濱工業(yè)大學(xué))自主研究課題(SKLRS201302C)

TP301.6;TP391

A

1001-2257(2014)07-0064-04

占正鋒(1990-)，男，福建福安人，碩士研究生，研究方向?yàn)闄C(jī)器人視覺。