基于CUDA技術(shù)圖像旋轉(zhuǎn)

1、武漢理工大學(xué)學(xué)士學(xué)位論文目 錄摘 要IAbstractII1 緒論12 CUDA技術(shù)概況42.1 CUDA架構(gòu)42.2 CUDA的硬件模型52.3 CUDA軟件環(huán)境和編程模型62.4 CUDA開(kāi)發(fā)平臺(tái)的搭建102.4.1軟件安裝和環(huán)境配置102.4.2創(chuàng)建工程113 基于CUDA的圖像旋轉(zhuǎn)實(shí)現(xiàn)143.1 實(shí)現(xiàn)圖像旋轉(zhuǎn)的方案選擇143.1.1適用于SPMD計(jì)算的圖像變換算法143.1.2 紋理技術(shù)153.2 基于CUDA的數(shù)字圖像幾何變換方案183.3 基于CUDA的數(shù)字圖像幾何變換的方案實(shí)現(xiàn)193.3.1 CUDA初始化193.3.2 紋理參數(shù)的拾取和坐標(biāo)變換193.3.3 數(shù)據(jù)傳送至GPU并

2、由GPU完成并行化核心過(guò)程203.3.4 數(shù)據(jù)回傳203.4 實(shí)驗(yàn)結(jié)果和分析214 結(jié)論與展望23參考文獻(xiàn)2426摘 要CUDA 是GPU 通過(guò)并發(fā)執(zhí)行多個(gè)線程以實(shí)現(xiàn)大規(guī)?？焖俨⑿杏?jì)算能力的技術(shù)，它能使對(duì)GPU 編程變得更容易。介紹了CUDA 基本特性及主要編程模型，在此基礎(chǔ)上,提出并實(shí)現(xiàn)了基于NVIDIA CUDA技術(shù)的圖像快速幾何變換。采用CUDA紋理技術(shù)，并把CUDA 技術(shù)的快速并行計(jì)算能力應(yīng)用到數(shù)字圖像幾何變換中，解決了基于CPU 的傳統(tǒng)圖像幾何變換運(yùn)算效率低下的問(wèn)題。實(shí)驗(yàn)結(jié)果證明使用CUDA 技術(shù)，隨著處理圖像尺寸的增加，對(duì)數(shù)字圖像幾何變換處理效率最高能夠提高到近100 倍。關(guān)鍵字：

3、CUDA; 并行化; 數(shù)字圖像; GPU 編程;紋理；AbstractCUDA is a type of technology that performs general purposes fast and parallel computation by running tens of thousands of threads concurrently. It makes users develop general GPU programs easily. This paper analyzes the distinct features of CUDA and summarizes the

4、general program mode of CUDA. Furthermore,it proposes and implements an image geometric transform process by the CUDA. By using the texture technology, and by applying the rapid parallel computing ability of CUDA to transform the geometric image, solve the problem of the inefficient image transform

5、based on CPU. The experimental results show that, as the image size increase, application of CUDA on image geometric transform can reach to more than 100x.Keywords: CUDA; parallelization; digital image; GPU programming;Texture;1 緒論l 研究背景在某些如車牌識(shí)別系統(tǒng)等對(duì)實(shí)時(shí)性要求較高的有大量圖像處理工作的系統(tǒng)開(kāi)發(fā)中，對(duì)從攝像設(shè)備中獲取的原始圖像的預(yù)處理所耗費(fèi)的時(shí)間占了很

6、大比重，如何最大限度地減少圖像預(yù)處理所耗費(fèi)的時(shí)間成為提高這類實(shí)時(shí)系統(tǒng)性能的關(guān)鍵。圖像幾何變換是預(yù)處理的一個(gè)重要方面，并行化是解決這類問(wèn)題的一個(gè)思路。近年來(lái)計(jì)算機(jī)圖形處理器(Graphic ProcessingUnit,GPU)高速發(fā)展，極大地提高了計(jì)算機(jī)圖形處理的速度和圖形質(zhì)量,同時(shí)也具有高度的可程序化能力，這使得通過(guò)GPU 實(shí)現(xiàn)數(shù)字圖像幾何變換并行化成為可能。由于顯示芯片通常具有相當(dāng)高的內(nèi)存帶寬，以及大量的執(zhí)行單元，繪制流水線的高速度和并行性為圖形處理以外的通用計(jì)算提供了良好的運(yùn)行平臺(tái),這使得基于GPU 的通用計(jì)算成為近幾年人們關(guān)注的一個(gè)研究熱點(diǎn)。 并行計(jì)算已成為突破摩爾定理局限性的重要研究

7、方向，而GPU強(qiáng)大的并行計(jì)算能力也因此吸引了全球廣泛的研究興趣。然后，在時(shí)間通用并行計(jì)算時(shí)，GPU計(jì)算模式存在著一些限制。首先，GPU的設(shè)計(jì)初衷是為了加速程序中的圖形繪制運(yùn)算，因此開(kāi)發(fā)人員需要通過(guò)OpenGL或者API來(lái)訪問(wèn)GPU,這不僅要求開(kāi)發(fā)人員掌握一定的圖形編程知識(shí)，而且要想方設(shè)法將通用計(jì)算問(wèn)題轉(zhuǎn)為圖形計(jì)算問(wèn)題。其次，GPU與多核CPU在計(jì)算構(gòu)架上有著很大不同，GPU更注重于數(shù)據(jù)并行計(jì)算，即在不同的數(shù)據(jù)上并行執(zhí)行相同的計(jì)算，而對(duì)并行計(jì)算中的互斥性，同步性以及原子性等方面支持不足，這些因素都限制了GPU在通用并行計(jì)算中的引用范圍。CUDA架構(gòu)的出現(xiàn)解決了上述問(wèn)題。CUDA 是NVIDIA&

8、#174; 公司的并行計(jì)算架構(gòu)。該架構(gòu)通過(guò)利用 GPU 的處理能力，可大幅提升計(jì)算性能。目前為止基于 CUDA的GPU銷量已達(dá)數(shù)以百萬(wàn)計(jì)，軟件開(kāi)發(fā)商、科學(xué)家以及研究人員正在各個(gè)領(lǐng)域中運(yùn)用CUDA，其中包括圖像與視頻處理、計(jì)算生物學(xué)和化學(xué)、流體力學(xué)模擬、CT圖像再現(xiàn)、地震分析以及光線追蹤等等。計(jì)算正在從CPU中央處理向CPU與GPU協(xié)同處理的方向發(fā)展。為了實(shí)現(xiàn)這一新型計(jì)算模式，NVIDIA發(fā)明了CUDA 并行計(jì)算架構(gòu)。該架構(gòu)現(xiàn)在正運(yùn)用于 Tesla® 、Quadro® 以及 GeForce® GPU 上。對(duì)應(yīng)用程序開(kāi)發(fā)商來(lái)說(shuō)，CUDA 架構(gòu)擁有龐大的用戶群。在科學(xué)研

