主機與設備

第四章 GPU 理論

4.2 GPU 模型架構

4.2.1 主機與設備

在 CUDA 的模型之中，將 CPU 當成主機端(host 端)，並將 GPU 作為協助主機端處理運算的元件稱之為設備端(device 端)。然而因兩種處理器特性上差異，將 CPU 用處理邏輯性較強的事件處理以及串行運算，GPU 則執行高度線程化的平行處理。一旦在確定程式架構上何處可以進行平行化演算，

便可以將這部分的計算工作交給 GPU 進行，然而在 GPU 上的 CUDA 平行運算的函數稱為 kernel，即在顯示卡上運行的程式碼區段。如圖 4-2 所示，

一個完整的 CUDA 程式架構，是透過一系列的設備端 kernel 函數的平行處理，及主機端的串行處理步驟所組成。在圖 4-2 中可以看出，一個 kernel 函數中存在有兩個層面的平行，即 Grid 中的 Block 間平行，以及 Block 中的 Thread 平行，而兩層平行架構也是 CUDA 主要的賣點之一。

圖 4- 2 CUDA 架構模型

4.2.2 線程結構

CUDA 的平行化模型是將核心(kernel)交由一組格網(grid)執行，再將網格切成數個區塊(block)，然後每個區塊再分成數個執行緒(thread)，依次分發進行平行運算，如圖 4-3 所示。

GPU 是以多處理器(SM , stream multiprocessor)為群組，每個多處理器有 8 個單精度浮點數的串流處理器(SP , stream processor) 負責運算這些串

圖 4- 4 GPU 計算單元概略圖

在 CUDA 之中的 kernel 函數實際上是以 Block 為單位執行的，在同一個 Block 之中數據是共享的，因此他們必須被配置到同一個 SM 之中，而 Block 中的每一個 Thread 則在被配置到一個 SP 上進行運作。

對於 GPU 記憶體頻寬相對於 CPU 大的許多，則是由於顯示卡中記憶體顆粒部分直接焊在顯示卡的 PCB 板上，故信號傳遞的完整性較高並且可擁有較高之工作頻率。而在 GPU 之記憶體架構中可以分為六大主要記憶體，

如圖 4-5 所示。在程式撰寫完畢後，可以透過不同的記憶體類型來針對程式進行優化，藉此提高程式計算效率，本研究主要使用的記憶體為 Global Memory 部分。

圖 4- 5 GPU 內記憶體架構圖 4.3 簡易程式撰寫及效率評估

本節以簡易的向量加法為例子，使用隨機向量測試運算效能，選取三個計算模式進行一系列解說。

4.3.1 單一區塊，單一執行緒

可用來測試單一串流處理器 (SP，stream processor) 之計算效能，因僅使用單一個串流處理器，其程式碼寫法與 C 語言寫法一致。程式的運作情況如圖 4-6 示意。

34 SP 輪流在這些執行緒中執行，其程式碼寫法將用到 threadIdx.x(區塊中的執行緒索引)、blockDim.x(區塊中的執行緒數量)。假設有八個執行緒，在同一

