崔 文，李 強(qiáng)，劉曉春，李 由

(1. 中國(guó)人民解放軍96941部隊(duì)， 北京 102208； 2. 國(guó)防科技大學(xué) 空天科學(xué)學(xué)院， 湖南 長(zhǎng)沙 410073；3. 中國(guó)航天員科研訓(xùn)練中心人因工程重點(diǎn)實(shí)驗(yàn)室， 北京 100194 )

紅外成像末制導(dǎo)[1]是精確制導(dǎo)武器的一種重要制導(dǎo)方式，采用前視紅外末制導(dǎo)的導(dǎo)彈在進(jìn)行攻擊檢查時(shí)可用圖像匹配[2-8]實(shí)時(shí)推演判斷識(shí)別點(diǎn)可跟蹤性，以特定姿態(tài)仿真導(dǎo)彈導(dǎo)引至目標(biāo)點(diǎn)的可行彈道。霧靄氣象條件下，為保障命中精度，紅外導(dǎo)引頭一般需要對(duì)實(shí)時(shí)圖進(jìn)行預(yù)處理而后進(jìn)行匹配運(yùn)算。為仿真紅外導(dǎo)引頭實(shí)時(shí)匹配導(dǎo)引過(guò)程，完成目標(biāo)實(shí)時(shí)跟蹤性判斷，本文針對(duì)限制對(duì)比度自適應(yīng)直方圖均衡(Contrast-Limit Adaptive Histogram Equalization, CLAHE[9])和歸一化互相關(guān)(Normalization Cross Correlation, NCC)算法開(kāi)展加速策略研究，提出基于統(tǒng)一計(jì)算設(shè)備架構(gòu)(Compute Unified Device Architecture, CUDA)的霧靄圖像實(shí)時(shí)處理加速方案，程序優(yōu)化后，可實(shí)現(xiàn)目標(biāo)可跟蹤性實(shí)時(shí)判斷。

1 軟硬件設(shè)計(jì)考慮

在實(shí)時(shí)圖像處理算法設(shè)計(jì)領(lǐng)域，學(xué)者們研究的熱點(diǎn)集中在數(shù)字信號(hào)處理器(Digital Signal Processor, DSP)、現(xiàn)場(chǎng)可編程門(mén)陣列(Field Programmable Gate Array, FPGA)和多核中央處理器(Central Processing Unit, CPU)算法設(shè)計(jì)與應(yīng)用。陳正剛等[10]研究了基于DSP與FPGA的視頻跟蹤系統(tǒng)硬件設(shè)計(jì)和NCC方法快速實(shí)現(xiàn)，將一幀圖像的跟蹤算法計(jì)算時(shí)間降至20 ms以內(nèi)。肖漢等[11]采用圖形處理器(Graphics Processing Unit, GPU)并行化方法研究圖像匹配問(wèn)題，獲得7倍于CPU實(shí)現(xiàn)的圖像匹配運(yùn)算速度，雖然在影像數(shù)據(jù)較大時(shí)，由于未對(duì)主機(jī)和設(shè)備之間的數(shù)據(jù)傳輸進(jìn)行優(yōu)化，實(shí)時(shí)度明顯降低，但研究給后繼者以啟迪，推動(dòng)著GPU并行化應(yīng)用加速發(fā)展。宋驥等[12]改進(jìn)優(yōu)化了基于GPU并行計(jì)算的圖像快速匹配，改進(jìn)優(yōu)化后的算法對(duì)圖像尺寸有較強(qiáng)適應(yīng)性。實(shí)時(shí)圖像處理硬件設(shè)計(jì)方面，曾有研究人員[13-14]指出，對(duì)于數(shù)據(jù)處理能力要求很高的應(yīng)用領(lǐng)域，GPU具有CPU甚至高端DSP不可比擬的性能優(yōu)勢(shì)。以上研究表明，選用合適的硬件和軟件實(shí)現(xiàn)對(duì)實(shí)時(shí)圖像處理速度至關(guān)重要，是值得設(shè)計(jì)人員重點(diǎn)考慮的因素?；谶壿嬮T(mén)和觸發(fā)器等數(shù)字電路進(jìn)行并行任務(wù)處理的FPGA，實(shí)時(shí)性很高，但造價(jià)昂貴、編程速度慢，并不適合做實(shí)時(shí)圖像處理算法驗(yàn)證實(shí)驗(yàn)。鑒于上述研究成果和工程經(jīng)驗(yàn)，本文初步確定采用GPU+CPU的硬件組成形式進(jìn)行實(shí)驗(yàn)。

