離線下載 vertex it領域 硬件 體系結構 占用 oid 過大 crc

第一章 為什麽需要並行程序

CUDA，全稱是Compute Unified Device Architecture，一般翻譯成中文為計算統一設備架構。筆者以為這樣的名字會讓人對CUDA感到很迷惑，CUDA到底是什麽呢？筆者用自己的大白話來說下自己對CUDA的理解，CUDA就是一個基於GPU（Graphics processing unit）（目前是單指Nvidia公司的）的通用並行計算平臺。這裏有3個關鍵字，GPU，通用計算和並行！

關於GPU，相信它是什麽，不用多說，不過關於CUDA的硬件架構後面會有一些分析，因為要寫出高質量的CUDA程序，不了解底層運行機制是不夠的，這就是侯捷講的“勿在高臺築浮沙”，這裏說一點題外話，筆者經常在同行老朋友面前很得意的吹噓：“你們這些家夥寫的程序都是在CPU

上跑的，太土了！俺的可是在GPU上跑的哦”，所以嘛，學習CUDA是很有前途的！

第二個重要概念是通用計算，這個主要是與以前的GPU只做渲染流水線相對應的，由於筆者比較年輕，沒有經歷過師兄他們走過的黑暗歲月，呵呵！傳統的GPU架構需要按一個標準的流水線編程，要經過vertex processor，fragment processor和pixel operation，這會使編程變得困難和不容易控制。早期的GPGPU（General Purpose Computing on GPU通用計算為目的的GPU）也需要按照那樣的標準的流水線編程，而CUDA的出現對我們學軟件的來講，才標誌著真正意義上GPU通用計算的到來！CUDA

從技術本身來講，僅僅是C/C++的一個小超集（這個後面會詳細分析），從軟件工程的角度看，降低了學習成本，加之其靈活性，必將在工程應用中發揮巨大的作用！在美國幾乎所有頂級學府如哈弗，MIT等都有CUDA實驗室，美國NASA的很多大型項目在多年前都已經使用CUDA技術。基本上，可以說CUDA提供了一個無論是在軟件編程還是硬件處理都適合通用計算的平臺架構。

在學習CUDA編程以前，我們必須深刻認識下CUDA程序中最核心的概念——並行。在操作系統的課程裏面，我們會遇到另外一個名詞——並發。在單核時代或者說對於普通的CPU運行程序，這兩個詞的意思是一模一樣的！但我們應該清楚知道兩者的區別，並發指復用，同一個時刻只有一個指令在處理器上執行；而並行則是在同一時刻多個指令同時運行，在GPU

上處理器的稱謂換成了個多處理器，每個多處理器上有8個名叫CUDA core的處理單元。在本科答辯的時候，有個老師問我，是不是可以說1個多處理器上有8個核（註：他不做CUDA的）。筆者以為作為初學者可以這樣看，但如果想深入了解和掌握高級優化手段就必須明白CPU上的多核不等同於CUDA core，這個問題會在後面章節逐步說明。這裏花點筆墨來說明一下並行與並發的區別，主要是讓讀者明白CPU並行程序和GPU並行程序在底層運行機制的區別，同時初學者往往遇到EmuDebug運行正確，而實際運行卻錯誤的情況。這裏最根本的原因還在於EmuDebug運行模式下是串行而非並行。雖然有很多牛人覺得EmuDebug沒有意義，但我還是建議初學者遇到結果錯誤的時候先保證EmuDebug是正確的，因為絕大多數情況，保證EmuDebug是正確的說明程序邏輯沒有問題（當然也有例外，後面會講到）。

知道了並行的概念，那麽為什麽我們需要並行程序，根據摩爾定律：集成電路上可容納的晶體管數目，約每隔18個月便會增加一倍，性能也將提升一倍，當價格不變時；或者說，每一美元所能買到的電腦性能，將每隔18個月翻兩倍以上，如圖1。他的核心思想就是硬件會越來越越快，可是隨著單核CPU的時鐘頻率不斷增加，也落入了收益遞減規律裏面了，其中一個主要的性能瓶頸是存儲器延遲。因此我們有了新的摩爾定律：No longer get faster, just wider（未來的計算機不會更快，而是更“寬”）。新的需求也給我們這個做軟件的帶來了挑戰和機遇：必須重新設計算法，To be aggressively parallel!筆者慶幸自己在不到兩年的時光發表10來篇EI論文也得益於這樣的變革機遇！在並行處理這個領域，並行體系結構、並行軟件和並行算法三者缺一不可，而其中並行算法則是核心和瓶頸技術，也是我們從事軟件行業的人應該肩負的使命。

在結束這個開場白的時候，我們來比較學術化的看看什麽是並行算法。並行算法是指在各種並行計算機上求解問題和處理數據的算法，其本質是把多任務映射到多處理器中執行，或將現實的多維問題映射到具有特定的拓撲結構的多處理器上求解。一定要牢記的是並行算法的實現強烈的依賴於計算機硬件和軟件環境，這是我們軟件專業出身的人最容易忽視的，筆者個人比較特別的經歷就是在系統開發時的直接合作師兄是硬件工程師，我們經常被對方弄的很崩潰。但自己現在能夠寫出比較高質量的CUDA程序也得益於這些對硬件問題的崩潰和思考！

圖1 CPU chip進化速度

1.1 並行算法的目標

計算需求是永無止境的，可以說高性能計算是計算機科學研究中的“日不落”課題。並行計算是其中最有效的手段。作為軟件編程人員，設計編寫並行算法是最為核心的工作任務。筆者想從3個基本概念：時間重疊、資源重復和資源共享，來讓讀者初步直觀的認識一下並行算法的總體設計目標。

