GPU并行計算的CUDA架構淺析

吳輝 羅清海 彭文武

摘要：本文闡述了GPU并行運算的一種主流架構——CUDA架構，包括CUDA編程模型、程序的運行模式、線程架構、存儲器結構、指令結構等。

關鍵詞：GPU；CUDA架構；并行計算

中圖分類號：G642.41 文獻標志碼：A 文章編號：1674-9324（2019）06-0277-02

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

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

一、CUDA編程模型

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

二、CUDA程序的運行模式

a.開辟主機端存儲器（主存）空間并進行初始化。

b.開辟設備端存儲器（顯存）空間。

c.將已初始化的主機端存儲器的數據通過PCI-E數據總線復制并傳遞到設備端存儲器上的開辟空間內。

d.GPU進行并行計算。

e.將GPU上已處理完成的數據拷貝回主機端。

f.主機端進行數據處理。

三、GPU的線程架構

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

四、CUDA存儲器架構

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

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

五、CUDA指令結構

可拓展的多線程流多處理器（SMs）架構是CUDA架構的核心，如果運行在主機端的CUDA程序調用kernel核函數后，線程格內塊枚舉并分發到處于空閑狀態的流多處理器上執行，SM設計為能同時控制數百的線程進行運算，因此，流多處理器采用單指令多線程（SIMT）方式來同時操控如此多的線程。

六、結論

本文所述的GPU并行計算架構為NVIDIA的CUDA計算平臺，這是當前GPU并行計算應用及研究的主流架構，本文通過介紹該架構的主要結構及運行模式，為其在各個領域的應用提供前瞻性的參考。

參考文獻：

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

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