1. CPU 并行

1.1 概述

之所以先從 CPU 聊起而不是直接切入正題講 GPU, 是因為并行思想在CPU和GPU上是相通的, 而首先了解 CPU 的模式, 也更有助于之后熟悉 GPU 與 CPU 之間的交互膏燕。

因為 CPU 首先得取得數(shù)據(jù), 才能進行運算, 所以很多時候, 限制我們程序運行速度的并非是 CPU 核的處理速度, 而是數(shù)據(jù)訪問的速度, 下面羅列了一些硬件設備的訪問速度 (1Gbps = 125 MB/s):

網(wǎng)卡 (NIC), 1Gbps
機械硬盤(HDD) , 1-2Gbps, 如果連接了 SATA3 的話, 最高可達到 6Gbps. 其實對于 HDD 而言, 最耗時的部分不是傳送數(shù)據(jù)的時間, 而是尋找數(shù)據(jù)的時間, 因為 HDD 內部構造就像老式的唱片機, 一個旋轉頭和一個圓盤, 而定位某個數(shù)據(jù)需要將旋轉頭移動到數(shù)據(jù)儲存在圓盤上的位置, 這個過程是較為耗時的, 所以如果數(shù)據(jù)是不規(guī)則儲存的, 即分散儲存, 那么找到所有數(shù)據(jù)的時間還需要更長, 因此雖然 SATA3 最高可支持 6Gbps 的速率, 但是 HDD 是很難達到這個最高速率的
USB 2.0 的速率是 0.48Gbps, 而 USB3.0 支持 5Gbps, USB3.1 支持 10Gbps
固態(tài)硬盤(SDD) 使用 SATA3 的話可支持 4-5Gbps
DDR4 內存和 CPU(Core i7) 之間的傳輸速率是 160-480Gbps
GPU 內部的內存?zhèn)鬏斔俾试?800-8000Gbps

市面上賣的 CPU 大多都是 x核2x線程, 譬如 8核16線程, 10核20線程, 這是因為這些核都采用了 Hyper-Threading 技術, 該技術可以讓一個 microprocessor(核) 表現(xiàn)出兩個分離的 processor 的形式, 不過這兩個 processor 需要共享相同的核資源, 譬如 cache memory 等探熔。

但是如果你打開當前系統(tǒng)的資源管理器, 你可能會看到當前系統(tǒng)中存在大概 1000+ 個線程, 但是你的 CPU 是 10 核 20 線程, 很顯然是不合理的樱调。這是因為資源管理器中顯示的線程不是正在活動的線程, 而其中絕大部分都在睡眠狀態(tài), 譬如某個網(wǎng)絡監(jiān)聽窗口, 一直在沉睡, 直到某個網(wǎng)絡包到來, 系統(tǒng)喚醒了該線程, 線程處理完這個網(wǎng)絡包后繼續(xù)睡眠专酗。此處的10核20線程指的是能同時執(zhí)行(execution) 的線程, 而某個線程被啟動(launch), 但是在后臺睡眠而不被執(zhí)行也是可以的。

程序可以被分為以下三類:

core-intensive (核資源密集型)
memory-intensive (內存資源密集型)
I/O-intensive (I/O 資源密集型)

一個 4GHz 的單核 CPU 的能耗大概是一個 1GHz 的單核 CPU 的 16 倍, 而一個 4GHz 的單核 CPU 的能耗與一個雙核的每個核 3GHz 的CPU能耗相當屡律。

1.2 CPU 并行化程序設計

需要明確的點是:

當我們需要 launch 多個線程的時候, 確保大量的計算放在線程中, 而不要放在多線程執(zhí)行完并同步后的代碼中
launch 多個線程是有 overhead 的, 因為需要給每個線程分配資源, 句柄等
在 c 語言中, 一旦一個線程被成功 launch, 系統(tǒng)會給其分配一個 handler(句柄), 一個虛擬CPU (10核20線程相當于有20個虛擬CPU), 一個stack(棧)區(qū), 并將 handler 返回給launch這個線程的主線程
總共有兩種類型的核, 一種是 in-order(inO) 的, 另一種是 out-of-order(ooO) 的吱抚。
- 其中 inO 意味著其嚴格按照二進制碼中的順序來執(zhí)行代碼, 而 ooO 則是當前哪個 operand(操作元) 可以執(zhí)行, 就執(zhí)行哪個, 譬如對于一系列指令, inO 只能等一條指令執(zhí)行完了后才能執(zhí)行后一條指令, 而 ooO 可以在等待當前這條指令執(zhí)行完之前, 先著手執(zhí)行其后面的指令。
- 很顯然, 在指令執(zhí)行調度上 ooO 是比 inO 高效的, 然而, 因為 inO 較為簡單, 所以制造 inO 核所需要的基礎 chips(芯片) 更小, 所以一個 inO 核中可以放置更多的基礎芯片, 因為基礎芯片的數(shù)量增加了, 所以 inO 可以以更高的時鐘頻率來工作
- 同時, 因為 inO 更為簡單, 所以耗能較低
如果一個線程需要使用大量的核的資源, 那么這個線程就被稱為 thick 線程, 反之被稱為 thin 線程, 如果一個線程被設計為盡可能少的需要核的資源, 那么由這種線程組成的多線程程序, 能最大地提升多線程的效率, 這也是為什么微軟設計 windows 系統(tǒng)時, 確保其每個系統(tǒng)線程所占用的核的資源最少, 這樣一來就不會影響用戶的應用程序的執(zhí)行

1.3 內核與內存

CPU架構以及 L1, L2, L3 緩存

1.jpg

上圖是 i7-5930K 的 CPU 結構, 下圖是每個核中詳細的結構。

2.jpg

芯片組和DRAM

