吳輝 羅清海 彭文武

摘要：本文闡述了GPU并行運(yùn)算的一種主流架構(gòu)——CUDA架構(gòu)，包括CUDA編程模型、程序的運(yùn)行模式、線程架構(gòu)、存儲(chǔ)器結(jié)構(gòu)、指令結(jié)構(gòu)等。

關(guān)鍵詞：GPU；CUDA架構(gòu)；并行計(jì)算

中圖分類號(hào)：G642.41 文獻(xiàn)標(biāo)志碼：A 文章編號(hào)：1674-9324（2019）06-0277-02

從20世紀(jì)90年代，第一款圖形處理器誕生到現(xiàn)在，摩爾定律預(yù)測(cè)到由于制造工藝和設(shè)計(jì)的進(jìn)步，單一芯片內(nèi)部集成的晶體管數(shù)量大約每隔一年都會(huì)翻一翻，但是隨著工藝進(jìn)步的空間日趨縮小，制造工藝反而成了一種限制，CPU上的晶體管不能無限增加，所以怎么把任務(wù)更均衡地分配給GPU，實(shí)現(xiàn)CPU-GPU的負(fù)載均衡，是當(dāng)代計(jì)算機(jī)技術(shù)的一個(gè)重要研究方向。

而CUDA是英偉達(dá)公司2006年推出的通用并行計(jì)算架構(gòu)，其提供直接的硬件訪問接口，使得程序開發(fā)人員不必再依賴應(yīng)用程序編程接口（Application Programming Interface，API）來實(shí)現(xiàn)GPU的交互通信，完整地解決GPU設(shè)備的通用并行計(jì)算的程序開發(fā)及實(shí)現(xiàn)的問題，CUDA平臺(tái)主要包含三個(gè)部分：開發(fā)庫(kù)CUFFT（離散快速傅立葉變換）和CUBLAS（離散基本線性計(jì)算）等、運(yùn)行期環(huán)境（宿主代碼和設(shè)備代碼）和驅(qū)動(dòng)。

一、CUDA編程模型

CUDA的編程模型為中央處理器CPU與圖形處理器GPU聯(lián)合運(yùn)行的異構(gòu)編程平臺(tái)，即也就是在該平臺(tái)上，程序同時(shí)使用了CPU和GPU，其中CPU作為主機(jī)端（Host），GPU作為設(shè)備端（Device）或協(xié)處理器，CPU利用強(qiáng)大的邏輯門判斷功能，能夠快速解釋計(jì)算機(jī)指令，分配數(shù)據(jù)處理事物和串行計(jì)算，并指揮設(shè)備端的運(yùn)作，而GPU的邏輯門判斷功能簡(jiǎn)單，但是算術(shù)邏輯單元數(shù)量眾多，數(shù)據(jù)計(jì)算能力十分強(qiáng)大，負(fù)責(zé)處理可以高度并行化的任務(wù)。在Nvida顯卡平臺(tái)，CPU和GPU之間通過PCI-E數(shù)據(jù)總線進(jìn)行數(shù)據(jù)傳遞，以保證數(shù)據(jù)傳輸擁有足夠的帶寬。具體執(zhí)行模式是主機(jī)端（CPU）運(yùn)行的程序段通過kernel函數(shù)來控制設(shè)備端（GPU）的程序的運(yùn)行；此外，它們擁有各自互相獨(dú)立的內(nèi)存空間：顯存和內(nèi)存；另外CUDA對(duì)儲(chǔ)存空間的調(diào)配方式也不同，CUDA平臺(tái)使用調(diào)配函數(shù)CUDA API來調(diào)配設(shè)備端（GPU）上的顯存進(jìn)行空間開辟、空間釋放、初始化顯存等動(dòng)作，對(duì)主機(jī)端（CPU）上的內(nèi)存管理，則使用C語(yǔ)言集成的內(nèi)存管理函數(shù)來讀寫。

二、CUDA程序的運(yùn)行模式

a.開辟主機(jī)端存儲(chǔ)器（主存）空間并進(jìn)行初始化。

b.開辟設(shè)備端存儲(chǔ)器（顯存）空間。

c.將已初始化的主機(jī)端存儲(chǔ)器的數(shù)據(jù)通過PCI-E數(shù)據(jù)總線復(fù)制并傳遞到設(shè)備端存儲(chǔ)器上的開辟空間內(nèi)。

d.GPU進(jìn)行并行計(jì)算。

e.將GPU上已處理完成的數(shù)據(jù)拷貝回主機(jī)端。

f.主機(jī)端進(jìn)行數(shù)據(jù)處理。

三、GPU的線程架構(gòu)

GPU的線程架構(gòu)是一種邏輯架構(gòu)，這樣使開發(fā)人員不必深入地學(xué)習(xí)每一個(gè)GPU的內(nèi)部物理結(jié)構(gòu)，從而從繁雜的計(jì)算機(jī)圖形學(xué)工作中脫離出來，事實(shí)上，該線程架構(gòu)包含了由淺入深的三個(gè)部分，由大到小依次為線程格（grid）、線程塊（block）和線程（thread）。一個(gè)kernel函數(shù)映射到GPU上后，稱為一個(gè)網(wǎng)格（grid）；一個(gè)線程格由若干個(gè)線程塊構(gòu)成，線程塊與線程塊之間通過共享存儲(chǔ)器進(jìn)行數(shù)據(jù)傳輸，并把數(shù)據(jù)發(fā)射到流多處理器（Streaming Multiprocessor，SM）上執(zhí)行，這些線程塊由數(shù)量和排列方式能夠組成一維和二維結(jié)構(gòu)，且每個(gè)線程塊都用一個(gè)唯一的坐標(biāo)（blockIdx）進(jìn)行區(qū)別編號(hào)；同時(shí)，每個(gè)線程塊都包含了若干個(gè)線程（thread），線程在組織和結(jié)構(gòu)上與線程塊相類似，但是線程能組成三維結(jié)構(gòu)。線程塊中的線程通過編號(hào)變量threadIdx.x、threadIdx.y、threadIdx.z中的x、y、z來使用線程，并構(gòu)成在線程塊中的一維、二維和三維中的索引標(biāo)識(shí)。

