第1章 并行編程概覽

對絕大多數(shù)人而言，編程語言只是一個工具，講究簡單高效。科學(xué)家的主要精力應(yīng)該用在科研創(chuàng)新活動上，編程工作僅僅是用來驗(yàn)證創(chuàng)新的理論，編程水平再高也不可能獲得諾貝爾獎。對學(xué)生和企業(yè)程序員而言，技術(shù)無窮盡，永遠(yuǎn)學(xué)不完，不用即忘，應(yīng)該認(rèn)清技術(shù)發(fā)展方向，學(xué)習(xí)有前途的技術(shù)，不浪費(fèi)青春年華。

OpenACC語言專為超級計算機(jī)設(shè)計，因此讀者需要了解超級計算機(jī)的技術(shù)演進(jìn)方向，特別是主流加速器的體系架構(gòu)、編程模型，看清OpenACC的應(yīng)用場景，有的放矢。普通讀者雖然不會用到大型機(jī)群，但小型機(jī)群甚至單臺服務(wù)器、普通顯卡的計算模式都是相同的。

最近幾年的著名超級計算機(jī)（見附錄A）均采用加速器作為主要計算部件，可預(yù)見未來幾年的上層應(yīng)用仍將圍繞加速器展開。

1.1 加速器產(chǎn)品

超級計算機(jī)的加速器歷史上有多種，本節(jié)只介紹當(dāng)前流行的兩種：英偉達(dá)GPU和英特爾融核處理器。加速器的物理形態(tài)是PCIe板卡，樣子大致如圖1.1所示，圖1.2是拆掉外殼后的樣子，正中央的是GPU芯片，芯片周圍的小黑塊是顯存顆粒，金黃色的邊緣處是與PCIe連接的金手指，通過PCIe插槽與CPU相連。圖1.3中的機(jī)架式服務(wù)器左下部裝有4塊GPU卡，圖1.4是服務(wù)器的主板俯視圖，箭頭處就是4個PCIe插槽。

圖1.1 英偉達(dá)Tesla GPU

圖1.2 英偉達(dá)GPU內(nèi)部構(gòu)造

圖1.3 GPU服務(wù)器

圖1.4 GPU服務(wù)器的主板

從邏輯關(guān)系上來看，目前市場在售的英特爾服務(wù)器CPU中已經(jīng)集成了內(nèi)存控制器和PCIe通道。2顆CPU通過1個或2個QPI接口連接，CPU通過內(nèi)置的內(nèi)存控制器連接DDR3或DDR4內(nèi)存條，GPU加速器通過PCIe總線與CPU相連（圖1.5）。2016年6月主流的英特爾至強(qiáng)E5-2600 v3和E5-2600 v4 CPU每顆擁有40個PCIe Lan，而每塊GPU的接口需要16個PCIe Lan，因此每顆CPU上最多全速掛載2塊GPU。有些服務(wù)器的PCIe ×16插槽上只安排×8，甚至×4的信息速率，可以掛載更多的GPU。

圖1.5 加速器在服務(wù)器上的邏輯位置

英特爾融核處理器在服務(wù)器上的位置與GPU一樣，不再贅述。

1.1.1 英偉達(dá)GPU

通用計算GPU是英偉達(dá)公司發(fā)明的，每隔2～3年就會升級硬件架構(gòu)。通用計算GPU的第一代架構(gòu)代號是G80，每8個核心封裝在一起稱為一個流式多處理器（Streaming Multiprocessor, SM）,2個流多處理器共用一塊L1緩存，所有的核心共用L2緩存，任意核心都能訪問芯片外部的GPU顯存（圖1.6中的FB）。從圖1.6中可以看出，G80架構(gòu)GPU產(chǎn)品最多可以有128個核心。

圖1.6 通用計算GPU的G80架構(gòu)

接下來的架構(gòu)代號是GT200，然后是Fermi。Fermi架構(gòu)（圖1.7）最多包含14個流式多處理器，每個流式多處理器包含32個核心。6個DRAM接口，1個PCIe接口（圖1.16中的Host Interface）。

圖1.7 Fermi架構(gòu)圖

流式多處理器內(nèi)部組件也相當(dāng)多（圖1.8）:1個指令緩存（Instruction Cache）,2組Warp調(diào)度器（Warp Scheduler）、分發(fā)單元（Dispatch Unit），一堆寄存器；最重要的是32個核心，每個核心擁有一個單精度浮點(diǎn)單元和一個整數(shù)單元；16個Load/Store單元；4個特殊功能單元（Special Function Unit），負(fù)責(zé)計算雙精度浮點(diǎn)數(shù)、三角函數(shù)、超越函數(shù)、倒數(shù)、平方根等。共享內(nèi)存+L1緩存共64KB，可以靈活配置。

圖1.8 Fermi架構(gòu)的流式多處理器

接下來是Kepler架構(gòu)（圖1.9），這一代架構(gòu)最大可以包含15個流式多處理器（Streamong Multiprocessor extreme, SMX)和6個64位內(nèi)存控制器。不同的產(chǎn)品型號將使用不同的配置，可以包含13或14個SMX，多種參數(shù)都有升級和更改。Kepler架構(gòu)的主要設(shè)計目標(biāo)是提高用電效率，臺積電的28nm制造工藝在降低功耗方面起著重要的作用。Kepler架構(gòu)還提高了雙精度計算能力。