芯片組和CPU的構建模塊主要是MOS晶體管昙篙，但DRAM的構建模塊是存儲電荷的極小電容器腊状。CPU 通過將 MOS 晶體管縮小使 CPU 中能添加更多的晶體管, 而 DRAM通過將更多的電容器轉入連續(xù)的區(qū)域來獲得更大的儲存。
由于電荷存儲在極小的電容器中苔可，因此在一定時間后（例如50 ms）會消耗（即泄漏）電荷缴挖。由于這種泄漏，必須讀取數(shù)據(jù)并將其放回DRAM（即刷新）焚辅∮澄荩考慮到刷新的缺點，允許一次一個字節(jié)地訪問數(shù)據(jù)是沒有意義的同蜻。因此棚点，每次訪問數(shù)據(jù)最好是大塊大塊地存取（例如湾蔓，每次4KB）瘫析。
DRAM一次訪問一行，每行是DRAM內存的最小可訪問量∧穑現(xiàn)代DRAM中的行大約為2-8 KB贬循。要訪問某行，需要一定的時間（延遲）桃序。但是杖虾，一旦訪問該行（由DRAM內部進入行），該行實際上可以自由訪問媒熊。CPU訪問DRAM的延遲大約為200-400個周期奇适，而訪問同一行中的后續(xù)元素只需幾個周期。即獲取一行數(shù)據(jù)有很長的延遲芦鳍，但一旦讀取嚷往，對該行的訪問速度非常快怜校。除了可以訪問行之外间影，DRAM還具有各種其他延遲注竿，例如行到行延遲等茄茁。這些參數(shù)中的每一個都由存儲器接口標準指定，這些標準由協(xié)議定義巩割。
內存帶寬的增加應允許更多的硬件線程更快地將數(shù)據(jù)帶入內核, 從而當更多數(shù)量的已啟動軟件線程同時需要來自DRAM主內存的數(shù)據(jù)時, 不會造成阻塞, 避免內存帶寬飽和裙顽。

L1, L2, L3 緩存和執(zhí)行單元

L1, L2, L3 緩存是 SRAM(static random access memory), 而不是 DRAM, 其訪問速度遠遠大于 DRAM。
L1 緩存
- 總共有 64KB, 其中 32KB 用來儲存數(shù)據(jù), 即上圖中的 L1D, 剩下 32KB 用來儲存指令, 即 L1I, L1I 中儲存的是最常用的一些指令
- 訪問 L1 緩存的速度很快 (4 周期的 load-to-use 延遲)
- 每個核都有獨有的 L1 緩存, 因為每個核共有兩個線程, 所以這兩個線程需要共享 L1 緩存, 而這兩個線程之間可以彼此之間可以通過 L1 緩存收/發(fā) 數(shù)據(jù)
L2 緩存
- 總共有 256KB, 其不像 L1 緩存將數(shù)據(jù)和指令分開, 而是將指令和數(shù)據(jù)儲存在一起
- 訪問 L2 的緩存速度較塊 ( 11-12 周期的 load-to-use 延遲 )
- 當系統(tǒng)決定某個數(shù)據(jù)或某個指令不再常用時, 會將其從 L1 緩存中移除, 因為 L2 緩存比 L1 緩存大, 所以從 L1緩存中移除的數(shù)據(jù)或指令可能仍然存在與 L2 緩存中
- 每個核都有其獨有的 L2 緩存
L3 緩存
- 總共有 15MB
- 訪問 L3 緩存比 DRAM 塊, 但是比 L2 緩存稍慢 (約等于 22 周期的 load-to-use 延遲)
- 每個核都有其獨有的 L3 緩存
L1, L2, 和 L3 中數(shù)據(jù)的出入完全由CPU控制宣谈，不受程序員的控制愈犹。但是，通過將數(shù)據(jù)操作保持在小循環(huán)中，程序員可以很大程度上影響高速緩沖存儲器的效率漩怎。
數(shù)據(jù)會首先進入 L3 緩存, 然后從 L3 緩存進入 L2 緩存, 最后再從 L2 緩存進入 L1 緩存, 所以, 為了在最大程度上利用緩存提高效率, 我們應該盡可能做到:
1. 每個線程重復訪問32 KB數(shù)據(jù)區(qū)域
2. 嘗試在可能的情況下將更廣泛的訪問限制為256 KB勋颖，
3. 在考慮所有已啟動的線程時, 嘗試在L3 (例如，15 MB) 儲存累積的數(shù)據(jù)
4. 如果必須超過L3的大小勋锤，請確保在超過此區(qū)域之前大量使用 L3
執(zhí)行單元
- 執(zhí)行單元分為兩類：ALU（算術邏輯單元）負責整數(shù)運算饭玲，邏輯運算如OR，AND叁执，XOR等.FPU（浮點單元）負責浮點（FP）操作茄厘，如FP ADD和FP MUL（乘法）。除（整數(shù)或FP）除法比加法和乘法更復雜谈宛，因此有一個單獨的除法單位次哈。但是整數(shù)除法顯著地比整數(shù)乘法要慢。所有這些執(zhí)行單元都由一個核中的兩個線程共享吆录。
- 在每一代中窑滞，更復雜的計算單元可用作共享執(zhí)行單元。但是径筏，多個單元被合并用于可能由兩個線程 (例如ALU) 執(zhí)行的常見操作, 不過每一代的確切細節(jié)可能會改變葛假。但在過去3到3年的CPU設計中，ALU-FPU功能分離從未改變過滋恬。
- 計算兩個線程生成的地址聊训，以將兩個線程中的數(shù)據(jù)寫回內存。對于地址計算恢氯，加載和存儲地址生成單元（LAGU和SAGU）由兩個線程共享带斑，以及正確排序目標存儲器地址 (MOB) 的單元。
- 指令只被預取并解碼一次并傳遞到所有者線程勋拟。因此勋磕，兩個線程共享預取器和解碼器。
- 我們在寫代碼時需要考慮某個計算所消耗的資源, 譬如計算 √(x2+y2) 就是計算密集型的指令, 其需調用的運算單元包括
  - 兩次 FP-MUL (浮點相乘)
  - 一次 FP-ADD (浮點相加)
  - 一次開方操作 (用的應該是卡馬克算法)

3.jpg

假設我們現(xiàn)在要左右翻轉一張圖片, 即上圖的狗, 我們初步的想法是將該 2MB 的圖片整個儲存在全局內存中, 然后一行行迭代, 每迭代一行則將該行左右相對的像素點交換, 那么按照這種邏輯設計出來的程序如下:

...
for(row=ts; row<=te; row++) {
  col=0;
  while(col<ip.Hpixels*3/2){ 
    // example: Swap pixel[42][0] , pixel[42][3199] 
    pix.B = TheImage[row][col];   // TheImage 使用的是全局內存(DRAM), 儲存了圖片數(shù)據(jù)
    pix.G = TheImage[row][col+1]; 
    pix.R = TheImage[row][col+2]; 
    TheImage[row][col] = TheImage[row][ip.Hpixels*3-(col+3)]; 
    ...

上述程序如果想要翻轉一個像素, 其需要訪問6次內存(TheImage), 前三次讀取該像素點的三通道的值, 后三次一一將讀取的值放到對應的翻轉后的位置上, 因為圖片的一行總共有3200個像素點, 所以翻轉一行總共需要 3200×6=19200 次內存訪問

很顯然, 上述程序設計十分低效, 因為:

我們每次都讀取全局內存, 而且每次讀取的內容都不一樣, 這樣系統(tǒng)根本沒辦法使用 L1, L2, L3 緩存將常用的數(shù)據(jù)緩存下來用以加速
當我們訪問 DRAM 時, 盡可能地一次存取較大并且連續(xù)的塊, 譬如 1KB, 4KB 之類大小的塊, 而不是一個字節(jié)一個字節(jié)地存取, 那樣相當耗時, 而上述程序讀取 TheImage 時, 都是以像素點通道為單位讀取, 一個像素點通道一個字節(jié), 讀取粒度太小

那么改進后的代碼為:

unsigned char Buffer[16384]; // This is the buffer to use to get the entire row 
...
for(row=ts; row<=te; row++) {
// bulk copy from DRAM to cache
  memcpy((void *) Buffer, (void *) TheImage[row], (size_t) ip.Hbytes); col=0;
  while(col<ip.Hpixels*3/2){
     pix.B = Buffer[col];
     pix.G = Buffer[col+1];
     pix.R = Buffer[col+2];
     Buffer[col] = Buffer[ip.Hpixels*3-(col+3)];
     Buffer[col+1] = Buffer[ip.Hpixels*3-(col+2)];
     Buffer[col+2] = Buffer[ip.Hpixels*3-(col+1)];
     Buffer[ip.Hpixels*3-(col+3)] = pix.B;
     Buffer[ip.Hpixels*3-(col+2)] = pix.G;
     Buffer[ip.Hpixels*3-(col+1)] = pix.R;
     col+=3;
}
// bulk copy back from cache to DRAM
memcpy((void *) TheImage[row], (void *) Buffer, (size_t) ip.Hbytes); 
...

上述代碼的優(yōu)點是:

通過一個 Buffer, 每次將一行的像素數(shù)據(jù)緩存到 L1, L2, L3 緩存中, 這樣一來有兩個好處, 一是不再需要訪問內存, 并提高了效率, 而且 buffer 相當是連續(xù)的儲存, 不像之前訪問 TheImage 那樣離散地訪問像素點
因為所處理的圖片是 22MB 的, 假設我們的 L3 緩存是 15MB, 那么在程序運行的過程中, L3 緩存會不斷的將不常用的數(shù)據(jù)刪除, 并讀取常用的像素點, 如此周而復始

上述說法有一點不準確, 即不是我們主動地將其放入 L1, L2, L3 緩存中, 在我們創(chuàng)建 Buffer 的時候, 其還是在 DRAM 中的, 只是操作系統(tǒng)發(fā)現(xiàn) Buffer 比較常用, 因此主動地將其放入 L1, L2, L3 緩存中

程序設計所用內存

4.jpg

當我們執(zhí)行一個程序時, 它需要幾個內存空間

一個棧內存區(qū)
- 其用來儲存?zhèn)魅牒瘮?shù)的的參數(shù)變量和從函數(shù)返回的參數(shù)變量及指針, 在函數(shù)中定義的一些基本類型的變量和對象的引用變量都在函數(shù)的棧內存中分配敢靡。當在一段代碼塊定義一個變量時挂滓，編譯器就在棧中為這個變量分配內存空間，當超過變量的作用域后啸胧，編譯器會自動釋放掉為該變量所分配的內存空間赶站，該內存空間可以立即被另作他用。
- 棧的優(yōu)勢是纺念，存取速度比堆要快贝椿，僅次于寄存器，棧數(shù)據(jù)可以共享陷谱。但缺點是烙博，存在棧中的數(shù)據(jù)大小與生存期必須是確定的，缺乏靈活性。棧中主要存放一些基本類型的變量（,int, short, long, byte, float, double, boolean, char）和對象句柄渣窜。棧有一個很重要的特殊性铺根，就是存在棧中的數(shù)據(jù)可以共享。
- 在上下文切換的時候(context switch), 原有線程的棧區(qū)會被保存下來
一個堆內存區(qū)
- 在標準C語言上乔宿，使用malloc等內存分配函數(shù)獲取內存即是從堆中分配內存, 從堆中分配的內存需要手動釋放夷都，如果不釋放，而系統(tǒng)內存管理器又不自動回收這些堆內存的話, 那就一直被占用予颤。如果一直申請堆內存囤官，而不釋放，內存會越來越少蛤虐，很明顯的結果是系統(tǒng)變慢或者申請不到新的堆內存党饮。而過度的申請堆內存 (譬如在函數(shù)中申請一個1G的數(shù)組), 會導致堆被壓爆
- 我們掌握堆內存的權柄就是返回的指針，一旦丟掉了指針驳庭，我們便無法主動釋放它刑顺。這便是內存泄露。而如果在函數(shù)中申請一個數(shù)組饲常，在函數(shù)體外調用使用這塊堆內存蹲堂，結果是未定義的。我們知道在c/c++ 中定義的數(shù)組大小必需要事先定義好贝淤，他們通常是分配在靜態(tài)內存空間或者是在棧內存空間內的柒竞，但是在實際工作中，我們有時候卻需要動態(tài)的為數(shù)組分配大小播聪，這時就要用到堆內存分配的概念朽基。
- 在堆內存分配時首先應該知道操作系統(tǒng)有一個記錄空閑內存地址的鏈表，當系統(tǒng)收到程序的申請時离陶，會遍歷該鏈表稼虎，尋找第一個空間大于所申請空間的堆結點，然后將該結點從空閑結點鏈表中刪除招刨，并將該結點的空間分配給程序霎俩，另外，對于大多數(shù)系統(tǒng)沉眶，會在這塊內存空間中的首地址處記錄本次分配的大小打却，這樣，代碼中的delete語句才能正確的釋放本內存空間沦寂。另外学密，由于找到的堆結點的大小不一定正好等于申請的大小淘衙，系統(tǒng)會自動的將多余的那部分重新放入空閑鏈表中传藏。
- 堆內存是向高地址擴展的數(shù)據(jù)結構，是不連續(xù)的內存區(qū)域。這是由于系統(tǒng)是用鏈表來存儲的空閑內存地址的毯侦，自然是不連續(xù)的哭靖，而鏈表的遍歷方向是由低地址向高地址。堆內存的大小受限于計算機系統(tǒng)中有效的虛擬內存侈离。
- 由此可見试幽，堆內存獲得的空間比較靈活，也比較大卦碾。堆內存是由new分配的內存铺坞，一般速度比較慢，而且容易產生內存碎片,不過用起來最方便
一個儲存代碼的區(qū)域
- 儲存程序代碼和在程序中定義的 constant 變量, 這個區(qū)域不會被修改, 包括這個區(qū)域內儲存的常量變量

我們需要知道的是, 一個程序設計為 12 個線程執(zhí)行并不意味著速度加快了 12 倍, 在不同線程之間切換以及上下文切換, 或多線程多線程搶占內存帶寬造成的阻塞等等, 會產生一個 Parallelization Overhead, 我們可以通過下述公式計算 Parallelization Overhead

$Parallelization Overhead=1- \frac{單個線程執(zhí)行時間}{N個線程總共執(zhí)行時間\times N}$

物理內存與虛擬內存

系統(tǒng)將所請求的內存大小看做一系列的 pages 的組合, 即頁, 每個頁是 4KB, 如果一個用戶請求 1MB 的內存大小, 那么其請求 256 頁
而系統(tǒng)會將當前常用的頁放置在內存中, 不常用的放在磁盤中, 而僅僅會提供給程序一個虛擬地址, 這個虛擬地址可能在內存中, 也可能在磁盤中
因為虛擬內存的緣故, 所以假設我們的物理內存是 8GB, 但我們卻可以分配超過 8GB 的虛擬內存
malloc() 會分配一個虛擬內存地址, 該地址指向的位置既可能在磁盤中, 也可能在內存中

鎖頁內存與可分頁內存

鎖頁內存是分配的內存地址就是物理內存內的地址, 而不像虛擬內存一樣可能在磁盤中, 不過若分配過多的鎖頁內存, 會導致其他程序可用的內存減小洲胖。
- 假設物理內存是 8GB, 而虛擬內存是 64GB, 為某個程序分配了 2GB 的鎖頁內存后, 此時可用物理內存還剩 6GB, 不過虛擬內存還是 64GB, 雖然虛擬內存還是 64GB, 但是其靈活性下降, 因為可用的真實物理內存從 8GB 下降到 6GB
malloc()將分配標準的济榨，可分頁的主機內存。而cudaHostAlloc()將分配頁鎖定的主機內存绿映。頁鎖定的主機內存也稱為固定內存或不可分頁內存擒滑，它的重要屬性就是：操作系統(tǒng)將不會對這塊內存分頁并交換到磁盤上，從而確保了該內存始終駐留在物理內存中叉弦。因此丐一，操作系統(tǒng)能夠安全的使用應用程序訪問該內存的物理地址，因為這塊內存將不會被破壞或者重新定位淹冰。
當使用可分頁內存進行從CPU到GPU的復制時库车，復制操作將執(zhí)行兩遍，第一遍從可分頁內存復制到一塊"臨時的"鎖頁內存樱拴，然后再從這個鎖頁內存復制到GPU上凝颇。因此，當在GPU和主機間復制數(shù)據(jù)時疹鳄，這種差異會使也鎖定主機內存的性能比標準可分頁內存的性能要高大約2倍拧略。
cudaHostAlloc 和 malloc 操作都會在系統(tǒng)內存中分配一塊區(qū)域, 但其區(qū)別是, cudaHostAlloc 除了像 malloc 分配內存外, 還額外有一個鎖頁操作, 而這鎖頁內存會消耗額外的時間
固定內存是一把雙刃劍，當使用固定內存時瘪弓，將失去虛擬內存的功能垫蛆。特別是，應用程序中使用每個鎖頁內存時都需要分配物理內存腺怯，因為這些內存不能交換到磁盤上袱饭。這意味著，與使用標準的malloc調用相比呛占，系統(tǒng)將更快的耗盡內存虑乖。因此，建議僅對 cudaMemcpy() 調用中的源內存或者目標內存晾虑，才使用鎖頁內存疹味，并且在不需要的時候立即釋放仅叫。
鎖頁內存允許GPU上的DMA控制器請求主機傳輸，而不需要CPU主機處理器的參與
CPU仍然可以訪問上述鎖頁內存糙捺，但是此內存是不能移動或換頁到磁盤上的
在GPU上分配的內存默認都是鎖頁內存诫咱，這只是因為GPU不支持將內存交換到磁盤上
在主機上分配鎖頁內存有以下兩種方式：
- 使用特殊的cudaHostAlloc函數(shù)，對用的釋放內存使用cudaFreeHost函數(shù)進行內存釋放
- 使用常規(guī)的malloc函數(shù)洪灯，然后將其注冊為(cudaHostRegister)鎖頁內存,注冊為鎖頁內存只是設置一些內部標志位以確保內存不被換出坎缭，并告訴CUDA驅動程序，該內存為鎖頁內存签钩，可以直接使用而不需要使用臨時緩沖區(qū)

使用鎖頁內存需要注意以下幾點：

鎖頁操作會消耗額外的時間, 隨著所分配的內存區(qū)域的增大, 其耗時也增長, 所以不能分配太多掏呼，太多的話會降低系統(tǒng)整體性能
鎖頁內存和顯存之間的拷貝速度是6G/s，普通的內存和顯存之間的拷貝速度是3G/s（顯存之間的拷貝速度是30G/s,CPU之間的速度是10G/s）
使用cudaHostAlloc函數(shù)分配內存铅檩，其內的內容需要從普通內存拷貝到鎖頁內存中哄尔，因此這種拷貝會帶來額外的CPU內存拷貝時間開銷，CPU需要把數(shù)據(jù)從可分頁內存拷貝到鎖頁柠并，但是采用cudaHostRegister把普通內存改為鎖頁內存岭接，則不會帶來額外的cpu內存拷貝時間開銷，因為cudaHostAlloc的做法是先分配鎖頁內存臼予，這時里面是沒有數(shù)據(jù)的鸣戴，那么需要將一般的內存拷貝過來，而對于cudaHostRegister內存粘拾，他是之間就使用malloc分配好的窄锅，cudaHostRegister只是設置一些內部標志位以確保其不被換出，相當于只是更改了一些標志位缰雇，就不存在前面說的數(shù)據(jù)拷貝
在某些設備上入偷，設備存儲器和主機鎖頁存儲器之間的數(shù)據(jù)拷貝和內核函數(shù)可以并發(fā)執(zhí)行
在某些設備上，可以將主機的鎖頁內存映射到設備地址空間械哟，減少主機和設備之間的數(shù)據(jù)拷貝疏之，要訪問數(shù)據(jù)的時候不是像上面那那樣將數(shù)據(jù)拷貝過來，而是直接通過主機總線到主機上訪問暇咆，使用cudaHostAlloc分配時傳入cudaHostAllocMapped锋爪，或者使用cudaHostRegister時傳入cudaHostRegisterMapped標簽
默認情況下，鎖頁內存是可以緩存的爸业。在使用cudaHostAlloc分配時傳入cudaHostAllocWriteCombined標簽其骄，將其標定為寫結合，這意味著該內存沒有一級二級緩存扯旷，這樣有利用主機寫該內存拯爽，而如果主機讀取的話，速度將會極其慢钧忽，所以這種情況下的內存應當只用于那些主機只寫的存儲器

1.4 線程管理與同步

5.jpg

如果當前 CPU 是 10核20線程的, 那么在某一時刻, 不能超過 20個線程出于 Running 狀態(tài)
當我們調用 pthread_create() 創(chuàng)建一個進程時, 操作系統(tǒng)則會去查看, 是否有足夠的資源來新建一個線程? 如果有的話, 那么一個 handle 會被分配給該線程, 并為其創(chuàng)建棧區(qū)和必要的內存空間, 此時, 這個新建立的線程就進入了 Runnable 隊列, 等待被調度
一旦某個在 Runnable 隊列中的線程準備執(zhí)行, 那么一個虛擬的 CPU 則會分配給它, 而其此時進入了 Running 狀態(tài)
當某個 Running 狀態(tài)的線程因為等待某個資源(譬如調用了 scanf 函數(shù), 等待用戶輸入), 而卡住時, 其會被重新放入到 Runnable 隊列中, 同時, 其在 Running 狀態(tài)得到的寄存器信息等等, 即其運行狀態(tài), 會被保存到一個區(qū)域, 當其等待的資源準備好后, 重新運行該線程, 并將其寄存器信息恢復, Running ==> Runnable 這個過程稱為 context switch(上下文切換)
當某線程等待的資源暫時無法得到或無法確定大概需要多久能得到時, 那么該線程則會被放到 Stopped 隊列中, 當該資源準備好后, 則該線程又會回到 Runnable 隊列中
如果一個線程執(zhí)性完成, 那么其就會進入 Terminated 狀態(tài), 被放入 Terminated 中的線程無法再回到 Runnable 隊列中來
假設一個 CPU 最多只能支持 20核, 如果某個程序launch了25個線程, 有可能其效率還是比launch 20個線程高, 這是因為 context switch(上下文切換) 的緣故

2. GPU 并行

2.1 GPU / CUDA 概述

在游戲領域, 3D 人物的建模都是用一個個小三角形拼接上的, 而不是以像素的形式, 對多個小三角形的操作, 能使人物做出多種多樣的動作, 而 GPU 在此處就是用來計算三角形平移, 旋轉之后的位置, 如下圖
- 6.jpg
  
  *
  
  E2C991444715105B658510DADAB9FF7D.jpg
- 而為了提高游戲的分辨率, 程序會將每個小三角形細分為更小的三角形
- 每個小三角形包含兩個屬性, 它的位置和它的紋理
在游戲領域應用的 GPU 與科學計算領域的 GPU 使用的不同是, 當通過 CUDA 調用 GPU 來進行科學計算的時候, 計算結果需要返回給 CPU, 但是如果用 GPU 用作玩游戲的話, GPU 的計算結果直接輸出到顯示器上, 也就不需要再返回到 CPU
因為 CPU 和 GPU 的硬件架構上的不同, 譬如 CPU 的內存是用 DDR4 而 GPU 的內存是用 GDDR5, 它們的 ISAs (Instruction set architecture) 應該完全不同, 但是 nvcc 編譯器解決了這一問題, 其編譯的 .cu 文件(類 c 語言), 將其中涉及 CPU 的代碼編譯為 CPU 指令, 而將涉及 GPU 的代碼編譯為了對應的 GPU 指令, 大大簡化了 CUDA 編程的難度毯炮。

2.1 GPU 通用架構

GPU 軟件架構

每 32 個線程組成一個線程束 (warp), 一個線程束可以被看做是程序執(zhí)行的一個最小單元, 而 GPU 也是基于線程束來執(zhí)行程序, 假如某個代碼只需要 12 個線程, 那么 GPU 也會啟動 32 個線程來執(zhí)行這個代碼, 其中的 12 個線程就被完全浪費了
但一般來說, 一個僅包含 32 個線程的 warp 實在是太小了, 因此我們就將多個線程束合在一起, 組成 block, 常用的 block 大小是 32, 64, 128, 256, 512, or 1024 threads/block
- 每個 warp 相當于是 code execution 的最小單元
- 每個 block 相當于是 code launch 的最小單元
而多個 block 在一起組成一個 grid, grid 可以使 1D, 2D或3D 的 blocks 的排列
- 雖然GPU可以用3D grid, 譬如 x, y , z 每一維度都可以 launch 假設200個block, 總共200^3=8000000 個block, 但不代表這 800000 可以在用 1D grid 的時候全部用一個維度 launch, 譬如對于 GT630 GPU 而言其 x 維度最大只支持 65535 個block
block 之間不能有資源上的相互依賴
假設一個 block 有 256 個線程, 不代表每個線程都會被即刻執(zhí)行, 其首先會被分為 ceil(256/32)=8 個 warp, 然后將這 8 個 warp 按從 0 到 7 編號, 然后按 warp0, warp1, …, warp7 的書序執(zhí)行
- 正常情況下, 在寫 CUDA 程序時, 我們是不用考慮 warp 的 ID 的, 僅當我們寫底層的 CUDA assembly language (PTX) 編碼時才需要考慮 warp ID
CPU 和 GPU 之間的任何交互都會通過 NRE (Nvidia Runtime Engine), NRE 存在于顯卡驅動中
- 當 NRE 檢測到代碼中有錯誤的地址訪問時, 會終止程序, 但是需要注意的是, 程序是運行在 CPU 中的, 所以僅僅操作系統(tǒng)才可以終止程序, 所以 NRE 會通知操作系統(tǒng), 然后操作系統(tǒng)將該程序終止
每一代的 Parallel Thread Execution (PTX) 和 Instruction Set Architecture (ISA) 都會改進
- PTX 是 Intermediate Representation (IR) 并且對于不同操作系統(tǒng)而言都是一樣的, 我們在用 nvcc 編譯 .cu 文件時, 可以加上 --ptx 來生成 PTX代碼, 也可加上 --cubin 來生成 CUDA 的二進制文件, CUBIN 針對不同系統(tǒng)是不同的

GPU 硬件架構

7.jpg

Giga Thread Scheduler (GTS) 模塊用來管理安排 block, 即每當一個 SM 完成當成工作后, GTS 會給其安排一個新的 block
一個 GPU 有越多的 SM, 其速度就越快
雖然每個 SM(Streaming Multiprocessor) 一次只能執(zhí)行一個 block, 但其可以一次接收多個 block, 并將其放入一個隊列中緩存起來
假設 launch 了 166656 個 block, 但是一個 Pascal 顯卡的 SM 最多接收 32 個 block, 而該顯卡有 60 個 SM, 32*60=1920, 遠遠小于 launch 的 166656 個block, 那么 GTS 就讓剩下的 block 就排在隊列之中等待之前的 SM 執(zhí)行完畢后空出來, 同時需要注意的是, 分配給每個 SM 的 32 個 block 中, 一次只能執(zhí)行一個 block, 其他的 31 個 block 都需要等待, 同時 GTS 也會將 blockDim, blockId 以及 gridDim 等參數(shù)隨著 block 傳入到 SM 中, 用來幫助計算每個 thread 的 ID
每一個 block 都會得到一個當前執(zhí)行的 CUDA 命令代碼的二進制版本, 即 CUBIN, 而每個 SM 在收到一個 block 時會將其包含的 CUBIN 緩存在 SM 中的指令緩存中
假設現(xiàn)在有 6 個 SM, 那么加載 blocks 的順序就是:
- Block0→SM0, Block1→SM1, Block2→SM2, Block3→SM3, Block4→SM4, Block5→SM5
- Block6→SM0, Block7→SM1, Block8→SM2, Block9→SM3, Block10→SM4, Block11→SM5
- 那么每個 SM 中緩存 Block 的隊列即為:
- SM0 =? [ Block0, Block6, Block12, Block18, Block24, Block30, Block36, Block42 ]
- SM1 =? [ Block1, Block7, Block13, Block19, Block25, Block31, Block37, Block43 ]
- …...
- SM5 =? [ Block5, Block11, Block17, Block23, Block29, Block35, Block41, Block47 ]
當 SM 中當前執(zhí)行的 block 因為訪問某些資源而阻塞時, SM 可以選擇執(zhí)行隊列中的其他的 block, 以避免 SM 空閑下來, 因此, 對于 SM1而言, 有可能出現(xiàn)的情況是, Block7 率先完成了工作, 此時 GTS 就會將下一個 Block 傳入到 SM1 中原來 Block7 所在的位置, 因為下一個需要執(zhí)行的是 Block48 (Block47 已經傳給 SM5 了), 那么此時 SM1 的隊列就是
- SM1 =? [ Block1, Block48, Block13, Block19, Block25, Block31, Block37, Block43 ]
當 GTS 給一個 SM 分配一個 block 后, 其就與這個 block 沒關系了, 就該輪到該 SM 給這個 block 分配 thread ID, 緩存資源, 寄存器等

2.2 GPU 不同系列架構

Fermi 架構

8.jpg

9.jpg

SFU 是 Special Function Unit, 其用來執(zhí)行一些特殊的函數(shù), 譬如 sin(), cos(), log() 等
LD/ST 是 Load/Store, 對于內存數(shù)據(jù)的讀寫請求都被放置在該模塊的隊列中, 當讀寫完成后, 該請求從隊列中移除
每個 Core 都包含一個 Floating Point (FP) 和一個 Integer (INT) 執(zhí)行單元用來執(zhí)行浮點或整型的指令
L1 緩存用來儲存那些常用的數(shù)據(jù), 其與 Shared Memory 共用 64 KB 的大小, 64 KB 被分為 (16 KB+48 KB) 或 (48 KB+16 KB)

Kepler 架構

10.jpg

11.jpg

SM 在 Kepler 架構中被稱為 SMX
DPU 是 Double Precision Unit, 即可以更高效地計算 double 類型數(shù)據(jù)

Maxwell 架構

12.jpg

13.jpg

SM 在 Maxwell 中被稱為 SMM
一個 SMM 包含四個 sub-uint, 其中每個 sub-unit 包含 32 個 core, 這四個 sub-unit 共用一個 Instruction Cache, 而每個 sub-unit 獨占一個 Instruction Buffer

Pascal 架構

14.jpg

15.jpg

High Bandwidth Memory (HBM2), 其通過使用4096-bit 內存總線帶寬, 最高可支持 720 GBps 傳輸速率
GTS 在此處稱為 GIGA THREAD ENGINE (GTE)

模塊介紹

FPU (Floating Point Unit)
- FPU 可以用來計算 double 類型數(shù)據(jù), 其通過不斷地循環(huán)來計算
- 對于那些沒有 DPU 模塊的 GPU, 在使用 FPU 來計算 double 類型數(shù)據(jù)時 (譬如 24 位或 32 位的double), 其耗時是計算單精度浮點類型數(shù)據(jù)的 24 或 32 倍
DPU (Double Precision Unit)
- DPU 往往會比 Core 在物理上大一些, 因為隨著尾數(shù)的增加, 乘法器的大小呈平方擴大, 即一個單浮點類型需要 23 位尾數(shù), 而一個雙精度浮點需要 52 位尾數(shù), 故與 FPU 相比一個 GPU 大概是其四倍的大小 (注意 FPU 是包含在 Core 中的)
SFU (Special Function Unit)
- 其用來計算一些特殊的運算, 如 sin(), cos(), exp(), log(), sqrt() 等
RF (Register FIle)
- 假設現(xiàn)在程序中有如下變量
  - double R, G, B
  - unsigned int ThrPerBlk, MYbid, MYtid, MYgtid,
  - unsigned int BlkPerRow, RowBytes, MYrow, MYcol,
  - unsigned int MYsrcIndex, MYpixIndex
- 假設每個寄存器是 32-bit 的, 那么每個 double 類型都需要 2 個寄存器, 而每個 unsigned int 需要 1 個, 因此為了儲存這些變量總共需要耗費 16 個寄存器, 而編譯器可能還需要消耗一些寄存器儲存臨時運算的值, 那么假設該程序總共需要耗費 24 個寄存器
- 假設我們每個 block 有 128 個線程, 那么每個 block 就需要 24*128=3072≈3K 個寄存器, 而假設 Pascal 架構的 GPU 每個 SM 的隊列中最多可容納 32 個 block, 那么總共需要 32*3K=96K 個寄存器, 而若每個 SM 只有 32K 個寄存器(總共 128KB), 因此對于 GTE 而言, 因為寄存器(RF)數(shù)量的限制, 其最多只能向每個 SM 中加載 10 個 block
- 所以在設計 CUDA 程序時, 盡可能地限制 RF 的使用
LDST (Load/Store Queues)
- 該模塊用來在 core 和 memory 之間傳輸數(shù)據(jù), 當一個 core 請求從 memory 中讀取或寫入數(shù)據(jù)時, 該請求則會被儲存在 LD/ST 中, 并等待其完成, 在等待其完成的過程中, 另外一個 warp 會被啟動并執(zhí)行
L1 和 Texture Cache
- L1 緩存是硬件控制的緩存, 即代碼無法控制 L1 緩存中儲存的數(shù)據(jù), L1 緩存用以儲存常用的數(shù)據(jù)
- Texture 緩存用來儲存之前提到的游戲建模中小三角形的紋理
Shared Memory
- 共享內存是代碼控制的儲存, 通過 CUDA 編程, 可以顯式地規(guī)定哪部分數(shù)據(jù)儲存在共享內存中
- 當使用 Shared Memory 的時候要注意內存大小, 不然有可能比沒用 Shared Memory 速度還慢, 因為當 shared memory 用完的時候, 沒有資源供新的 block 執(zhí)行, 故新的 block 會被阻塞, 而對于那些沒有用 shared memory 的, 雖然沒有 shared memory 稍顯耗時, 但其 block 不會被阻塞, 所以最終可能速度更快
Constant Cache
- 其用來保存不可變的值, 即常量
- 其只被寫入一次(即初始值), 但可被多次讀取
Instruction Cache
- 其用來保存當前 SM 所執(zhí)行 Block 中的指令, 每個 block 中都包括其所需執(zhí)行的指令
Instruction Buffer
- 其用來保存當前 SM 的局部指令, 即從 Instruction Cache 中復制來的指令, Instruction Buffer 相對于 Instruction Cache 的關系就像 L1 相對于 L2 的關系
Warp Schedulers
- 當每個 block 被傳遞給 SM 時, warp scheduler 用來將每個 block 轉化成線程束的形式, 即假設我們啟動了 256 threads/block 那么就會轉化成 8 warps/block, 即轉化為如下形式
  - schedule warp0: gridDim.x=166656, blockDim.x=256, blockIdx.x=0
  - schedule warp1: gridDim.x=166656, blockDim.x=256, blockIdx.x=0
  - …...
  - schedule warp7: gridDim.x=166656, blockDim.x=256, blockIdx.x=0
- 需要注意的是此處只是 schedule 了, 而不是 dispatch, 所以此處沒有分配 thread ID,僅當所有資源都準備好時, 才會 dispatch
Dispatch Units
- 該模塊在每個 warp 已經被 schedule 后, 并且所有的資源都準備完畢時啟動
- 其會給每個線程分配 threadIdx.x, threadIdx.y, 和 threadIdx.z, 即如下形式
  - gridDim.x=166656, blockDim.x=256, blockIdx.x=0, threadIdx.x=0…32
  - 32 個 core, 每個 core 負責執(zhí)行一個線程
GLOBAL MEMORY
- 全局內存與 DDR4 內存一樣, 一次讀取一塊內存的效率比一個比特一個比特讀取的效率更高

2.3 CPU / GPU 傳輸

一些定義
- latency(延遲), 第一個包從發(fā)出到收到的時間差
- throughput(吞吐量), 在一段時間內的平均傳輸速率, 即傳輸?shù)目偭砍钥倳r間
- bandwidth(帶寬), 最大的 throughput(吞吐量)
- UpStream 帶寬是 CPU->GPU 的帶寬, 而 DownStream 帶寬是 GPU->CPU 的帶寬, 而 PCIe 的好處是支持 UpStream和DownStream 的同時傳輸
當數(shù)據(jù)通過 PCIe 傳輸時, 無論是從 CPU->GPU 還是 GPU->CPU, CPU 始終通過虛擬內存頁來參與到這個過程中
僅當 CPU, GPU 以及主板都支持 PCIe 的某個特定版本時, 譬如 PCIe 3.0, 才能充分發(fā)揮 PCIe 的傳輸速率, 如果其中某一個不支持, 那么則無法使用 PCIe 3.0
??