首先，時間重疊是指多個處理過程在時間上相互錯開，輪流重疊地使用同一套硬件設備的各個部分。這個概念可從計算復雜度的角度來理解，一個算法的復雜度可表示為空間復雜度和時間復雜度。從算法樹的結構來看，通常的串行算法樹“深而窄”，因為串行算法的本質是為一維問題設計的。而並行算法的目標則是盡可能減少時間復雜度，通常是通過以空間換時間的方式實現的，即增加空間復雜度。典型的時間重疊就是流水線處理。雖然CUDA平臺上單個GPU暫時是不能設計流水線算法，但它也提供了異步訪問以及Fermi架構的雙kernel調度等時間重疊的處理方式。

除了時間重疊外，資源重復也可以實現將時間復雜度轉化為空間復雜度。資源重復是指設置多個相同的處理器，同時從事處理工作，以數量取勝的方式提高處理速度。

這裏給一個經典的以空間換時間的並行算法例子，尋找一個數列最大值。

我們知道在數列中尋找最大值的一般算法復雜度為O（n）但如果采用並行算法則可為O（1），假定有2n個處理單元，算法如下：

1. int i=Px //範圍0～n-1
2. int j=Py //範圍0～n-1
4. arrB[i] = 1
5. _syncthreads() //線程同步函數
6. if arrA[i] < arrA[j]
7. arrB[i] = 0
8. _syncthreads()
9. if arrB[i] == 1
10. print arrA[i]

這個算法並行的思路分析：是很簡單，當A[i]不是最大值時，B[i]標識為0，結合圖2，相信很容易想明白。看到這個比較神奇的並行算法idea，會否激起你的挑戰欲望呢！

圖2

註意，這裏給出的例子只是一個理論算法，要實現它必須假定處理系統是CRCW（Concurrent Read Concurrent Write（同時讀同時寫））系統而且處理器足夠多（2n個），所以很難在實際中應用，當然也不是在CUDA平臺上解決這一問題的理想算法。但這個經典算法體現了並行算法設計目標與傳統串行程序的本質區別，同時它也揭示了並行算法的理論加速比，關於加速比在章節1.3會有詳細說明。

由此，我們可以看到並行算法樹采用的是與串行算法樹截然不同的“淺而寬”的結構，即每個時刻可容納的計算量相應增加，使整個算法的執行步驟數目盡可能接近問題的關鍵路徑長度，也可以這樣說，通過增加每個時刻步的算法復雜度來減少整體的時間復雜度，從而達到把時間復雜度轉化為空間復雜度的目的。

最後由於數據在處理的過程中涉及到讀取和寫入操作，尤其是CUDA平臺下存在多種訪存類型，如主機端到設備端，GM（global memory全局存儲器）到SM（share memory共享存儲器）等等，資源共享也是並行處理技術重要概念之一，原始的定義是多個處理器按照一定的規則對同一數據資源進行訪問，也包括處理器負載的均衡，即多個服務請求如何映射到多個處理器。筆者以為在CUDA平臺下就是如何通過資源共享來減小訪存延遲，提高存儲器的帶寬（見5.4黃三）。

1.2 並行算法設計的基本方法

在前面的小節裏，我們已經從不同角度認識了什麽是並行處理，並行程序的大體模樣和設計理念。下面筆者將結合一些實例從總體上介紹下並行算法設計的基本方法。並行算法的一般設計方法大致包含以下3個途徑：

（1） 檢測和開發現有串行算法中固有並行性而直接將其並行化

（2） 從問題本身特征出發，設計一個新的並行算法

（3） 修改已有的並行算法使其可求解另一類相似問題

以上三個途徑是最正統的說法，筆者以為凡事從事並行算法設計的同仁們都應當對它們“死去活來”。記住它們，因為它們是設計的鑰匙；靈活運用，因為每一條都只是一個思想而不是一個技巧。筆者在本科畢業設計答辯時還給出了第四把鑰匙：充分利用CUDA平臺的各種技術特性。它不是一個設計準則而是所有設計的基礎，在本書裏面，筆者最強調的是是並行算法的實現強烈的依賴於計算機硬件和軟件環境。為了讓大家更形象化，筆者貼出自己的一頁答辯PPT ，見圖3。

首先，我們來看看第一個路徑，是很長的標準化學術化說法。思想很簡單，就是原先一個處理器一次處理一個，現在多個處理器同時處理多個。但註意，要做到這一點，需要有一個基礎就是計算無關性，有很顯式的，有半遮面的，也有隱式的。首先，我們還是不厭其煩的看看最簡單的一個例子。

for(int i=0;i<256;++i)

X[i]=i;

這個for循環，大家可以看到，計算X[i]與任何的X[i+n]沒有任何的依賴關系，所以可以很輕易的寫出並行處理程序。這裏因為假定大家還沒有CUDA基礎，所以給出一個直觀的偽代碼：

//設置256線程

//讓第i個線程取得值i

//讓第i個線程將取得的值i寫到X[i]這個位置

圖3

上面這個例子是不是太easy了，那麽稍稍加點難度，來看看下面這個例子：

X[0]=0.05;

for(i=1;i<256;i++)

X[i]=X[i-1]+0.05;

從直觀上看，與第一個例子不同，X[i]的得到是基於X[i-1]，這個說法對嗎？筆者以為也對也不對，按傳統的理解，這個程序的執行是相互依賴的。但稍稍想想這個猶抱琵琶半遮面的計算無關性就浮現於眼前了，並行方式如下：

//設置256線程