圖1.9 Kepler架構(gòu)

流式多處理器SMX（圖1.10）包含192個單精度核心、64個雙精度單元、32個特殊功能單元（SFU）和32個加載/存儲單元（LD/ST）、3個Warp調(diào)度器、6個分發(fā)器、一大堆寄存器。每個核心由1個浮點(diǎn)計算單元和1個整數(shù)算術(shù)邏輯單元組成，支持融加（FMA）運(yùn)算。每個SMX擁有64KB的片上存儲器，可配置為48KB的共享存儲器和16KB的L1緩存，或配置為16KB的共享內(nèi)存和48KB的L1緩存。

圖1.10 kepler架構(gòu)中的流式多處理SMX

這里簡要介紹市面的主力GPU產(chǎn)品型號，見表1.1。

表1.1 2014～2016年市場主力GPU產(chǎn)品規(guī)格

2016年4月，英偉達(dá)在GPU技術(shù)會議（GPU Technology Conference）上發(fā)布了新一代Pascal架構(gòu)和旗艦產(chǎn)品Tesla P100，引入一些新特性。

? 極致性能：為高性能計算、深度學(xué)習(xí)等計算領(lǐng)域設(shè)計。雙精度浮點(diǎn)峰值5.3 Tflops，單精度浮點(diǎn)峰值10.6 Tflops，專為深度學(xué)習(xí)設(shè)計的半精度浮點(diǎn)峰值達(dá)到驚人的21.2 Tflops。按雙精度峰值對比，Tesla P100是同期主力高端CPU英特爾至強(qiáng)E5-2680 v4的10倍。按照深度學(xué)習(xí)應(yīng)用性能對比，Tesla P100是E5-2680 v4的20倍。

? NVLink：為應(yīng)用擴(kuò)展性全新設(shè)計的高速、高帶寬互連協(xié)議。

? HBM2：快速、大容量、高效片上堆疊式內(nèi)存。

? 統(tǒng)一內(nèi)存：用統(tǒng)一的代碼來管理主機(jī)CPU內(nèi)存和GPU內(nèi)存，方便開發(fā)代碼。

Pascal最強(qiáng)勁的是GP100硬件架構(gòu)（圖1.11），包含60個Pascal流式多處理器和8個512位內(nèi)存控制器（共4096位）。每個流式多處理器擁有64個CUDA核心和4個紋理單元。GP100共有3840個單精度核和240個紋理單元。每個內(nèi)存控制器附帶512KB的L2緩存，每個HBM2堆疊內(nèi)存由一對內(nèi)存控制器管理。L2緩存共計4096KB。Tesla P100共配置有56個流式多處理器。

圖1.11 GP100架構(gòu)全景圖

GP100的第6代流式多處理器架構(gòu)提升了CUDA核心的利用率和能源效率，可以運(yùn)行在更高的頻率上。每個流多處理器包含64個單精度（FP32）CUDA核心，分成兩部分，每部分32個；含有1個指令緩沖器，1個Warp調(diào)度器，2個分發(fā)單元。與前面的Kepler、Maxwell流式多處理器相比，每個核心分到的寄存器數(shù)量增多，從而可以運(yùn)行更多的線程。由于增加了流式多處理器，共享內(nèi)存總量也相應(yīng)增加，合并帶寬也翻番。每個流式多處理器中更多的共享內(nèi)存、寄存器、Warp，使代碼可以更高效地執(zhí)行。Pascal架構(gòu)中的FP32 CUDA核心新增一項(xiàng)功能：既能處理16位精度又能處理32位精度的指令和數(shù)據(jù)，F(xiàn)P16操作的吞吐率是FP32的2倍。

如圖1.12所示，每個流式多處理器擁有32個雙精度浮點(diǎn)單元，是單精度核心的一半。共享內(nèi)存獨(dú)占64KB，不再與L1緩存相互調(diào)配。紋理和L1緩存共用一塊緩存，可以靈活配置。

圖1.12 Pascal GP100流式多處理器

Tesla P100的高帶寬內(nèi)存（High Bandwidth Memory 2, HBM2）跟前面幾代相比有顯著變化：從GPU芯片封裝的外部移到了內(nèi)部（圖1.13），安裝在同一塊基板上，距離更近，傳輸更快。由于在豎直方向進(jìn)行堆疊，占用的芯片面積更小，16GB的容量原來需要十幾個GDDR5顆粒，現(xiàn)在只需要4個HBM2顆粒。

HBM2內(nèi)存帶來的另一個好處是原生支持錯誤校驗(yàn)碼（Error Correcting Code, ECC）功能。有些應(yīng)用會用到很多塊GPU，計算時間也很長，對數(shù)據(jù)差錯十分敏感，一旦中間步驟有微小誤差，后續(xù)計算就會將誤差迅速放大，污染最終結(jié)果。ECC技術(shù)能夠檢測并糾正1位差錯。

圖1.13 Tesla P100正視圖

先前使用的GDDR5顯存內(nèi)部不提供ECC功能，僅能偵測GDDR5總線上的錯誤，內(nèi)存控制器中的錯誤或DRAM自身的錯誤都不能偵測。Kepler架構(gòu)GPU的ECC功能實(shí)現(xiàn)方法是劃出6.25%的容量專門用來存放校驗(yàn)數(shù)據(jù)，從而內(nèi)存帶寬將損失12%～15%。

