吳輝 羅清海 彭文武

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

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

中圖分類(lèi)號(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ú)限增加，所以怎么把任務(wù)更均衡地分配給GPU，實(shí)現(xiàn)CPU-GPU的負(fù)載均衡，是當(dāng)代計(jì)算機(jī)技術(shù)的一個(gè)重要研究方向。

而CUDA是英偉達(dá)公司2006年推出的通用并行計(jì)算架構(gòu)，其提供直接的硬件訪(fǎng)問(wèn)接口，使得程序開(kāi)發(fā)人員不必再依賴(lài)應(yīng)用程序編程接口（Application Programming Interface，API）來(lái)實(shí)現(xiàn)GPU的交互通信，完整地解決GPU設(shè)備的通用并行計(jì)算的程序開(kāi)發(fā)及實(shí)現(xiàn)的問(wèn)題，CUDA平臺(tái)主要包含三個(gè)部分：開(kāi)發(fā)庫(kù)CUFFT（離散快速傅立葉變換）和CUBLAS（離散基本線(xiàn)性計(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)大的邏輯門(mén)判斷功能，能夠快速解釋計(jì)算機(jī)指令，分配數(shù)據(jù)處理事物和串行計(jì)算，并指揮設(shè)備端的運(yùn)作，而GPU的邏輯門(mén)判斷功能簡(jiǎn)單，但是算術(shù)邏輯單元數(shù)量眾多，數(shù)據(jù)計(jì)算能力十分強(qiáng)大，負(fù)責(zé)處理可以高度并行化的任務(wù)。在Nvida顯卡平臺(tái)，CPU和GPU之間通過(guò)PCI-E數(shù)據(jù)總線(xiàn)進(jìn)行數(shù)據(jù)傳遞，以保證數(shù)據(jù)傳輸擁有足夠的帶寬。具體執(zhí)行模式是主機(jī)端（CPU）運(yùn)行的程序段通過(guò)kernel函數(shù)來(lái)控制設(shè)備端（GPU）的程序的運(yùn)行；此外，它們擁有各自互相獨(dú)立的內(nèi)存空間：顯存和內(nèi)存；另外CUDA對(duì)儲(chǔ)存空間的調(diào)配方式也不同，CUDA平臺(tái)使用調(diào)配函數(shù)CUDA API來(lái)調(diào)配設(shè)備端（GPU）上的顯存進(jìn)行空間開(kāi)辟、空間釋放、初始化顯存等動(dòng)作，對(duì)主機(jī)端（CPU）上的內(nèi)存管理，則使用C語(yǔ)言集成的內(nèi)存管理函數(shù)來(lái)讀寫(xiě)。

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

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

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

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

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

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

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

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