//讓第i個線程計算（i+1）*0.05

//讓第i個線程將計算結果寫到X[i]這個位置

是不是也很容易呢！好吧，我們把這個例子在推廣一下下，應該也就沒有難度了,至少有了前面的熱身不會太崩潰吧：

X[0]=b[0];

for(int i=1;i<256;++i)

{

for(int j=0;j<i;++j)< span="">

{

b[i] += X[i-j-1]*a[i][j];

}

X[i] = b[i];

}

這是一個線性遞推公式，用數學符號表示如下：

（1）

打了這麽多鋪墊，就不用筆者再說了吧！要並行化就是要將其展開，其實一點都不難，用筆劃劃就找到規律了，好吧，筆者來寫寫展開式。大家註意看看a和b的下標規律，是否想起了，咱們高一上學期數學的最後一章呢！等差數列，如果找到這個規律了，大家就會有感覺了！縱向看看我們的b，是不是階梯形，呵呵！在橫向看看b的系數a，好，我們給出兩個定義式來說明下這個規律。

定義總是讓人覺得學術化，但數學有時候是最直觀的表達，由我們的展開式，可以得到

（2）

到了這裏，這個例子的計算無關性也找到了，最直接的並行化處理（肯定不是最優，比如有大量的冗余計算可以去除等優化方法）可以這樣描述如下：

//設置256線程

//讓第i個線程計算公式（2）

//讓第i個線程將計算結果寫到X[i]這個位置

通過這三個難度等級不同的例子，可以看到挖掘現有算法中的計算無關性，雖然很直接但未必很簡單，而且串行程序的固有並行性出現的方式各式各樣，不一而足，需要我們的數學素養（如姜伯駒院士所講，數學素養是國民素質的重要元素，筆者以為，然也！），同時設計經驗也是我們的重要財富。

下一把鑰匙是從問題本身特征出發，設計一個新的並行算法，這類情況主要是指問題本身不具備計算無關性，或者叠代或者數據相關。這裏以CUDA SDK下的直方圖生成算法為例。

傳統的圖像直方圖算法是通過逐點統計實現的，基本程序算法如下：

for(int i = 0; i < BIN_COUNT; i++)

histogram[i] = 0;

for(int i = 0; i < dataNum; i++)

histogram[image[i]]++;

由於要確定的是整個圖像的灰度分布情況，所以即使在並行處理環境中，不進行叠代運算也是無法生成圖像的直方圖的。所以最顯然的算法就是分而治之，然後歸並，也就是讓每一個線程生成一個局部區域的直方圖，然後加在一起形成全局直方圖。如果不考慮性能，這樣做是完全OK的，但可以很負責任的說，在CUDA平臺下這樣做是不好的，甚至是不行的。我的一個學生在做圖像紋理分析時曾經很無奈的說，無法在CUDA平臺下生成和差直方圖（這個可能需要有點計算機視覺的背景才容易明白，因為對於圖像的和差直方圖可能會需要512個針腳）。

回到這個問題上來，問題到底出在哪裏了呢？因為還沒有完全介紹CUDA，這裏只稍稍分析下，讓某一個線程去處理其中一塊形成一個局部直方圖，最後歸並成一個全局直方圖。對於CUDA平臺這個算法是非常不適宜的，因為GPU平臺的存儲器最優的是共享存儲器而非全局存儲器（見第五章黃三）。在2.0以下的CUDA設備上，共享存儲器的大小為16KB，同時一個有效工作的線程塊應該包含128～256個線程以獲得相對高效的運行性能。那麽如果采用一個線程直接對應於一個子直方圖的話，一個明顯的限制就是16KB平均到最大一個包含192條線程的塊，每個線程只有85bytes，因此這種方法最大情況共享存儲區能適合每線程子直方圖可達64單字節針腳數。更為嚴重的是，單個執行線程所處理的數據大小也受限於字節計數器所規定的255字節，也就是說單個線程不能生成針腳數大於64的子直方圖。從硬件上看，也就是受維護單個線程上下文存儲器資源的限制。那麽SDK下解決這個問題的方法是一個線程單幹不可以，就依靠群體嘛，所謂團結就是力量，即以一個warp塊（32個線程）輸出一個局部區域的子直方圖。多線程維護同一個直方圖，在不使用原子操作的情況下（因為1.1以下設備的共享存儲器不支持原子操作，筆者以為即使使用現在的高端設備，原子操作越少越好）是非常棘手的問題。我們來看看SDK的解決方案：

__device__ void addData256(volatile unsigned int *s_WarpHist,unsigned int data,unsigned int threadTag)

{

unsigned int count;

do{

count = s_WarpHist[data] & 0x07FFFFFFU;

count = threadTag | (count + 1);

s_WarpHist[data] = count;

}while(s_WarpHist[data] != count);

}

這個代碼，除了關鍵字__device__（暫時不需要管它），其它全是普通的C代碼，應該是無障礙理解吧！

算法解決訪問沖突的技巧是很巧妙的，這裏圖像數據值的取值範圍是[0，255]，每個warp線程根據輸入的圖像數據使數組s_WarpHist[]相應位置增加1。算法為了解決warp內的線程沖突，通過最末一個執行寫入的線程將直方圖標識，標識為直方圖頻數變量的高5位，因為warp的大小為32，所以只需要5位。這樣的話，對於每一個線程讀取當前直方圖針腳值高5位就被替代為當前線程的標識符。如果線程間要寫入的不是同一針腳位置，就沒有額外的操作，但當兩個或多個線程在寫同一位置時硬件會執行共享存儲區寫結合，使線程中被標識的頻數被接受而拒絕所有其它未定的線程修改。在寫嘗試後每個線程讀取來自同一個共享內存的位置。而那個可以寫回自己修改的頻數的線程，則退出循環並空閑等待warp中其余的線程。Warp繼續它的執行，直到所有線程退出循環。由於每個warp使用自己的子直方圖並且warp線程總是同步的，所以不依賴於warp的調度順序，這個循環執行的叠代次數不會超過32次，而且只有在warp內所有線程讀到的灰度值相同時才會發生。這個解決訪問沖突的算法主要包含三個步驟