由于HBM2自帶ECC功能，Tesla P100就不再損失內(nèi)存容量和內(nèi)存帶寬。Tesla P100的寄存器、共享內(nèi)存、L1緩存、L2緩存和HBM2內(nèi)存都能偵測2位錯誤并糾正1位錯誤。

1.1.2 英特爾至強(qiáng)融核處理器

2012年，英特爾推出至強(qiáng)Phi融核處理器，代號Knights Corner，融核處理器又稱為MIC（Many Integrated Core）。物理形態(tài)上跟GPU一樣是PCIe板卡（圖1.14），分為被動散熱和主動散熱（帶風(fēng)扇），具體分為6個型號（表1.2），市面上常用的型號是5110P和7120。

圖1.14 Intel Xeon Phi協(xié)處理器

表1.2 英特爾融核處理器技術(shù)規(guī)格

此代架構(gòu)MIC擁有最多63個核心（圖1.15），每個核心有512KB的L2緩存，所有核心使用雙向環(huán)互連。每個核心有一個512位向量處理單元，支持4個硬件線程，線程并發(fā)隱藏延時。

圖1.15 Knights Corner架構(gòu)

MIC上編程模式有4種（圖1.16）：主機(jī)模式（Multicore Only）、卸載模式（Multicore Hosted with Many-core Offload）、對稱模式（Symmetric）、原生模式（Many-Core Only）。主機(jī)模式中只使用主機(jī)CPU，不在MIC上運(yùn)行代碼；卸載模式跟GPU的模式類似，將部分代碼卸載到MIC上運(yùn)行；對稱模式將CPU核心和MIC核心一視同仁，但是兩種核心的運(yùn)算能力不一樣，極易導(dǎo)致負(fù)載不均；原生模式只使用MIC運(yùn)行程序，CPU空閑。實(shí)際上，常用的還是卸載模式和原生模式。

圖1.16 MIC的運(yùn)行模式

此代MIC的理論性能是同時期主流CPU的3.2～3.45倍，實(shí)測性能2.2～2.9倍。因此，使用MIC編寫運(yùn)行程序不要期待過高的加速比。

第二代MIC的代號為Knights Landing，開發(fā)版于2016年4月開始發(fā)貨，已經(jīng)不是PCIe板卡，而是可以獨(dú)立運(yùn)行的CPU。

Knights Landing由36片組成（圖1.17），每片包含2個CPU核心，每個核心包含2個VPU（Vector Processing Units，向量處理器單元）。物理形態(tài)上，不再是協(xié)處理器，而是真正的中央處理器，與通常的CPU一樣使用。每個CPU核心可以運(yùn)行4個線程，芯片的亂序執(zhí)行也有實(shí)質(zhì)性提升，每個線程的性能提高約3倍，整個處理器的理論峰值約為3Tflops。芯片上有16GB MCDRAM（Multi-Channel DRAM）內(nèi)存，6個內(nèi)存通道支持高達(dá)384GB DDR4外部內(nèi)存。與Knights Corner中的環(huán)狀拓?fù)浣Y(jié)構(gòu)不同，新款處理器中使用了網(wǎng)狀的拓?fù)浣Y(jié)構(gòu)。盡管圖1.17顯示為6×7，但是由于封裝方面的一些問題，布局上可能會調(diào)整為4×9。

圖1.17 英特爾至強(qiáng)融核處理器Knights Landing

與2016年2月發(fā)布的英特爾至強(qiáng)系列CPU的主力型號E5-2680 v4（主頻2.4GHz,14核心，理論峰值537.6Gflops）相比，Knights Landing的理論峰值提高至約5.6倍。

1.2 并行編程語言

在并行計算發(fā)展史上出現(xiàn)過多種并行編程語言，至今仍在使用的只剩幾種，它們各有特色。

（1）Pthreads

20世紀(jì)70年代，貝爾實(shí)驗(yàn)室發(fā)明了UNIX，并于20世紀(jì)80年代向美國各大高校分發(fā)V7版的源碼以做研究。加利福尼亞大學(xué)伯克利分校在V7的基礎(chǔ)上開發(fā)了BSD UNIX。后來很多商業(yè)廠家意識到UNIX的價值也紛紛以貝爾實(shí)驗(yàn)室的System V或BSD為基礎(chǔ)來開發(fā)自己的UNIX，較著名的有Sun OS、AIX、VMS。隨著操作系統(tǒng)的增多，應(yīng)用程序的適配性工作越來越繁重。為了提高UNIX環(huán)境下應(yīng)用程序的可遷移性，電氣和電子工程師協(xié)會（Institute of Electrical and Electronics Engineers, IEEE）設(shè)計了POSIX標(biāo)準(zhǔn)。然而，POSIX并不局限于UNIX，許多其他的操作系統(tǒng)也支持POSIX標(biāo)準(zhǔn)。POSIX.1已經(jīng)被國際標(biāo)準(zhǔn)化組織所接受，POSIX已發(fā)展成為一個非常龐大的標(biāo)準(zhǔn)族，一直處在發(fā)展之中。