9、究領(lǐng)域，CUDA受到狂熱追捧。例如，CUDA能夠加快AMBER 這款分子動(dòng)力學(xué)模擬程序的速度。全球有6萬(wàn)余名學(xué)術(shù)界和制藥公司的科研人員使用該程序來(lái)加速新藥開(kāi)發(fā)。在金融市場(chǎng)，Numerix和CompatibL 已宣布在一款對(duì)手風(fēng)險(xiǎn)應(yīng)用程序中支持 CUDA，而且因此實(shí)現(xiàn)了18倍速度提升。在 GPU 計(jì)算領(lǐng)域中，Tesla GPU的大幅增長(zhǎng)說(shuō)明了CUDA 正被人們廣泛采用。 目前，全球財(cái)富五百?gòu)?qiáng)企業(yè)已經(jīng)安裝了700多個(gè)GPU集群，從能源領(lǐng)域中的斯倫貝謝和雪佛龍到銀行業(yè)中的法國(guó)巴黎銀行，這些企業(yè)的范圍十分廣泛。l CUDA的優(yōu)勢(shì)CUDA(Computer Unified DeviceArchitect

10、ure)是NVIDIA 的GPU 編程模型，它的編程模式是單程序多重?cái)?shù)據(jù)(Single Program MultipleData,SPMD) ,即多個(gè)并發(fā)的線程執(zhí)行單一的程序來(lái)處理多重?cái)?shù)據(jù)1。本文提出的基于CUDA 的圖像并行幾何變換并不僅是將傳統(tǒng)的基于CPU 的圖像幾何變換算法簡(jiǎn)單機(jī)械移植到具有并行計(jì)算能力的GPU 中。通過(guò)對(duì)傳統(tǒng)的數(shù)字圖像幾何變換的分析,盡量減少原算法中的冗余乘法運(yùn)算，然后利用CUDA 技術(shù)并發(fā)執(zhí)行的優(yōu)點(diǎn),并發(fā)執(zhí)行多個(gè)線程對(duì)圖像的像素點(diǎn)進(jìn)行同步幾何變換，因此大幅提高了變換效率?；贑UDA計(jì)算的優(yōu)點(diǎn)：GPU通常具有更大的內(nèi)存帶寬。GPU具有更多的執(zhí)行單元。GPU較同等級(jí)的C

11、PU具有價(jià)格優(yōu)勢(shì)。在CUDA上實(shí)現(xiàn)的應(yīng)用具有一個(gè)共同的特征，就是個(gè)CUDA線程所處理的數(shù)據(jù)相互之間獨(dú)立性很高，線程間通信很少或者沒(méi)有。每個(gè)線程所處理的數(shù)據(jù)被稱為單位數(shù)據(jù)，是一段內(nèi)存空間，可以是1bit，或者1Byte，或者一個(gè)數(shù)據(jù)或者字符串。我們將這種單位數(shù)據(jù)間能夠被并行處理的性質(zhì)稱為數(shù)據(jù)并行性。處理的數(shù)據(jù)具有數(shù)據(jù)并行性的應(yīng)用稱為數(shù)據(jù)并行應(yīng)用。l 研究?jī)?nèi)容和目的本文介紹了CUDA基本特性以及主要編程模型，在此基礎(chǔ)上，提出并實(shí)現(xiàn)了基于NVIDIA CUDA技術(shù)的數(shù)字圖像快速幾何變換，采用CUDA的紋理技術(shù)進(jìn)行數(shù)字圖像的幾何變換，代替原變換中大量乘法運(yùn)算，并把CUDA技術(shù)的快速并行計(jì)算能力應(yīng)用到數(shù)

12、字圖像幾何變換中，解決了基于CPU的傳統(tǒng)圖像幾何變換運(yùn)算效率低下的問(wèn)題。l CUDA的應(yīng)用領(lǐng)域由于GPU的特點(diǎn)是處理密集型數(shù)據(jù)和并行數(shù)據(jù)計(jì)算，因此CUDA非常適合需要大規(guī)模并行計(jì)算的領(lǐng)域。目前CUDA平臺(tái)支持C語(yǔ)言和FORTRAN的開(kāi)發(fā)，CUDA平臺(tái)以后預(yù)計(jì)將進(jìn)一步支持Python、C+、Java等語(yǔ)言。CUDA技術(shù)已經(jīng)成功應(yīng)用在下列涉及大規(guī)模密集型數(shù)據(jù)計(jì)算的領(lǐng)域：金融、電信、證券分析；能源勘探開(kāi)發(fā)；搜索引擎的應(yīng)用；數(shù)據(jù)庫(kù)和數(shù)據(jù)挖掘；醫(yī)藥工程等。在未來(lái)，隨著CUDA平臺(tái)的繼續(xù)發(fā)展，CUDA技術(shù)將應(yīng)用于更多的領(lǐng)域。1. CUDA應(yīng)用舉例：實(shí)時(shí)的裸眼立體醫(yī)療成像系統(tǒng)裸眼立體成像技術(shù)是成像技術(shù)中一個(gè)

13、很有趣的技術(shù)應(yīng)用，裸眼立體成像技術(shù)不需要使用者使用特殊的眼鏡就能讓使用者看到三維的立體圖像。裸眼立體成像技術(shù)可以和其它領(lǐng)域的技術(shù)結(jié)合，實(shí)現(xiàn)其它領(lǐng)域內(nèi)的重要應(yīng)用。特別是這種技術(shù)在娛樂(lè)方面有著巨大的應(yīng)用潛力。這種技術(shù)涉及到特別巨大的計(jì)算量，東京大學(xué)的Takeyoshi教授認(rèn)為NVDIA的CUDA平臺(tái)可成功實(shí)現(xiàn)裸眼立體成像技術(shù)應(yīng)用于醫(yī)療成像。并且東京大學(xué)的研究小組在2000年時(shí)已經(jīng)成功開(kāi)發(fā)出一種應(yīng)用裸眼立體成像的醫(yī)療成像系統(tǒng)，該系統(tǒng)使用CT和MRI技術(shù)實(shí)時(shí)掃描得到的活體截面圖作為被體紋理，通過(guò)進(jìn)一步體繪制實(shí)現(xiàn)為裸眼立體成像的三維圖像。系統(tǒng)涉及的計(jì)算量極其巨大，并且需要在極短的時(shí)間下高精度地實(shí)現(xiàn)這樣的

