莫德林，戴晨光，張振超，胡玲

（信息工程大學(xué)地理空間信息學(xué)院，鄭州450001）

一種基于OpenACC的遙感影像正射糾正快速實(shí)現(xiàn)方法

莫德林，戴晨光，張振超，胡玲

（信息工程大學(xué)地理空間信息學(xué)院，鄭州450001）

利用CUDA語(yǔ)言移植舊程序時(shí)需要重新設(shè)計(jì)算法，花費(fèi)較多的時(shí)間，效率不高。針對(duì)這一問(wèn)題，本文在分析正射糾正算法并行性的基礎(chǔ)上，提出一種基于OpenACC的遙感影像正射糾正快速實(shí)現(xiàn)方法，并與基于CUDA的正射糾正方法進(jìn)行對(duì)比。通過(guò)正射糾正實(shí)驗(yàn)表明，OpenACC能通過(guò)對(duì)源代碼的較小改動(dòng)將其移植到GPU中，獲得一定的加速比，其可移植性好，代碼開(kāi)發(fā)效率較高。

OpenACC；CUDA；正射糾正；加速比

1 引言

多核的CPU和眾核的GPU（Graphic Processing Unit，即圖形處理器）已經(jīng)成為目前大多數(shù)計(jì)算機(jī)中最重要的兩種處理器，科研人員積極尋找專(zhuān)業(yè)領(lǐng)域與GPU并行處理的結(jié)合點(diǎn)，對(duì)GPU并行計(jì)算進(jìn)行了有益的探索。數(shù)字正射影像是地球空間數(shù)據(jù)框架的一個(gè)基礎(chǔ)數(shù)據(jù)層［1］，被視為快速成圖與更新的重要手段。因此，提高數(shù)字正射影像的生產(chǎn)效率具有重要的作用。楊靖宇［2］提出一種基于CUDA（Compute Unified Device Architecture，統(tǒng)一計(jì)算設(shè)備架構(gòu)）的遙感影像正射糾正GPU-CPU協(xié)同處理方法，實(shí)現(xiàn)了重采樣操作的GPU細(xì)粒度并行化。然而，初次接觸GPU的科研人員往往需要較長(zhǎng)的時(shí)間才能完全掌握CUDA程序的編寫(xiě)與優(yōu)化［3］。為此，PGI、Cray和NVIDIA3家公司聯(lián)合創(chuàng)立OpenACC應(yīng)用編程接口標(biāo)準(zhǔn)。OpenACC使得科研人員能夠更加自由地將時(shí)間投入到自己的研究中，而又能在更短的時(shí)間內(nèi)使程序得到加速。本文實(shí)現(xiàn)了一種基于OpenACC的遙感影像正射糾正方法，通過(guò)將計(jì)算密集區(qū)域加載到GPU上進(jìn)行并行計(jì)算，提高了算法的執(zhí)行效率。

2 OpenACC介紹

OpenACC借鑒OpenMP的導(dǎo)語(yǔ)（directive）模式，在原始代碼上添加導(dǎo)語(yǔ)，告訴編譯器將哪些代碼塊加載到加速器上執(zhí)行、如何在主機(jī)與加速器之間移動(dòng)數(shù)據(jù)［4］。如果編譯器不支持OpenACC標(biāo)準(zhǔn)或支持選項(xiàng)沒(méi)有打開(kāi)，編譯器將忽略所有的OpenACC導(dǎo)語(yǔ)，編譯出的程序只在CPU上運(yùn)行。OpenACC程序不必對(duì)原始代碼作過(guò)多的修改，可以讓相同的代碼在多核CPU、GPU或任何編譯器支持的其他類(lèi)型的并行硬件上運(yùn)行［5］。

在C/C++語(yǔ)言中，用語(yǔ)言本身提供的#pragam機(jī)制來(lái)引入OpenACC導(dǎo)語(yǔ)。OpenACC導(dǎo)語(yǔ)在C/C++語(yǔ)言中的語(yǔ)法是：