四、CUDA存儲(chǔ)器架構(gòu)

CUDA存儲(chǔ)器分為GPU片內(nèi)內(nèi)存：寄存器（register）、共享內(nèi)存（shared memory，SM）、本地內(nèi)存（local memory，LM）；板載內(nèi)存：常量?jī)?nèi)存（constant memory，CM）、紋理內(nèi)存（texture memory，TM）、全局內(nèi)存（global memory，GM）。

共享內(nèi)存位于圖形處理器內(nèi)部，每個(gè)流多處理器SM能使用的共享內(nèi)存分為16KB、64K等幾種，如果要獲取較高的存儲(chǔ)器帶寬，則流多處理器上的共享內(nèi)存一般被分割合并成一些16位或者32位的Bank，線程訪問Bank的速度與寄存器的訪問速度在理論上能達(dá)到一致，實(shí)際上由于幾個(gè)線程讀寫相同的Bank時(shí)產(chǎn)生沖突，而造成存取延時(shí)，因此共享寄存器的延遲有30—40個(gè)時(shí)鐘。本地內(nèi)存用來存儲(chǔ)寄存器溢出的數(shù)據(jù)，通常是可能消耗了大量寄存器空間的大數(shù)組或大結(jié)構(gòu)、無法確定大小的數(shù)組和內(nèi)核使用的寄存器溢出的任何變量。另外，由于本地內(nèi)存存在于板載內(nèi)存空間，所以所有的本地內(nèi)存的讀寫訪問延遲與全局內(nèi)存一樣搞，帶寬和全局內(nèi)存一樣低。全局內(nèi)存位于設(shè)備存儲(chǔ)器中，儲(chǔ)存空間最大，但距離運(yùn)算執(zhí)行單元流處理器最遠(yuǎn)，存取速度最慢的存儲(chǔ)器，只有約177GB/s，訪問延時(shí)高達(dá)400—600個(gè)時(shí)鐘中期。設(shè)備存儲(chǔ)器通過32，64或128字節(jié)的存儲(chǔ)器通信訪問，當(dāng)一個(gè)束需要讀取全局內(nèi)存時(shí)，它會(huì)采用合并訪問（coalesced access）的形式，將束內(nèi)線程的存儲(chǔ)器的訪問一次或者多次完成，以提高訪問效率。整合內(nèi)存是指，若干線程的內(nèi)存請(qǐng)求在硬件上被合并和整合成一次多數(shù)據(jù)的數(shù)據(jù)傳輸，從而提高全局內(nèi)存的存取速度，這是實(shí)現(xiàn)高性能GPGPU編程極為重要的手段。常量?jī)?nèi)存是一段位于設(shè)備存儲(chǔ)器上的只讀地址空間，特點(diǎn)是對(duì)其訪問可以獲得緩存加速，如果從常量?jī)?nèi)存中獲得所需要的常量時(shí)，只要一個(gè)時(shí)鐘周期即可返回，大大提高了程序運(yùn)行的速度，但是常量?jī)?nèi)存的容量一般只有64K大小，故只存儲(chǔ)可能被頻繁讀取的只讀參數(shù)。常量?jī)?nèi)存在函數(shù)體外使用_constant_變量類型限制符進(jìn)行存放，能夠被GPU的所有線程訪問。本文的程序的權(quán)系數(shù)、流體的參數(shù)、系數(shù)矩陣等存儲(chǔ)于常量?jī)?nèi)存空間。紋理內(nèi)存的位置在設(shè)備存儲(chǔ)器上，能夠?qū)⒉糠秩謨?nèi)存的數(shù)據(jù)緩存到紋理緩存中，但是紋理內(nèi)存也是一段只讀的內(nèi)存空間。

五、CUDA指令結(jié)構(gòu)

可拓展的多線程流多處理器（SMs）架構(gòu)是CUDA架構(gòu)的核心，如果運(yùn)行在主機(jī)端的CUDA程序調(diào)用kernel核函數(shù)后，線程格內(nèi)塊枚舉并分發(fā)到處于空閑狀態(tài)的流多處理器上執(zhí)行，SM設(shè)計(jì)為能同時(shí)控制數(shù)百的線程進(jìn)行運(yùn)算，因此，流多處理器采用單指令多線程（SIMT）方式來同時(shí)操控如此多的線程。

六、結(jié)論

本文所述的GPU并行計(jì)算架構(gòu)為NVIDIA的CUDA計(jì)算平臺(tái)，這是當(dāng)前GPU并行計(jì)算應(yīng)用及研究的主流架構(gòu)，本文通過介紹該架構(gòu)的主要結(jié)構(gòu)及運(yùn)行模式，為其在各個(gè)領(lǐng)域的應(yīng)用提供前瞻性的參考。

參考文獻(xiàn)：

[1]雷德川，陳浩，王遠(yuǎn)，張成鑫，陳云斌，胡棟材.基于CUDA的多GPU加速SART迭代重建算法[J].強(qiáng)激光與粒子束，2013，（09）：2418-2422.

[2]岳俊，鄒進(jìn)貴，何豫航.基于CPU與GPU/CUDA的數(shù)字圖像處理程序的性能比較[J].地理空間信息，2012，（04）：45-47+180.