POSIX線程（POSIX threads, Pthreads），是線程的POSIX標(biāo)準(zhǔn)。該標(biāo)準(zhǔn)定義了創(chuàng)建和操縱線程的一整套接口。在類UNIX操作系統(tǒng)（UNIX、Linux、Mac OS X等）中，都使用Pthreads作為操作系統(tǒng)的線程。Pthreads用來開發(fā)與操作系統(tǒng)緊密相關(guān)的應(yīng)用程序，管理粒度很細(xì)，例如線程的創(chuàng)建與銷毀、線程鎖、線程屬性、線程優(yōu)先級、線程間通信等瑣碎操作均需要程序員安排。對科學(xué)與工程類計算程序來說，程序員的精力應(yīng)集中在業(yè)務(wù)模型和代碼算法上，不應(yīng)浪費(fèi)在底層代碼細(xì)節(jié)上。

雖然Pthreads不適合編寫高性能計算程序，但它多線程并發(fā)的設(shè)計理念啟發(fā)了其他并行語言。

（2）OpenMP

OpenMP是由一些大型IT廠商和一些學(xué)術(shù)機(jī)構(gòu)組成的非盈利組織，官網(wǎng)是www.openmp.org。永久成員包括AMD、CAPS-Entreprise、Convey Computer、Cray、Fujitsu、HP、IBM、Intel、NEC、NVIDIA、Oracle Corporation、Red Hat、ST Microelectronics、Texas Instruments；正式成員是對OpenMP標(biāo)準(zhǔn)感興趣，但不生產(chǎn)銷售相關(guān)產(chǎn)品的組織，例如ANL、ASC/LLNL、BSC、cOMPunity、EPCC、LANL、NASA、ORNL、RWTH Aachen University、SNL-Sandia National Lab、Texas Advanced Computing Center、University of Houston。

計算熱點(diǎn)都是在循環(huán)上，OpenMP的并行化思路是將循環(huán)的迭代步分?jǐn)偟蕉鄠€線程上，每個線程只承擔(dān)一部分計算任務(wù)，循環(huán)運(yùn)行的墻上時間（從開始到結(jié)束的自然流逝時間）自然也就減少了。分割方法是在循環(huán)上面添加一些預(yù)處理標(biāo)記（圖1.18），編譯器識別到這些標(biāo)記以后，將關(guān)聯(lián)的循環(huán)翻譯成并行代碼，然后再與剩余的串行代碼合并起來編譯、鏈接成可執(zhí)行程序。

圖1.18 OpenMP的并行化模式

使用OpenMP編譯并行程序時，程序員需要先保證串行代碼正確，找出熱點(diǎn)循環(huán)，然后在循環(huán)上添加OpenMP預(yù)處理標(biāo)記。打開編譯器的openmp選項(xiàng)可以編譯成并行版本，否則編譯器將忽略預(yù)處理器標(biāo)記，仍然編譯成串行版本。既不破壞原有代碼，開發(fā)速度又快，省時省力。

（3）CUDA

CUDA（Compute Unified Device Architecture）是英偉達(dá)公司設(shè)計的GPU并行編程語言，一經(jīng)推出就引發(fā)了GPU通用計算研究熱潮。CUDA是閉源的，只能運(yùn)行在英偉達(dá)的產(chǎn)品上。CUDA C/C++是對C/C++語言的擴(kuò)展，添加了一些數(shù)據(jù)類型、庫函數(shù)，并定義一種新的函數(shù)調(diào)用形式。CUDA起初支持C和少量C++特性，后來逐漸提高對C++的支持度。從CUDA 3.0開始與PGI合作支持Fortran。CUDA語言還可以細(xì)分為CUDA C/C++、CUDA Fortran，本書成稿時的CUDA C/C++最新版本是7.5,8.0版本即將發(fā)布。CUDA Fortran沒有版本號，只是隨著PGI編譯器的升級而增加新特性。1.3節(jié)會介紹CUDA C/C++的編程模型和一些示例代碼，CUDA Fortran的詳細(xì)情況可以參考官網(wǎng)https://developer.nvidia.com/cuda-fortran，網(wǎng)絡(luò)上有英偉達(dá)工程師撰寫的圖書《Best Practices for Efficient CUDA Fortran Programming》和中文翻譯版《CUDA-Fortran高效編程實(shí)踐》，此處不展開介紹。

（4）OpenCL

OpenCL（Open Computing Language，開放計算語言）是一個面向異構(gòu)系統(tǒng)并行編程的免費(fèi)標(biāo)準(zhǔn)，支持多種多樣的設(shè)備，包括但不限于CPU、GPU、數(shù)字信號處理器（DSP）。OpenCL的優(yōu)勢是一套代碼多處運(yùn)行，只要為新的設(shè)備重新編譯代碼就可以運(yùn)行，移植方便。

OpenCL由蘋果公司首先提出，隨后Khronos Group成立相關(guān)工作組，以蘋果草案為基礎(chǔ)，聯(lián)合業(yè)界各大企業(yè)共同完成了標(biāo)準(zhǔn)制定工作，工作組的成員來自各行各業(yè)，且都是各自領(lǐng)域的領(lǐng)導(dǎo)者，成員名單請參見官網(wǎng)www.khronos.org/opencl/。

（5）OpenACC

這里不多說，后面會全面、詳細(xì)講解。