#pragma acc directive-name［clause［［,］clause］…］new-line

OpenACC導(dǎo)語(yǔ)包含四個(gè)部分［5］：

①數(shù)據(jù)管理

數(shù)據(jù)管理導(dǎo)語(yǔ)包括data構(gòu)件和update導(dǎo)語(yǔ)。data構(gòu)件指明的標(biāo)量、數(shù)組和子數(shù)組都會(huì)在加速器內(nèi)存上開(kāi)辟空間。進(jìn)入本區(qū)域時(shí)，數(shù)據(jù)被從主機(jī)復(fù)制到加速器內(nèi)存，離開(kāi)本區(qū)域時(shí)，數(shù)據(jù)被從設(shè)備復(fù)制到主機(jī)內(nèi)存。update導(dǎo)語(yǔ)用在顯式或隱式數(shù)據(jù)區(qū)域中，使加速器內(nèi)存中數(shù)組的值和主機(jī)內(nèi)存中相應(yīng)數(shù)組的值相互傳遞。

②工作管理

工作管理導(dǎo)語(yǔ)包括parallel構(gòu)件、kernel構(gòu)件和loop構(gòu)件。OpenACC提供兩個(gè)計(jì)算構(gòu)件（parallel和kernels），并能詳細(xì)指定并行方式（gang、worker、vector的值，對(duì)應(yīng)于CUDA的grid、block、thread的值），通過(guò)編譯器將循環(huán)轉(zhuǎn)換為高效的低層級(jí)語(yǔ)言代碼（CUDA/OpenCL）。loop構(gòu)件可以描述執(zhí)行這個(gè)循環(huán)的并行類(lèi)型，還可以聲明循環(huán)的私有變量、數(shù)組和歸約操作。

③其他語(yǔ)法

其他語(yǔ)法包括cache構(gòu)件、host_data構(gòu)件、wait導(dǎo)語(yǔ)和declare構(gòu)件。Cache構(gòu)件指定哪些數(shù)組元素或子數(shù)組需要為循環(huán)體而預(yù)取到最高層級(jí)的緩存中，host_data構(gòu)件使加速器上數(shù)據(jù)的地址在主機(jī)上可用，wait導(dǎo)語(yǔ)使主機(jī)等待一個(gè)異步活的完成declare構(gòu)件指定的變量或數(shù)組需要在加速器內(nèi)存上開(kāi)辟空間。

④運(yùn)行時(shí)例程

OpenACC提供多個(gè)運(yùn)行時(shí)例程，可以設(shè)置加速器設(shè)備的參數(shù)，如例程acc_set_device_type告訴運(yùn)行時(shí)環(huán)境使用哪種類(lèi)型的設(shè)備來(lái)執(zhí)行加速器parallel區(qū)域和kernels區(qū)域。

3 基于OpenACC的正射糾正

數(shù)字微分糾正中，一般是利用反解公式求解坐標(biāo)變換系數(shù)，計(jì)算對(duì)應(yīng)像元的坐標(biāo)，然后進(jìn)行灰度重采樣，最后將重采樣后的灰度值賦值給糾正后的像元。影像灰度重采樣操作是典型的計(jì)算密集型模塊，其處理流程相對(duì)固定，每個(gè)數(shù)據(jù)點(diǎn)上的計(jì)算形式相同，數(shù)據(jù)點(diǎn)之間相互獨(dú)立，具有內(nèi)在的并行性，因此非常適合GPU并行處理。由于遙感影像的數(shù)據(jù)量一般都比較大，相對(duì)于坐標(biāo)變換系數(shù)的求解，重采樣操作將耗費(fèi)更多的時(shí)間，所以正射糾正并行化的重點(diǎn)應(yīng)該是重采樣操作的并行化［3］。正射糾正的流程如圖1所示。