GPU的線(xiàn)程架構(gòu)是一種邏輯架構(gòu)，這樣使開(kāi)發(fā)人員不必深入地學(xué)習(xí)每一個(gè)GPU的內(nèi)部物理結(jié)構(gòu)，從而從繁雜的計(jì)算機(jī)圖形學(xué)工作中脫離出來(lái)，事實(shí)上，該線(xiàn)程架構(gòu)包含了由淺入深的三個(gè)部分，由大到小依次為線(xiàn)程格（grid）、線(xiàn)程塊（block）和線(xiàn)程（thread）。一個(gè)kernel函數(shù)映射到GPU上后，稱(chēng)為一個(gè)網(wǎng)格（grid）；一個(gè)線(xiàn)程格由若干個(gè)線(xiàn)程塊構(gòu)成，線(xiàn)程塊與線(xiàn)程塊之間通過(guò)共享存儲(chǔ)器進(jìn)行數(shù)據(jù)傳輸，并把數(shù)據(jù)發(fā)射到流多處理器（Streaming Multiprocessor，SM）上執(zhí)行，這些線(xiàn)程塊由數(shù)量和排列方式能夠組成一維和二維結(jié)構(gòu)，且每個(gè)線(xiàn)程塊都用一個(gè)唯一的坐標(biāo)（blockIdx）進(jìn)行區(qū)別編號(hào)；同時(shí)，每個(gè)線(xiàn)程塊都包含了若干個(gè)線(xiàn)程（thread），線(xiàn)程在組織和結(jié)構(gòu)上與線(xiàn)程塊相類(lèi)似，但是線(xiàn)程能組成三維結(jié)構(gòu)。線(xiàn)程塊中的線(xiàn)程通過(guò)編號(hào)變量threadIdx.x、threadIdx.y、threadIdx.z中的x、y、z來(lái)使用線(xiàn)程，并構(gòu)成在線(xiàn)程塊中的一維、二維和三維中的索引標(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，線(xiàn)程訪(fǎng)問(wèn)Bank的速度與寄存器的訪(fǎng)問(wèn)速度在理論上能達(dá)到一致，實(shí)際上由于幾個(gè)線(xiàn)程讀寫(xiě)相同的Bank時(shí)產(chǎn)生沖突，而造成存取延時(shí)，因此共享寄存器的延遲有30—40個(gè)時(shí)鐘。本地內(nèi)存用來(lái)存儲(chǔ)寄存器溢出的數(shù)據(jù)，通常是可能消耗了大量寄存器空間的大數(shù)組或大結(jié)構(gòu)、無(wú)法確定大小的數(shù)組和內(nèi)核使用的寄存器溢出的任何變量。另外，由于本地內(nèi)存存在于板載內(nèi)存空間，所以所有的本地內(nèi)存的讀寫(xiě)訪(fǎng)問(wèn)延遲與全局內(nèi)存一樣搞，帶寬和全局內(nèi)存一樣低。全局內(nèi)存位于設(shè)備存儲(chǔ)器中，儲(chǔ)存空間最大，但距離運(yùn)算執(zhí)行單元流處理器最遠(yuǎn)，存取速度最慢的存儲(chǔ)器，只有約177GB/s，訪(fǎng)問(wèn)延時(shí)高達(dá)400—600個(gè)時(shí)鐘中期。設(shè)備存儲(chǔ)器通過(guò)32，64或128字節(jié)的存儲(chǔ)器通信訪(fǎng)問(wèn)，當(dāng)一個(gè)束需要讀取全局內(nèi)存時(shí)，它會(huì)采用合并訪(fǎng)問(wèn)（coalesced access）的形式，將束內(nèi)線(xiàn)程的存儲(chǔ)器的訪(fǎng)問(wèn)一次或者多次完成，以提高訪(fǎng)問(wèn)效率。整合內(nèi)存是指，若干線(xiàn)程的內(nèi)存請(qǐng)求在硬件上被合并和整合成一次多數(shù)據(jù)的數(shù)據(jù)傳輸，從而提高全局內(nèi)存的存取速度，這是實(shí)現(xiàn)高性能GPGPU編程極為重要的手段。常量?jī)?nèi)存是一段位于設(shè)備存儲(chǔ)器上的只讀地址空間，特點(diǎn)是對(duì)其訪(fǎng)問(wèn)可以獲得緩存加速，如果從常量?jī)?nèi)存中獲得所需要的常量時(shí)，只要一個(gè)時(shí)鐘周期即可返回，大大提高了程序運(yùn)行的速度，但是常量?jī)?nèi)存的容量一般只有64K大小，故只存儲(chǔ)可能被頻繁讀取的只讀參數(shù)。常量?jī)?nèi)存在函數(shù)體外使用_constant_變量類(lèi)型限制符進(jìn)行存放，能夠被GPU的所有線(xiàn)程訪(fǎng)問(wèn)。本文的程序的權(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)

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

六、結(jié)論

本文所述的GPU并行計(jì)算架構(gòu)為NVIDIA的CUDA計(jì)算平臺(tái)，這是當(dāng)前GPU并行計(jì)算應(yīng)用及研究的主流架構(gòu)，本文通過(guò)介紹該架構(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.