1.3 CUDA C

本節(jié)簡要介紹CUDA C編程的相關(guān)概念，使讀者能夠看懂OpenACC編譯過程中出現(xiàn)的CUDA內(nèi)置變量，理解并行線程的組織方式。如果讀者已有CUDA編程經(jīng)驗(yàn)，請?zhí)^。

CPU用得好好的，為什么要費(fèi)心費(fèi)力地改寫程序去到GPU上運(yùn)行呢？只有一個理由：跑得更快。小幅的性能提升吸引力不夠，必須有大幅提升才值得采購新設(shè)備、學(xué)習(xí)新工具、設(shè)計新算法。從圖1.19可以看出，在雙精度浮點(diǎn)峰值和內(nèi)存帶寬這兩個關(guān)鍵指標(biāo)上，GPU的性能都達(dá)到同時期主力型號CPU的5～7倍。如果利用得當(dāng)，可以預(yù)期獲得5～7的性能提升。以前只在CPU上運(yùn)行，計算方法的數(shù)學(xué)理論和程序代碼實(shí)現(xiàn)已經(jīng)迭代發(fā)展多年，花很大力氣才能提速10%～20%，提速50%已經(jīng)很厲害了。簡單粗暴地更換硬件設(shè)備就能立刻提速幾倍，全世界的科學(xué)家、工程師一擁而上，GPU加速的應(yīng)用遍地開花。注意，評價GPU應(yīng)用性能的時候，至少要和2顆中高端CPU相對，并且兩種代碼都優(yōu)化到最好。任何超過硬件潛能的加速結(jié)果都是有問題的。

圖1.19 同時期主力CPU與GPU的性能對比

那么問題來了。GPU的芯片面積與CPU差不多，價格也接近，為什么性能這么強(qiáng)悍呢？圖1.20是CPU和GPU芯片的組成示意圖，左邊是一個單核超標(biāo)量CPU,4個算術(shù)邏輯單元（ALU）承擔(dān)著全部計算任務(wù)，卻只占用一小部分芯片面積。“控制”是指分支預(yù)測、亂序執(zhí)行等功能，占用芯片面積大而且很費(fèi)電。服務(wù)器CPU通常有三級緩存，占用的芯片面積最大，有的型號甚至高達(dá)70%。ALU、控制、緩存都在CPU內(nèi)部，大量內(nèi)存條插在主板上，與CPU通過排線相連。GPU中絕大部分芯片面積都是計算核心（4行緊挨著的小方塊，每行12個），帶陰影的水平小塊是控制單元，控制單元上面的水平條是緩存。

圖1.20 CPU（左）和GPU（右）的芯片面積占用情況

通用CPU對追蹤鏈表這樣擁有復(fù)雜邏輯控制的程序運(yùn)行得很好，但大規(guī)模的科學(xué)與工程計算程序的流程控制都比較簡單，CPU的長處難以施展。為了解釋GPU如何獲得極高的性能，需要先了解一下CPU中的控制、緩存、多線程的作用。

ALU承擔(dān)最終的計算工作，越多越好。“控制”的目標(biāo)是預(yù)取到正確的指令和數(shù)據(jù)以保證流水線不中斷，挖掘指令流里的并行度，讓盡量多的部件都在忙碌工作，從而提高性能。緩存的作用是為了填補(bǔ)CPU頻率與內(nèi)存條頻率的差距、減小CPU與內(nèi)存條之間數(shù)據(jù)延時。目前中高端CPU的頻率在2.0～3.2GHz，而內(nèi)存條的頻率還處于1600MHz、1866MHz、2133MHz，內(nèi)存條供應(yīng)、承接數(shù)據(jù)的速度趕不上CPU處理數(shù)據(jù)的速度。由于ALU到主板上內(nèi)存條的路徑較長，延時高，而如果需要的數(shù)據(jù)已經(jīng)在緩存中，那么就能有效降低延時，提高數(shù)據(jù)處理速度。緩存沒有命中怎么辦？只能到內(nèi)存條上取，延時高。為了進(jìn)一步降低延時，英特爾CPU有超線程功能，開啟后，一個CPU物理核心就變成了兩個邏輯核心，兩個邏輯核心分時間片輪流占用物理核心資源。當(dāng)然了，按時間片切換是有代價的：換出時要保留正在運(yùn)行的程序的現(xiàn)場，換入時再恢復(fù)現(xiàn)場以便接著上次繼續(xù)運(yùn)行。在緩存命中率比較低的情況下，超線程功能能夠提高性能。

GPU天生是為并行計算設(shè)計的：處理圖像的大量像素，像素之間相互獨(dú)立，可以同時計算，而且沒有復(fù)雜的流程跳轉(zhuǎn)控制。正如圖1.19所示，GPU的大部分芯片面積都是計算核心，緩存和控制單元很小，那么它是怎么解決分支預(yù)測、亂序執(zhí)行、數(shù)據(jù)供應(yīng)速度、存取數(shù)據(jù)延時這些問題的呢？