針對(duì)實(shí)時(shí)圖像處理并行加速，有文獻(xiàn)[15]指出，基于開(kāi)源計(jì)算機(jī)視覺(jué)庫(kù)OpenCV進(jìn)行數(shù)字圖像處理、模式識(shí)別、運(yùn)動(dòng)跟蹤等多核CPU算法開(kāi)發(fā)可以提升運(yùn)算速度；文獻(xiàn)[16]指出，采用高效的GPU并行計(jì)算算法，構(gòu)建CPU/GPU異構(gòu)計(jì)算平臺(tái)已解決眾多模擬計(jì)算加速問(wèn)題；文獻(xiàn)[17]指出，在目標(biāo)跟蹤點(diǎn)搜索及跟蹤應(yīng)用領(lǐng)域，基于CUDA并行程序設(shè)計(jì)的GPU編程加快了運(yùn)算速度。也就是說(shuō)，多核CPU并行算法、GPU并行算法都可用于圖像處理加速。更好更優(yōu)的算法決定了更少的計(jì)算時(shí)間，是實(shí)現(xiàn)實(shí)時(shí)性的關(guān)鍵，也是本文要著重考慮的問(wèn)題。

2 實(shí)時(shí)處理算法

2.1 問(wèn)題描述

圖1 霧靄圖像實(shí)時(shí)處理算法計(jì)算流程圖Fig.1 Real-time foggy image processing flowchart

在進(jìn)行加速策略研究之前，先對(duì)加速的對(duì)象——霧靄圖像實(shí)時(shí)處理算法進(jìn)行簡(jiǎn)單描述。該算法以CLAHE算法和模板圖像匹配算法為核心要素，計(jì)算流程見(jiàn)圖1?；诜謮K處理的自適應(yīng)直方圖均衡算法可對(duì)霧靄天氣條件下的圖像進(jìn)行去霧增強(qiáng)，并通過(guò)限制對(duì)比度閾值去除圖像噪聲影響，均衡效果較經(jīng)典直方圖優(yōu)秀且不會(huì)過(guò)度放大噪聲。圖像匹配就是以“圖”搜圖，較為成熟的有基于模板的匹配和基于特征的匹配兩種[1]。基于模板的匹配是從實(shí)時(shí)圖中搜索比對(duì)出模板圖，計(jì)算量大[18]，對(duì)旋轉(zhuǎn)、形變、遮擋比較敏感，但匹配準(zhǔn)確度高，需要考慮硬件加速；基于特征的匹配是從實(shí)時(shí)圖中搜索到特征量，計(jì)算量相對(duì)較小，對(duì)灰度、形變及遮擋的適應(yīng)性較好，匹配精度較高。常用的圖像匹配計(jì)算公式有絕對(duì)平均誤差(Mean Absolute Deviation, MAD)、絕對(duì)差和(Sum of Absolute Differences, SAD)、最小均方誤差(Mean Squared Error, MSE)和NCC(如式(1)所示)。對(duì)于本文研究的目標(biāo)實(shí)時(shí)跟蹤性判斷來(lái)說(shuō)，需要考慮相互匹配的圖像間亮度變化可能，NCC能更好消除這種影響，準(zhǔn)確度更高。

(1)

要解決的問(wèn)題就集中到如何在毫秒間完成圖像匹配，即算法實(shí)時(shí)化加速問(wèn)題。

2.2 加速方案一：OpenMP多核CPU并行

OpenMP多核CPU并行加速方案是采用多線程技術(shù)進(jìn)行并行計(jì)算加速。特點(diǎn)是CPU所有的核共享一個(gè)內(nèi)存，不同核執(zhí)行不同線程，在同一內(nèi)存的不同部分操作多數(shù)據(jù)，加速的目的是將各線程工作負(fù)載均衡分配，CPU整體計(jì)算速度達(dá)到最高。常見(jiàn)的For循環(huán)語(yǔ)句采用圖2方式實(shí)現(xiàn)代碼并行化；不同代碼段并行執(zhí)行時(shí)，采用OpenMP Section語(yǔ)句(見(jiàn)圖3)實(shí)現(xiàn)。