14、計(jì)算。2. CUDA應(yīng)用舉例：VMD/NAMD分子動(dòng)力學(xué)NAMD是一個(gè)重要的醫(yī)學(xué)應(yīng)用程序類別。分子動(dòng)力學(xué)是一種用于生物學(xué)研究的基本工具。之前，伊力諾依州立大學(xué)的研究人員在它們的網(wǎng)站上發(fā)布了一款GPU加速版的視覺(jué)分子動(dòng)力學(xué)工具,而最近，又發(fā)布了一款納米級(jí)分子動(dòng)力學(xué)工具。該工具可以從他們的網(wǎng)站上免費(fèi)下載。NAMD MD是將并行計(jì)算應(yīng)用于生物學(xué)中的基本工具的一個(gè)典型案例.3. CUDA應(yīng)用舉例：視頻轉(zhuǎn)換加速利用CUDA實(shí)現(xiàn)GPU計(jì)算來(lái)為應(yīng)用程序提速，Badaboom就是很好的一例，這是一款CUDA開(kāi)發(fā)的視頻轉(zhuǎn)換軟件，可以把mpeg2的視頻轉(zhuǎn)換為ipod或者iphone這樣的所使用的H.264視頻格式

15、。選取一段碼率較高的MPEG2視頻，可以看到GTX 280的處理速度達(dá)到了80FPS左右，如果碼率較小，還可以達(dá)到100FPS以上甚至更高。328MB的MPEG2視頻轉(zhuǎn)換成17.4MB的iPhone可用的MP4視頻（640*365），只用了37s。而同樣平臺(tái)下用CPU進(jìn)行計(jì)算，得到的結(jié)果是耗時(shí)107s，幾乎是用GTX 280轉(zhuǎn)換耗時(shí)的三倍。4. CUDA應(yīng)用舉例：FoldinghomeFoldinghome是美國(guó)史丹佛大學(xué)推動(dòng)的分散式運(yùn)算計(jì)劃，目的在于使用聯(lián)網(wǎng)式的計(jì)算方式和大量的分布式計(jì)算能力來(lái)模擬蛋白質(zhì)折疊的過(guò)程，并指引對(duì)由折疊引起的疾病的一系列研究。NVIDIA的CUDA顯卡也加入了Fol

16、dinghome計(jì)劃，目前全球有8000萬(wàn)塊CUDA顯卡，平均擁有100GFLOPS的浮點(diǎn)運(yùn)算能力，如果這其中有0.1%參與FoldingHome，就能夠?yàn)樵撚?jì)劃帶來(lái)7PFLOPS的運(yùn)算能力。2 CUDA技術(shù)概況2.1 CUDA架構(gòu)GPU在速度上之所以會(huì)有如此巨大的變革主要是因?yàn)樗旧淼脑O(shè)計(jì)原理。GPU使用更多的晶體管來(lái)用于數(shù)據(jù)處理，其數(shù)量遠(yuǎn)大于數(shù)據(jù)緩存和流控制。GPU特別適合處理那些并行數(shù)據(jù)計(jì)算一即同一程序并行地在許多數(shù)據(jù)元素上執(zhí)行。這樣一來(lái)，它對(duì)于流控制的要求就不是很高，而且運(yùn)算將代替大數(shù)據(jù)緩存來(lái)隱藏內(nèi)存存取延遲。數(shù)據(jù)元素通過(guò)數(shù)據(jù)并行處理機(jī)制映射到每個(gè)并行處理線程上。許多處理大數(shù)據(jù)集的應(yīng)用

17、程序，如數(shù)組，可以通過(guò)使用一個(gè)數(shù)據(jù)并行處理模型來(lái)加速計(jì)算。在3D渲染中，像素和頂點(diǎn)集可以映射到并行線程。同樣，圖形圖像處理程序，如視頻編解碼，圖像縮放和模式識(shí)別等，可以映射塊和像素到并行處理線程上。事實(shí)上，除了圖像渲染和處理領(lǐng)域，其他領(lǐng)域的許多計(jì)算也可以通過(guò)數(shù)據(jù)并行處理得以加速，例如信號(hào)處理，物理仿真，金融工程以及生物信息學(xué)。然而，直到現(xiàn)在，GPU要實(shí)現(xiàn)這樣的應(yīng)用還是存在許多困難的：1)GPU只能通過(guò)一個(gè)圖形的API來(lái)編程，這不僅加重了學(xué)習(xí)負(fù)擔(dān)更造2)由于DRAM內(nèi)存帶寬，一些程序會(huì)遇到瓶頸。3)無(wú)法在DRAM上進(jìn)行通用寫操作。CUDA(Compute Unified Device Archi

18、tecture，計(jì)算機(jī)統(tǒng)一設(shè)備架構(gòu))是一種新的處理和管理GPU計(jì)算的硬件和軟件架構(gòu)，它將GPU視作一個(gè)數(shù)據(jù)并行計(jì)算沒(méi)備，并且無(wú)需把這些計(jì)算映射到圖形API。GeForce 8，Tesla解決方案和一些Quadro解決方案已經(jīng)提供了CUDA。操作系統(tǒng)的多任務(wù)機(jī)制可以同時(shí)管理CUDA訪問(wèn)GPU和圖形程序的運(yùn)行時(shí)刻。CUDA在軟件方面組成有：一個(gè)硬件驅(qū)動(dòng)，一個(gè)應(yīng)用程序接口和運(yùn)行時(shí)，兩個(gè)通用算術(shù)庫(kù)一CUFFT和CUBI。AS。CUDA改進(jìn)了DRAM的讀寫靈活性，使得GPU與CPU的機(jī)制相吻合。另一方面，CUDA提供了片上共享內(nèi)存，使得線程之間可以共享數(shù)據(jù)。應(yīng)用程序可以利用共享內(nèi)存來(lái)減少DRAM的數(shù)據(jù)傳

19、送，更少的依賴DRAM的內(nèi)存帶寬。在之前的圖形處理架構(gòu)中，計(jì)算資源劃分為頂點(diǎn)著色器和像素著色器，而CUDA架構(gòu)則不同，它包含了一個(gè)統(tǒng)一的著色器流水線，使得執(zhí)行通用計(jì)算的程序能夠?qū)π酒系拿總€(gè)數(shù)學(xué)邏輯單元（ALU）進(jìn)行排列。由于NVIDIA希望使新的圖形處理器能適用于通用計(jì)算，因此在實(shí)現(xiàn)這些ALU時(shí)都確保它們能夠滿足IEEE單精度浮點(diǎn)數(shù)學(xué)運(yùn)算的需求。此外，GPU上的執(zhí)行單元不僅能任意地讀/寫內(nèi)存，同事還能訪問(wèn)由軟件管理的緩存，也稱為共享內(nèi)存。CUDA架構(gòu)的所有這些功能都是為了使GPU不僅能執(zhí)行傳統(tǒng)的圖形計(jì)算，還能高效的執(zhí)行通用計(jì)算。NVIDIA并不局限于通過(guò)幾成CUDA架構(gòu)的硬件來(lái)為消費(fèi)者同時(shí)提