圖1中虛線框內(nèi)是并行計(jì)算區(qū)域。要將并行計(jì)算區(qū)域加載到加速器上進(jìn)行計(jì)算，必須先設(shè)定環(huán)境以及將數(shù)據(jù)復(fù)制到加速器內(nèi)存上。

圖1 正射糾正流程圖

基于OpenACC的正射糾正步驟如下：

①環(huán)境設(shè)定；

//設(shè)定執(zhí)行并行計(jì)算區(qū)域的設(shè)備類(lèi)型

void acc_set_device_type(acc_device_t);

②使用data構(gòu)件將原始影像、結(jié)果影像和DEM數(shù)據(jù)復(fù)制進(jìn)顯存；

#pragma acc data copyin(SourceImage［0:width*height］)

③使用kernels指令定義核心；

#pragma acc kernels present_or_copyin (SourceImage［0:width*height］){

④使用loop指令定義外層循環(huán)；

#pragma acc loop independent

for(….){

⑤使用loop指令定義內(nèi)層循環(huán)；

#pragma acc loop independent

for(…){

//計(jì)算坐標(biāo)變換系數(shù)、灰度重采樣以及灰度賦值，同時(shí)把結(jié)果放入目標(biāo)內(nèi)存

}}}

⑥將結(jié)果復(fù)制到主機(jī)端，釋放空間。

//#pragma acc update host();

void acc_shutdown(acc_device_t);//斷開(kāi)程序與加速器設(shè)備的連接

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

本文使用的實(shí)驗(yàn)平臺(tái)CPU為Intel(R)Core(TM) i5，內(nèi)存大小為4GB，顯卡為NVIDIA Quadro FX 3700。實(shí)驗(yàn)平臺(tái)的軟件開(kāi)發(fā)環(huán)境為：Windows7 32位操作系統(tǒng)，PGI編譯器；NVIDIA公司提供的CUDA 5.0版本的開(kāi)發(fā)包。

本文使用的實(shí)驗(yàn)數(shù)據(jù)為2009年獲取的河南登封某地區(qū)DMC航攝影像，地面采樣間隔為0.25米。正射糾正所用的DEM數(shù)據(jù)為影像多視匹配獲得的點(diǎn)云數(shù)據(jù)經(jīng)過(guò)濾波獲得，將其內(nèi)插為0.25米規(guī)則格網(wǎng)數(shù)據(jù)。表1為不同范圍大小的影像基于CUDA與基于OpenACC的正射糾正時(shí)間對(duì)比結(jié)果。

由表1可見(jiàn)，基于CUDA和基于OpenACC的并行算法與基于CPU的串行算法相比，都能獲得一定的加速比，且加速比隨著數(shù)據(jù)量的增大而增大。在數(shù)據(jù)量較小時(shí)，加速效果并不明顯。這是因?yàn)榛贑UDA和基于OpenACC的并行計(jì)算過(guò)程中，要在主機(jī)和加速器之間進(jìn)行數(shù)據(jù)傳遞，這將耗費(fèi)一定的時(shí)間資源，如果數(shù)據(jù)量較小，則在加速器上運(yùn)算時(shí)間與數(shù)據(jù)傳遞時(shí)間之比也小，加速效果則不明顯。當(dāng)數(shù)據(jù)量增大，在加速器上運(yùn)算滿載時(shí)，運(yùn)算時(shí)間與數(shù)據(jù)傳遞時(shí)間之比合理，則能達(dá)到最大的加速比。