#pragma omp parallel for schedule(dynamic)

for(i=0;i

{

可并行部分代碼；

}

圖2 OpenMP并行化For循環(huán)實(shí)現(xiàn)方式
Fig.2 For loop′s OpenMP parallelization

#pragma omp parallel sections
{
# pragma omp section
{
可并行代碼
}
# pragma omp section
{
可并行代碼
}
}
圖3 不同段并行執(zhí)行的OpenMP Section語(yǔ)句
Fig.3 Different section parallelization in OpenMP

2.3 加速方案二：高效GPU并行算法

GPU并行加速方案是通過(guò)使用NVIDIA公司的CUDA編程模型實(shí)現(xiàn)“block間粗粒度并行”和“thread間細(xì)粒度并行”雙層并行加速。

加速方案用CUDA實(shí)現(xiàn)時(shí)，由CUDA編程指定CPU指令和內(nèi)存分配，GPU完成具體并行計(jì)算。具體說(shuō)就是，用CPU可編譯執(zhí)行的代碼實(shí)現(xiàn)并行數(shù)據(jù)加載內(nèi)存，數(shù)據(jù)映射到GPU，然后GPU開(kāi)展并行計(jì)算、共享存儲(chǔ)器塊內(nèi)數(shù)據(jù)，定制運(yùn)行的眾線程數(shù)量，完成核函數(shù)(用_global_前綴標(biāo)明)并行計(jì)算，完成后再把控制權(quán)交回CPU，數(shù)據(jù)映射回CPU，最終進(jìn)行數(shù)據(jù)存取和文件讀寫(xiě)，由GPU可編譯執(zhí)行的代碼和CPU可編譯執(zhí)行代碼共同完成全算法加速。加速算法實(shí)現(xiàn)時(shí)，要注意劃分的線程塊內(nèi)warp數(shù)目合理，每個(gè)流處理器都可交替執(zhí)行內(nèi)部warp，提高整體計(jì)算效率；也要根據(jù)NVIDIA GPU設(shè)備的流處理器個(gè)數(shù)，并行使用合理的線程塊block數(shù)，提升流多處理器(Streaming Multiprocessor, SM)利用率。

2.4 CUDA實(shí)現(xiàn)并行加速策略

GPU算法開(kāi)銷(xiāo)常見(jiàn)的有調(diào)度開(kāi)銷(xiāo)、GPU內(nèi)核函數(shù)啟動(dòng)開(kāi)銷(xiāo)和內(nèi)核函數(shù)之間數(shù)據(jù)傳輸開(kāi)銷(xiāo)。算法優(yōu)化主要考慮循環(huán)融合和內(nèi)核融合。循環(huán)融合是盡量將明顯獨(dú)立的循環(huán)進(jìn)行歸并，減少總體迭代的次數(shù)；內(nèi)核融合是將數(shù)據(jù)重用、內(nèi)存合并、線程間數(shù)據(jù)共享。霧靄圖像實(shí)時(shí)處理算法加速策略中的歸一化互相關(guān)部分，考慮申請(qǐng)共享內(nèi)存進(jìn)行CUDA并行實(shí)現(xiàn)，使得每個(gè)線程塊block、每個(gè)線程thread都能訪問(wèn)和操作這塊共享內(nèi)存。由于共享內(nèi)存是基于存儲(chǔ)體切換架構(gòu)的，無(wú)論多少個(gè)線程發(fā)起操作，每個(gè)存儲(chǔ)體每個(gè)周期只執(zhí)行一次操作，即如果有線程束中的一個(gè)線程訪問(wèn)片上共享內(nèi)存，所有的線程都能在這同一指令周期內(nèi)同時(shí)執(zhí)行訪問(wèn)共享內(nèi)存，無(wú)須順序訪問(wèn)，因而能有效減少圖像匹配算法像元讀取訪問(wèn)次數(shù)，降低數(shù)據(jù)訪問(wèn)開(kāi)銷(xiāo)。

圖4 實(shí)時(shí)圖內(nèi)block劃分Fig.4 Block division in real-time figure

需要說(shuō)明的是，匹配循環(huán)次數(shù)也對(duì)算法效率產(chǎn)生影響。由于匹配循環(huán)次數(shù)等于[(實(shí)時(shí)圖寬度-模板圖寬度)/ 線程塊Y方向大小, (實(shí)時(shí)圖長(zhǎng)度-模板圖長(zhǎng)度)/ 線程塊X方向大小]，匹配循環(huán)次數(shù)越少， CUDA內(nèi)核函數(shù)1次可使用核心數(shù)越多，算法開(kāi)銷(xiāo)也越小。當(dāng)內(nèi)核函數(shù)計(jì)算需被拆分成多次，無(wú)法一次完成時(shí)，加速效果將大大降低。以上性能加速均基于線程塊大小的平方與匹配循環(huán)次數(shù)的乘積小于GPU設(shè)備標(biāo)稱核心數(shù)目。實(shí)驗(yàn)使用的GPU單個(gè)線程塊block管理的最大線程數(shù)量為1024，BLOCKSIZE必須小于等于32。固定模板圖為64像素×64像素，模板圖線程塊大小BLOCKSIZE可從典型值8、16、32中取一種；固定模板圖為32像素×32像素，模板圖線程塊大小BLOCKSIZE可從典型值8、16中取一種。又鑒于實(shí)驗(yàn)所用GPU設(shè)備標(biāo)稱核心數(shù)目，線程塊大小、匹配循環(huán)次數(shù)最好按表1所列配對(duì)方式選取，具體的實(shí)驗(yàn)結(jié)果參見(jiàn)下一部分內(nèi)容。

表1 主要設(shè)計(jì)參數(shù)配對(duì)表

CLAHE部分的CUDA優(yōu)化加速是分別對(duì)直方圖生成、直方圖裁剪、直方圖再均衡環(huán)節(jié)進(jìn)行，將圖像各個(gè)分組箱分塊內(nèi)的像元灰階數(shù)、分組箱分塊內(nèi)超出限定灰度閾值范圍的灰階數(shù)目使用共享存儲(chǔ)器實(shí)現(xiàn)數(shù)據(jù)共享，實(shí)現(xiàn)圖像像素點(diǎn)For循環(huán)累加計(jì)數(shù)運(yùn)算CUDA并行化，達(dá)到降低算法開(kāi)銷(xiāo)、優(yōu)化加速的目的。

3 結(jié)果與分析

針對(duì)上述兩種加速方案開(kāi)展對(duì)比實(shí)驗(yàn)。針對(duì)圖5(a)所示的霧靄圖像進(jìn)行去霧處理，處理后的圖像增強(qiáng)效果見(jiàn)圖5(b)。近處樓房、遠(yuǎn)處樓房、天空均得到對(duì)比度增強(qiáng)，去霧效果顯著。進(jìn)行方案一加速和方案二加速后，圖像匹配歸一化算法部分GPU加速后的性能指標(biāo)、OpenMP多核CPU加速后的性能指標(biāo)分別見(jiàn)表2、表3；CLAHE算法部分GPU加速后性能指標(biāo)、OpenMP多核CPU加速后性能指標(biāo)分別見(jiàn)表4、表5。

(a) 原始霧靄圖像(a) Origin foggy image

(b) 去霧靄效果圖(b) Foggy image processing effect show圖5 霧靄圖像處理效果Fig.5 Foggy image processing effect

表2 NCC算法GPU加速后性能指標(biāo)

表3 NCC算法OpenMP多核CPU加速后性能指標(biāo)Tab.3 Performance indicators of NCC algorithm with multicore CPU OpenMP acceleration

表4CLAHE算法GPU加速后性能指標(biāo)

Tab.4 Performance indicators of CLAHE with GPU acceleration

表5 CLAHE算法OpenMP多核CPU加速后性能指標(biāo)Tab.5 Performance indicators of CLAHE with multicore CPU OpenMP acceleration

查看表2中加速比為1.37和1.86的兩個(gè)數(shù)據(jù)發(fā)現(xiàn)，選用的線程塊尺寸大小不合適(圖像大小不能整除線程塊尺寸大小)會(huì)造成加速比過(guò)低，導(dǎo)致效果接近表3中OpenMP多核CPU加速效果。為提高加速比，在GPU CUDA加速實(shí)現(xiàn)策略中避免選擇此種線程塊尺寸。剔除兩數(shù)據(jù)后的表2，圖像匹配 GPU優(yōu)化加速后較未使用加速策略前計(jì)算時(shí)間縮短至4.85%～17.11%，計(jì)算性能提高5～20倍,加速效果理想。對(duì)比表4、表5可以發(fā)現(xiàn)，CLAHE算法GPU優(yōu)化加速后時(shí)間可縮短至10.17%～14.32%，計(jì)算性能提高6～9倍。分析表2中圖像大小與加速比關(guān)系可發(fā)現(xiàn)，使用CUDA并行實(shí)現(xiàn)NCC加速時(shí)，選定線程塊大小后，存在實(shí)時(shí)圖與模板圖最優(yōu)尺寸配對(duì)，使加速比達(dá)到最大。

透過(guò)數(shù)據(jù)本身，進(jìn)一步分析可得到：線程塊大小設(shè)計(jì)對(duì)加速性能有很大影響，這進(jìn)一步驗(yàn)證了優(yōu)化設(shè)計(jì)的原理；當(dāng)圖像大小是線程塊尺寸大小的整數(shù)倍時(shí)，相比隨意的線程塊大小尺寸，計(jì)算時(shí)間更短，這是因?yàn)闇p少了總迭代次數(shù)，降低了算法開(kāi)銷(xiāo)；同時(shí)線程塊越小，GPU使用的核心數(shù)越多，加速性能越明顯。對(duì)于圖像匹配N(xiāo)CC算法，實(shí)時(shí)圖為64像素×64像素、模板圖為32像素×32像素、線程塊大小為8像素×8像素時(shí)，CUDA使用核心數(shù)為4×4×8=128<384(GPU核心數(shù))，并且中間成果數(shù)據(jù)融合效率最高，因而加速性能最高。當(dāng)實(shí)時(shí)圖為128像素×128像素、模板圖為64像素×64像素、線程塊大小為8像素×8像素時(shí)，CUDA使用核心數(shù)為8×8×8=512>384(GPU核心數(shù))，內(nèi)核函數(shù)計(jì)算被拆分成多次計(jì)算，加速性能有所減弱，但因分塊形式對(duì)于數(shù)據(jù)融合較為高效，也能最終獲得9.33倍的加速效果。