20、供計(jì)算功能和圖形功能。機(jī)關(guān)NVIDIA的芯片中增加了許多功能來(lái)加速計(jì)算，但任然只能通過(guò)OpenGL或者DirectX來(lái)訪問(wèn)這些功能。這不僅要求用戶任然要將他們的計(jì)算任務(wù)偽裝為圖形問(wèn)題，而且還需要使用面向圖形的著色語(yǔ)言來(lái)編寫計(jì)算代碼。為了盡可能地吸引更多的開(kāi)發(fā)人員，NVIDIA采取了工業(yè)標(biāo)準(zhǔn)的C語(yǔ)言，并且增加了一小部分關(guān)鍵字來(lái)支持CUDA架構(gòu)的特殊功能，在發(fā)布了了GeForce 8800 GTX之后的幾個(gè)月，NVIDIA公布了一款編譯器來(lái)編譯CUDA C語(yǔ)言，這樣，CUDA C就稱為了第一款專門有GPU公司設(shè)計(jì)的編程語(yǔ)言，用于在GPU上編寫通用計(jì)算。除了專門設(shè)計(jì)一種語(yǔ)言來(lái)為GPU編寫代碼以后，N

21、VIDIA還提供了專門的硬件驅(qū)動(dòng)程序來(lái)發(fā)揮CUDA架構(gòu)的大規(guī)模計(jì)算功能。現(xiàn)在，用戶不需要了解OpenGL或者DirectX圖形編程結(jié)構(gòu)，也不需要將通用計(jì)算問(wèn)題偽裝為圖形計(jì)算問(wèn)題。2.2 CUDA的硬件模型CUDA的硬件結(jié)構(gòu)模型如圖2.1，包括多個(gè)SIMD多處理器，每個(gè)多處理器具有4種類型的芯片內(nèi)存儲(chǔ)器：本地寄存器，并行數(shù)據(jù)告訴緩存，只讀常量高速緩存，只讀紋理高速緩存，后三種存儲(chǔ)器可以通過(guò)主機(jī)讀或者寫，并永久存在于響應(yīng)應(yīng)用程序的內(nèi)核啟動(dòng)中。圖2.1 CUDA的硬件結(jié)構(gòu)這個(gè)模型有以下幾個(gè)特點(diǎn)：1) 全局讀寫存儲(chǔ)器：GPU 能夠從任意位置和存儲(chǔ)器中獲取數(shù)據(jù)，也可以將數(shù)據(jù)寫入任何存儲(chǔ)器中，這幾乎和CP

22、U 一樣具有靈活性。2) 線程共享存儲(chǔ)器：它可以使得在同一個(gè)處理器中的線程快速獲取數(shù)據(jù)，避免從全局變量存儲(chǔ)器中存取。共享存儲(chǔ)器的訪問(wèn)速度幾乎寄存器一樣快，從這里獲取數(shù)據(jù)僅僅需要4 個(gè)時(shí)鐘周期，然而全局變量存儲(chǔ)器的訪問(wèn)花費(fèi)為400-600 個(gè)時(shí)鐘周期。3) 線程同步：處在同一個(gè)線程組的線程組能夠進(jìn)行同步操作，因此他們可以通迅和協(xié)作以解決一些復(fù)雜的問(wèn)題。2.3 CUDA軟件環(huán)境和編程模型CUDA的軟件體系由CUDA庫(kù)函數(shù)、運(yùn)行時(shí)API、驅(qū)動(dòng)組成，如圖2.2所示。在CUDA的軟件層面，NVIDIA C編譯器是其中的核心。CUDA程序是GPU和CPU的混合代碼， 它首先由NVIDIA C編譯器進(jìn)行編譯

23、。經(jīng)過(guò)編譯后，GPU和CPU的代碼將被分離，GPU代碼被編譯成GPU計(jì)算的機(jī)器碼，而CPU的C代碼輸出由標(biāo)準(zhǔn)的C編譯器進(jìn)行編譯。因此一個(gè)完整的CUDA軟件開(kāi)發(fā)環(huán)境還需要有一個(gè)面向CPU的C編譯器。CUDA可以支持多種運(yùn)行在Windows XP和Linux操作系統(tǒng)下的C開(kāi)發(fā)系統(tǒng)諸如Microsof Visual C+等。右圖是NVIDIAC編譯器結(jié)構(gòu)。其中，EDG將CPU和GPU的代碼分離；Open64生成GPU PTX(ParallelThread eXecution) 匯編碼。CUDA運(yùn)行需要CUDA runtime driver的支持，而Profiler則可以提供GPU和CPU kerne

24、l調(diào)用和內(nèi)存拷貝的時(shí)序分析，從而對(duì)性能進(jìn)行評(píng)估并且發(fā)現(xiàn)潛在的性能上的問(wèn)題。CUDA應(yīng)用程序CUDA庫(kù)函數(shù)CUDA運(yùn)行APICUDA驅(qū)動(dòng)GPU 圖2.2 CUDA軟件體系除了編譯器外，NVIDIA提供了一些非常實(shí)用的函數(shù)庫(kù)。目前有兩個(gè)數(shù)字計(jì)算庫(kù)包含在已經(jīng)發(fā)布的軟件包里面，分別是CUDA FFT和CUDA BLAS子程序庫(kù)。CUDA FFT是快速傅里葉變換的子程序庫(kù)，快色傅里葉變換時(shí)信號(hào)處理之類應(yīng)用的基本算法，BLAS是基本線性代數(shù)的子程序庫(kù)。CUDA FFT和BLAS都是針對(duì)GPU高度優(yōu)化的高性能數(shù)學(xué)函數(shù)庫(kù)，在CUDA程序中可以方便調(diào)用，節(jié)省大量的代碼編寫時(shí)間。CUDA程序構(gòu)架分為兩部分：Hos