第一步count = s_WarpHist[data] & 0x07FFFFFFU;是將當前該針腳的頻數放入count變量中，並將高位清零。

第二步增加相應針腳頻數，並寫入線程標識符，這裏特別說明的是threadTag = threadIdx.x << (32 - WARP_LOG_SIZE);也就是將線程號經過移位操作放到高5位作為標識符。

第三步則是將增加後的頻數寫回到數組s_WarpHist[]。然後判斷當前寫入的頻數是不是最後修改頻數的線程所寫入的，是就接受，不是則通過循環重新執行上述操作。

在完成了子直方圖計算後就是歸並每一個子直方圖，這個過程應該分兩步首先歸並各個warp的子直方圖到塊直方圖，然後將塊直方圖歸並為全局直方圖。要理解為什麽這個算法要費這麽大的力氣，需要掌握後面的內容，在？.？會有詳細的分析。但筆者這裏給出一個實例是想說明，問題本身不具備計算無關性，無法直接並行化的時候，可以針對問題本身的特征結合並行處理平臺的技術特性來設計一個新的並行算法。

另外一個重要的設計原則，其實就是學會舊瓶裝新酒，或者文言一點，它山之石，可以攻玉。這裏就分析介紹下並行算法中最經典的例子，Fan-in算法。

所謂Fan-in算法求和就是利用多處理器實現樹型求和，圖1.4是一個樹型求和的實例：

圖1.4 Fan-in算法處理實例

可以看到該算法是利用多處理器的並行計算能力，將計算時間縮短，也可以很清楚的看到並行處理程序的特色“淺而寬”。這個算法如果結合CUDA平臺的技術特性，會有多種變化，在後面會有詳細分析。這裏對算法本身做一點延伸介紹，有一種重要的並行設計策略——加速級聯策略。

加速級聯策略的核心思想是將一個最優但不是最快的和一個最快但不最優的級聯起來。具體做法

（1） 開始使用最優算法，直到求解問題規模減少到某一閾值

（2） 接著使用快而非最優的算法，繼續完成問題的求解

將Fan-in算法進行一下加速級聯，算法如下：

假設，N個數a₁，a₂，…，a_n，其中N=P²

第一步：a₁+a₂，…，a_p+1+a_p+2，…,a_N-P+1+a_N-P+2

第二步：(a₁+a₂)+a₃,…，(a_p+1+a_p+2)+a_P+3,…,(a_N-P+1+a_N-P+2)+a_N-P+3

第P-1步：(a₁+a₂+…+a_P-1)+a_p,…,(a_P+1+a_P+2+…+a_2P-1)+a_2P,…,(a_N-P+1+a_N-P+2+…+a_N-1)+a_N

第P步到第P-1+log₂P步按Fan-in算法處理。

如果我們對以上兩種算法做簡要的理論分析，原始的Fan-in算法使用的處理線程個數為N/2，加速比為S_f和效率E_f（定義見1.3）分別為：

而使用加速級聯之後的算法所使用的線程個數為 ，加速比S_z和效率E_z分別為

從上面的式子可以看到，傳統Fan-in算法的加速比高於加速級聯算法，但後者為常數效率的並行算法，也就是效率不隨問題規模和處理器個數的改變而改變，也就是說當問題規模較小，線程切換延時也比較小的時候采用Fan-in算法，程序運行時間更快，當超過CUDA的執行函數（後面知道就是kernel）承擔線程極限，函數切換延時較大時則應采用加速級聯獲得的執行效率更高。

這個地方簡單的介紹了經典的Fan-in算法和加速級聯策略，這些都是常見的並行算法設計技術。在本節中通過幾個例子對並行算法設計的基本方法進行了說明，可以用一句口號來講，一個中心，三項基本原則。

所謂一個中心，就是並行處理平臺，在這裏就是CUDA的各項技術特性，這是一切設計的基礎。三項基本原則，挖掘計算無關性，舊瓶裝新酒和獨立自主創新，既有各自的技術含量，又離不開CUDA平臺這個中心。下面1.3會繼續並行算法設計的另一個重要方面——性能度量，接下來就會從此刻的多核時代（1.4），引出我們故事的主角CUDA（1.5）。

1.3 並行算法性能的度量

應該說在計算機科學裏面，沒有任何一個領域比高性能計算更加需要一個性能度量的標準了。因為它因計算需求而存在，它的價值也就體現在那一點一點的性能提升上面。對於並行程序更是這樣，如果不是為了在便宜的處理器上得到高性價比的性能，何必要費弄個並行呢！所以想做CUDA的同誌們一定要時刻記住，性能就是生產力！所以本節就稍稍分析下並行算法的性能度量。

對於一個給定問題，在設計完成一個並行處理算法之後，對並行處理算法性能的評價一般包括算法運行時間、算法並行度、算法成本、加速比和效率等方面。

（1）並行算法運行時間