16.jpg
- 圖中的內存總線是 L3 與 DDR4 之間的連接部分, 68 GB/s, 即 CPU 緩存與內存的傳輸速率
- PCIe 3.0 是 PCI EXPRESS BUSS, 傳輸速率為 16 GB/s
- GPU 內部的儲存, 被稱作 GPU global memory, 即圖中的 GDDR5, 其與 GPU L2 緩存交互的速率是 336 GB/s , 需要注意的是 GPU 最多只有 L2, 沒有 L3 緩存, 即 GPU 中的 L2 緩存與 CPU 中的 L3 緩存都是最后一層緩存 Last Level Cache (LLC), LLC 直接與內存交互
17.jpg
- 上圖是不同內存的帶寬

2.4 限制 SM 速度的因素

Shared Memory
- 如果程序被設計為需要大量依賴共享內存, 那么當共享內存耗盡時, SM 就會被阻塞, 而不能在不同 block 之間切換執(zhí)行
Register File
- 使用的寄存器的數(shù)量與代碼中定義的變量數(shù)量有關, 其用來保存 kernel 中變量的值
- 假設每個 kernel 中可使用的寄存器是 255 個, 如果使用的寄存器超出了這個部分, 則其就會用內存當做寄存器, 而內存當做寄存器很顯然會慢, 而每個 block 中最多有 32K 或 64K 個寄存器, 而對每個 block 允許的寄存器的數(shù)量的限制, 直接影響了每個 block 中所能容納的 thread 數(shù)量
  - if your kernels require 240 registers, launching 512 threads/block will require 512 × 240 = 122,880 = 120 K registers inside the register ?le. However, even if your Compute Capability allows 64 K registers, this is still more than the 120 K you need. This means that you cannot launch more than 256 threads/block before hitting the register ?le limitation, which would require 256 × 240 = 61,440 = 60 K registers
  - 假設當前的 kernel 需要 240 個寄存器, 那么啟動 512 threads/block 就需要 512 × 240 = 122,880 = 120 K 個在RF中的寄存器, 然而, 假設根據(jù)我們當前的 compute capability, 我們只能允許 64K 個寄存器, 很顯然是不夠的, 因此我們不被允許啟動 512 threads/block, 而最多只能啟動 256 threads/block, 因為 256 × 240 = 61,440 = 60 K < 64 K