25、t和Device。一般而言，Host指的是CPU，Device指的是GPU。在CUDA程序構(gòu)架中，主程序還是由CPU來(lái)執(zhí)行，而當(dāng)遇到數(shù)據(jù)并行處理的部分，CUDA就會(huì)將程序編譯成GPU能執(zhí)行的程序，并傳送到GPU。而這個(gè)程序在CUDA里叫做內(nèi)核。CPU主要負(fù)責(zé)要進(jìn)行邏輯性強(qiáng)的事務(wù)處理和串行運(yùn)算，GPU則專注于執(zhí)行高度線程化的并行處理任務(wù)。CPU、GPU各自擁有相互獨(dú)立的存儲(chǔ)器地址空間：主機(jī)端的內(nèi)存和設(shè)備端的顯存。CUDA對(duì)內(nèi)存操作與一般的C程序基本相同，但增加了一種可以被并行執(zhí)行的步驟。運(yùn)行在GPU上的CUDA并行計(jì)算函數(shù)稱為kernel（內(nèi)核函數(shù)）。一個(gè)kernel函數(shù)并不是一個(gè)完整的程序，而

26、是整個(gè)CUDA程序中的一個(gè)可以被并行執(zhí)行的步驟。CPU串行代碼完成的工作包括在kernel啟動(dòng)前進(jìn)行數(shù)據(jù)準(zhǔn)備和設(shè)備初始化的工作，以及在kernel之間進(jìn)行一些串行計(jì)算。理想情況是，CPU串行代碼的作用應(yīng)該只是清理上一個(gè)內(nèi)核函數(shù)，并啟動(dòng)下一個(gè)內(nèi)核函數(shù)。這樣，就可以在設(shè)備上完成盡可能多的工作，減少主機(jī)與設(shè)備之間的數(shù)據(jù)傳輸。在實(shí)際運(yùn)行時(shí)，CUDA會(huì)產(chǎn)生許多在GPU上執(zhí)行的線程，每一個(gè)線程都會(huì)去執(zhí)行內(nèi)核這個(gè)程序，雖然程序是同一份，但是因?yàn)樗饕煌?，而取得不同的?shù)據(jù)來(lái)進(jìn)行計(jì)算。在GPU中要執(zhí)行的線程，根據(jù)最有效的數(shù)據(jù)共享來(lái)創(chuàng)建塊(Block)，其類型有一維、二維或三維。在同一個(gè)塊里的線程，使用同一個(gè)共享

27、內(nèi)存。另外，塊的大小是有限制的，所以不能把所有的線程都塞到同一個(gè)塊里(一般GPGPU程序線程數(shù)目是很多的)；這時(shí)，可以用同樣維度和大小的塊(需同一個(gè)內(nèi)核)，來(lái)組成一個(gè)網(wǎng)格(Grid)做批處理(在同一個(gè)網(wǎng)格里的塊，GPU只需編澤一次)。具體框架如圖2.2所示。圖2.3 CUDA編程模型CUDA對(duì)C語(yǔ)言進(jìn)行了一些擴(kuò)展，允許定義稱為內(nèi)核（kernel）的C函數(shù)。調(diào)用內(nèi)核函數(shù)時(shí)，N個(gè)CUDA線程將內(nèi)核函數(shù)并行執(zhí)行N次。定義好使用_global_限定符聲明內(nèi)核函數(shù)，使用<<<>>>語(yǔ)法設(shè)定調(diào)用的線程數(shù)。我們看下面的CUDA程序例子：內(nèi)核函數(shù)的對(duì)應(yīng)執(zhí)行線程都內(nèi)置一個(gè)的線

28、程ID，所有的線程ID不會(huì)相同。我們可以使用內(nèi)置的threadIdx變量進(jìn)行訪問(wèn)。程序代碼實(shí)現(xiàn)把大小為N的向量A和B相加，結(jié)果存在向量C。執(zhí)行核函數(shù)VecAdd()的每一個(gè)線程都處理一次相加運(yùn)算。_global_void VecAdd(int*A,int*B,int*C)int i=threadIdx.x;Ci=Ai+Bi;int main()VecAdd<<<1,N>>>(A,B,C);l 線程層次結(jié)構(gòu)可以使用多維的索引對(duì)線程進(jìn)行標(biāo)識(shí)，組成多維的線程塊（thread block），核函數(shù)內(nèi)置的threadIdx為一個(gè)有3分量的向量。如下面的CUDA代碼示例

29、，大小為NxN矩陣A和B相加，結(jié)果存在矩陣C中。_global_void MatAdd(int ANN,int BNN,int CNN)int i=threadIdx.x;int j=threadIdx.y;Cij=Aij+Bij;int main()dim3 dimBlock(N,N);MatAdd<<<1,dimBlock>>>(A,B,C);同個(gè)塊內(nèi)線程可同步，使用共享存儲(chǔ)器（shared memory）實(shí)現(xiàn)共享數(shù)據(jù)。通過(guò)調(diào)用內(nèi)建函數(shù)_syncthreads(）實(shí)現(xiàn)在核函數(shù)中同步。同一個(gè)塊的所有的線程須都在同一處理器內(nèi)核。一個(gè)核函數(shù)可由多個(gè)相同規(guī)格的線

30、程塊執(zhí)行，組成一個(gè)多維度的線程塊網(wǎng)格Grid，Grid的維度由<<<>>>語(yǔ)法設(shè)置。網(wǎng)格中的塊可由多維索引標(biāo)識(shí)，內(nèi)核函數(shù)可通過(guò)blockIdx中訪問(wèn)，可以通過(guò)blockDim訪問(wèn)線程塊的維度。我們可以用多維的網(wǎng)格實(shí)現(xiàn)上面的程序。_global_void matAdd(int ANN,int BNN,int CNN)int i=blockIdx.x*blockDim.x+threadIdx.x;int j=blockIdx.y*blockDim.y+threadIdx.y;if(i<N&&j<N)Cij=Aij+Bij;int ma

31、in()dim3 dimBlock(16,16);dim3 dimGrid(N+dimBlock.x1)/dimBlock.x,(N+dimBlock.y1)/dimBlock.y);matAdd<<<dimGrid,dimBlock>>>(A,B,C)l 存儲(chǔ)器層次結(jié)構(gòu)CUDA有一個(gè)多層次的存儲(chǔ)器結(jié)構(gòu)。一個(gè)線程有一個(gè)局部存儲(chǔ)器（local memory）。每一個(gè)線程塊有一個(gè)塊內(nèi)所有線程共有的共享存儲(chǔ)器（shared memory）。所有線程共有一個(gè)全局存儲(chǔ)器（global memory）和兩個(gè)只讀存儲(chǔ)器：常量存儲(chǔ)器（constant memory）和紋理存