這個標準是最直接又最簡單的，首先來看看學術化定義，並行算法的運行時間是算法在並行計算機上求解一個問題所需的時間，多處理器就是最早開始到最晚結尾。在CUDA平臺上一般有兩種主要的時間測量方式，主機端測量和設備端測量（測量方法見5.1）。當然，還有采用CPU的計時器測量（很不精確，筆者已經將其徹底拋棄）。運行時間是我們衡量並行算法是否有用的最重要的指標，是第一生產力！

（2）算法的並行度

關於並行度，個人以為它更加是一種並行算法設計的評價。所謂算法的並行度，就是指該算法中可並行執行的操作數。若處理器資源無限，則算法的並行度可理解為可用來做並行運算的處理器個數。但在實際的算法中並行度不會是固定不變的，只能限制在一定的執行步驟或時間範圍內。並行度刻畫的是一個並行算法的並行程度，反映了軟件並行性與硬件並行性的匹配程度。

這就引出了一個與並行度相關的主要概念是粒度，一般而言，大的粒度意味著能獨立並行運行的是大任務，算法的並行度就小，反之算法的並行度就大。對於CUDA而言，由於線程是輕量級線程，所以在盡可能的情況下讓線程的並行度達到GPU最大的活動量（後面會解釋活動線程個數的概念），也就是達到處理器的最大吞吐量。當然有時候為了減少冗余計算和有效利用存儲器帶寬，也可以使用粗粒度的並行方式，畢竟我們實際的情況是資源有限。

最後給出一個形式化定義，算法的平均並行度定義為假設並行算法可以在m個並行步驟內完成，第i步時算法的並行度為DOP(i)，則算法的平均並行度為

這個定義可以用來評估你算法的一個整體資源使用情況，找到你並行算法的處理瓶頸之處。並行度小的地方，必然存在資源的空閑。我們後面會知道在Tesla架構下一個kernel函數占用整個GPU處理資源，如果並行度小，就會讓處理資源空閑。相反並行度過大的時候，若不考慮CUDA資源的承載能力，就反而會讓程序性能下降甚至執行失敗，比如寄存器和共享存儲器這樣的稀缺資源。

（3）並行算法的成本

經濟學上最經常出現的一個名詞就是成本，當我們踏上並行處理這個靠效益吃飯的領域時，就不能再把自己當作象牙塔裏闊佬了。我們要對節能，減排，經濟，節約有深刻的認識。那麽什麽是並行算法的成本？下面的定義很直觀，它就是並行算法的運行時間與並行算法所需的處理器個數的乘積，即

也就是說，成本等於最壞情況下求解某一問題時總的執行步數，其中包含了硬件和軟件代價。對於我們CUDA程序而言（這裏指單個GPU），每次執行占用的是處理器的全部資源，也就是說P是定值。我們的並行算法成本就是運行時間為主導，我們的目標就是在給定資源平臺下讓我們的時間降下來，成本下降就是效益啊！

（4）加速比，效率與加速比模型

筆者以為加速比是CUDA門外漢（比如說我們的客戶）最關心的，我早年經常拿一些數據去唬人，GPU比CPU加速多少。但其實這個說法很不科學，也沒有意義。由前面的並行算法成本就知道，運行時間才是評價和產生效益的關鍵，我的老板們都是很精明的，現在關心的都是你這個算法幀速率多少？還有沒有可以優化的地方？那為什麽我還要列出這個性能指標呢！因為它的理論模型對我們設計和評估並行算法還是很有借鑒意義的。

首先，加速比是衡量並行處理算法最傳統的評價標準，體現了在並行計算機上運行並行算法求解實際問題所獲得的效益。

加速比被定義為

其中T_l是最優串行算法在單處理器上的運行時間；T_p是並行算法在P個處理器上所需的時間。由這個定義就可以得到理論的最大加速比，這是我們自己給自己設定目標，以及在談判桌上可以用到的。事實上，對於一個問題，如果按照某種條件，保持每個處理器的計算規模一定，並行算法的加速比S_p與處理器個數P成正比，則稱該並行算法在該條件下，在該並行計算機上具有線性加速比。若在某些條件下，S_p＞P，則稱該並行算法在該條件下，在該並行計算機上具有超線性加速比。

但上面對加速比的定義比較適合用於理論分析，一般理論分析也比較困難，所以在實際應用中往往使用下式：

一般的人只了解加速比這個概念，但其實與之密不可分的重要性能指標並行處理效率卻容易被忽視。

首先，並行處理效率的定義是

其中P是處理器的個數。

我們來稍稍分析下，加速比其實是一個忽視了處理資源的評價，舉個例子同一個CUDA程序在GT9600上（8個多處理器）加速40倍，在GTS 250上（16個多處理器）是不是就會80倍呢？由前面的分析可以知道，我們所期望的是處理效率高的線性加速比或超線性加速比，那麽就不能脫離當前的計算資源來評價加速效果。

理解了加速比和效率這兩個定義，後頭去看看前面關於加速級聯策略的分析，就會很清楚這些評價指標對設計的重要意義。最後再來說說兩個經典的並行加速比模型，了解它們的目的在於認識到並行算法如何突破一些設計瓶頸的思路。

首先，並行加速比模型作為一個度量並行處理性能的參數，用來表示並行求解一個實際問題所獲得的性能，即相對單處理器上的串行處理而言使用並行處理所獲得的性能。從問題規模的角度出發可將其分為固定規模問題和可變規模問題的加速比模型。

①Amdahl模型

其中f是串行所占比例，N為並行所占比例