CUDA程序經(jīng)過(guò)優(yōu)化后，其加速比較OpenACC大。其原因是CUDA程序經(jīng)過(guò)存儲(chǔ)器優(yōu)化后，可大大減少存儲(chǔ)訪問(wèn)的時(shí)間。在OpenACC中，數(shù)據(jù)傳遞與訪問(wèn)是隱式進(jìn)行的，無(wú)法使用類(lèi)似CUDA的共享內(nèi)存，所以其訪存時(shí)間要比CUDA程序長(zhǎng)。然而，從代碼開(kāi)發(fā)的效率方面看，OpenACC程序只需要在串行代碼中增加OpenACC指令，就能將指定的并行區(qū)域加載到GPU上，取得10倍左右的加速比。這種方式保留了代碼的通用性，不破壞原代碼，開(kāi)發(fā)速度快，既可并行執(zhí)行又可恢復(fù)串行執(zhí)行。而且，在硬件升級(jí)時(shí)，重新編譯一次代碼即可，不必手工修改代碼。更為重要的是，隨著編譯器的進(jìn)一步優(yōu)化和硬件技術(shù)的發(fā)展，OpenACC與CUDA在底層技術(shù)實(shí)現(xiàn)上的差距將會(huì)越來(lái)越小［6］，而且支持OpenACC指令的設(shè)備將不僅僅限于CUDA設(shè)備，還將擴(kuò)展到其他更多廠商的硬件加速器，從而提高OpenACC程序的可移植性。

表1 GPU粗粒度并行下灰度重采樣時(shí)間結(jié)果表

5 結(jié)束語(yǔ)

本文以正射糾正算法為對(duì)象，通過(guò)在串行代碼的基礎(chǔ)上加入OpenACC指令，將坐標(biāo)計(jì)算、灰度重采樣計(jì)算以及灰度賦值區(qū)域加載到GPU上，得到較高的加速比，提高了算法的執(zhí)行效率。這種方法不僅效率高，且不改變?cè)瓉?lái)的代碼結(jié)構(gòu)，大大增強(qiáng)了代碼的可移植性，對(duì)其他的影像處理算法有一定的參考價(jià)值。

［1］耿則勛，張保明，范大昭.數(shù)字?jǐn)z影測(cè)量學(xué)［M］.北京：測(cè)繪出版社，2010.8.

［2］楊靖宇，張永生，李正國(guó)等.遙感影像正射糾正的GPU-CPU協(xié)同處理研究［J］.武漢大學(xué)學(xué)報(bào)（信息科學(xué)版），2011，36（9）：1043-1046.

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

［4］Tetsuya Hoshino,Naoya Maruyama,Satoshi Matsuoka. CUDA vs OpenACC:Performance Case Studies with Kernel Benchmarks and a Memory-Bound CFD Application［C］. 13th IEEE/ACM International Symposium on Cluster, Cloud,and Grid Computing,2013(136-143).

［5］"The OpenACCApplication Programming Interface,Version 1.0,"November 2011.

［6］曾文權(quán)，胡玉貴，何擁軍等.一種基于OPENACC的GPU加速實(shí)現(xiàn)高斯模糊算法［J］.計(jì)算機(jī)技術(shù)與發(fā)展，2013，23（7）：147-151.

A Fast Implementation M ethod of Remote Sensing Image Ortho-rectification Based on OpenACC

MO De-lin,DAIChen-guang,ZHANG Zhen-chao,HU Ling
(Institude of Geospatial Information,Information Engineering University,Zhengzhou 450001,China)

Algorithms need to be redesigned when using CUDA to transplant old programs.This will cost a long period of time,which leads to a low efficiency.Aiming at solving this problem,a fast implementationmethod of remote sensing image ortho-rectification based on OpenACC is proposed,and compared with ortho-rectification based on CUDA in this paper.Ortho-rectification experiments show that OpenACC can transplant program to the GPU with small changes,and get a speed up ratio.The portability is good,and the efficiency of code development is high.

OpenACC;CUDA;Ortho-rectification;Accelerated ratio

P231

B

10.3969/j.issn.1001-0270.2014.02.21

2014-01-20

莫德林（1988-），男，廣西貴港人，碩士研究生，主要從事遙感圖像處理的研究。