Threads / Block
- 如果該值太小, 即每個 block 過小, 那么我們就需要更多的 block 完成特定任務, 但是每個 SM 一次啟動的 block 數(shù)量是有上限的
- 如果該值太大, 那么有可能會造成浪費, 譬如我們總共需要 5220 個線程, 而我們的 block 大小是 512 threads / block, 那么我們就需要 11 個 block, 11 * 512 - 5220=412, 即我們浪費了 412 個線程
Warp
- 每個 SM 中可容納的 warp 數(shù)量, 假設該值是 64, 那么一個 warp 中有 32 個線程, 那么也就意味著每個 block 最多可容納 32*64=2048 個線程
我們將上述情況匯總, 以一個例子來看一下:
- 現(xiàn)在有以下條件
  1. 假設 block 大小是 320 threads/block (10 warps/block)
  2. 假設 GTS 給當前 SM 分配了 6 個 block
  3. 每個 SM 中可容納 64 個 warp
  4. 假設每個 block 都需要 10 KB 的 shared memory 大小
  5. 當前 SM 可允許的最大 shared memory 大小是 48 KB
- 根據(jù)前三個條件, 因為 SM 分配了 6 個 block, 故總共啟動了 6*10=60 個 warp, 也就是說每個 SM 浪費了 64-60=4 個warp, 此時的占用率是 60/64≈94%
- 根據(jù)后兩個條件, 因為 shared memory 的限制, 我們最多只能啟動 4 個 block, 因為 4*10=40<48KB, 此時占用率為 4*10(warps/block) / 64 ≈ 63%
- 假如我們將 block 的大小改為 512 threads/block (16 warps/block), 此時每個 block 依然只消耗 10 KB 的 shared memory, 那么以上條件不變的話, 我們啟動 4 個 block, 就可以達到 4*16/64=100% 的占用率