由此模型可以看出無論處理器數目如何，加速比都不能超過1/f，這就意味著在當前問題的計算需求下，無論你如何增加處理器的個數都無法繼續增大加速比。例如在我們遇到一些無法展開的叠代運算時，無論你的GPU上有多少個多處理器都無能為力。但不要過於悲觀，此模型忽視了問題計算規模的變化，強調的是通過並行處理來縮短求解問題的時間。再看看下一個模型之前，我們要明白，Amdahl模型告訴我們在問題規模一定的情況下，我們要盡可能的讓計算負載分布到所有的計算資源上去，通過最大化並行提高加速比。

②Gustafson模型

與Amdahl模型不同，此模型說明了應該隨處理器數目的增加而增加問題的規模，強調的是在同樣的時間內，通過並行處理能運行多大的運算量，即通過運行時間來限制問題規模的增長程度。剛剛那個很悲觀的例子，就可以采用增加運算規模來提高加速比，事實上在實時系統上這是有現實意義的。比如CUDA上PCI-Express是一個突出的性能制約因素，但通過大塊數據傳輸可以提高處理效率；對於一些並行度很小的處理環節通過增加數據規模來提高加速性能，比如原先一個並行度小的kernel只處理一幀，現在合並處理多幀，這也可以提高處理效率和整體加速比。

雖然Amdahl模型和Gustafson模型只是是從不同角度去看待並行處理，但卻告訴我們並行算法在設計和評價中的一些重要思想。比如要最大化利用計算資源，但不要盲目增加處理器；在計算規模可變的情況下，算法的串行瓶頸可以通過數據規模的增加而忽略等。

1.4 多核時代的各路諸侯

隨著新摩爾定律的產生，我們已經進入到了一個多核時代。本人所在學校的外面不遠處有個網吧，取名為四核時代。網吧都如此了，何況在我們這些專業的IT領域裏面。本章的內容第一目的在於科普教育，讓我們一起看看當今的多核技術發展，然後筆者希望通過一些比較讓我們一起來看看多核GPU到底有些什麽與眾不同的地方。

首先，在多核處理器產生以前的並行處理主要以並行計算機和借助網絡實現的大規模集群或分布式並行為主的兩大派系。20世紀60年代初期，由於晶體管以及磁芯存儲器的出現，處理單元的面積小型化，存儲器也更小巧和廉價。這一時期出現了規模不大的共享存儲多處理系統，就是我們所謂的大型機，典型代表IBM 360。與之相對應的是通常由數百數千甚至更多的處理器（機）組成的超級計算機，比如我國的天河-1A，速度全球第一，比第二名的美國國家實驗室的計算機快30%，速度達到每秒2.5千萬億次運算。這兩大類可以說是高性能計算領域的巨無霸。

另外一類可與之爭鋒的是借助網絡實現的大規模集群或分布式並行計算。由於網絡技術的高速發展，出現了以獨立計算機連接組成的分布式並行處理系統——集群計算機。一個集群系統中的計算機節點可以是在一起的，也可以是物理上分離的，他們對於用戶和應用程序來講是透明的，如果只有一個單一的系統一樣。這樣就可以提供高性價比的服務，用來解決大型計算問題。隨著技術的進一步發展，近年來產生了利用互聯網上的計算機的 CPU的閑置處理能力來解決大型計算問題的分布式計算系統——網格計算。網格計算把一個需要非常巨大的計算能力才能解決的問題分成許多小的部分，然後把這些部分分配給許多計算機進行處理，最後把這些計算結果綜合起來得到最終結果。相信大家都使用過迅雷，如果你用過離線下載，你會看到雲端下載。這就是當今比較流行的雲計算產物。雲計算是網格計算、分布式計算、並行計算、網絡存儲、虛擬化等多種傳統計算機技術和網絡技術發展融合的產物。雲計算在當今被稱為是一種劃時代的技術，因為它將數量龐大的廉價計算機放進資源池中，用軟件容錯來降低硬件成本，通過將雲計算設施部署在各種能節省成本的區域，通過規模化的共享使用來提高資源利用率。國外代表性雲計算平臺提供商達到了驚人的10-40倍的性能價格比提升。不管有多麽新潮的技術和名稱，這一派系的高性能運算是以計算機網絡資源為依托。

前面的兩大門派，走的是高端路線，很難飛入尋常百姓家。多核處理器的出現打破了由大型機和網絡集群式分布式系統的壟斷，同時也使得並行算法和並行編程技術成為程序員不得不會的本領。

所謂多核處理器是指在一枚處理器中集成兩個或多個完整的計算引擎。多核技術的產生源於工程師們認識到，僅僅提高單核芯片的速度會產生過多熱量且無法帶來相應的性能改善，同時單靠提高單核芯片速度的性價比也令人難以接受，速度稍快的處理器價格要高很多。作為處理器技術發展的先驅龍頭企業，英特爾在1971推出的全球第一顆通用型微處理器4004，由2300個晶體管構成。這時，戈登摩爾提出了後來被業界奉為信條的“摩爾定律”——每過18個月，芯片上可以集成的晶體管數目將增加一倍。但到了2005年，當主頻接近4GHz時，英特爾和AMD發現，速度也會遇到自己的極限：那就是單純的主頻提升，已經無法明顯提升系統整體性能。以英特爾公司的奔騰系列為例，按照當時的預測，奔騰4在該架構下，最終可以把主頻提高到10GHz。但由於流水線過長，使得單位頻率效能低下，加上由於緩存的增加和漏電流控制不利造成功耗大幅度增加，3.6GHz奔騰4芯片在性能上反而還不如早些時推出的產品。所以，該系列只達到3.8G，就戛然而止。所以戈登摩爾本人似乎也依稀看到了“主頻為王”這條路的盡頭——2005年4月，他曾公開表示，引領半導體市場接近40年的“摩爾定律”，在未來10年至20年內可能失效。

