用CUDA加速多波束測深系統(tǒng)波束形成-精品文檔

1、用CUDA加速多波束測深系統(tǒng)波束形成1 引言多波束測深系統(tǒng)是一種高效的水底地形測量系統(tǒng)。 設(shè)備通過 向水底發(fā)射一次信號， 接收信號的回波來進行處理， 得到測量平 臺兩側(cè)一條帶上幾百個采樣點的深度數(shù)據(jù)。 波束形成是多波束測 深系統(tǒng)信號處理中一個非常重要的任務(wù)， 往往耗費大量的計算資 源，通常對這類數(shù)據(jù)的處理采用 DSP來實現(xiàn)。然而這種DSP實時 處理的方式具有投入大， 研發(fā)周期長， 板卡多，升級困難的缺點。 由英偉達公司提出的 CUDA(Compute Unified Device Architecture )計算架構(gòu)，能使用標準的 C來進行GPU®程，對 大規(guī)模數(shù)據(jù)并行計算，不僅運行

2、速度快，并行化程度高，還具有 開發(fā)周期短、無需硬件設(shè)計調(diào)試、便于升級的優(yōu)點。本文通過CUDA技術(shù)對多波束測深系統(tǒng)的波束形成進行并行 優(yōu)化處理， 大大提高了信號處理速度， 從而在移動工作站上實現(xiàn) 了多波束測深數(shù)據(jù)的實時波束形成處理。2 多波束測深系統(tǒng)波束形成考慮多波束測深系統(tǒng)中基陣接收位于遠場的情況， 即信號波 前到達基陣時可看作平面波。在遠場情況下對基帶解調(diào)后的信 號，采用基帶延遲加相移波束形成。 對于波束形成需要的時延補 償，其整數(shù)點采樣部分，用時延實現(xiàn)，剩余的小數(shù)采樣部分，用 相移實現(xiàn)。在遠場平面波的假設(shè)條件下， 以任意陣型M元陣，入射單頻 脈沖信號為例。 假設(shè)第 i 號陣元上的接收信號經(jīng)

3、過基帶解調(diào)后為 xi (t)，其中i=1 , 2，M它的希爾波特變換(Hilbert )( t )。設(shè)信號原始頻率為 f0 ，接收信號經(jīng)過基帶解調(diào)后頻率為 fc，采樣頻率為fs。若要形成e方向的預(yù)成波束，根據(jù)常規(guī)波 束形成的時延計算方法， 可以計算出， 第 i 號陣元需要補償?shù)臅r 延，假定為T i ( e )，其中i=1 , 2，M。采用基帶信號的 延遲加相移波束形成方法，則 e 方向的預(yù)成波束輸出為：(1)其中： ( 2)( 3)ai 為陣元的權(quán)值， ni ( e )為需要延時的整數(shù)采樣點， T i (e )為剩下的小數(shù)點采樣對應(yīng)的時延。3 CUDA實現(xiàn)3.1 CUDA 編程模型CUDA 用

4、了“單指令多線程”(Si ngle In structio nMultiple Thread , SIMT)的執(zhí)行模型，只需一條指令，GPU的內(nèi)部機制會保證其所有配置的線程都會執(zhí)行同樣的運算。 CUDA 的線程模型分為三層，從上到下依次為 Grid ， Block ， Thread。 Thread (線程)是最小的單元，每個 Thread可以映射到一個流 處理器上，通過合理安排 threads 的工作可以提高程序的并行 化，提升程序性能。 kernel (內(nèi)核函數(shù))是一個只能運行在 GPU 上的函數(shù)，其調(diào)用格式為 _global_ kernel_function>>(paraml,

5、 )，其中 num_blocks 禾口 num_threads 是 dim3 整形 向量類型的變量， 分別表示網(wǎng)格和塊的維度， 兩者決定了內(nèi)核函 數(shù)執(zhí)行時所需要的線程數(shù)。3.2 CUDA 程序波束形成設(shè)計對于N通道的基帶解調(diào)數(shù)據(jù)，數(shù)據(jù)以N*M的2維數(shù)組排列，X維表示時間序列，Y維表示通道號，對 方向進行波束形成：定義dim3 threads_rect (x，y)，表示一個塊內(nèi)的線程數(shù)， 定義dim3 blocks_rect ( x，y，z)，表示線程塊的數(shù)量。其中 threads_rect.y* blocks_rect.y=N ， blocks_rect.x=( M+threads_rect.

6、x-1 ) /threads_rect.x ，一個波束時 z=1 ， 對應(yīng)的CUDA內(nèi)部kernel分配圖如圖1所示。其中 X維的索引 IDx= () +threadIdx.x ；Y維的索引 IDy= ()+threadldx.y 。3.3 性能優(yōu)化3.3.1 線程數(shù)的選擇表 1 顯示了在不同計算能力的設(shè)備上， 每個線程塊開啟不同 線程數(shù)時的設(shè)備利用率。通過表1可以看出，當線程數(shù)為128512之間時，在不同 的設(shè)備上都有比較好的利用率。 但是高利用率并不一定意味著高 性能，CUDA1以線程束(Warp)的方式訪問全局內(nèi)存，線程束 的大小通常為 32，每次以 Half-Warp 的 16 個獨立