2.5 設計 CUDA 代碼時的一些注意事項

通過如下代碼我們在 main 函數(shù)中定義了一個 buffer 變量, 并給其分配了 GPU 上的內存
- unsigned char *buffer;
- cudaMalloc((void**)&buffer, SIZE);
- 你可能會感到困惑, 從 buffer 的定義方式上來看, 其很顯然是一個 CPU 上的指針變量, 指向 GPU 上的某塊內存區(qū)域, 這說法看起來很讓人困惑, 因為在常理中, CPU 上的指針變量指向的應該是 CPU 的內存地址
- 我們首先來看一個指針變量是什么, 一個指針變量無非就是一個 64位的integer , 這個 integer 代表了內存上某塊地址, 正因為指針變量是一個 integer, 所以其可以進行加減操作, 通過加減操作可以指向當前地址前后的不同地址
- 那么對于這個指針變量本身, 其只知道自己儲存的是一個 64位的integer, 而無從了解其是內存中的還是GPU內存中的地址, 而判斷這個的是 nvcc 編譯器, 其通過 buffer 這個 64位integer 的值, 判斷出其是屬于內存還是GPU內存, 然后根據(jù)其所屬地, 分配出一塊空間
GPU 對于 a*(b+c) 有特殊的運算操作符, 但是對于 a+b+c 沒有, 這就可能導致 a+b+c 反而比 a*(b+c) 還慢
OR, AND, 移位操作 (>>) 等在 GPU 中計算速度很快, 可以在代碼中盡可能使用, 譬如如下代碼
- //NOW: A=[B1,R0,G0,B0] B=[G2,B2,R1,G1] C=[R3,G3,B3,R2], 其中 A, B, C 都是 32 bit 變量, 這三個變量每個都包含 4 個內部值, 每個值都是 8 bit
- // D=[B2,R3,G3,B3]
- D = (C >> 8) | ((B << 8) & 0xFF000000); //其中 0xFF000000 轉化為二進制為 11111111,00000000,00000000,00000000
- // E=[G1,B1,R2,G2]
- E = (B << 24) | (B >> 24) | ((A >> 8) & 0x00FF0000) | ((C << 8) & 0x0000FF00); // 0x0000FF00 轉化為二進制為 00000000,00000000,11111111,00000000
- // F=[R0,G0,B0,R1]
- F=((A << 8) & 0xFFFF0000) | ((A >> 16) & 0x0000FF00) | ((B >> 8) & 0x000000FF);
線程束分化 (thread divergence), 當一個 warp 中, 因為數(shù)據(jù)不同, 不同線程運行同一份代碼時, 根據(jù) if 而給出不同的 TRUE/FALSE 結果, 并因為結果不同而執(zhí)行不同的操作時, 就會導致線程束分化, 線程束分化會影響并行效率, 因為僅當一個 warp 中所有線程都在做相同的事時, 效率才最高

本文主要內容來自書籍: GPU Parallel Program Development Using CUDA (Tolga Soyata)
[http://www.hds.bme.hu/~fhegedus/00%20-%20Numerics/B2018%20GPU%20Parallel%20Program%20Development%20Using%20CUDA.pdf]

CPU / GPU原理與 CUDA