32、儲(chǔ)器（texture memory）。多層次的存儲(chǔ)器結(jié)構(gòu)適于不同的用途。l 宿主和設(shè)備CUDA線程在獨(dú)立設(shè)備（device）上執(zhí)行，C程序其它部分在CPU（宿主）上執(zhí)行。串行代碼在宿主上執(zhí)行，而并行代碼在設(shè)備上執(zhí)行。通過(guò)CUDA編程時(shí)，將GPU看作可以并行執(zhí)行非常多個(gè)線程的計(jì)算設(shè)備(compute device)。它作為主CPU的協(xié)處理器或者主機(jī)（host）來(lái)運(yùn)作：換句話說(shuō)，在主機(jī)上運(yùn)行的應(yīng)用程序中數(shù)據(jù)并行的、計(jì)算密集的部分卸載到此設(shè)備上。經(jīng)過(guò)了CUDA對(duì)線程、線程塊的定義和管理，在支持CUDA的GPU內(nèi)部實(shí)際上已經(jīng)成為了一個(gè)迷你網(wǎng)格計(jì)算系統(tǒng)。在內(nèi)存訪問(wèn)方面，整個(gè)GPU可以支配的存儲(chǔ)空間被分成

33、了寄存器（Register）、全局內(nèi)存（External DRAM）、共享內(nèi)存（Parallel Data Cache）三大部分。其中寄存器和共享內(nèi)存集成在GPU內(nèi)部，擁有極高的速度，但容量很小。共享內(nèi)存可以被同個(gè)線程塊內(nèi)的線程所共享，而全局內(nèi)存則是我們熟知的顯存，它在GPU外部，容量很大但速度較慢。經(jīng)過(guò)多個(gè)級(jí)別的內(nèi)存訪問(wèn)結(jié)構(gòu)設(shè)計(jì)，CUDA已經(jīng)可以提供讓人滿意的內(nèi)存訪問(wèn)機(jī)制，而不是像傳統(tǒng)GPGPU那樣需要開(kāi)發(fā)者自行定義。圖2.4為CPU與GPU通過(guò)Globle Memory交換數(shù)據(jù)。GridBlock(1,0)Block(0,0)Shared MemoryShared MemoryRegist

34、RegistRegistRegistersThread(1,0)Thread(0,0)Thread(1,0)Thread(0,0)HostGlobel Memory 圖2.4 CPU和GPU通過(guò)Global Memory交流數(shù)據(jù)2.4 CUDA開(kāi)發(fā)平臺(tái)的搭建2.4.1軟件安裝和環(huán)境配置本次開(kāi)發(fā)是在window7 32位上的，硬件條件為NVIDIA GeForce GT 550M顯卡下，使用到的軟件為Microsoft Visual Studio 2010和CUDA5.0。1）Visual Studio 2010及Visual Assist X安裝先安裝Visual Studio 2010 后再

35、安裝助手Visual Assist X。這里Visual Assist X 對(duì)于使用CUDA 作并行計(jì)算不是必須的，但為了使程序編寫更為方便，這里 推薦安裝。2）安裝cuda5.0 CUDA Toolkit的默認(rèn)安裝目錄為：C:Program FilesNVIDIA GPU Computing ToolkitCUDAv5.0 CUDA SDK的默認(rèn)安裝目錄為：C:ProgramDataNVIDIA CorporationCUDA Samplesv5.03）配置環(huán)境變量安裝完成Toolkit 和SDK 后，已自動(dòng)配置好系統(tǒng)環(huán)境變量。保險(xiǎn)起見(jiàn)，手動(dòng)配置環(huán)境變量。在系統(tǒng)環(huán)境變量中新建如下項(xiàng)： CUD

36、A_SDK_PATH = C:ProgramDataNVIDIA CorporationCUDA Samplesv5.0common CUDA_PATH = C:Program FilesNVIDIA GPU Computing ToolkitCUDAv5.0 CUDA_LIB_PATH = %CUDA_PATH%libWin32 CUDA_BIN_PATH = %CUDA_PATH%bin CUDA_SDK_LIB_PATH = %CUDA_SDK_PATH%commonlibWin32 CUDA_SDK_BIN_PATH = %CUDA_SDK_PATH%binWin322.4.2創(chuàng)建工程

37、在Visual Studio 2010 菜單選擇“file|new|project（文件|新建|工程）”，在打開(kāi)的新建項(xiàng)目窗口的“已安裝的模板”一欄中選擇“NVIDIA|CUDA”，類型選擇為“CUDA 5.0 Runtime”，見(jiàn)下圖3.2。圖2.5 創(chuàng)建工程在“名稱”中輸入工程名后，點(diǎn)擊確定。可對(duì)系統(tǒng)提供的kernel.cu 示例進(jìn)行編譯運(yùn)行，運(yùn)行結(jié)果如下圖3.3。圖2.6 kernel.cu程序運(yùn)行圖工程建立完畢，接下來(lái)，我們需要添加.cu文件來(lái)編寫代碼，右擊項(xiàng)目名稱，選擇“添加-新建項(xiàng)”，如圖2.7所示：圖2.7 添加.cu文件選擇CUDA C/C+ File,輸入項(xiàng)名稱，點(diǎn)擊提提那家

38、，.cu文件添加完畢，我們可以在文件中添加代碼，代碼添加完畢之后，點(diǎn)“CTRL+F5”編譯。3 基于CUDA的圖像旋轉(zhuǎn)實(shí)現(xiàn)3.1 實(shí)現(xiàn)圖像旋轉(zhuǎn)的方案選擇3.1.1適用于SPMD計(jì)算的圖像變換算法圖像幾何變換主要包括圖像的平移，縮放，旋轉(zhuǎn)及扭曲等。平移及縮放變換處理較為簡(jiǎn)單，只需要做加法處理即能得到最終結(jié)果，因此在變換算法上沒(méi)有大的提升空間。圖像的幾何變換可以用一個(gè)線性方程組來(lái)描述： x'=a0x+a1y+a2y'=b0x+b1y+b2 式（1）圖像的幾何變換過(guò)程實(shí)質(zhì)就是一個(gè)確定線性方程組系數(shù)值的過(guò)程。由于圖像的像素點(diǎn)的位置是離散的，經(jīng)過(guò)式(1)變換以后新的位置為P'(x

39、 ', y ')不一定正好落在一個(gè)像素上,即所謂的空穴現(xiàn)象7。采用逆向變換的方法可以解決這一問(wèn)題。若變換過(guò)程P 采用直接求式(1)逆變換方法處理新圖像的每一個(gè)像素點(diǎn)，每個(gè)點(diǎn)將會(huì)進(jìn)行4 次乘法運(yùn)算，那么對(duì)于變換之后的大小為M *N 的圖像，將會(huì)產(chǎn)生4*M *N 次乘法運(yùn)算。考慮采用位置增量來(lái)代替大量乘法運(yùn)算，以進(jìn)一步提高變換效率。將一大小為M *N 圖像任意點(diǎn)(x, y)以點(diǎn)0 0 (x , y )為中心旋轉(zhuǎn) ，變換后的位置為： x'=(x-x0)cos-(y-y0)siny'=(y-y0)cos+(x-x0)sin 式(2)為了避免出現(xiàn)提到的所謂空穴點(diǎn)，采用逆變