void add(float *a, float* b, float* c, int n) {

__global__ void Gadd(float *a, float* b, float* c, int n) {

Gadd <<< Grid, Block >>> (a, b, c, n);

}

void add(float *a, float* b, float* c, int n) {

__global__ void Gadd(float *a, float* b, float* c, int n) {

for( int k = threadIdx.x; k<n; k+=blockDim.x) c[k] = a[k] + b[k];

Gadd <<< Grid, Block >>> (a, b, c, n);

}

4.3.3 多區塊，多執行緒

當迴圈的數目小於 gridDim 與 blockDim 之乘積，前提為資料前後無相依性，便可以運用網格及區塊設定代替迴圈的部分，每一個執行緒只計算一次，即將整個 kernel 函數當成一平行化過的大迴圈。 (註:每個格網最大的區塊數量 gridDim 為 65535 個，每個區塊最大的執行緒數量為 512 個)，

程式的運作情況如圖 4-8 示意。

圖 4- 8 多區塊，多執行緒程式碼編寫部分

Block=512 代表了 512 個執行緒

Grid=n/Block+1 代表 n/Block+1 個區塊記憶體單元執行緒區塊(thread)

t⁰

共n個

void add(float *a, float* b, float* c, int n) {

__global__ void Gadd(float *a, float* b, float* c, int n) {

int k = blockIdx.x * blockDim.x + threadIdx.x;

c[k] = a[k] + b[k];

Gadd <<< Grid, Block >>> (a, b, c, n);

}

中央處理器(CPU) Intel® Core™2 Duo Processor E7400 (3M Cache, 2.80 GHz, 1066 MHz FSB) 記憶體(RAM) DDR2-800 4G

顯示卡(GPU) GTS250 256 MB GDDR3

處理核心數共 128 個單一時脈 1.8GHz 系統(OS) Windows Xp Sp3

編譯器(Copmiler) Visual Studio 2005 掛載 CUDA 版本 3.0 編譯數值格式精度採用單倍精度

(1) 單一區塊，單一執行緒

主機端運算時間 : 0.5 秒設備端運算時間 : 36.89 秒

(2) 單一區塊，多執行緒

主機端運算時間 : 0.5 秒設備端運算時間 : 0.078 秒

(3) 多區塊，多執行緒

主機端運算時間 : 0.5 秒設備端運算時間 : 0.0016 秒

透上述模擬可以發現，在單執行緒單區塊中，即 1 個 SP 單元進行運算，

GPU 的運算效率可以說是相當低落，CPU 整體運算效能為 GPU 的 7.4 倍左右。隨後以單一區塊多執行緒，即 1 個 MP 單元進行運算，GPU 的運算效率已有初步成效大約為 CPU 的 6.4 倍，隨後透過多區塊，多執行緒進行運算，GPU 運算效能為 CPU 之 31.3 倍左右，已是相當可觀。

4.3 透過 GPU 解 Navier-Stokes 相關演算法

在本研究中將流線方程式的解法稍做改寫，以及針對數值四則運算過程中採用了簡約算法，達到平行化運算的概念。

在運用 GPU 進行計算時，在主機端與設備端兩端進行資料交換時所花費之時間，與計算數據量的多寡有關，為了避免花費多餘的資料傳輸時間，

盡量將大量的數據運算在 GPU 上進行，並且盡可能的將設備端傳輸至主機端的資料量少量化。因此在最後以式(3-24)判定是否滿足收斂條件時，將運算數據由設備端拷貝至主機端來進行運算，故在此透過簡約算法(reduction algorithm) 將數值在設備端上先行運算，以減少傳輸至主機端的數據量。

4.4.1 流線方程式

此方程式以 Jacobi 法進行離散，在計算次數上，一次便運行兩次，以減少資料的交換，再運行兩次計算後，便透過式(3-24)判斷是否滿足收斂條件。

程式碼如下：

4.4.2 簡約算法(reduction algorithm)

傳統的樹狀簡約算法如圖 4-9 所示，假若有八個元素做累加動作，則先兩兩相加，重複動作到累加完畢為止。

圖 4- 9 樹狀計算過程

GPU 的簡約計算類似傳統樹狀簡約計算，但是其運算過程有些改變，

運算過程如圖 4-10。

圖 4- 10 GPU 簡約計算過程

Result

解說如下：

 在第一次循環(Stride=1)時，僅有 Thread ID=0、2、4…14 線程的運算，即程式碼中 idate[i]與其往後 1 個單位的元素進行加總動作。

 在第二次循環(Stride=2)時，僅有 Thread ID=0、4、8…線程的運算，

即程式碼中 idate[i]與其往後 2 個單位的元素進行加總動作。

 以此類推。

 直到最後一次時假若為(Stride=8)，僅有 Thread ID=0 此線程進行運算，即程式碼中 idate[i]與其往後 8 個單位的元素進行加總動作，其和即為輸入之 28 組數據之和。

程式碼如下：

第五章模擬驗證與結果討論

本章首先以 Ghia et al. (1982) 穴流模擬結果進行模擬驗證，再調整穴流模擬範圍之長寬比，以 Cheng (2006) 模擬結果進行驗證，隨後探討 GPU 加速之成效。

5.1 模擬驗證與分析

5.1.1 模擬設備

本研究採用單倍精度進行模擬，Kuznik (2009) 指出運用單倍精度已足夠應付目前多數的數值演算。而選用單精度做為數值儲存格式，因其在 GPU 目前所開發的編譯器架構上較為成熟並且有較高的運算效能，雙精度數值格式下 GPU 上運算成效較為低落，並且在有限的記體容量限制之下，採用單倍精度能夠存取較多數值。根據 IEEE-754 浮點數之定義可得知單倍精度下，有效數字可達 7 位，即精確度到小數點後第 6 位，在小數點後第 6 位

中央處理器(CPU) Intel® Core™2 Duo Processor E7400 (3M Cache, 2.80 GHz, 1066 MHz FSB) 記憶體(RAM) DDR2-800 4G

顯示卡(GPU) GTX 480 1536MB GDDR5

處理核心數共 480 個單一時脈 1.4GHz 系統(OS) Windows XP Service pack 3

編譯器(Compiler) Visual Studio 2005 掛載 CUDA 版本 3.0 編譯數值格式精度採用單倍精度

會產生數值的震盪，實則為電腦儲存格式之問題。因此在流線模擬成果繪製上僅輸出到10⁻⁷。

5.1.2 方型穴流之模擬

拉蓋方穴流(lip-driven cavity flow)的幾何形狀如圖 5-1 所示，為一長寬比為一的矩型，當流體以均勻速度 U 由上壁面的左方流至右方，將在方穴裡面形成渦流，其餘三壁面邊界速度均為零，而透過不同的速度大小可以得到不同雷諾數下的渦度及流線模擬情況。

圖 5- 1 方穴流示意圖

根據 Ghia et al. (1982) 的研究，本研究分別於雷諾數 100、400、1000、

3200、5000、7500 及 10000 進行模擬，採用 Ghia et al. (1982)模擬之網格數數大小，分別在雷諾數 100、400、1000 及 3200 部分採用網格數為 129×129，

雷諾數 5000、7500 及 10000 部分網格數採用 257×257。模擬至穩定狀態 (steady state)，其流線與渦度模擬結果分別如圖 5-2 至圖 5-15 所示。

根據模擬結果可以發現到當雷諾數改變時將使流場結構產生變化，雷諾數越高則流場形狀越複雜，所形成的渦漩數量也會變多。在雷諾數 1000 以下時僅在下壁面左右兩側各產生一渦漩，隨後在雷諾數達到 3200，在左壁面上方又在產生一渦漩，而後雷諾數提高到 5000 之際，左上角之渦漩也更往中心點移動，而在中央渦漩部分隨著雷諾數增加而逐漸往幾何中心移動(x=1/2，y=1/2)。本研究模擬之流線與渦度趨勢圖與 Ghia 所模擬之成果有一樣的趨勢，而渦漩中心點也大置上位於相同的位置。

隨後根據 Ghia et al. (1982)之水平與垂直中心線上速度計算結果，做為比較之基準如圖 5-16 至圖 5-29 所示，由圖中可以發現雷諾數增加下，速度分布曲度變化隨之更加明顯。本文研究之模擬結果，與 Ghia 模擬結果之 u、

v 速度相差無幾，而在將程式碼改編成 CUDA 演算之模式下，所求得之模擬成果亦相當準確。

圖 5- 2 雷諾數 100 之流線比較圖

圖 5- 3 雷諾數 100 之渦度比較圖