前言
《并行編程》系列是學(xué)習(xí)《Intro to Parallel Programming》過程中所做的筆記記錄以及個(gè)人一些所思所想。
GPU 與 CPU
衡量一個(gè)高性能處理器的時(shí)候坎拐，采用兩個(gè)指標(biāo)痘儡。
- 執(zhí)行時(shí)間（Latency）：執(zhí)行一項(xiàng)任務(wù)所花時(shí)間侯勉，采用時(shí)間單位猜极。
- 吞吐量（Throughput）：?jiǎn)挝粫r(shí)間完成的任務(wù)量缨叫。
而非常遺憾的是，這兩項(xiàng)指標(biāo)并不總是一致的缅糟，它們通常是矛盾的挺智。比如說：
A地和B地相距4500 KM祷愉，從A到B可以有兩種選擇窗宦。一種方法是開跑車，車上乘坐2個(gè)人二鳄，以200 KM/H的速度開到B地赴涵；另一種方法是乘坐客車，車上乘坐著40個(gè)人订讼，以50 KM/H的速度開到B地髓窜。
| 方案 | Latency(hour) | Throughput(people/hour) |
|---|---|---|
| 開跑車 | 4500 / 200 = 22.5 | 2 / 22.5 = 0.0889 |
| 客車 | 4500 / 50 = 90 | 40 / 90 = 0.444 |
雖然這個(gè)例子并不是很合理，但是它展示了 Latency 和 Throughput 的計(jì)算方式欺殿。
傳統(tǒng)的 CPU 設(shè)計(jì)就是嘗試去最優(yōu)化執(zhí)行時(shí)間寄纵，使其在每一項(xiàng)任務(wù)上的處理時(shí)間都能夠達(dá)到最優(yōu)。而 GPU 的設(shè)計(jì)與 CPU 不同脖苏，它的目標(biāo)最大化吞吐量程拭。因?yàn)樵谟?jì)算機(jī)圖形學(xué)中，我們更加關(guān)心每秒能處理的像素?cái)?shù)量棍潘，而不是每個(gè)像素需要花多少時(shí)間處理恃鞋，甚至只要每秒能處理的像素?cái)?shù)量只要能增加，即便單個(gè)像素處理的時(shí)間需要增加兩倍也是可以接受的亦歉。
GPU 設(shè)計(jì)原則
- GPU 有許多簡(jiǎn)單的計(jì)算單元恤浪，它們組合在一塊可以執(zhí)行大量的計(jì)算，GPU 通常會(huì)犧牲更多的控制力來換取更強(qiáng)大的計(jì)算能力肴楷。
- GPU 采用顯式并行編程模型(課程一大重點(diǎn))水由，即編程的時(shí)候就是按多處理器的思路進(jìn)行，而不是假設(shè)只有單個(gè)處理器赛蔫，然后將程序交給編譯器映射到多個(gè)處理上砂客。
- GPU 設(shè)計(jì)的優(yōu)化目標(biāo)是 Throughput 而不是 Latency ，它可以接受單個(gè)任務(wù)執(zhí)行時(shí)間延長(zhǎng)濒募，只要每秒能處理的任務(wù)總數(shù)能增加鞭盟，也因此 GPU 適用于以 Throughput 為最重要衡量指標(biāo)的應(yīng)用程序中。
CUDA 編程模型
異構(gòu)型計(jì)算機(jī)擁有兩種不同的處理器瑰剃，它們是 CPU 和 GPU齿诉。如果只是簡(jiǎn)單地寫一個(gè) C 程序，那么它只使用到了 CPU，而如果想要使用 GPU 就要借助 CUDA粤剧。CUDA 編程模型允許我們通過一個(gè)程序同時(shí)對(duì)兩個(gè)處理器進(jìn)行編程歇竟，另外雖然 CUDA 支持多門編程語言，但是本課程中主要使用 C 語言抵恋。
CUDA 中普通 C 語言部分的程序會(huì)運(yùn)行在 CPU （也稱為"HOST"）中焕议，而另外一部分將在 GPU （相對(duì)于"HOST"被稱為"DEVICE"）中運(yùn)行。然后 CUDA 編譯器會(huì)將 CPU 部分的代碼和 GPU 部分的代碼分開編譯弧关，為每個(gè)處理器生成各自的編譯結(jié)果盅安。
CUDA 將 GPU 當(dāng)做 CPU 的協(xié)處理器（co-processor）來對(duì)待，并且假設(shè) HOST 和 DEVICE 各自擁有獨(dú)立的內(nèi)存用于存儲(chǔ)數(shù)據(jù)世囊，GPU 通常采用高性能的內(nèi)存塊來作為內(nèi)存别瞭。當(dāng)談到 GPU 和 CPU 的關(guān)系時(shí)，CPU 則處于主導(dǎo)地位株憾。CPU 負(fù)責(zé)運(yùn)行主程序蝙寨，并向 GPU 發(fā)送控操作指令。
操作內(nèi)容包含有：
- 將數(shù)據(jù)從 CPU 內(nèi)存中移動(dòng)到 GPU 內(nèi)存中嗤瞎。
- 將數(shù)據(jù)從 GPU 內(nèi)存中移動(dòng)到 CPU 內(nèi)存中墙歪。
- 向 GPU 中的內(nèi)存申請(qǐng)空間。
- 調(diào)用 GPU 中的程序贝奇，以并行的方式進(jìn)行運(yùn)算虹菲，這些程序也稱為內(nèi)核，所以 HOST 能夠啟動(dòng) DEIVCE 中的內(nèi)核弃秆。
操作1和2涉及的命令是 cudaMemcpy 届惋，操作3涉及的指令是 cudaMalloc。
CUDA 程序流程
一個(gè)典型的 CUDA 程序流程是：
- CPU 為 GPU 申請(qǐng)存儲(chǔ)內(nèi)存空間（cudaMalloc）菠赚。
- CPU 將輸入數(shù)據(jù)復(fù)制到 GPU 內(nèi)存中（cudaMemcpy）脑豹。
- CPU 啟動(dòng) GPU 內(nèi)核處理數(shù)據(jù)（Kernel launch）。
- CPU 將結(jié)果從 GPU 內(nèi)存中復(fù)制回來（cudaMemcpy）衡查。
容易發(fā)現(xiàn)瘩欺，步驟2與4屬于數(shù)據(jù)傳輸?shù)倪^程。在程序中我們通常都希望能盡量減少數(shù)據(jù)傳輸所消耗的時(shí)間拌牲，而使更多時(shí)間花在計(jì)算上俱饿。所以對(duì)于 I/O 密集型的程序，便不適用于 CUDA 或者 GPU 編程塌忽。事實(shí)上拍埠，成功的 GPU 程序在計(jì)算時(shí)間與傳輸通信時(shí)間的比率上通常具有較高的值。
GPU 的優(yōu)點(diǎn)
GPU 擅長(zhǎng)處理以下兩個(gè)事項(xiàng)：
- 高效地啟動(dòng)大量的線程
- 并行地運(yùn)行大量的線程
舉個(gè)例子土居，比如說要對(duì)一個(gè)長(zhǎng)64的數(shù)組進(jìn)行求平方運(yùn)算枣购。
CPU 的做法
首先是只運(yùn)行于 CPU 中的做法嬉探。
程序中對(duì)數(shù)組進(jìn)行遍歷，然后依次對(duì)每一個(gè)元素都執(zhí)行相同的乘積操作棉圈。這些操作是在一個(gè)線程中串行執(zhí)行的涩堤，所以該線程將會(huì)執(zhí)行循環(huán)64次。
注：此處的線程指的是執(zhí)行完整代碼的一條獨(dú)立路徑分瘾。
GPU 的做法
理論知識(shí)
之前介紹過胎围，CUDA 的代碼需要分成兩部分，一部分運(yùn)行于 CPU德召，一部分運(yùn)行于 GPU白魂。GPU 部分所要實(shí)現(xiàn)的邏輯很簡(jiǎn)單，這里是使得輸出等于輸入的平方氏捞，但是這部分并沒有說明并行運(yùn)算的程度（或者是線程數(shù)量）碧聪。事實(shí)上冒版，指明并行運(yùn)算程度的任務(wù)將交給 CPU 進(jìn)行液茎。所以 CPU 需要為 GPU 分配內(nèi)存空間，再將數(shù)據(jù)復(fù)制到 GPU 內(nèi)存中辞嗡，然后再啟動(dòng) GPU 計(jì)算平方數(shù)的內(nèi)核（此處聲明了64個(gè)線程）捆等。
同時(shí)創(chuàng)建64個(gè)線程用于執(zhí)行平方運(yùn)算的好處是，每個(gè)線程都有一個(gè)唯一的線程索引续室，所有就可以將數(shù)組的第 n 個(gè)元素分配給第 n 個(gè)線程進(jìn)行處理栋烤。
代碼實(shí)踐
定義 GPU 內(nèi)核代碼。
__global__ void square(float *d_out, float *d_in){
// 獲取線程索引挺狰，將線程索引也作為數(shù)組的元素索引
int idx = threadIdx.x;
float f = d_in[idx];
d_out[idx] = f * f;
}
然后通過內(nèi)核啟動(dòng)語句配置并啟動(dòng)內(nèi)核明郭。
...
const int ARRAY_SIZE = 64;
...
square<<<1, ARRAY_SIZE>>>(d_out, d_in);
...
所以，這里啟動(dòng)了一個(gè)含有64個(gè)線程的塊丰泊，每個(gè)線程各自負(fù)責(zé)計(jì)算數(shù)組中的一個(gè)元素薯定。
配置啟動(dòng)內(nèi)核的參數(shù)
啟動(dòng)內(nèi)核的語句形式如：
kernel_function<<<blocks_number, thread_per_block>>>(d_out, d_in)
kernel_function 是自定義的內(nèi)核函數(shù)名稱。<<<....>>> 則是 CUDA 定義的特殊符號(hào)瞳购，其中接受兩個(gè)參數(shù)话侄，分別用于啟動(dòng)的塊的數(shù)量以及每個(gè)塊的線程數(shù)。
例如学赛，SQUARE<<<1, 64>>>(d_out, d_in) 語句啟動(dòng)了一個(gè)內(nèi)核年堆，并指定了內(nèi)核具有一個(gè)塊，每個(gè)塊存在64個(gè)線程盏浇。
如果需要使用到更多的計(jì)算資源变丧，那么便可以通過 <<<...>>> 的參數(shù)進(jìn)行配置。其中關(guān)于內(nèi)核的配置具有兩個(gè)特點(diǎn)需要知道：
- 一個(gè)內(nèi)核可以同時(shí)運(yùn)行多個(gè)塊绢掰。
- 每個(gè)塊可以運(yùn)行多個(gè)線程痒蓬，不過線程具有上限译蒂，通常而言在較舊的 GPU 中這個(gè)上限為 512，在較新的 GPU 中上限是 1024谊却。
所以當(dāng)我們想要啟動(dòng)128個(gè)線程計(jì)算128個(gè)數(shù)的平方時(shí)柔昼，代碼可以改為 SQUARE<<<1, 128>>>(...) 。那么如果想要想要啟動(dòng)1280個(gè)線程呢炎辨？這時(shí)便有多種策略捕透。比如 SQUARE<<<5, 256>>> 或者 SQUARE<<<10, 128>>> ，但是需要注意不能寫成 SQUARE<<<1, 1280>>> 碴萧，因?yàn)檫@樣超出了最大線程限制乙嘀。
但是當(dāng)前的塊和線程都是一維的，而如果我們需要處理二維或者是三維結(jié)構(gòu)的數(shù)據(jù)破喻，這樣顯然就不方便了虎谢，所以 CUDA 也支持二維和三維的塊與線程布局方式。
借助 dim3(x, y, z) 函數(shù)可以創(chuàng)建指定維度的布局方式曹质。默認(rèn)情況下婴噩，每個(gè)維度的值都為1，所以 dim3(w, 1, 1) == dim3(w) == w 羽德。
多維內(nèi)核
啟動(dòng)內(nèi)核的最一般形式是：
kernel<<<dim3(bx, by, dz), dim3(tx, ty, tz), shmem>>>(...)
dim3(bx, by, bz) 指定了塊的維度几莽。dim3(tx, ty, tz) 指定了每個(gè)塊的線程維度。shmem 這個(gè)參數(shù)不太常用宅静，默認(rèn)為 0章蚣，它以字節(jié)為單位指定了每個(gè)線程塊分配的共享內(nèi)存量，關(guān)于該參數(shù)的具體使用方法姨夹，后續(xù)會(huì)介紹到纤垂。
上圖展示了一個(gè)二維的布局，其中 thread 是最基本的單位磷账，若干個(gè) thread 組成了一個(gè) block峭沦，而若干個(gè) block 組合成一個(gè) grid。CUDA 還提供了多個(gè)屬性用于實(shí)現(xiàn)獲取線程索引够颠、塊索引和塊大小等熙侍。
-
threadIdx：獲取塊內(nèi)的線程索引，最多具有三個(gè)維度x, y, z履磨。 -
blockDim：獲取塊大小蛉抓，最多具有三個(gè)維度x, y, z。 -
blockIdx：獲取塊索引剃诅，與線程索引一樣巷送，最多具有三個(gè)維度。 -
gridDim：獲取 grid 大小矛辕，也是最多具有三個(gè)維度x, y, z笑跛。
注意事項(xiàng)：這邊有一個(gè)在編程處理圖像的時(shí)候經(jīng)常容易犯錯(cuò)的坑付魔。在處理圖像數(shù)據(jù)的時(shí)候，通常會(huì)將
grid of blocks定義為與圖像大小一致飞蹂，然后每個(gè)block中的thread數(shù)定義為1几苍。 這樣對(duì)于一個(gè)分辨率為n*m的圖像，則使用了n*m個(gè)塊陈哑，每個(gè)塊中只有一個(gè)線程負(fù)責(zé)處理當(dāng)前像素值妻坝。但是需要注意的是，在定義gridSize的時(shí)候需要定義成dim3(m, n)惊窖。但是在編程時(shí)出于習(xí)慣刽宪，我們很可能會(huì)寫成dim3(n, m)，然而這樣是錯(cuò)誤的 界酒。這是因?yàn)?dim3()接受的三個(gè)參數(shù)依次對(duì)應(yīng)了x, y, z的取值范圍圣拄，所以在使用圖像的寬高參數(shù)指定dim3參數(shù)的時(shí)候應(yīng)該是先制定寬度再指定高度。否則處理之后的圖像很可能出現(xiàn)有一半是黑色的情況毁欣。