40、換方式，因此將式(2)做逆變換得： x=x'cos+y'sin+x0y=y'cos-x'sin+y0 式(3) 圖3.1 任意兩點(diǎn)位置關(guān)系如圖3.1所示，設(shè)p '和q 分別表示變換后新圖像中兩點(diǎn)，其位置關(guān)系可用式(4)表示 xq'=xp'+ xyq'=yp'+y 式(4)將式（4）帶入式（3），最終得到： xq=xp+xcos+ysinyq=yp+ycos-xsin 式(5)設(shè) p '和q '為同一行的兩個(gè)像素點(diǎn)，則y = 0。因此在同一行中任意一點(diǎn)q '與行首點(diǎn) 0 X 在原圖像的位置偏移量8，x

41、 ，y為：x = 0*cos Xq'= X01*cos Xq'= X1n-1*cos Xq'= Xn 式(6) y =0*sin Xq'=X0-1*sin Xq'=X1-(n-1)sin Xq'=Xn 式(7)由上式可以看出， 任意同一行中點(diǎn)q' , 1, 2,., n q n = N 的 X 方向及Y 方向上的位置偏移量與其對(duì)應(yīng)的坐標(biāo)具有線性增量關(guān)系。同一列像素點(diǎn)位置偏移關(guān)系類似。這樣，只要計(jì)算旋轉(zhuǎn)圖像中原點(diǎn)在原圖中的位置，其余像素點(diǎn)均可通過(guò)上式求出其位置偏移量，進(jìn)而計(jì)算出在原圖像中的位置。由此，對(duì)于變換之后的大小為的M *N 圖像，只

42、需要計(jì)算原點(diǎn)位置時(shí)進(jìn)行4 次乘法操作，其余點(diǎn)均可以加法操作代替乘法操作。由采用位置偏移量的方法將圖像變換需要4*M *N 次乘法運(yùn)算變成只需要2*M *N 次加法運(yùn)算，在效率上很大提高。至此，M *N 次加法運(yùn)算成為圖像變換算法中最耗時(shí)的部分。注意到每個(gè)像素點(diǎn)在變換過(guò)程中都是執(zhí)行相同的加法操作指令，且操作的數(shù)據(jù)具有線性關(guān)系，非常符合并行化的特點(diǎn)，可以進(jìn)行并行運(yùn)算。設(shè)X 為輸入圖像，Y 為輸出圖像，P為圖像變換過(guò)程，|表示并行執(zhí)行，則圖像變換并行化可以表示為：Y = P(X)=P1(X)|P2(X)|Pn(X)3.1.2 紋理技術(shù)(1) 紋理屬性紋理可以在線性內(nèi)存或是CUDA數(shù)組(紋理內(nèi)存)的任

43、何區(qū)域。所以紋理拾取也就對(duì)存在與線性內(nèi)存或CUDA數(shù)組中的紋理讀取數(shù)據(jù)。共用運(yùn)行組件(既可以運(yùn)行在host又可以運(yùn)行在設(shè)備)中給出了紋理類型texture。紋理拾取的第一個(gè)參數(shù)就是紋理參考，紋理參考定義要拾取哪部分紋理內(nèi)存，它必須通過(guò)宿主運(yùn)行時(shí)(只運(yùn)行在宿主上)函數(shù)綁定到一些內(nèi)存區(qū)域(稱為紋理(texture)，然后才能供內(nèi)核使用。下面就來(lái)看看紋理參考的不可變屬性和可變屬性。l 不可變屬性紋理參考的聲明是texture texRef;Type、Dim和ReadMode都是不可變屬性。其中：Type指定拾取紋理時(shí)返回的數(shù)據(jù)類型;Type限制為基本的整數(shù)和浮點(diǎn)數(shù)類型，以及1-、2-和4-分量的向量

44、類型之一。Dim指定紋理參考的維度，等于1、2或3;Dim是可選參數(shù)，缺省值為1;ReadMode等于cudaReadModeNormalizedFloat或cudaReadModeElementType。如果;ReadMode為cudaReadModeNormalizedFloat，且Type為16-位或8-位整數(shù)類型，則其值實(shí)際返回值為浮點(diǎn)數(shù)類型，即根據(jù)原整數(shù)類型的全范圍進(jìn)行歸一化處理，結(jié)果被映射到 0.0, 1.0區(qū)間(對(duì)于無(wú)符號(hào)整數(shù))或-1.0, 1.0為區(qū)間(對(duì)于有符號(hào)整數(shù));例如，值0xff的無(wú)符號(hào)8-位紋理元素返回值為1;如果ReadMode為cudaReadModeElemen

45、tType，則不執(zhí)行任何轉(zhuǎn)換;ReadMode是可選參數(shù)，默認(rèn)為cudaReadModeElementType.l 可變屬性紋理參考的可變屬性包括尋址模式、紋理過(guò)濾和紋理坐標(biāo)是否歸一化。這些屬性都是執(zhí)行宿主運(yùn)行時(shí)改變的，如cudaBindTextureToArray()執(zhí)行時(shí)在硬件紋理單元(texture units)時(shí)。其中：addressMode，即如何處理超出范圍的紋理坐標(biāo)。尋址模式是大小為2的數(shù)組，數(shù)組的第一個(gè)和第二個(gè)元素分別指定第一個(gè)第二個(gè)紋理坐標(biāo)的尋址模式：當(dāng)尋址模式等于cudaAddressModeClamp時(shí)，超出范圍的紋理坐標(biāo)將使用clamp尋址：非歸一化紋理坐標(biāo)的情況下，小

46、于0的值設(shè)置為0，大于等于N的值設(shè)置為N-1;歸一化紋理坐標(biāo)的情況下，小于0.0或大于1.0的值設(shè)置到區(qū)間0.0，1.0)中。當(dāng)尋址模式是cudaAddressModeWarp時(shí)，超出范圍的紋理坐標(biāo)使用warp模式：warp尋址僅使用紋理坐標(biāo)的小數(shù)部分：如1.25當(dāng)作0.25處理，-1.25當(dāng)作0.75處理。默認(rèn)情況下是cudaAddressModeClamp。cudaAddressModeWarp只支持歸一化的紋理坐標(biāo)。filterMode，過(guò)濾模式即當(dāng)拾取紋理時(shí)，如何基于輸入紋理坐標(biāo)來(lái)計(jì)算返回的值。filterMode等于cudaFilterModePoint或cudaFilterMode