4 結(jié)論

1)充分利用共享寄存器的CUDA加速算法設(shè)計(jì)，可以降低GPU設(shè)備讀取數(shù)據(jù)次數(shù)，實(shí)現(xiàn)并行線程間數(shù)據(jù)共享，實(shí)現(xiàn)算法加速。單幅圖像CUDA加速算法線程塊尺寸設(shè)計(jì)對(duì)加速性能有較大影響，當(dāng)模板圖圖像大小是線程塊尺寸大小整數(shù)倍時(shí)，相比隨意的線程塊大小尺寸，計(jì)算時(shí)間更短；對(duì)于性能指標(biāo)固定的GPU，選定線程塊大小后，能實(shí)時(shí)處理的實(shí)時(shí)圖和模板圖存在最優(yōu)尺寸大小，可通過(guò)仿真實(shí)驗(yàn)方法進(jìn)行測(cè)試標(biāo)定。

2)從總體時(shí)間性能指標(biāo)方面來(lái)看，基于CUDA的霧靄圖像GPU并行處理算法計(jì)算速度較未進(jìn)行CUDA加速優(yōu)化前提高5～20倍，且GPU使用的核心數(shù)越多，加速性能越明顯；64像素×64像素實(shí)時(shí)圖與32像素×32像素模板圖匹配的計(jì)算時(shí)間僅為5.57 ms，達(dá)到實(shí)時(shí)仿真計(jì)算的要求，能夠用于目標(biāo)實(shí)時(shí)跟蹤性判斷檢測(cè)。

3)在解決單幅圖像并行處理需求方面，使用OpenMP實(shí)現(xiàn)多核CPU并行加速效果不如使用GPU CUDA加速明顯，OpenMP算法開(kāi)支時(shí)間較大，實(shí)時(shí)性較差。相比之下，OpenMP多核CPU并行更適合多幅圖像同時(shí)進(jìn)行處理，減少總體作業(yè)開(kāi)支時(shí)間。