GPU的設(shè)計目標(biāo)是大批量的簡單計算，沒有復(fù)雜的跳轉(zhuǎn)，因此直接取消分支預(yù)測、亂序執(zhí)行等高級功能。更進(jìn)一步，多個計算核心（例如32個）共用一個控制單元再次削減控制單元占用的芯片面積。這樣做的效果就是：發(fā)射一條指令，例如加法，32個計算核心步調(diào)一致地做加法，只是每個計算核心操作不同的數(shù)據(jù)。如果只讓第1個計算核心做加法，那么在第1個計算核心做加法運(yùn)算的時候，剩余的計算核心空閑等待。這種情形下資源浪費(fèi)，性能低下，要盡量避免。讓大量計算核心空轉(zhuǎn)的應(yīng)用程序不適合GPU，用CPU計算性能更好。

計算核心與顯存之間的頻率差異如何填補(bǔ)？特別簡單，降低計算核心的頻率?？紤]到芯片功耗與頻率的平方近似成正比，降低頻率不但能解決數(shù)據(jù)供應(yīng)速度問題，而且能降低GPU的功耗，一舉兩得。從表1.1可以看出GPU產(chǎn)品的頻率在562～875MHz，遠(yuǎn)低于主力CPU的2.0GHz～3.2GHz。

最重要是延時，GPU的緩存那么小，怎么解決訪問顯存的巨大延時呢？答案是多線程，每個計算核心分?jǐn)?0個以上的線程。執(zhí)行每條指令之前都要從就緒隊(duì)列中挑選出一組線程，每組線程每次只執(zhí)行一條指令，執(zhí)行完畢立即到后面排隊(duì)。如果恰巧碰上了延時較多的訪存操作，那么該線程進(jìn)入等待隊(duì)列，訪存操作完成后再轉(zhuǎn)入就緒隊(duì)列。只要線程足夠多，計算核心總是在忙碌，隱藏了訪存延時。有人立刻會問，這么頻繁地切換線程、保存現(xiàn)場、恢復(fù)現(xiàn)場也需消耗不少時間吧，會不會得不償失呢？實(shí)際上GPU線程切換瞬間完成，這是因?yàn)槊總€線程都有一份獨(dú)占資源（例如寄存器），不需要保存、恢復(fù)現(xiàn)場，線程切換只是計算核心使用權(quán)的轉(zhuǎn)移。

1.3.1 線程組織方式

一塊GPU上有幾千個核心，每個核心都能運(yùn)行10個以上線程，可見線程數(shù)量龐大，需要按照一定結(jié)構(gòu)組織起來，方便使用和管理。所有的線程合在一起稱為一個網(wǎng)格（grid），網(wǎng)格再剖分成線程塊（block），線程塊包含若干線程。圖1.21中的線程按照二維形式組織，網(wǎng)格包含2×3個線程塊，每個線程塊又包含3×4個線程。實(shí)際上，線程還可以按照一維、三維形式組織。

圖1.21 線程網(wǎng)格與線程塊

既然線程能夠以不同的形式組織起來，那么每個線程都要有一個唯一的編號。為此CUDA C引入了一個新的數(shù)據(jù)類型dim3。dim3相當(dāng)于一個結(jié)構(gòu)體，3個成員分別為：

        unsigned int x;
        unsigned int y;
        unsigned int z;

dim3類型變量的3個成員的默認(rèn)值都是1。網(wǎng)格尺寸用內(nèi)置變量gridDim表示，gridDim. x、gridDim.y、gridDim.z分別表示x、y、z方向上的線程塊數(shù)量；網(wǎng)格中每個線程塊的編號用內(nèi)置變量blockIdx表示，blockIdx.x、blockIdx.y、blockIdx.z分別表示當(dāng)前線程塊在x、y、z方向上的編號，從0開始編號；線程塊的尺寸用內(nèi)置變量blockDim表示，blockDim.x、blockDim.y、blockDim.z分別表示當(dāng)前線程塊在x、y、z方向上擁有的線程數(shù)量；任意一個線程塊內(nèi)的線程編號用內(nèi)置變量threadIdx來表示，threadIdx.x、threadIdx.y、threadIdx.z分別表示當(dāng)前線程在x、y、z方向上的編號，從0開始編號。以圖1.21中的網(wǎng)格、線程塊（1,1）、線程塊（1,2）為例，這些內(nèi)置變量的值如表1.3：

表1.3 內(nèi)置變量的取值

1.3.2 運(yùn)行過程

在GPU編程話語體系里，稱CPU為主機(jī)，稱GPU為設(shè)備。圖1.22演示了CUDA C程序的執(zhí)行過程：在帶有設(shè)備的計算機(jī)上，與C語言程序一樣，從主機(jī)開始執(zhí)行，主機(jī)上執(zhí)行串行代碼，并為設(shè)備上的并行計算做準(zhǔn)備，包括數(shù)據(jù)初始化、開辟設(shè)備內(nèi)存、將數(shù)據(jù)復(fù)制到設(shè)備內(nèi)存中。準(zhǔn)備工作完成之后，在主機(jī)上以特殊形式調(diào)用一個在設(shè)備上執(zhí)行的函數(shù)（稱為內(nèi)核，調(diào)用時比C函數(shù)多了一對三尖號），然后設(shè)備執(zhí)行內(nèi)核中的并行代碼。內(nèi)核代碼執(zhí)行完以后，控制權(quán)交還主機(jī)，主機(jī)從設(shè)備上取回內(nèi)核的并行計算結(jié)果，程序繼續(xù)向下執(zhí)行。圖1.22中只畫出一個內(nèi)核，實(shí)際上一個CUDA程序可以包含多個內(nèi)核。