7、線程對全局內(nèi) 存中的一段連續(xù)數(shù)據(jù)進行訪問，可以最大化 IO 的效率。因此我 們線程塊開啟的線程數(shù)設(shè)置為 16x16,對應(yīng)threads_rect.x=16 , threads_rect.y=16 。3.3.2 常量內(nèi)存在CUDA中，常量內(nèi)存其實只是全局內(nèi)存的一種虛擬地址形 式，但是它具有高速緩存， 訪問速度要比沒有緩存的全局存儲器 快，因此把系統(tǒng)中所有只讀常量保存在常量內(nèi)存中可以提高性 能。3.3.3 鎖頁內(nèi)存 為了對某一數(shù)據(jù)集進行計算，需要把數(shù)據(jù)從主機傳輸?shù)?GPU設(shè)備上，計算完成后再把數(shù)據(jù)結(jié)果傳輸回主機，主機與GPL設(shè)備的數(shù)據(jù)傳輸是依靠 PCI-E 總線來實現(xiàn)的。 PCI-E 傳輸實際上只

8、能 使用基于DMA的傳輸方式，在不使用鎖頁內(nèi)存時，驅(qū)動程序會在 后臺先分配一塊鎖頁內(nèi)存， 實行一個從常規(guī)內(nèi)存到鎖頁內(nèi)存的主 機端復(fù)制操作，這些時間額外消耗了寶貴的 CPU周期。直接使用 cudaHostAlloc 或 cudaMallocHost 分配鎖頁內(nèi)存，可以增加傳 輸帶寬，提高通信效率。3.3.4 異步傳輸在主機與GPU設(shè)備的數(shù)據(jù)傳輸過程中，GPU勺會出現(xiàn)大量空 閑，這樣導(dǎo)致GPU執(zhí)行效率不高。因此需要通過 CUDA流和 cudaMemcpyAsync來使CPU和GPU都盡量處于滿負荷狀態(tài)，提高GPI使用效率，掩蓋傳輸帶來的延遲。4 性能測試4.1 實驗環(huán)境實驗在移動工作站上進行，CP

9、U為In tel Core i7-4600M,主頻 2.9GHz，2 核 4 線程，內(nèi)存 16GB GPU為 GeForce GT745M 384個流處理器，2GB顯存；軟件支持環(huán)境為： Win7-64位操作 系統(tǒng)、MicrosoftVS 2010、NVIDIA GPU Computing SDK 7.5。4.2 GPU性能測試選取 128通道一個量程掃描周期的數(shù)據(jù)進行 256個波束形成 速度比較。不同通道點數(shù)的CPU和GPU加速度對比，左邊為消耗的時間， 右邊為加速比。4.3 實驗分析從表3和圖2可以看出，相對于 CPU GPU寸多波束測深系 統(tǒng)的波束形成計算具有明顯的加速作用， 并且隨著每通道點數(shù)的 增加，CPU的計算速度大幅下降，而 GPU勺耗時卻平穩(wěn)增長，到 每通道8192點數(shù)