多核心CPU解決方案的出現，標誌著新摩爾時代的來臨。事實上早在上世紀90年代末，就有眾多業界人士呼籲用單芯片多處理器技術來替代復雜性較高的單線程CPU。IBM、惠普、Sun等高端服務器廠商，更是相繼推出了多核服務器CPU。不過由於當時的技術並不成熟，造價也太高，並未引起大眾廣泛的註意。直到AMD率先推出64位處理器後，英特爾才想起利用“多核”這一武器進行“帝國反擊戰”。2005年4月，英特爾推出簡單封裝雙核的奔騰D和奔騰四至尊版840。AMD在之後也發布了雙核皓龍(Opteron)和速龍(Athlon) 64 X2和處理器。盡管如此，多核時代的元年應該是公元2006年。2006年7月23日，英特爾基於酷睿(Core)架構的處理器正式發布。2006年11月，又推出面向服務器、工作站和高端個人電腦的至強(Xeon)5300和酷睿雙核和四核至尊版系列處理器。與上一代臺式機處理器相比，酷睿2雙核處理器在性能方面提高40%，功耗反而降低40%。多核CPU由於其在傳統系統上重要的地位，無可爭議的開啟了當今的多核時代。

而與多核CPU差不多同時出現的就是本書的關鍵角色——多核GPU。GPU是相對於CPU的一個概念，傳統的GPU是一個專門的圖形的處理器，也是顯卡的“心臟”。NVIDIA公司在1999年發布GeForce256圖形處理芯片時首先提出GPU的概念。GPU使顯卡減少了對CPU的依賴，並進行部分原本CPU的工作，尤其是在3D圖形處理時。GPU所采用的核心技術有硬體T&L、立方環境材質貼圖和頂點混合、紋理壓縮和凹凸映射貼圖、雙重紋理四像素256位渲染引擎等，而硬體T&L技術可以說是GPU的標誌。這也可以說是GPU脫離中央領導，自立門戶的前奏。CUDA技術的前身GPGPU，即通用計算圖形處理器。如筆者前面的介紹，CUDA事實上也是一種GPGPU技術，但為了與早期的GPGPU技術相區別和突出CUDA是真正意義上的的通用計算，筆者將兩者視為不同。事實上，GPU的性能提升速度大大超過了CPU所遵照的摩爾定律，而且可編程性和功能都大大擴展，支持越來越復雜的運算。2006年，隨著支持DirectX 10的GPU的發布，基於GPU的通用計算開始普及。Nvidia公司於2007年正式發布了CUDA這一通用計算平臺。它是第一種不需要借助圖形學API就可以使用以C/C++為基礎的超集的通用計算的開發環境和軟件體系。

從前面的介紹可以看到，在新的多核時代，多核處理器成為時代最強音。因為它們可以集成到超級計算機，集群系統與網絡分布式計算系統，也可以廣泛運用於個人PC，公共機等普通消費市場。這一時代的兩大代表就是多核CPU和多核GPU，首先看看圖1.5.

圖1.5 CPU VS GPU

結構決定上層，有一種說法，CPU線程為重量級線程，GPU線程為輕量級線程。我們從結構對比來領會一下。首先，GPU上更多的資源放在了ALU（運算邏輯單元），這就意味著GPU擁有強大的計算能力。而在CPU上擁有更完整的控制單元和緩存空間，這就意味著CPU更適合處理復雜分支，不規則數據結構，不可預測存取，遞歸等情況。事實上，CUDA不支持遞歸程序。也正是因為CPU這樣的結構，CPU線程的啟動、切換、通信、同步操作開銷巨大，和GPU有1000：1的關系。

總體來說，GPU 是專為計算密集型、高度並行化的計算而設計。因此，GPU 的設計能使更多晶體管用於數據處理，而非數據緩存和流控制。更具體地說，GPU 專用於解決可表示為數據並行計算的問題，即在許多數據元素上並行執行的程序，具有極高的計算密度（數學運算與存儲器運算的比率）。由於所有數據元素都執行相同的程序，也就對精密流控制的要求不高；又由於有許多數據元素執行相同處理且具有較高的計算密度，因而可通過計算隱藏存儲器訪問延遲，而不必使用較大的數據緩存。這也是就GPU和CPU在結構設計上區別的根源。

值得註意的是，由於這兩大處理器各有自己的優缺點，CPU+GPU的並行處理方式不僅可以滿足更高的計算需求，節約成本和功耗，還可以將高性能計算廣泛普及，引領新的程序設計方法革命。

1.5 高性能計算利劍之CUDA

首先，筆者很喜歡NVIDIA的CUDA海報——高性能計算的利器，見圖1.6。2007年6月，NVIDIA推出了 CUDA。CUDA是一種將GPU作為數據並行處理設備的軟硬件體系。它也是一種 GPGPU 的技術，但相對傳統的GPGPU開發方法更為靈活方便，它是透過它的C語言的函數庫和一些 CUDA延伸的語法來編寫，因此不需用到 OpenGL 或 Direct3D，也不需要傳統的圖形函數庫，也不會被傳統的render pipeline設計束縛。

CUDA架構使GPU能夠解決復雜的計算問題。它包含了CUDA指令集架構（ISA）以及GPU內部的並行計算引擎。開發人員現在可以使用C語言來為CUDA?架構編寫程序，C語言是應用最廣泛的一種高級編程語言。所編寫出的程序於是就可以在支持CUDA?的處理器上以超高性能運行。而且還支持其它語言，包括FORTRAN以及C++等。更為重要的是，NVIDIA CUDA技術是當今世界上第一種針對NVIDIA GPU 的C語言環境，該技術充分挖掘出NVIDIA GPU巨大的計算能力。憑借NVIDIA CUDA技術，開發人員能夠利用NVIDIA GPU（圖形處理器）攻克極其復雜的密集型計算難題。