圖1.22 CUDA程序運(yùn)行過程

下面以實(shí)際例子演示CUDA C代碼的編寫方法和執(zhí)行過程。兩個長度為N的向量a和b對應(yīng)元素相加，將結(jié)果存入向量c。從圖1.23可以看出，N個加法操作之間沒有依賴關(guān)系，可以并行計算。實(shí)現(xiàn)代碼見例1.1。

圖1.23 兩個向量的對應(yīng)元素相加

【例1.1】addvec.cu：向量并行相加。

      1  #include<stdio.h>
      2  #define N 64
      3
      4  __global__ void add( int ＊a_d, int ＊b_d, int ＊c_d ) {
      5      int tid = blockIdx.x ＊ blockDim.x + threadIdx.x;
      6      if (tid < N) c_d[tid] = a_d[tid] + b_d[tid];
      7  }
      8  int main()
      9  {
      10    int a[N], b[N], c[N];
      11    int ＊a_d, ＊b_d, ＊c_d;
      12    cudaMalloc((void＊＊)&a_d, N ＊ sizeof(int));
      13    cudaMalloc((void＊＊)&b_d, N ＊ sizeof(int));
      14    cudaMalloc((void＊＊)&c_d, N ＊ sizeof(int));
      15    for(int i=0; i<N; i++)
      16    {
      17      a[i] = 1;
      18      b[i] = 2;
      19    }
      20    cudaMemcpy(a_d, a, N＊sizeof(int), cudaMemcpyHostToDevice);
      21    cudaMemcpy(b_d, b, N＊sizeof(int), cudaMemcpyHostToDevice);
      22    dim3 block(32,1,1), grid;
      23    grid.x = (N+block.x-1)/block.x;
      24    add<<<grid, block>>>(a_d, b_d, c_d);
      25    cudaMemcpy(c, c_d, N＊sizeof(int), cudaMemcpyDeviceToHost);
      26
      27    for(int i=0; i<N; i++)
      28      printf("%2d +%2d =%2d\n", a[i], b[i], c[i]);
      29    cudaFree( a_d );
      30    cudaFree( b_d );
      31    cudaFree( c_d );
      32  return 0;
      33  }

例1.1中第10行定義3個主機(jī)向量a、b、c，第11行定義3個指針用于存放設(shè)備向量，第12～14行為3個設(shè)備向量分配設(shè)備內(nèi)存空間。第15～19行的循環(huán)為主機(jī)向量a、b賦初值，第20～21行使用內(nèi)置函數(shù)cudaMemcpy將主機(jī)向量a和b中的元素值復(fù)制到設(shè)備向量a_d和b_d之中，即從主機(jī)內(nèi)存復(fù)制到設(shè)備內(nèi)存。第22行定義了2個dim3變量block和grid。block用于指定每個線程塊的形狀：一維，x方向長度為32; grid用于指定線程網(wǎng)格的形狀：一維，x方向的尺寸用block.x和N計算出來，以適應(yīng)N不能被32整除的情形。至此，準(zhǔn)備工作完畢。

第24行從主機(jī)調(diào)用內(nèi)核add，三尖號＜＜＜＞＞＞里的參數(shù)稱為執(zhí)行配置，第1個參數(shù)指定線程網(wǎng)格的形狀，第2個參數(shù)指定線程塊的形狀，緊跟著的圓括號里面是和C函數(shù)一樣的實(shí)參。執(zhí)行配置參數(shù)要求啟動2個線程塊共64個線程來執(zhí)行內(nèi)核add。內(nèi)核add在設(shè)備上運(yùn)行，它將設(shè)備向量a_d和b_d并行相加，結(jié)果存入設(shè)備向量c_d。內(nèi)核add的定義在第4～7行，第4行上的修飾符__global__表示該函數(shù)需要在主機(jī)上調(diào)用且在設(shè)備上執(zhí)行。第5行計算線程的全局編號，N為64，每個線程塊有32個線程，因此網(wǎng)格中有2個線程塊。在每個線程塊中，線程的本地編號threadIdx.x分別是0,1,2, …,31, blockDim.x的值為32，所以執(zhí)行內(nèi)核的64個線程的tid分別為0,1,2, ...,63，見圖1.24。第6行也被64個線程同時執(zhí)行，每個線程執(zhí)行1次加法，共同完成兩個向量的對應(yīng)相加。

圖1.24 線程在線程塊內(nèi)的編號和全局編號

第25行將設(shè)備上的計算結(jié)果復(fù)制回主機(jī)內(nèi)存，即把向量c_d的元素值復(fù)制到向量c中。第27～28行輸出計算結(jié)果以便檢驗(yàn)正確性，可以預(yù)見是64行1+2=3。第29～31行釋放設(shè)備內(nèi)存。

在已經(jīng)部署CUDA C開發(fā)工具的Linux環(huán)境上編譯、運(yùn)行：

        $ nvcc -o addvec.exe addvec.cu
        $ ./addvec.exe
          1 + 2 = 3
          1 + 2 = 3
          1 + 2 = 3
        【共64行，后面省略】

1.3.3 內(nèi)存層級