47、Linear;如果它為cudaFilterModePoint，則返回值是紋理坐標(biāo)最接近輸入紋理坐標(biāo)的紋理元素;如果它為cudaFilterModeLinear，則返回值是紋理坐標(biāo)最接近輸入紋理坐標(biāo)的2個(gè)(對(duì)于1D紋理)、4個(gè)(對(duì)于1個(gè)2D紋理)或8給(對(duì)于3D紋理)的紋理元素的線性插值。cudaFilterModeLinear只對(duì)返回值是浮點(diǎn)類型有效。normalize，指定紋理坐標(biāo)是否是歸一化的;如果其值非0，則紋理中的所有元素都使用區(qū)間0,1，而非區(qū)間0,width-1、0,height-1或0,depth-1中的紋理坐標(biāo)來(lái)尋址，其中width、height和depth是紋理大小。在線性內(nèi)

48、存中分配的紋理：1) 維度只能等于1;2）不支持紋理過(guò)濾;3）只能使用非歸一化的整數(shù)紋理坐標(biāo)尋址;4）尋址模式單一：越界的紋理訪問(wèn)將返回0值。(2) 紋理拾取函數(shù)紋理拾取函數(shù)是設(shè)備運(yùn)行時(shí)函數(shù)。紋理存儲(chǔ)的區(qū)間不同，拾取的方式也不同。從線性內(nèi)存中拾取，使用的函數(shù)是tex1Dfetch()的函數(shù)簇。如：template<class Type>Type tex1Dfetch(texture<Type, 1, cudaReadModeElementType> texRef, int x); float tex1

49、Dfetch(texture<unsigned char, 1, cudaReadModeNormalizedFloat> texRef,int x);這些函數(shù)用紋理坐標(biāo)x拾取綁定到紋理參考texRef 的線性內(nèi)存的區(qū)域。不支持紋理過(guò)濾和尋址模式。對(duì)于整數(shù)型，這些函數(shù)將會(huì)將整數(shù)型轉(zhuǎn)化為單精度浮點(diǎn)型。除了這些函數(shù)，還支持2-和4-分量向量的拾取。如：float4 tex1Dfetch(texture<uchar4, 1, cudaReadModeNormalizedFloat> texRef,int x

50、);用紋理坐標(biāo)x拾取綁定到紋理參考的線性內(nèi)存的區(qū)域。從CUDA數(shù)組中拾取，是用函數(shù)tex1D()，tex2D()，tex3D()：template<class Type, enum cudaTextureReadMode readMode>Type tex1D(texture<Type, 1, readMode> texRef,float x); template<class Type, enum cudaTextureReadMode readMode>Ty

51、pe tex2D(texture<Type, 2, readMode> texRef,float x, float y); template<class Type, enum cudaTextureReadMode readMode>Typetex3D(texture<Type, 3, readMode> texRef,float x, float y, float z);這些函數(shù)用用紋理坐標(biāo)x，y，

52、z拾取綁定到紋理參考 texRef 的CUDA數(shù)組。紋理參考的可變和不可變屬性決定了如何理解坐標(biāo)，在紋理拾取和返回值將做怎樣的處理。(3) 拾取紋理內(nèi)存與讀取全局或常量?jī)?nèi)存相比的優(yōu)點(diǎn)1) 有高速緩存，如果CUDA數(shù)組中的紋理在片上的高速緩存中，則可以潛在的獲得較高帶寬2) 不受訪問(wèn)模式的約束。全局或常量?jī)?nèi)存讀取必須遵循相應(yīng)訪存模式才能獲得好的性能。如全局內(nèi)存在單個(gè)指令中將32-位、64-位或128-位從全局內(nèi)存讀取到寄存器，單個(gè)指令讀取的位數(shù)要盡量多;另外每個(gè)半warp中同時(shí)訪問(wèn)全局內(nèi)存地址的每個(gè)線程應(yīng)該進(jìn)行排列，以便內(nèi)存訪問(wèn)可以合并到單個(gè)鄰近的、對(duì)齊的內(nèi)存訪問(wèn)中。3) 尋址計(jì)算的延遲隱藏的更

53、好，有時(shí)候會(huì)改善應(yīng)用程序執(zhí)行隨機(jī)訪問(wèn)數(shù)據(jù)的性能。4) 打包的數(shù)據(jù)可以在單個(gè)操作中廣播到多個(gè)獨(dú)立變量中5) 8-位和16-位整數(shù)輸入數(shù)據(jù)可以有選擇地轉(zhuǎn)化為0.0,1.0或-1.0,1.0區(qū)間內(nèi)的32位浮點(diǎn)值6) 如果訪問(wèn)的是CUDA數(shù)組還有其他的功能，過(guò)濾、歸一化紋理坐標(biāo)、尋址模式3.2 基于CUDA的數(shù)字圖像幾何變換方案實(shí)現(xiàn)基于CUDA的數(shù)字圖像幾何變換的實(shí)現(xiàn)流程圖如圖3.2所示。CUDA初始化紋理參數(shù)拾取坐標(biāo)變換數(shù)據(jù)傳至GPUGPU完成并行核心過(guò)程數(shù)據(jù)回傳圖3.2 基于CUDA的數(shù)字圖像幾何變換流程圖 由上述流程圖，我們可以看出，本次實(shí)現(xiàn)基于CUDA的數(shù)字圖像幾何變換主要分為六個(gè)部分：CUD

54、A初始化，紋理參數(shù)拾取，坐標(biāo)變換，數(shù)據(jù)傳送至GPU，由GPU完成并行化核心過(guò)程，數(shù)據(jù)回傳。3.3 基于CUDA的數(shù)字圖像幾何變換的方案實(shí)現(xiàn)3.3.1 CUDA初始化首先，初始化CUDA，代碼如附錄1，結(jié)果如圖4.2.圖3.3 初始化CUDA3.3.2 紋理參數(shù)的拾取和坐標(biāo)變換通過(guò)紋理拾取函數(shù)，進(jìn)行紋理參數(shù)的拾取，并對(duì)圖片進(jìn)行紋理坐標(biāo)的計(jì)算和歸一化紋理左邊的變換。/ Constantsconst float angle = 0.5f; / angle to rotate image by (in radians)/ Texture reference for 2D float texturete

55、xture<float, 2, cudaReadModeElementType> tex;/ Auto-Verification Codebool testResult = true;/! Transform an image using texture lookups/! param outputData output data in global memory_global_ void transformKernel(float *outputData, int width, int height, float theta)/ calculate normalized texture coordinates unsigned int x = blockIdx.x*blockDim.x + threadIdx.x; unsigned int y = blockIdx.y*blockDim.y + threadIdx.y; float u = x / (float) width; float v = y / (float) height;/ transform coordinates u -=