圖1.6 高性能計算之利器CUDA

作為並行計算平臺的CUDA，首先是它通過采用統一處理架構，更加有效的利用了以往分布在頂點渲染器和像素渲染器的計算資源。這樣既提高了GPU通用計算能力，又讓開發人員無須用匯編或者高級著色器語言編寫shader程序，然後再通過圖形學API執行。

CUDA技術另一個重要的方面是片內共享存儲器的引入，這樣就使得GPU可以完成隨機寫入和線程間通信。在Fermi架構下，又加入了L1和L2緩存，使得GPU對數據讀寫的方式更為靈活，例如合並訪問規則的放寬等。

從程序設計的角度來看，線程組的層次結構是一個很重要的核心概念（詳見3.1.1）。它將指導程序員如何將問題分解為更小的片段，以便各個線程塊通過協作的方法並行的解決。這樣的分解通過允許線程在解決各子問題時協作的形式，也保持了程序語言的問題表達能力，同時保持了程序的伸縮性。各個線程塊處理的子問題，安排在任何的CUDA core上進行處理。由此，編譯後的 CUDA 程序可以在任何數量的處理器內核上執行（需要編譯為相應的機器碼，詳見Nvcc），只有運行時系統需要了解物理多處理器數量。而這對程序員是透明的。

再來看看CUDA的強大計算能力，GPU的浮點運算能力是CPU的十數倍，如圖1.7。有了這個數據對比，然後我們來看看CUDA到底給我們帶來了什麽。

圖1.7 GPU與CPU浮點運算能力比較

第一個利用GPU的科學領域是天體物理學，在N體天體物理學研究中，天體物理學的研究人員利用GPU在一臺普通的PC上模擬2百萬個粒子，依靠GPU強大的並行計算能力，研究人員們做到了這一點。據性能分析報告知道，在N=16384時，GeForce 8800 GTX每秒可計算100多億次的相互作用，每秒執行38個積分時間間隔。

在Manifold 8地理信息項目研究中，CUDA作為第一款標配應用軟件用於地理信息系統（GIS）處理。該軟件可制作出一幅地圖，並疊加上人口信息，如該區域居民的年齡、住房類型、公路的數量等等 所有描述居住區的信息。規劃人員使用GPU可以正確設計道路、房屋以及各種服務的位置，打造更加高效的城市。

在成像技術中，東京大學信息科學與技術研究生院機械信息系的Takeyoshi Dohi教授與他的同事研究了NVDIA的CUDA並行計算平臺之後認為，醫療成像是CUDA這種平臺非常有前途的應用領域之一。自2000年以來，這所大學的研究小組已經開發出一種系統，通過CT或MRI掃描實時獲得的活體截面圖被視為體紋理，這種系統不僅能夠通過體繪制再現為三維圖像，還可作為立體視頻顯示，供IV系統使用。該系統為實時、立體、活體成像帶來了革命性的變化。但是，它的計算量極其龐大，僅體繪制本身就會帶來極高的處理工作量，況且此後還需要進一步處理來實現立體成像。對於每一個圖像幀，都有眾多角度需同時顯示。將此乘以視頻中的幀數，您會看到令人震驚的龐大計算數量，且必須在很短的時間內高度精確地完成這樣的計算。在2001年的研究中，使用了一臺Pentium III 800 MHz PC來處理一些512 x 512解析度的圖片，實時體繪制和立體再現要花費10秒鐘以上的時間才能生成一幀。為了加速處理，研究小組起初嘗試使用配備60塊CPU的UltraSPARC III 900 MHz機器，這是當時性能最高的計算機。但可以得到的最佳結果也不過是每秒鐘五幀。從實用的角度考慮，這樣的速度還不夠快。隨後，研究人員使用NVIDIA GPU GeForce 8800 GTX開發了一個原型系統。在使用CUDA的GPU上運行2001年研究所用的數據集時，性能提升到每秒13至14幀。UltraSPARC系統的成本高達數千萬日元，是GPU的上百倍，而GPU卻交付了幾乎等同於其三倍的性能，研究人員為此感到十分驚訝。同時，根據小組的研究，NVIDIA的GPU比最新的多核CPU至少要快70倍。測試顯示，對於較大規模的體紋理數據，GPU的性能會更為突出。

我們通過這些不勝枚舉的應用實例可以看到CUDA平臺提供的強大計算能力給我們的生活帶來的改變。筆者已經經歷了由CUDA 1.0，1.1，2.3，3.1到現在的4.0的N多次改版，每一次改版都會引入一些新的技術特性，而且CUDA平臺的功能也日趨完善，對新的硬件平臺的特性也提供了良好的支持。筆者相信CUDA作為一把高性能計算的利劍，必將為解決科學研究和工程應用中的通用計算瓶頸建立卓越的功勛！多核時代，CUDA已經劍起了風塵，君等準備好了沒有！

好了，到目前為止，我們已經充分的了解了為什麽需要並行程序，知道了並行算法的設計目標，基本方法與性能度量指標，也清楚了我們處在一個怎樣的多核時代。對我們要學習的CUDA技術也有了一個概括性的基本認識。做好準備，下一章開始，讓我們完成第一個CUDA程序，You can do it！

深入淺出說CUDA程序設計（一）