從1.1.1節(jié)的硬件架構(gòu)圖中已經(jīng)看到，GPU中有多種內(nèi)存：處于芯片外部的全局內(nèi)存（Global Memory)，芯片內(nèi)部的共享內(nèi)存（Shared Meory）、寄存器（Register）、紋理內(nèi)存、常量內(nèi)存、L1緩存、L2緩存。每種內(nèi)存都有不同的特性，有不同的使用技巧。對開發(fā)CUDA程序最重要的三種內(nèi)存分別是寄存器、共享內(nèi)存和全局內(nèi)存。

如圖1.25所示，每個線程都有自己專用的寄存器，從內(nèi)核開始時，一旦擁有某個寄存器的使用權(quán)，就一直獨(dú)占，直到內(nèi)核結(jié)束才釋放，從而線程之間無法通過寄存器交換數(shù)據(jù)。雖然有大量的寄存器，但也有大量的線程，平均下來每個線程只能分配到幾十個至幾百個寄存器，復(fù)雜程序仍然要控制線程消耗的寄存器數(shù)量。每個線程塊都能分配一塊共享內(nèi)存，本塊內(nèi)的線程可以訪問這塊共享內(nèi)存的任意位置，因此可以用共享內(nèi)存來交換數(shù)據(jù)。一個線程塊不能訪問其他線程塊的共享內(nèi)存，因而線程塊之間不能用共享內(nèi)存交換數(shù)據(jù)。共享內(nèi)存容量比寄存器要大，例如Tesla P100的每個流式多處理器擁有64KB共享內(nèi)存，每個線程塊最多可以擁有32KB。所有的線程塊、線程網(wǎng)格都能訪問全局內(nèi)存，只要不顯式地釋放或者程序結(jié)束，全局內(nèi)存中的數(shù)據(jù)會一直存在，因此可以用于線程塊之間、線程網(wǎng)格之間的數(shù)據(jù)交換。全局內(nèi)存更大，以GB為單位。

不同內(nèi)存的訪問延時差別很大，寄存延時最小，共享內(nèi)存次之，全局內(nèi)存最大。對Pascal之前的架構(gòu)，全局內(nèi)存與GPU芯片相互分離，通過板卡上的排線相連，訪問延時達(dá)到幾百個時鐘周期。Pascal架構(gòu)中，全局內(nèi)存與GPU芯片距離很近，延時應(yīng)該有大幅減小，具體數(shù)值還需要等待官方公布。

圖1.25 GPU的內(nèi)存層級

不同構(gòu)件下的內(nèi)存層級多少都有些變化，要想使CUDA程序達(dá)到最好性能，必須做針對性優(yōu)化。

1.3.4 性能優(yōu)化技術(shù)

CUDA程序編寫容易，調(diào)優(yōu)不易。程序員能夠掌控很多事情，包括但不限于分配全局內(nèi)存：全局內(nèi)存中的數(shù)據(jù)對齊、維數(shù)，為每個線程塊分配的共享內(nèi)存大小，將哪些數(shù)據(jù)以什么樣的組織方式放入共享內(nèi)存，哪些數(shù)據(jù)放入紋理內(nèi)存，哪些數(shù)據(jù)放入常量內(nèi)存，線程網(wǎng)格如何劃分，線程塊是一維、二維還是三維，線程塊每個維度的大小是多少，線程與數(shù)據(jù)元素的對應(yīng)關(guān)系，不同線程訪問的數(shù)據(jù)是否有沖突，不同線程同時訪問的數(shù)據(jù)是否會走相同的通道；單個內(nèi)核是否能夠用滿資源，如何同時運(yùn)行多個內(nèi)核以提高設(shè)備利用率，有幾個數(shù)據(jù)復(fù)制引擎，如何安排異步隊(duì)列來重疊數(shù)據(jù)的來往傳輸，如何重疊數(shù)據(jù)傳輸與計算，如何填補(bǔ)PCIe帶寬與全局內(nèi)存帶寬之間的差異，數(shù)據(jù)復(fù)制操作是否需要錨定主機(jī)內(nèi)存；計算密度夠不夠大，計算核心要等待數(shù)據(jù)多久，一個Warp內(nèi)的線程的流程分支有多少，多少個線程才能隱藏延時；GPU上的算術(shù)指令與CPU上對應(yīng)指令的差異，雙精度操作、單精度操作、半精度操作、三角函數(shù)等特殊操作的計算資源分配。

管事多，操心就多。每個問題都有相應(yīng)的優(yōu)化方法和一定的約束條件，具體技巧請參考英偉達(dá)官方文檔《CUDA C BEST PRACTICES GUIDE》。需要注意，不同架構(gòu)下的優(yōu)化技術(shù)會有一些差別。

影響最大的優(yōu)化技巧是主機(jī)與設(shè)備間的數(shù)據(jù)傳輸。從圖1.4可以看出，設(shè)備與主機(jī)通過PCIe×16通道相連，在采用2016年發(fā)布的最新CPU的服務(wù)器上，PCIe 3.0×16的理論帶寬為16GB/s，與表1.1中幾百GB/s的顯存（全局內(nèi)存）帶寬差別可達(dá)30倍，與Tesla P100的差別會更大。因此，應(yīng)盡量減少主機(jī)與設(shè)備間的數(shù)據(jù)傳輸量與傳輸次數(shù)。