CUDA平行計算 | 執行緒模型與記憶體模型
文章目錄
前言
CUDA（Compute Unified Device Architecture）是顯示卡廠商NVIDIA推出的通用平行計算平臺和程式設計模型，它利用NVIDIA GPU中的平行計算引擎能更有效地解決複雜的計算問題。通過使用CUDA，開發人員可以像在CPU上那樣直接訪問GPU裝置的虛擬指令集和儲存裝置,大大提高了GPU演算法或程式的開發效率。CUDA平臺可以通過CUDA加速庫、編譯器指令、應用程式設計介面以及行業標準程式語言的擴充套件（包括C、C++、Fortan、Python）來使用。
在CUDA學習的過程中，CUDA的執行緒模型和記憶體結構是必須要掌握的基礎知識，這決定了你是否能夠寫一個完整的CUDA程式。CUDA演算法開發通常有兩個步驟：演算法初稿、效能優化。如果不掌握執行緒模型和記憶體模型，演算法初稿就無法完成，更談不上效能。而效能優化更是對這兩個方面有著較高的要求，需要不僅掌握，而且深入掌握。
CUDA執行緒模型（如何組織執行緒）
在CUDA執行緒模型中，執行緒（Thread）是GPU的最小執行單元，能夠完成一個最小的邏輯意義的操作，每個執行緒都有自己的指令地址計數器和暫存器狀態，利用自身的資料執行當前的指令。而執行緒束（Warp）則是GPU的基本執行單元，包含32個執行緒，GPU每次呼叫執行緒都是以執行緒束為單位的，在一個執行緒束中，所有的執行緒按照單指令多執行緒（SIMT）方式執行，即所有執行緒執行相同的指令。多個執行緒束位於一個最高維度為3的執行緒塊（Block）中，同一個執行緒塊中的所有執行緒，都可以使用共享記憶體來進行通訊、同步。執行緒塊又通過一個最高維度為3的網格（Grid）來管理。CUDA的執行緒結構圖如圖1所示：
圖1 CUDA執行緒結構
舉個最簡單常用的例子，對於二維影像，我們可以一個執行緒對應一個畫素，一個執行緒塊對應一塊區域，一個Grid對應整個影像。
CUDA記憶體模型（瞭解不同記憶體優缺點，合理使用）
CUDA記憶體模型中，有兩種型別的儲存器：不可程式設計儲存器和可程式設計儲存器，前者並不對開發人員開放存取介面，包括一級快取和二級快取；後者可以顯式地控制資料在記憶體空間中的存取，包括暫存器、共享記憶體、本地記憶體、常量記憶體、紋理記憶體以及全域性記憶體。可程式設計儲存器的結構如圖2所示：
圖2 CUDA記憶體結構
它們的特點如下：
暫存器：暫存器是GPU上執行速度最快的記憶體空間，分配於每個執行緒中，儲存執行緒核函式中宣告的沒有其它修飾符的自變數。暫存器對於每個執行緒來說是私有的，數量也是有限的，在Fermi架構中，每個執行緒最多有63個暫存器，Kepler架構將限制數量擴充套件至256個，線上程核函式中使用較少的暫存器可以使SM中有更多的常駐執行緒塊，增加使用率和效能。而如果一個核函式使用了超過了硬體限制數量的暫存器，將會用本地記憶體代替多出的暫存器，降低演算法效能。
本地記憶體：執行緒核函式中原則上應該儲存在暫存器中但由於某些原因（如暫存器數量使用超出限制或變數佔用記憶體過大）而無法進入暫存器空間的變數將溢位到本地記憶體中，本地記憶體訪問符合高效記憶體訪問要求，對於計算能力在2.0以上的GPU來說，本地記憶體資料儲存在每個SM的一級快取和每個裝置的二級快取中。
共享記憶體：共享記憶體是片上記憶體，與本地記憶體和全域性記憶體相比具有更高的頻寬和更低的延遲。共享記憶體由執行緒塊分配，宣告週期伴隨著整個執行緒塊，執行緒塊中的每個執行緒都可以共享其儲存空間，這也是共享記憶體的意義所在，當一個執行緒塊執行完畢，其分配的共享記憶體將被釋放且重新分配給其他執行緒塊。
共享記憶體是同一個執行緒塊的執行緒之間通訊的基本方式，一個塊內的執行緒可以通過使用共享記憶體而相互合作。常用的合作方式是在計算前，將全域性記憶體讀進共享記憶體中，而讀取的方式是每個執行緒負責讀取某一個位置的資料，讀完之後塊內的所有執行緒都能夠使用整個共享記憶體中的資料。因為共享記憶體的延遲低頻寬高，所以這種方式比直接讀取全域性記憶體要高效得多。需要注意的是讀取全域性記憶體至共享記憶體時要注意同步（執行緒塊內只有一個執行緒束的情況除外），在CUDA C中同步的方式是使用執行緒同步函式__syncthreads()來設立一個執行障礙點，即同一個執行緒塊中的所有執行緒必須等待其他所有執行緒執行到這個函式處才能往下執行，這樣可以確保需要的全域性記憶體被全部載入共享記憶體，避免潛在的資料衝突。
核函式中儲存在共享記憶體的變數通過修飾符__shared__修飾。
__shared__ int shared_memory[];
常量記憶體：常量記憶體駐留在裝置記憶體中，並在每個SM專用的常量快取中快取,常量記憶體變數使用修飾符__constant__修飾，必須在全域性空間內和所有核函式之外進行宣告，對同一編譯單元中的所有執行緒核函式可見，核函式只能對常量進行讀操作。
紋理記憶體：紋理記憶體駐留在裝置記憶體中，並在每個SM的只讀快取中快取，是一種通過指定的制度快取訪問的全域性記憶體，是對二維空間區域性性的優化，訪問二維資料時可以達到最優效能。
全域性記憶體：全域性記憶體常駐於裝置記憶體中，是GPU中最大、延遲最高、最常使用的記憶體，它貫穿於應用程式的整個生命週期。全域性記憶體通過快取來實現資料存取，對全域性記憶體進行訪問時，必須注意記憶體訪問的兩個特性：對齊記憶體訪問和合並記憶體訪問。
當訪問記憶體的第一個地址是快取粒度的偶數倍時（二級快取為32位元組，一級快取為128位元組），滿足合併訪問，可以獲得更高的訪問效率，而非對齊訪問則會造成頻寬浪費。
當一個執行緒束中全部的32個執行緒訪問一個連續的記憶體塊時，滿足合併記憶體訪問，可以最大化全域性記憶體吞吐量。這是由於GPU可通過一次定址和一次讀寫指令對連續的32位元組（二級快取）或128位元組（一級快取）進行一次讀取，如果滿足合併訪問，則一個執行緒束通過一次定址即可完成訪問，效率非常高；反之如果不滿足合併訪問，則最壞的情況需要32次定址才能完成訪問，這是對記憶體頻寬的一種極大的浪費。圖3描述了合併訪問與非合併訪問的兩種方式。
圖3 合併訪問示意圖
從個人專案經驗來說，
1. 全域性記憶體是不得不用，但是要儘量減少讀寫次數，因為它真的很慢。對於新手，演算法初稿可以全部使用全域性記憶體，可以將演算法最快實現，後面再去優化。對於老手，肯定是在演算法初稿就會考慮更快的讀寫方式了。
2. 共享記憶體是優化演算法記憶體讀寫效率的利器，是減少全域性記憶體讀寫次數的不二選擇，這得益於同一個執行緒塊中的執行緒共享同一塊共享記憶體，所以最常見的思路就是將需要重複讀寫的全域性記憶體一次性讀入共享記憶體中反覆使用。
3. 暫存器是速度最快的，奈何一個kernel所能使用的數量實在有限，所以一般情況下要使用最大化，將能用的都用上，這就要考效能監測工具來監測每個執行緒的使用量了，如NVIDIA Nsight Visual Studio Edition，效能優化不可缺工具之一 link。
當然也有一些演算法，使用超量的暫存器，給每個執行緒增加運算元（記憶體操作、運算操作）來減少延遲Latency（記憶體延遲、運算延遲），可以達到更高的效能。詳細請參考 link。因為在CUDA運算裡，延遲確實非常恐怖，尤其是全域性記憶體讀寫延遲，高達400 ~ 600個時鐘週期，也就是說從你發讀寫指令到指令執行完，需要400 ~ 600時鐘週期。所以通常通過開闢大量的執行緒以及執行緒內增加併發指令來隱藏延遲。