說明

最近在學習CUDA，感覺看完就忘，於是這裡寫一個導讀，整理一下重點

主要內容來源於NVIDIA的官方文檔《CUDA C Programming Guide》，結合了另一本書《CUDA並行程序設計 GPU編程指南》的知識。 因此在翻譯總結官方文檔的同時，會加一些評註，不一定對，望大家討論指出。

另外，我才不會老老實實的翻譯文檔，因此細節還是需要從文檔里看的。

看完兩份文檔總的來說，感覺《CUDA C Programming Guide》這本書作為一份官方文檔，知識細碎且全面，且是針對最新的Maxwell、Pascal、Volta架構的闡述。但相對來說不夠深入，且有關程序設計方面所述甚少。

而《CUDA並行程序設計 GPU編程指南》這本書，講解的比較深入，不僅闡述了NVIDIA GPU的特性，並且在程序設計方面有比較深入的見解。美中不足的是該書是針對老舊的Tesla、Fermi架構GPU，沒有涉及到新架構的新特性。

Chapter 1 簡介

1.1 從圖形處理到通用並行計算

GPU是能夠高度並行化、具有很多處理器核心的器件，具有很強的計算能力和內存帶寬。下圖是CPU和GPU在浮點運算上的性能對比發展趨勢。

NVIDIA GPU和 Intel CPU 浮點計算能力對比

可以看到，NVIDIA的GPU在浮點運算能力上，吊打了Intel的CPU。其原因來自於CPU和GPU結構上的差異。

如下圖所示，CPU僅僅具有有限的核心數量。相比於GPU，CPU的核心屬於「少而精」的存在，核心數雖然很少，但是每個核心的性能很強，適合處理具有很多分支的複雜的邏輯。近些年來，CPU中集成了一些並行指令集，如SSE、AVX等，其中AVX可以同時處理256位(32個位元組)，可以大大加速並行計算。但是相比於GPU，還是小巫見大巫。

CPU 與 GPU 內部結構對比

GPU的設計理念與CPU不同，GPU具有龐大的核心數。以TITANX為例，流處理器（等同於CPU的核心）達到3072個之多。這意味著相比於8核的CPU處理器，TITANX可以同時並行處理384倍的任務。但是GPU的單個核心不如CPU的核心強大。因此相對於CPU，GPU更適合處理高度並行化的任務。

從另一個視角來看上圖，在CPU晶元中，運算單元(ALU)所佔的比例較小，CPU中更多的矽片被用來製作控制單元和緩存，以完成複雜的邏輯；而GPU的運算單元使用的矽片面積比例要大於CPU，以完成高強度的計算。因此CPU的側重點在於邏輯控制，而GPU的側重點在於計算。

深度學習，尤其是卷積神經網路中，有很多可以高度並行化的向量運算與矩陣運算。因此使用GPU進行深度學習運算，遠比CPU快速。

另外還需注意的是，由於GPU的核心數量太多，即使GPU的內存（也稱作顯存）優於CPU的內存，其內存仍然是瓶頸。因此，GPU希望程序是 計算密集型 而不是 內存密集型。

1.2 CUDA

CUDA(Compute Unified Device Architecture)，是NVIDIA推出的通用並行計算平台和編程模型。CUDA是在底層API的基礎上，封裝了一層，使得程序員可以使用C語言來方便的編程。

CUDA還支持C++/Python等更高級的語言編程；此外，NVIDIA還提供了CuDNN、TensorRT、NPP等更高級的庫函數。

各代顯卡、CUDA、上層庫之間的關係

從上圖中也可以看出各個系列的GPU屬於哪些架構、什麼定位。例如GeForece 1000系列，就是使用Pascal架構的消費顯卡。

1.3 可擴展的編程模型

CUDA的編程模型，使得同一個CUDA程序，可以在不同的顯卡上運行。

CUDA編程模型

如上圖所示，CUDA程序一般會創建一些線程塊(Block)，線程塊會被調度到空閑的流處理器簇(SM)上去。當線程塊執行完畢後，線程塊會退出SM，釋放出SM的資源，以供其他待執行線程塊調度進去。

因此，無論是只有2個SM的GPU，還是有4個SM的GPU，這些線程塊都會被調度執行，只不過是執行的時間有長有短。因此，同樣的程序，可以在具有不同SM數量上的GPU運行。

1.4 Document Structure

略

Chapter 2 編程模型

2.2 線程層級

在講解內核函數前，先講解一下線程層級，不然有點難講。

CUDA編程是一個多線程編程，數個線程(Thread)組成一個線程塊(Block)，所有線程塊組成一個線程網格(Grid)，如下圖所示：

CUDA線程層級

圖中的線程塊，以及線程塊中的線程，是按照2維的方式排布的。實際上，CUDA編程模型允許使用1維、2維、3維三種方式來排布。另外，即使線程塊使用的是1維排布，線程塊中的線程也不一定要按照1維排，而是可以任意排布。

目前的GPU限制一個線程塊中，最多可以安排1024個線程。

一個線程塊用多少線程，以及一個線程網格用多少線程塊，是程序員可以自由安排的。由於32個相鄰的線程會組成一個線程束(Thread Warp)，而一個線程束中的線程會運行同樣的指令。因此一般線程塊中線程的數量被安排為32的倍數，選用256是比較合適的。

在線程數定下來之後，一般根據數據的排布情況來確定線程塊的個數。

例如：一個數組的長度為4096，安排每個線程處理一個元素。如果安排一個線程塊為256個線程，則需要4096/256=16個線程塊。

2.1 內核函數(Kernels)

內核函數是CUDA 每個線程 執行的函數。CUDA使用擴展的C語言編寫內核函數，關鍵字為global。內核函數返回值只能是void。

下面是一段簡單的內核函數，用於求兩個數組的和：

// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
int i = threadIdx.x;
C[i] = A[i] + B[i];
}

int main()
{
...
// Kernel invocation with N threads
VecAdd<<<1, N>>>(A, B, C);
...
}

threadIdx.x是線程在所處線程塊中的X方向的ID。由於本例中是定義的1維排布，因此X方向ID即為線程的ID。

由於GPU中的每個線程都會執行相同的VecAdd函數，因此不同的線程需要使用自己獨有的ID來區分彼此，來獲取不同的數據。這就是SIMT的概念，即「相同指令，不同線程」。

在main()函數中，我們注意到，VecAdd函數的調用使用了<<<blockPerGrid, threadsPerBlock>>>關鍵字。這是調用內核函數所獨有的。程序員通過該關鍵字，制定網格中線程塊和線程的排布方式。排布方式與數據息息相關。

下面舉一個2維排布的例子，用於做矩陣加法：

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
float C[N][N])
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
if (i < N && j < N)
C[j][i] = A[j][i] + B[j][i];
}

int main()
{
...
// Kernel invocation
dim3 threadsPerBlock(16, 16);
dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
...
}

首先看主函數，當排布不使用1維時，需要使用dim3數據類型。該程序每個線程塊中線程為16x16排布，而線程塊的排布依賴於數據的多少。

在內核函數中，i代表x方向上的ID，j代表y方向上的ID。blockDim代表當前線程塊的尺寸。從程序中可以看到，x方向為行方向，y方向為列方向。（注意，這裡官方文檔裡面寫的有些錯誤）

每個線程讀取自己ID對應的數據A[j][i]和B[j][i]，並將結果寫回C[j][i]。其中A、B、C都存儲在GPU的全局內存上（後面會提及）

2.3 內存層級

同CPU一樣，GPU也有不同層級的內存。越靠近核心的內存速度越快，但容量越小；反之，越遠離核心的內存速度越慢，但容量較大。

CUDA內存層級

上圖是NVIDIA設備的硬體示意圖。

  • 最上方是主機端內存(host memory)，指的就是我們常說的內存。一般主機端內存通過PCI-E匯流排與設備端內存交換數據。數據交換的速度等於PCI-E匯流排的速度。
  • 全局內存(global memory) 、常量內存(constant memory)、紋理內存(texture memory)、本地內存(local memory)。都位於GPU板上，但不在片內。因此速度相對片內內存較慢。 常量內存和紋理內存對於GPU來說是只讀的。
  • GPU上有 L2 cache和 L1 cahce。其中L2 cache為所有流處理器簇(SM)共享，而L1 cache為每個SM內部共享。這裡的cache和CPU的cache一樣，程序員無法對cache顯式操縱。
  • 紋理緩存和常量緩存在SM內部共享，在早期1.x計算能力的時代，這兩種緩存是片上唯一的緩存，十分寶貴。而當Fermi架構出現後，普通的全局內存也具有了緩存，因此就不那麼突出了。
  • 共享內存(shared memory, SMEM) 具有和L1緩存同樣的速度，且可以被程序員顯式操縱，因此經常被用作存放一些需要反覆使用的數據。共享內存只能在SM內共享，且對於CUDA編程模型來說，即使線程塊被調度到了同一個SM內也無法互相訪問。
  • GPU的寄存器(registers) 和CPU不一樣，其空間非常巨大，以至於可以為每一個線程分配一塊獨立的寄存器空間。因此，不像CPU那樣切換進程時需要保存上下文，GPU只需要修改一下寄存器空間的指針即可繼續運行。所以巨大的寄存器空間，使得GPU上線程切換成為了一個幾乎無消耗的操作。 不過有一點需要注意，寄存器的空間不是無限大的。如果線程數過多，或一個線程使用的寄存器數量太多，多出來的數據會被保存到緩慢的本地內存上，影響程序速度，需要注意。

這裡想強調一下共享內存。

共享內存在物理上是一個個存儲體組成的。如果在訪問時沒有出現衝突，則可以實現高速的訪問。但如果出現了衝突(如對某一個存儲體的原子操作)，則不僅僅當前線程束會發生串列化，而且會導致其他線程束無法被調度(存疑，待考證)。

以上是從硬體的角度解讀了一下GPU的內存層級。從編程角度來看，CUDA的線程網格、線程塊、線程與各個內存見的關係如下圖：

線程層級與內存層級對應圖

  • 寄存器和本地內存綁定到了每個線程，其他線程無法訪問。
  • 同一個線程塊內的線程，可以訪問同一塊共享內存。注意，即使兩個線程塊被調度到了同一個SM上，他們的共享內存也是隔離開的，不能互相訪問。
  • 網格中的所有線程都可以自由讀寫全局內存。
  • 常量內存和紋理內存只能被CPU端修改，GPU內的線程只能讀取數據。

2.4 CPU/GPU混合編程

一種最簡單的CPU/GPU混合編程如下圖所示：

CPU/GPU混合編程

主機端(Host，即CPU)執行串列代碼，然後調用內核函數，讓設備端(Device，即GPU)執行並行代碼。如此交錯執行。

CPU和GPU的內存是獨立的。因此在運行內核函數前，主機端需要調用內存拷貝函數，將數據通過PCI-E匯流排拷貝到設備端。內核運行結束後，需要CPU再次調用內存拷貝函數，將數據拷回主機端內存。

另一種方式是使用統一編址，將設備端的內存和主機端內存編到一起。這樣主機就不需要顯式的調用函數將數據拷貝到設備端內存了。

除了CPU/GPU交錯執行代碼的方式外，還可以通過使用事件(event)和流(stream)等方式，讓CPU/GPU並行工作，提升整體的效率。

2.5 計算能力(Compute Capability)

所謂的計算能力(Compute Capability)，說白了就是GPU的版本號。有時也被稱作SM Version。

不同版本的GPU具有不同的特性，因此程序編寫也會有所差異。

計算能力為X.Y，其中主版本號X代表架構，各個架構如下表：

版本號與架構名稱

在CUDA的書籍文檔中，我們經常能看到"1.x"、"5.x"等這樣的字眼，代表第1代/第5代架構，也就是Tesla/Maxwell架構。

次版本號Y，代表在架構的基礎上，有一定改進，或者有一些新特性的引入。

最新的圖靈架構(Turing)，實際上計算能力是7.5，也就是說還是屬於Volta架構。

CUDA是軟體平台，其版本(CUDA7.5 CUDA8.0 CUDA10.0)與計算能力基本沒有關係。不過最新的CUDA一般會支持最新的架構。

從CUDA7.0起，Tesla架構不再被支持；從CUDA9.0起，Fermi架構不再被支持。

2.6 完整的例子

初學者的話，上面內容能看明白的都是勇士。這裡給一個例子，是我自己寫的BGR轉灰度圖的程序，希望能讓大家稍微明白一點CUDA程序如何寫。

/* main.cu */
#include <iostream>
#include <time.h>
#include "opencv2/highgui.hpp" //實際上在/usr/include下
#include "opencv2/opencv.hpp"
using namespace cv;
using namespace std;

//內核函數
__global__ void rgb2grayincuda(uchar3 * const d_in, unsigned char * const d_out,
uint imgheight, uint imgwidth)
{
const unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
const unsigned int idy = blockIdx.y * blockDim.y + threadIdx.y;

if (idx < imgwidth && idy < imgheight) //有的線程會跑到圖像外面去，不執行即可
{
uchar3 rgb = d_in[idy * imgwidth + idx];
d_out[idy * imgwidth + idx] = 0.299f * rgb.x + 0.587f * rgb.y + 0.114f * rgb.z;
}
}

//用於對比的CPU串列代碼
void rgb2grayincpu(unsigned char * const d_in, unsigned char * const d_out,
uint imgheight, uint imgwidth)
{
for(int i = 0; i < imgheight; i++)
{
for(int j = 0; j < imgwidth; j++)
{
d_out[i * imgwidth + j] = 0.299f * d_in[(i * imgwidth + j)*3]
+ 0.587f * d_in[(i * imgwidth + j)*3 + 1]
+ 0.114f * d_in[(i * imgwidth + j)*3 + 2];
}
}
}

int main(void)
{
Mat srcImage = imread("./test.jpg");
imshow("srcImage", srcImage);
waitKey(0);

const uint imgheight = srcImage.rows;
const uint imgwidth = srcImage.cols;

Mat grayImage(imgheight, imgwidth, CV_8UC1, Scalar(0));

uchar3 *d_in; //向量類型，3個uchar
unsigned char *d_out;

//首先分配GPU上的內存
cudaMalloc((void**)&d_in, imgheight*imgwidth*sizeof(uchar3));
cudaMalloc((void**)&d_out, imgheight*imgwidth*sizeof(unsigned char));

//將主機端數據拷貝到GPU上
cudaMemcpy(d_in, srcImage.data, imgheight*imgwidth*sizeof(uchar3), cudaMemcpyHostToDevice);

//每個線程處理一個像素
dim3 threadsPerBlock(32, 32);
dim3 blocksPerGrid((imgwidth + threadsPerBlock.x - 1) / threadsPerBlock.x,
(imgheight + threadsPerBlock.y - 1) / threadsPerBlock.y);

clock_t start, end;
start = clock();

//啟動內核
rgb2grayincuda<< <blocksPerGrid, threadsPerBlock>> >(d_in, d_out, imgheight, imgwidth);

//執行內核是一個非同步操作，因此需要同步以測量準確時間
cudaDeviceSynchronize();
end = clock();

printf("cuda exec time is %.8f
", (double)(end-start)/CLOCKS_PER_SEC);

//拷貝回來數據
cudaMemcpy(grayImage.data, d_out, imgheight*imgwidth*sizeof(unsigned char), cudaMemcpyDeviceToHost);

//釋放顯存
cudaFree(d_in);
cudaFree(d_out);

imshow("grayImage", grayImage);
waitKey(0);

return 0;

}

這裡我對比了CUDA、CPU、OPENCV三種實現方式的執行時間：

CUDA、CPU、OPENCV執行時間對比

速度：CUDA>OPENCV>CPU。其中OPENCV快於CPU的主要原因是OPENCV調用了並行運算指令，但慢於CUDA。

水水的介紹到此結束，剩下的我就要放飛自我了，看不懂別怪我

Chapter 3 編程介面

CUDA的編程介面由一系列C語言的擴展和運行庫(runtime library)組成。

C語言的擴展在第二章「編程模型」中有所提及，如內核函數、線程網格和線程塊等；

運行庫則是在CUDA Driver API的基礎上建立的。用戶可以直接在應用程序中跳過CUDA，直接調用CUDA Driver API，以便更底層地操作GPU，如操作GPU的上下文。不過對於大多數應用來說，使用CUDA提供的運行庫就足夠了。

本章講首先講解CUDA程序的編譯過程，之後會介紹CUDA運行庫，最後會介紹程序兼容性等問題。

3.1 使用NVCC編譯CUDA程序

CUDA程序使用NVCC編譯器。

NVCC提供了簡單方便的介面，能夠很好的同時處理主機端和設備端代碼。這裡將簡要介紹NVCC編譯CUDA程序的流程，更多信息請參考nvcc user manual。

3.1.1 編譯流程

3.1.1.1 離線編譯

NVCC進行離線編譯的操作流程是：

分離CUDA程序中的主機端代碼(host code)和設備端代碼(device code) 將設備端代碼編譯成一種虛擬彙編文件(名為PTX)，再接著編譯成二進位代碼(名為cubin) 將主機端代碼中含有"<<<>>>"的代碼(即內核調用)替換為CUDA運行庫中的函數調用代碼 之後NVCC會藉助其他編譯器(如gcc)將主機端代碼編譯出來 * 主機端代碼和設備端代碼被編譯好後，nvcc會將兩段代碼鏈接起來

3.1.1.2 在線編譯(JIT Compilation)

PTX是一個虛擬彙編文件。其形式雖然很像彙編，但裡面的每一條指令實際上是一個虛擬的指令，與機器碼無法對應。需要編譯器或設備驅動程序將其翻譯成對應平台的彙編/機器碼才能運行。

如果在編譯過程中，NVCC不將設備端代碼編譯為cubin文件，即二進位代碼，而是停在PTX代碼上。設備驅動(device driver)會負責在運行時，使用PTX代碼生成二進位代碼。這個過程被稱作在線編譯(JIT Compilation, Just-In-Time Compilation)。

在線編譯必然會使得程序啟動的時間延長，不過設備驅動程序會自動緩存編譯出來的二進位代碼(也被稱作compute cache)。

在線編譯一方面的優勢在於兼容性。另一方面的優勢在於，當設備驅動程序有關編譯的部分得到優化時，同樣的PTX編出來的cubin文件同樣會得到優化。也就是說，一段祖傳的PTX代碼，很有可能因為驅動程序不斷的優化，而躺著得到了優化。而如果直接離線編譯得到了cubin文件的話，則無法享受到這一優化。

3.1.2 二進位代碼的兼容性

二進位代碼cubin是受到GPU計算能力的限制的。在編譯時，需要使用-code來指定將代碼編譯到哪個計算能力平台上，如-code=sm_35代表生成的cubin代碼是運行在計算能力為3.5的平台上的。

二進位代碼若要兼容，首先架構得一致。不同架構上的二進位代碼不能互相兼容，如在Maxwell架構上編譯出來的代碼，不能在其他架構上運行。

其次，若執行平台的次版本號版本比編譯時指定的的次版本號高，則可以運行。例如如果在編譯時指定-code=sm_35，則在計算能力3.7的平台上也可以運行。反之則不可以。

另外需要說明的是，上述二進位代碼的兼容性原則只限於桌面款顯卡。

3.1.3 PTX代碼的兼容性

PTX代碼的兼容性遠強於二進位代碼。只要不涉及到不同架構上的特性差異，PTX可以在任何架構上運行。

不過PTX代碼在兩種情況下其兼容性會受限：

1. 若PTX代碼使用了較高級別架構的特有特性，則無法在較低架構上運行。例如若PTX代碼用到了計算能力3.0以上才能使用的Warp Shuffle特性，則無法在2.x或1.x平台上運行。 2. 若PTX在較低架構上生成，則雖然能夠在所有更高級別的架構上運行，但無法充分利用這些架構的硬體特性，造成性能無法最大化的問題。

在編譯時，可以通過-arch來指定生成的PTX代碼的版本，如-arch=compute_30。

3.1.4 應用程序兼容性

為了保證應用程序的兼容性，最好是將代碼編譯成PTX代碼，然後依靠各個計算能力的驅動程序在線編譯成對應平台的二進位代碼cubin。

除了使用-arch和-code來分別指定C->PTX和PTX->cubin的計算能力外，還可以用-gencode關鍵字來操作，如下例：

nvcc x.cu
-gencode arch=compute_35,code=sm_35
-gencode arch=compute_50,code=sm_50
-gencode arch=compute_60,code=compute_60,sm_60

使用上述編譯指令後，會生成3.5/5.0/6.0的cubin文件，以及6.0的PTX代碼。具體內容請參考nvcc user manual。

對於主機端代碼，會自動編譯，並在運行時決定調用哪一個版本的執行。對於上例，主機端代碼會編譯為：3.5/5.0/6.0的二進位文件，以及7.0的PTX文件。

另外，在程序中可以使用__CUDA_ARCH__宏來指定計算能力(只能用於修飾設備端代碼)。計算能力3.5在程序中對應的__CUDA_ARCH__為350。

有一點需要注意的是，7.0以前，都是以線程束為單位在調度，線程束內指令永遠是同步的。而Volta架構(計算能力7.x)引入了Independent Thread Scheduling，破壞了線程束內的隱式同步。因此，在不針對同步問題而修改代碼的前提下，要想使用7.x設備，可以指定-arch=compute_60 -code=sm_70，即將PTX編到Pascal架構下以禁用Independent Thread Scheduling特性。

另外，版本相關編譯指令有縮寫的情況，具體看手冊。

3.1.5 C/C++兼容性

對於主機端代碼，nvcc支持C++的全部特性；而對於設備端代碼，只支持C++的部分特性。具體查閱手冊。

3.1.6 32/64位兼容性

當且僅當主機端代碼按照64位編譯時，設備端代碼才能編譯為64位。當主機端代碼編譯為32位時，設備端代碼只能編譯成32位。即設備端代碼的位數和主機端永遠保持一致。

具體編譯成32/64位的哪一種，取決於nvcc本身的版本。32位nvcc會自動編出32位的代碼，不過可以使用-m64來編出64位代碼。對於64位編譯器亦然。

3.2 CUDA C 運行庫

運行庫實際上在cudart庫內，可以使靜態鏈接庫cudart.lib/libcudart.a，或者動態鏈接庫cudart.dll/cudart.so。

所有程序的入口都是cuda。

3.2.1 初始化

CUDA運行庫沒有顯式的初始化函數，在調用第一個函數時會自動初始化(設備和版本管理函數不行)。初始化時，會產生一個全局可見的設備上下文(device context)。

當主機端代碼調用了cudaDeviceReset()函數，則會銷毀掉這個上下文。注意，銷毀的上下文是主機端正在操縱的設備。如要更換，需要使用cudaSetDevice()來進行切換。

3.2.2 設備內存

CUDA運行庫提供了函數以分配/釋放設備端的內存，以及與主機端內存傳輸數據。

這裡的設備內存，指的是全局內存+常量內存+紋理內存。

設備內存有兩種分配模式：線性存儲(linear memory)、CUDA arrays。 其中CUDA arrays與紋理內存有關，本導讀略去不談。

線性內存是我們常用的內存方式，在GPU上用40位的地址線定址。線性內存可以用cudaMalloc()分配，用cudaFree()釋放，用cudaMemcpy()複製數據，用cudaMemset()賦值。

對於2D或3D數組，可以使用cudaMallocPitch()和cudaMalloc3D()來分配內存。這兩個函數會自動padding，以滿足內存對齊的要求，提高內存讀寫效率。內存對齊的問題，會在第五章里詳細闡述。

另外，如果要在設備內存中定義全局變數，則需要使用使用__constant__或__device__來修飾，並使用cudaMemcpyToSymbol()和cudaMemcpyFromSymbol()來讀寫。如下例：

__constant__ float constData[256];
float data[256];
cudaMemcpyToSymbol(constData, data, sizeof(data));
cudaMemcpyFromSymbol(data, constData, sizeof(data));

__device__ float devData;
float value = 3.14f;
cudaMemcpyToSymbol(devData, &value, sizeof(float));

__device__ float* devPointer;
float* ptr;
cudaMalloc(&ptr, 256 * sizeof(float));
cudaMemcpyToSymbol(devPointer, &ptr, sizeof(ptr));

實際上，當使用__constant__關鍵字時，是申請了一塊常量內存；而使用__device__時，是普通的全局內存。因此__device__申請的內存需要申請，而__constant__不用。不管是全局內存，還是常量內存，都是只讀的，需要用帶有Symbol的函數拷貝。

3.2.3 共享內存

不管是全局變數還是局部變數，都需要使用__shared__來修飾。不過需要注意的是，即使定義為全局變數，共享內存依舊只能被同一線程塊內的線程可見。

舉個例子，對於如下代碼，雖然是定義了一個全局的共享內存hist_shared，但實際上，在每一個線程塊被調度到SM上時，都會在SM的共享內存區開一塊內存。因此，每一個線程塊都有一個hist_shared，且之間無法互相訪問。

__shared__ unsigned int hist_shared[256]; //共享內存僅在線程塊內共享

__global__ void getGrayHistincuda_usesharemem(unsigned char * const grayData,
unsigned int * const hist,
uint imgheight,
uint imgwidth) //使用共享內存加速
{
const unsigned int idx = blockDim.x * blockIdx.x + threadIdx.x;
const unsigned int idy = blockDim.y * blockIdx.y + threadIdx.y;
const unsigned char inner_idx = threadIdx.y * blockDim.x + threadIdx.x;

hist_shared[inner_idx%256] = 0; //清空數據，由於每個塊的inner_idx可以超過256，所以這樣可以保證hist_shared被全部清零

__syncthreads(); //等待其他線程完成

if(idx < imgwidth && idy < imgheight)
{
const unsigned long pid = imgwidth * idy + idx;
const unsigned char value = grayData[pid];
atomicAdd(&(hist_shared[value]), 1);
}

__syncthreads();

if(threadIdx.y < 8) //每個線程塊將自己共享內存中的值合併到全局內存中去
{
atomicAdd(&(hist[inner_idx]), hist_shared[inner_idx]);
}

}

當然，共享內存的聲明放在內核函數裡面也是可以的，效果一致。

使用共享內存，可以獲得等同於L1 cache的訪存速度，其速度遠快於全局內存。

但是注意，並不是什麼時候都可以使用共享內存來獲取加速的。例如內核函數計算出來結果後，如果這個結果只需要傳輸回主機端，而不需要再次被用到時，直接寫回全局內存會比較快。如果先寫回共享內存，再寫回全局內存，反而會比較緩慢。

一般來講，當需要頻繁讀寫，或是有原子操作時，使用共享內存替代全局內存，會取得比較大的增益。

強調一下，共享內存只能為線程塊內的線程共享。如果需要整個線程網格中線程都能訪問，則需要全局內存或常量內存。

另外，共享內存是一個稀缺資源。有些架構可以通過配置，分配L1 cache和共享內存的比例。

3.2.4 鎖頁內存(Page-Locked Host Memory/Pinned Memory)

鎖頁內存指的是主機端上不會被換出到虛擬內存(位於硬碟)上的內存。

鎖頁內存的分配與釋放：

在CUDA程序中，使用cudaHostAlloc()，可以分配鎖頁內存，使用cudaFreeHost()來釋放鎖頁內存 或者使用cudaHostRegister()來將malloc()分配的內存指定為鎖頁內存

NVIDIA官方給出的鎖頁內存相對於普通的內存的的好處是：

使用鎖頁內存後，鎖頁內存與設備內存之間的數據傳輸，可以使用流的方式，和內核函數執行並行。 使用鎖頁內存後，可以將鎖頁內存映射到設備內存上。 對於使用前端匯流排*的系統，使用鎖頁內存可以提升主機端到設備端傳輸的帶寬； 如果將鎖頁內存指定為合併寫(write_combining)，則可以進一步提高帶寬。

另一本書對於鎖頁內存之所以快的解釋是：

如果主機端將數據放在鎖頁內存，則可以使用PCI-E的DMA與設備內存進行數據傳輸，而不需要CPU來搬運數據。 這也是為何使用了鎖頁內存後，可以使用流和內存映射，來讓CPU程序、數據傳輸和內核執行並行。 如果主機端將數據放在普通內存，則CUDA會先申請一塊鎖頁內存，然後將數據拷貝到鎖頁內存，再做後面的操作。 拷貝的過程浪費了一定時間。

注意，鎖頁內存在 non I/O coherent Tegra 設備上不支持

3.2.4.1 Portable Memory

NVIDIA官方文檔表示：上述所說的鎖頁內存的優點，只有在使用cudaHostAlloc()時，傳入cudaHostAllocPortable flag，或者在使用cudaHostRegister()時傳入cudaHostRegisterPortable flag，才能體現。否則鎖頁內存並不會有上述優點。

《GPU編程指南》一書中是這麼描述的：如果傳入了cudaHostAllocPortable flag，則鎖頁內存在所有的CUDA上下文中變成鎖頁的和可見的。如果需要在CUDA上下文之間或者主機處理器的線程之間傳遞指針，則必須使用這個標誌。

(好吧，從編程指南一書中確實沒出來用Portable的必要性，不是很明白)

3.2.4.2 合併寫內存(Write-Combining Memory)

鎖頁內存默認是使用緩存的。如果將flag cudaHostAllocWriteCombined 傳入到 cudaHostAlloc()，則可以將這塊鎖頁內存指定為合併寫內存。

合併寫內存不再使用主機端的L1&L2 cache，使得更多的cache可以供其他任務使用。

另外，對於通過PCI-E傳輸數據的情景，使用合併寫內存不會被snooped (是不是指的是不會被緩存管？不理解這個snooped什麼意思)，可以提升40%的傳輸性能。

此外需要注意的是，由於合併寫內存不使用緩存，因此讀入CPU核的操作會非常的慢。因此合併寫內存最好只用作向GPU傳數據的內存，而不是傳回數據的內存。

3.2.4.3 內存映射(Mapped Memory)

CUDA中的內存映射，指的是將CPU端的鎖頁內存，映射到GPU端。

通過向cudaHostAlloc()傳入cudaHostAllocMapped flag，或向cudaHostRegister()傳入cudaHostAllocMapped flag，來將一塊內存指定為向GPU映射的內存。

映射的內存有兩個地址，一個是CPU端訪問的地址，一個是GPU端訪問的地址。

CPU端的地址在調用malloc()或cudaHostAlloc()時就已經返回； GPU端的地址使用cudaHostGetDevicePointer()函數來獲取。

使用內存映射有以下好處：

使用內存映射，可以讓CPU/GPU之間的數據傳輸隱式執行，而不需要顯示的分配GPU內存並傳輸數據。 當設備端執行內核函數需要某一塊數據時，如果數據實際上在CPU端，會給出一個PCI-E傳輸請求(比全局內存還慢)，從主機端內存獲取數據。此時給出數據請求的線程會被換出，直到數據就位後再被換入。因此如果使用內存映射，需要使用足夠多的線程來隱藏PCI-E的傳輸延遲。 內存映射可以替代流，實現數據傳輸和內核執行的並行 有一點不是很確定：內存映射是否會在GPU端緩存數據；據我的記憶是不會緩存的，因此多次請求同一塊數據的話，會啟動多個PCI-E傳輸，效率很低 不清楚內存映射在GPU更新後，CPU端數據會何時更新。是在CPU訪問這些數據時？還是自動更新？

使用內存映射必須要注意的幾點：

由於映射的內存會被CPU和GPU兩方共享，因此程序需要注意數據同步問題 如果要使用內存映射，必須在其他CUDA函數執行前，執行cudaSetDeviceFlags()並傳入cudaDeviceMapHost，來使能設備的內存映射功能。否則cudaHostGetDevicePointer()函數會返回error。 如果設備本身不支持內存映射，則使用cudaHostGetDevicePointer()一定會返回error。可以通過查看設備的canMapHostMemory信息來確認。 如果使用原子操作(atomicXXX)，需要注意，主機端和設備端的同時操作是不原子的。

3.2.5 非同步並行執行

CUDA允許以下操作互相併行：

主機端計算 設備端計算(內核執行) 主機端to設備端傳數據 設備端to主機端傳數據 設備端內部傳數據 設備間傳數據(可通過PCI-E直接傳輸，不需要先傳到主機端再轉發，不過這一操作跟使用的操作系統有關)

3.2.5.1 主機端/設備端並行

設備端的如下操作，可以跟主機端並行：

內核啟動與執行(可以通過將CUDA_LAUNCH_BLOCKING設為1，來disable內核執行並行，debug使用) 設備端內部傳輸數據 64KB及以下的 host-to-device數據傳輸 使用流(帶有Async前綴的內存傳輸函數)或內存映射傳輸數據（不再受64KB的限制） * 設備端memset函數(cudaMemset())

其中第3、4條說明，在使用cudaMemcpy()時，如果數據小於等於64KB，其實傳輸相對於CPU是非同步的。 如果數據多於64KB，則CPU會阻塞到數據傳輸完成。 這時使用帶Async的內存傳輸函數，會釋放CPU資源。

使用Async傳輸函數，不僅可以和CPU並行，而且可以和內核執行並行。

需要注意的是，如果沒有使用鎖頁內存，即使使用了Async函數，內存傳輸也不是並行的(和CPU？還是GPU？)。

3.2.5.2 內核並行執行

計算能力2.x及以上的設備，支持多個內核函數同時執行。(可以通過檢查concurrentKernels來確定)

執行多個內核函數，需要主機端不同的線程啟動。如果一個線程依次啟動多個內核，則這些內核會串列執行。同一線程的內核函數返回時會觸發隱式的同步。

另外，多個內核函數必須位於同一個CUDA上下文(CUDA context)上。不同CUDA上下文上的內核不能並行。這意味著，啟動多個內核的多個線程必須使用相同的CUDA上下文。(如何傳遞CUDA上下文？)

3.2.5.3 數據傳輸和內核執行並行(需要使用鎖頁內存)

一些設備支持數據傳輸(主機端/設備端、設備端/設備端)和內核執行並行，可通過檢查asyncEngineCount來確認。

一些設備支持設備端內部數據傳輸和內核執行/數據傳輸並行，可通過檢查concurrentKernels來確認。

這一特性需要使用鎖頁內存。

3.2.5.4 數據並行傳輸(需要使用鎖頁內存)

計算能力2.x及以上的設備，支持數據傳入和傳出並行。

必須使用鎖頁內存。

3.2.5.5 流(streams)

在CUDA中，流(streams)指的是在GPU上一連串執行的命令。

不同的線程，可以向同一個流填入任務。

同一個流內的任務會按順序執行；同一設備上不同的流有可能並行，其執行順序不會有保證。

3.2.5.5.1 流的創建和銷毀

下述代碼是一個流的創建和銷毀的例子。該程序創建了兩個流，分配了兩個鎖頁內存傳輸數據，依次啟動了兩個內核，最後銷毀了這兩個流。

cudaStream_t stream[2];
for (int i = 0; i < 2; ++i)
cudaStreamCreate(&stream[i]);
float* hostPtr;
cudaMallocHost(&hostPtr, 2 * size);

for (int i = 0; i < 2; ++i) {
cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
size, cudaMemcpyHostToDevice, stream[i]);
MyKernel <<<100, 512, 0, stream[i]>>>
(outputDevPtr + i * size, inputDevPtr + i * size, size);
cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,
size, cudaMemcpyDeviceToHost, stream[i]);
}

for (int i = 0; i < 2; ++i)
cudaStreamDestroy(stream[i]);

從上例中可以看到，流的創建需要定義cudaStream_t結構，並調用cudaStreamCreate()來初始化。

流的銷毀需要調用cudaStreamDestroy()來實現。

當向流中添加內核函數任務時，<<<...>>>不再是<<<blocksPerGrid, threadsPerBlock>>>，而是<<<blocksPerGrid, threadsPerBlock, dynamic_shared_memory, stream>>>。

其中dynamic_shared_memory指的是動態共享內存的大小(回去翻書)； stream就是cudaStream_t結構。

當設備還在執行流中的任務，而用戶調用cudaStreamDestroy()函數時，函數會立刻執行(不會阻塞)。之後，當流中的任務完成後，與流相關的資源會自動釋放。

另外需要注意的是，上例中主機端線程、數據拷貝和內核執行完全非同步，因此在"拷貝回主機端"這一操作完成之前，主機端的內存數據是不正確的。必須在數據返回的一步做同步操作，方能保證數據是正確的。

(需要了解一下流內部是如何實現的，為什麼內核執行和內存拷貝能夠非同步且重疊？什麼樣的操作又不能重疊？)

3.2.5.5.2 默認流(Default Stream)

在調用內核函數時，不指定流或者將流指定為0，則代表使用了默認流(default stream)。

如果在編譯時使用了--default-stream per-thread，或是在include任何cuda頭文件前#define CUDA_API_PER_THREAD_DEFAULT_STREAM，則主機端的每一個線程都有自己專屬的默認流。

而如果在編譯時未指定相關flag，或指定--default-stream legacy，則默認流是一個特殊的流，稱作NULL stream。主機端的所有線程會共享這個NULL stream。NULL stream是一個同步流，所有命令會產生隱式的同步。

3.2.5.5.3 顯式同步(Explicit Synchronization)

可以使用如下函數進行顯式同步：

cudaDeviceSynchronize()：直到所有線程向設備端的所有流的所有已送入指令完成，才會退出阻塞。 cudaStreamSynchronize()：直到指定流的之前所有已送入指令完成，才會退出阻塞。此函數可以用作同步指定流，而其他流可以不受干擾地繼續運行。 cudaStreamWaitEvent()：需要stream和event作為輸入參數。在調用該函數之後的命令，需要等待該函數等待的事件(Event)發生後，才能執行。如果stream指定為0，則對於向所有stream加入的命令來說，只要加在了該函數之後，都會阻塞直到等待的時間發生方可執行。 (不知道我理解的對不對：如果是Event->內核1->WaitEvent->內核2，則內核1不用等到Event發生就可以執行，而內核2必須等到Event發生才能執行。還是說內核1其實只有等待Event發生後才會執行？) (如果多個線程向同一個流壓入了任務，然後線程0調用了cudaStreamWaitEvent()，則線程1會不會被阻塞？線程1壓入的任務會不會被阻塞？) cudaStreamQuery()：查詢流內所有壓入的指令(preceding commands)是否全部完成。

注意，同步函數慎用，因為有可能會產生速度的下降。

3.2.5.5.4 隱式同步(Implicit Synchronization)

一般來講，不同流內的命令可以並行。但是當任何一個流執行如下的命令時，情況例外，不能並行：

鎖頁內存的分配 設備端內存分配 設備端內存設置(memset) 設備內部拷貝 NULL stream內的命令 L1 cache/共享內存空間的重新分配

3.2.5.5.5 操作重疊(Overlapping Behavior)

操作的重疊程度，一方面取決於各個操作的順序，另一方面取決於設備支持重疊的程度(是否支持內核執行並行/數據傳輸與內核執行並行/數據傳輸並行)

3.2.5.5.6 回調函數(Callbacks)

可以使用cudaStreamAddCallback()函數，向流中添加callback。該callback會在流中之前所有的任務完成後被調用。如果stream參數設為0，則代表之前的所有stream的任務執行完後就調用該callback。

回調函數和cudaStreamWaitEvent()一樣，對於在加在callback之後的指令，必須等待callback執行完成後，才會繼續執行。

下例是一個使用回調的例子。該例中，兩個stream將數據拷回主機端後，會調用回調函數。

void CUDART_CB MyCallback(cudaStream_t stream, cudaError_t status, void *data){
printf("Inside callback %d
", (size_t)data);
}
...
for (size_t i = 0; i < 2; ++i) {
cudaMemcpyAsync(devPtrIn[i], hostPtr[i], size, cudaMemcpyHostToDevice, stream[i]);
MyKernel<<<100, 512, 0, stream[i]>>>(devPtrOut[i], devPtrIn[i], size);
cudaMemcpyAsync(hostPtr[i], devPtrOut[i], size, cudaMemcpyDeviceToHost, stream[i]);
cudaStreamAddCallback(stream[i], MyCallback, (void*)i, 0);
}

回調函數中不能直接或間接的執行CUDA函數，否則會因為等待自己完成而造成死鎖。 (原因尚不太明白)

3.2.5.5.7 流的優先順序(Stream Priorities)

可以通過cudaStreamCreateWithPriority()來在創建流時指定流的優先順序。可以指定的優先順序可由cudaDeviceGetStreamPriorityRange()來獲得。

運行時，高優先順序stream中的線程塊不能打斷正在執行的低優先順序stream的線程塊(即不是搶佔式的)。但是當低優先順序stream的線程塊退出SM時，高優先順序stream中的線程塊會被優先調度進SM。

3.2.5.6 事件(Event)

事件(Event)可以被壓入流中以監視流的運行情況，或者用於精確計時。

如果向stream 0壓入事件，則當壓入事件前向所有流壓入的任務完成後，事件才被觸發。

3.2.5.6.1 事件的創建和銷毀

cudaEvent_t start, stop; //創建
cudaEventCreate(&start);
cudaEventCreate(&stop);
...
cudaEventDestroy(start); //銷毀
cudaEventDestroy(stop);

3.2.5.6.2 計算時間

下例是一個使用Event計算時間的例子：

cudaEventRecord(start, 0); //記錄事件(將事件壓入流)，流0則代表所有流完成任務後事件才會被觸發
for (int i = 0; i < 2; ++i) {
cudaMemcpyAsync(inputDev + i * size, inputHost + i * size, size, cudaMemcpyHostToDevice, stream[i]);
MyKernel<<<100, 512, 0, stream[i]>>>(outputDev + i * size, inputDev + i * size, size);
cudaMemcpyAsync(outputHost + i * size, outputDev + i * size, size, cudaMemcpyDeviceToHost, stream[i]);
}
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop); //獲取兩個事件發生的時間差(ms)

3.2.6 多設備系統(Multi-Device System)

3.2.6.1 設備枚舉(Device Enumeration)

下例是如何枚舉設備，並獲取設備信息的例子：

int deviceCount;
cudaGetDeviceCount(&deviceCount); //獲取設備數量
int device;
for (device = 0; device < deviceCount; ++device) {
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, device);
printf("Device %d has compute capability %d.%d.
", device, deviceProp.major, deviceProp.minor);
}

3.2.6.2 設備選擇(Device Selection)

使用cudaSetDevice()選擇設備，當不選擇時，默認使用設備0。

注意，所有的內存分配、內核函數啟動、流和事件的創建等，都是針對當前選擇的設備的。

下例是一個設備選擇的例子：

size_t size = 1024 * sizeof(float);
cudaSetDevice(0); // Set device 0 as current
float* p0;
cudaMalloc(&p0, size); // Allocate memory on device 0
MyKernel<<<1000, 128>>>(p0); // Launch kernel on device 0
cudaSetDevice(1); // Set device 1 as current
float* p1;
cudaMalloc(&p1, size); // Allocate memory on device 1
MyKernel<<<1000, 128>>>(p1); // Launch kernel on device 1

3.2.6.3 (多設備下)流和事件的執行情況

下面將討論，如果對一個不屬於當前設備的流或事件進行操作，哪些操作會成功，哪些操作會失敗：

  • 內核啟動(will fail)：如果將內核壓入不屬於當前設備的流中，則內核會啟動失敗。也就是說，如果要向一個流中壓入內核，必須先切換到流所在的設備：

cudaSetDevice(0); // Set device 0 as current
cudaStream_t s0;
cudaStreamCreate(&s0); // Create stream s0 on device 0
MyKernel<<<100, 64, 0, s0>>>(); // Launch kernel on device 0 in s0
cudaSetDevice(1); // Set device 1 as current
cudaStream_t s1;
cudaStreamCreate(&s1); // Create stream s1 on device 1
MyKernel<<<100, 64, 0, s1>>>(); // Launch kernel on device 1 in s1

// This kernel launch will fail:
MyKernel<<<100, 64, 0, s0>>>(); // Launch kernel on device 1 in s0

  • 內存拷貝(will success)：如果對一個不屬於當前設備的流進行內存拷貝工作，內存拷貝會成功。
  • cudaEventRecord()(will fail)：必須現將設備上下文切換過去，在向流壓入事件。
  • cudaEventElapsedTime()(will fail)：計算時間差前，必須先切換設備。
  • cudaEventSynchronize() and cudaEventQuery()(will success)：即使處於不同的設備，事件同步和事件查詢依然有效。
  • cudaStreamWaitEvent()(will success)：比較特殊，即使函數輸入的流和事件不在同一個設備上，也能成功執行。也就是說，可以讓流等待另一個設備上(當然當前設備也可以)的事件。這個函數可以用作多個設備間的同步。

另外需要注意，每個設備都有自己的默認流。因此在沒有指定流的情況下，向不同設備分派的任務，實際上是壓入了各個設備的默認流，他們之間是並行執行的。

3.2.6.4 (設備間)對等內存訪問(Peer-to-Peer Memory Access)

計算能力2.0及以上的設備支持設備間對等內存訪問，這意味著兩個GPU之間的傳輸和訪問可以不經過主機端中轉，速度會有提升。查詢cudaDeviceCanAccessPeer()可以得知設備是否支持這一特性。(官方文檔說還需要一個條件：64位程序，存疑)

需要使用cudaDeviceEnablePeerAccess()來使能這一特性。

對等設備的的地址是統一編址的，可用同一個指針訪問，如下例：

cudaSetDevice(0); // Set device 0 as current
float* p0;
size_t size = 1024 * sizeof(float);
cudaMalloc(&p0, size); // Allocate memory on device 0
MyKernel<<<1000, 128>>>(p0); // Launch kernel on device 0
cudaSetDevice(1); // Set device 1 as current
cudaDeviceEnablePeerAccess(0, 0); // Enable peer-to-peer access with device 0

// Launch kernel on device 1
// This kernel launch can access memory on device 0 at address p0
MyKernel<<<1000, 128>>>(p0);

3.2.6.5 (設備間)對等內存拷貝(Peer-to-Peer Memory Copy)

對等設備的地址是統一編址的，可以使用cudaMemcpyPeer()、cudaMemcpyPeerAsync()、cudaMemcpy3DPeer、cudaMemcpy3DPeerAsync()來進行直接拷貝。無需先拷貝會主機端內存，再轉到另一塊卡上。如下例：

cudaSetDevice(0); // Set device 0 as current
float* p0;
size_t size = 1024 * sizeof(float);
cudaMalloc(&p0, size); // Allocate memory on device 0
cudaSetDevice(1);
float* p1;
cudaMalloc(&p1, size); // Allocate memory on device 1
cudaSetDevice(0); // Set Device 0 as Current
MyKernel<<<1000, 128>>>(p0); // Launch Kernel on Device 0
cudaSetDevice(1); // Set Device 1 as Current
cudaMemcpyPeer(p1, 1, p0, 0, size); // Copy p0 to p1
MyKernel<<<1000, 128>>>(p1); // Launch Kernel on Device 1

關於設備間的對等拷貝，如果使用的是NULL stream，則有如下性質：

如果拷貝的雙方中的任何一方，在設備拷貝前有任務未完成，則拷貝會被阻塞，直至任務完成。 只有拷貝結束後，兩者的後續任務才能繼續執行。

(使用的如果不是NULL Stream，又會怎樣呢？)

3.2.7 統一虛擬地址空間(Unified Virtual Address Space)

當程序是64位程序時，所有主機端內存，以及計算能力≥2.0的設備的內存是統一編址的。所有通過CUDA API分配的主機內存和設備內存，都在統一編址的範圍內，有自己的虛擬地址。因此：

  • 可以通過cudaPointerGetAttributes()，來確定指針所指的內存處在主機端還是設備端。
  • 進行拷貝時，可以將cudaMemcpy***()中的cudaMemcpyKind參數設置為cudaMemcpyDefault，去讓函數根據指針所處的位置自行判斷應該是從哪裡拷到哪裡。
  • 使用cudaHostAlloc()分配的鎖頁內存，自動是Portable的，所有支持統一虛擬編址的設備均可訪問。cudaHostAlloc()返回的指針，無需通過cudaHostGetDevicePointer()，就可以直接被設備端使用。

可以通過查詢unifiedAddressing來查看設備是否支持統一虛擬編址。

3.2.8 進程間通訊(Interprocess Communication)

線程間通訊，可以很方便的通過共享的變數來實現。然而進程間通訊不行。

為了在進程間共享設備端內存的指針或者事件，必須使用IPC(Inter Process Communication) API。IPC API只支持64位程序，並且要求設備計算能力≥2.0。

通過IPC中的cudaIpcGetMemHandle()，可以得到設備內存指針的IPC句柄。該句柄可以通過標準的IPC機制(interprocess shared memory or files)傳遞到另一個進程，再使用cudaIpcOpenMemHandle()解碼得到該進程可以使用的設備內存指針。

事件的共享也是如此。

3.2.9 錯誤檢查(Error Checking)

所有的runtime function都會返回一個error code，可通過檢查error code判斷是否出錯。

但是對於非同步函數，由於在執行前就會返回，因此返回的error code僅僅代表函數啟動時的錯誤(如參數校驗)；非同步函數不會返回運行時出現的錯誤。如果運行時出了錯，會被後面的某個函數捕獲並返回。

檢查非同步函數是否出錯的唯一方式，就是在非同步函數啟動後，進行同步。 如在非同步函數後，調用cudaDeviceSynchronize()，則非同步函數的錯誤會被cudaDeviceSynchronize()捕獲到。

事實上，除了runtime function會返回error code之外，每一個主機端線程都會有一個初始化為cudaSuccess的變數，用於指示錯誤。一旦發生了錯誤，該變數也會被設置為相應的error code。

該變數不會被直接調用，但可以被cudaPeekAtLastError()和cudaGetLastError()訪問到。不同的是，cudaGetLastError()在返回這一變數的同時，會把它重置為cudaSuccess。

內核函數不會返回值，因此只能通過cudaPeekAtLastError()或cudaGetLastError()來知悉調用內核是否有錯誤。

當然，為了排除錯誤出現在調用內核之前就有錯誤，可以先檢驗之前的錯誤變數是否為cudaSuccess。

另外需要注意的是，cudaStreamQuery()和cudaEventQuery()這類函數，有可能會返回cudaErrorNotReady。但這不被認為是錯誤，因此不會被cudaPeekAtLastError()和cudaGetLastError()捕獲到。

3.2.10 調用棧(Call Stack)

對於計算能力≥2.0的設備，可以通過cudaDeviceGetLimit()/cudaDeviceSetLimit()來查詢/設置調用棧的大小。

3.2.11 紋理內存

略

3.2.12 Graphics Interoperability

略

3.3 版本和兼容性(Versioning and Compatibility)

有兩個版本需要注意：計算能力，以及CUDA driver API的版本。其中計算能力及其兼容性在前面已有闡述。

CUDA driver API的版本定義在驅動的頭文件中的CUDA_VERSION宏內。可以在程序中將該宏調出，以檢查程序是否可以在目標設備上運行。

CUDA driver API不是向前兼容的。也就是說，針對新版本的CUDA driver API編譯的程序、插件、庫，並不能在舊版本的驅動上運行。

CUDA Driver API 兼容性

關於CUDA driver API，有幾點需要注意：

  • 由於一個系統只能裝一個版本的驅動。因此驅動版本要足夠高(至少程序所需的版本)，否則程序跑不起來。
  • 默認情況下，nvcc編譯程序時，庫和插件是靜態編譯的。靜態編譯不要求庫和插件的驅動版本和CUDA運行庫保持一致。但是動態鏈接則要求版本一致。

3.4 計算模式(Compute Mode)

NVIDIA的設備可以設置三種計算模式：

默認模式(Default Compute Mode)：多個主機端線程可以同時使用一個設備(通過調用cudaSetDevice()) 專屬進程模式(Exclusive-Process Compute Mode)：對於一個設備，只能由一個進程創建設備上下文。一旦創建成功後，該進程的所有線程都可以使用該設備，而其他進程則不行。 * 禁止模式(Prohibited Compute Mode)：無法對設備建立CUDA上下文。

正常情況下，如果程序沒有調用cudaSetDevice()，則會默認使用0號設備。但是如果0號設備被置成禁止模式，亦或是被其他進程所專屬，則會在其他設備上創建上下文並使用。 可以向cudaSetValidDevices()函數輸入一個設備列表，函數會在第一個可以使用的設備上創建上下文。

Pascal及以上架構，支持指令級的優先順序調度。不再是以線程塊為SM的最小調度單位，而是以指令為最小調度單位，且具有優先順序。 這意味著具有冗長kernel的線程塊不再會佔據太多的計算資源，或是發生timeout。但是這也有缺點：當多個進程創建了上下文時，以往基於線程塊的調度不會造成太多的上下文切換，但現在的指令級調度則會造成很多的上下文切換，降低效率。(注意跟GPU內線程的上下文切換不同，GPU內線程上下文切換幾乎不浪費時間，直接換一個指針就好)。因此最好設置為Exclusive-process，這樣只有一個進程會使用設備。(如果線程很多的話，效果不是一樣嗎？)

(講道理好好看看上下文是什麼鬼)

(這段需要再看一下Pascal架構的說明，看看指令級的調度是如何實現的，寄存器等又是如何分配的)

設備處於哪種計算模式，可通過檢查computeMode來查看。

3.5 模式切換(Mode Switches)

GPU會將一些內存專門分配給primary surface，用於刷新顯示設備。 一旦用戶將顯示模式切換，如增加解析度或增加彩色的位數，會造成primary surface所需的內存變多。此時系統會把原來分配給CUDA運算的內存，調撥給primary surface，從而造成CUDA runtime產生錯誤，並返回invalid context error。

(言外之意是說，跑cuda的時候不要切解析度？)

3.6 Tesla Compute Cluster Mode for Windows

略

Chapter 4 硬體架構

4.0 補充內容

這份官方文檔講的硬體內容太少了，從另一本書里補一點過來，可能內容有點老舊，見諒。

4.0.0 硬體結構

(在一本比較老舊的書上找到的結構圖)

下圖是一款比較老舊的顯卡(G80/GT200)的硬體結構圖：

NVIDIA顯卡硬體結構

從圖中可以看到，GPU由內存和一系列流處理器簇(Streaming Multiprocessors, SM)組成，不同GPU的具有不同的內存大小和SM數量。

多個GPU可以掛載在PCI-E匯流排上，可以跟主機端或其他GPU通信。

流處理器簇(SM)內部結構如下圖所示：

SM內部結構

SM由一系列流處理器(Streaming Processor，SP)、寄存器文件(Register File)、共享內存(Shared Memory, SMem)、SPU(特殊運算單元)，以及紋理/常量/L1緩存組成。L2 cache是由所有SM共享的。

  • 流處理器(SP)：GPU的ALU單元，每個SP運行調度器分配給它的一個線程。 在CUDA編程模型里，32個線程(稱為一個線程束)同時執行一套指令；但是實際調度時，是以半個線程束(16個線程)調度的。這對應了NVIDIA GPU的每個SM內，SP的數量都是16的整數倍。
  • 寄存器文件：SM內的寄存器文件很大，大到分配到SM每個線程塊內的每個線程，都可以擁有自己的寄存器空間。因此線程的上下文切換，實際上只需要換一下寄存器空間的指針即可，十分迅速。
  • 共享內存：程序可控制的高速緩存。
  • 特殊運算單元(SPU)：進行一些特殊運算。
  • 紋理/常量/L1緩存。

4.0.1 調度過程

線程塊調度

當主機端啟動內核時，會根據線程網格中的線程塊(thread block)所需的寄存器和共享內存，決定將線程塊調度到哪個SM上運行，或者等待調度(沒有SM有足夠的資源運行該線程塊)。 只要資源足夠，一個SM上可以同時運行多個線程塊。 當線程塊運行完畢時，線程塊會退出SM，以供其他線程塊被調度上去。

另外需要強調一點的是，如果線程塊需要的寄存器或共享內存太多，以至於SM連一個線程塊都無法滿足的時候，內核會啟動失敗。

CUDA Toolkit提供了 CUDA Occupancy Calculator以供分析。

線程束調度

當線程塊被調度到SM上後，具體調度到哪個SP上運行，是由SM內部的調度器執行的。 NVIDIA GPU在邏輯上以32個線程(線程束)，作為最小調度單位。但實際上在硬體方面，是以半個線程束(16個線程)調度的。只不過調度器在調度了前半個線程束後，會立刻調度後半個線程束。調度過程如下圖所示(計算能力2.0平台)：
CUDA線程束調度

當調度器選擇調度線程束0時，第一個時鐘周期會將半個線程束調度到第一個16xSP上，然後下一個周期會將另外半個線程束調度到另一個16xSP上。因此，調度一個完整的線程束0的指令，需要2個時鐘周期。

上圖所示的計算能力2.0的設備，具有兩個調度器，剛好能夠保證32個SP核可以連續工作。但前提是SM至少有2個線程束待調度。否則如果只有一個線程束的話，調度器0將會限制。

線程束有時會處於等待狀態(等待內存，或等待特殊計算的結果)。此時調度器會調度其它線程束到SP上運行。當等待的線程束等到了所需的元素後，會再次處於就位狀態，等待調度器調度運行。

不同計算能力的平台，具有不同的調度器數量和SP數量。

4.1 SIMT架構

NVIDIA GPU是典型的SIMT架構(Single-Instruction, Multiple-Thread Architecture，單指令多線程架構)。(半)線程束在同一時間內執行同樣的指令(相同的PC)，但每個線程有自己的數據空間(寄存器)；可以同時做同樣的事情，但是處理不同的數據。這樣可以有效節省指令帶寬(指令由線程束共享)，實現高效並行。

當線程束中的線程，必須要執行不同的條件分支時，滿足分支條件的線程會被激活並執行分支內的內容；不滿足分支條件的線程會接收同樣的指令，但不會被激活，不會實際執行，但也不能跳過去執行其他指令。換言之，當線程束中的線程遇到分支時，不論線程是否需要執行分支，都會消耗執行該分支的時間，因為線程束中的線程執行同樣的指令，是高度同步的。

舉個例子，假如程序是按如下方式編寫的：

...
if(條件)
{
操作1
}
else
{
操作2
}
...

如果是CPU，如果不滿足條件，則會直接跳轉去執行操作2，不會執行操作1。 但是對於GPU同一個線程束內的線程，即使不滿足條件，依舊會去跟其他線程一起去執行操作1。等到操作1執行完畢後，該線程又會跟其他線程一起執行操作2。這也就意味著，不管線程實際上需不需要執行某一分支，它都要跟其他線程一起跑完這一分支。

另外，GPU不支持分支預測和推測執行，只會老老實實的一條條執行各個分支里的內容。因此，對於CUDA程序來說，分支實際上是一種低效的行為。

但是有兩種情況是例外：

線程束中所有線程均只需要執行一個分支，如線程束中的所有線程滿足條件，則所有線程執行完操作1後，不會去執行操作2。 線程束中的半個線程束(線程0~15、線程16~31)同時滿足條件，則這半個線程束不會去執行操作2(因為硬體實際上是按照半個線程束調度的)。

但是從Volta架構開始，Independent Thread Scheduling被引入，線程束內的線程不再完全同步。每個線程都會有自己獨立的PC。遇到分支時，不再像之前的架構一樣，只有(半)線程束內的線程條件一致時，才會跳過分支；Volta架構的調度優化器會將線程束中的線程，按照分支條件是否滿足，重新組合成SIMT單元，從而跳過分支。

Volta架構的Independent Thread Scheduling無疑是高效的，但是這是一個跟舊架構完全不同的特性。在編寫舊架構的CUDA程序時，程序員會默認線程束內的線程一定會同步執行。Volta架構的新特性破壞了這一假設，無疑會給程序帶來一些問題，需要注意。

4.1.1 線程的激活與原子操作

在一個線程束中，參與當前指令的線程稱為active threads，不參與的被稱為inactive threads。造成線程inactive的原因有：

某個線程比線程束中其他線程先退出(應該只發生在半線程束同時滿足分支的情況，或是Volta架構上)。 在分支結構中，線程不滿足當前分支的條件 * 設置線程塊大小時，為了補齊32的倍數而創建的線程(如線程塊定義為31個線程，則會補1個線程，這個線程是inactive的)

當線程束中所有線程，向同一個內存地址寫數據時，不能保證哪個線程先寫，哪個線程後寫，即程序的正確性是無法保證的。 這時需要使用CUDA提供的原子操作(atomic)函數，如atomicAdd()。

原子操作可以保證程序的正確性，但是會造成線程束中線程的串列化(serialization)，執行時間比並行執行要長。

另外需要注意的是，即使沒使用原子操作，向同一個內存地址寫數據，一樣也會產生一定程度的串列化，串列化程度依架構而定。

4.2 硬體多線程

GPU線程束中的線程的上下文(寄存器&PC等)都存在片內空間龐大的寄存器文件中，直到線程束執行完畢才會被釋放(生命周期為整個線程束執行過程)。因此，不像CPU，GPU的線程束上下文切換十分迅速，沒有損耗。 線程束調度器(warp scheduler)會選擇ready狀態的線程束，將其調度到SP上執行。

Chapter 5 性能優化

5.1 性能優化概述

CUDA程序性能優化有三個原則：

最大化並行，以提升資源利用率 優化內存排布，以最大化內存吞吐 * 最大化指令吞吐

在性能優化前，需要先分析程序性能的瓶頸，再針對瓶頸優化，否則收益會很低。

分析程序瓶頸，可以使用CUDA profiler等工具。

5.2 最大化利用率(Maximize Utilization)

最大化利用率的方法就是並行。

5.2.1 應用級別並行(Application Level)

從程序最高層來看，應該儘可能讓主機端、設備端、PCI-E匯流排並行工作。對此可以使用非同步CUDA函數，以及流(Stream)來實現。

同步操作，以及內存的共享會影響程序的並行性。因此需要仔細設計演算法流程，盡量減少同步和內存共享。 如果一定需要同步和內存共享，盡量在線程塊內完成(線程塊同步——使用__syncthreads()涉及到的線程少，且可以通過SM內的共享內存共享數據。如果需要線程網格內同步，則需要兩個內核調用，且共享數據只能通過全局內存，速度慢)。

5.2.2 設備級別並行(Device Level)

可以通過流的方式，儘可能的讓多個內核並行，提升利用率。

5.2.3 處理器級別並行(Multiprocessor Level)

延遲(latency)指的是線程束(從上一個動作開始)到它處於ready狀態的時鐘數。 例如線程束先提交了一個內存訪問請求，然後等了400個時鐘周期，內存管理系統才返回數據，線程束可以繼續執行。這400個時鐘周期稱為延遲。

當一個線程束髮生延遲時，線程束調度器(warp scheduler)會將其他處於ready狀態的線程束調度到SP上。等到延遲結束後，再將該線程調度回SP繼續執行。這樣一來，前一個線程束的延遲，就被另一個線程束的執行所隱藏了。 這一過程被稱作延遲的隱藏(hidden latency)。

隱藏延遲是GPU編程的核心概念。由於GPU具有巨大的寄存器空間，線程的切換不存在損耗。因此，通過向GPU上分配足夠多的線程，可以讓這些線程延遲互相交錯，以起到隱藏延遲的作用，提高硬體利用率。

造成線程(束)產生延遲的原因有：

指令執行：不同指令有不同的執行延遲 內存請求：共享內存、全局內存、PCI-E(Mapped Memory)的讀寫請求 * 同步操作：如使用__syncthreads()後，先完成的線程(束)，會等待線程塊中其他線程(束)達到同步點。

通過配置線程網格、線程塊、寄存器和共享內存用量，讓SM可以運行儘可能多的線程束，以隱藏延遲。例如對於計算能力3.x的設備，為了完全隱藏全局內存讀取的延遲(200-400時鐘)，需要大概40個線程束。

舉個例子，設SM有32KB共享內存空間。程序每個線程需要32B共享內存，即一個線程束需要1KB共享內存，考慮下述兩種方案：

方案1：每個線程塊有16個線程束，則每個線程塊需要16KB共享內存。可以調度兩個線程塊到SM上。 方案2：每個線程塊有18個線程束，則每個線程塊需要18KB共享內存，則只能調度一個線程塊到SM上。

雖然方案2在一個線程塊上，有更多的線程束，但是實際上SM上運行的線程束減少了(32->18)。因此方案2隱藏延遲的能力弱於方案1，資源利用率較低。

此外，如果寄存器使用過多，超過了SM上的寄存器空間，則會使用本地內存作為寄存器。本地內存是存在在全局內存上的，速度很慢，會嚴重影響程序速度。因此需要嚴格考慮寄存器使用數量。 (這裡官方文檔和另一本書里說的有矛盾，難道是新的架構把本地內存給取消掉了？)

最後強調一點，線程塊中的線程數量，最好是32的整數倍。這樣，就不會有為了補齊線程束，而出現的永遠不會激活的線程。這些不激活的線程也會佔用SM的資源，降低資源利用率。

CUDA具有Occupancy Calculator，幫助程序員設計。

5.3 最大化內存吞吐(Maximize Memory Throughput)

最大化內存吞吐，主要手段就是少用低帶寬的內存。這意味著首先要儘可能減少主機端和設備端間的設備傳輸(PCI-E，特別慢)，其次要儘可能減少全局內存的讀寫(快於PCI-E，但是相對於片內內存來說，還是挺慢的)；儘可能的使用片內的內存(寄存器、cache、共享內存)。

這裡需要強調一下cache和共享內存的事情。

共享內存是程序可控的高速緩存。一般情況下，共享內存的使用流程為： 將數據從全局內存拷貝到共享內存，或初始化共享內存 進行一個同步操作，確保共享內存全部被賦值 利用共享內存的數據，運行程序 如果出現了共享內存的寫操作，一般需要進行一個同步操作，確保寫操作全部完成後再進行下面的操作 * 將數據寫回全局內存

這裡有一點要強調，只有在數據需要反覆讀寫的時候，共享內存才有意義。如果數據只會被讀一次，處理完後又寫回並不再處理。則直接從全局內存讀出->寄存器運行->寫回全局內存是最快的。在共享內存中轉反而是慢的。

緩存(L1/L2 cache)是程序員無法顯式編程的。但是如果了解緩存的特性的話，可以通過合適的程序設計，增加緩存命中率。

事實上，硬體控制的cache，擁有更好的數據局部性(locality)。

5.3.1 主機端和設備端間數據傳輸

由於PCI-E傳輸並不快，因此要盡量減少主機端和設備端間的數據傳輸： 一種方式是讓中間結果儘可能的在設備端產生，在設備端使用。 另一種方式是將很多小的數據，打包傳輸。 還有可以通過分配鎖頁內存來加快前端匯流排*系統的帶寬。

當使用內存映射時，需要注意，每次內存訪問都會啟動一次PCI-E傳輸。因此，盡量保證數據只被讀寫一次，且儘可能合併訪問以提升有效內存帶寬。

有些GPU設備，主機端和設備端內存，在物理上就是同一塊。這種情況下，主機端和設備端傳輸是不存在的。可通過標誌integrated來查看。

5.3.2 設備內存訪問

5.3.2.1 全局內存(global memory)

全局內存支持合併訪問，可以一次性傳輸連續的32、 64、 128位元組的數據。因此，在設計內核時，線程束內的線程盡量連續的訪問內存。

考慮如下兩個內核：

//假設gpuData是一個二維數組，尺寸為32x32
int gpuData[32][32]; //這樣是不合法的，因為這麼定義實際上是在主機端，還需要拷貝到設備端，這裡只是為了方便說明問題

__global__ void Kernel1(int gpuData[][32])
{
const int tid = blockIdx.x * blockDim.x + threadIdx.x;
int sum = 0;
for(int i = 0; i < 32; i++)
sum += gpuData[i][tid]; //行訪問
...
}

__global__ void Kernel2(int gpu[][32])
{
const int tid = blockIdx.x * blockDim.x + threadIdx.x;
int sum = 0;
for(int i = 0; i < 32; i++)
sum += gpuData[tid][i]; //列訪問
...
}

上例中，執行Kernel1的線程束中的線程，在一次循環中，32個線程依次訪問gpuData[0][0], gpuData[0][1], gpuData[0][2], ..., gpuData[0][31]。在內存中，這32個變數是連續存儲的，因此可以被合併訪問。這種訪問被稱為行訪問。

而Kernel2在一次循環中，讀取的變數為gpuData[0][0], gpuData[1][0], gpuData[2][0], ..., gpuData[31][0]。這32個變數是不連續的，需要進行32次內存請求。這種訪問被稱為列訪問。

上例中，列訪問之所以效率低，原因有二：

對於執行一次循環，行訪問只需要一個內存請求指令，而列訪問需要32個內存請求指令。從指令角度來講，行訪問的內存請求指令帶寬是列訪問的1/32。 全局內存的最大帶寬為一次取128Byte，但是內核每次只需要4個Byte的數據。這使得列訪問的內存帶寬為峰值帶寬的1/32。事實上，即使內核只需要4Byte，GPU也會取連續的32Byte，然後丟掉後面的28Byte，造成資源的浪費。但是緩存的引入(自計算能力2.x開始)，這一問題得到了緩解，28Byte會先放到緩存中，下次會命中。

因此，從上例中可以看到，好好安排內存排布，盡量使得內存訪問可以合併，可以加速全局內存的讀寫。

5.3.2.2 對齊(Alignment)

當變數的尺寸為1/2/4/8/16位元組時，變數會對齊。但如果不是的話，變數無法對齊，會產生額外的內存訪問。

C/C++內建的變數(int/float等)，以及CUDA支持的向量(float2/float4等)，是對齊的。

一些結構體可能會產生不對齊的情況，看下例：

struct struct1{
float x;
float y;
};

struct struct2{
float x;
float y;
float z;
};

struct struct3 __align__(16){
float x;
float y;
float z;
};

上例中，struct1是8位元組的結構體，自動會對齊； struct2具有12個位元組，無法對齊； struct3使用了__align__(16)關鍵字，顯式指定對齊到16。

使用各類malloc分配的設備內存，一定是256位元組對齊的。

5.3.2.3 本地內存(local memory)

當使用了自動變數(不明白，再看)時，有可能會將數據放到本地內存上：

Arrays for which it cannot determine that they are indexed with constant quantities 大的結構體，寄存器放不下 * 寄存器溢出(register spilling)，即內核使用的寄存器多於SM上可用的寄存器

通過看PTX代碼，可以看到標記為.local的變數，就是本地內存。

即使PTX代碼里沒有使用本地內存，在編譯到cubin代碼的過程中，仍然會使用本地內存，編譯器會報告lmem的使用情況。

前面多次強調過了，一旦使用了本地內存，其速度會非常慢。不過本地內存在存儲的時候，是按照32個線程連續存儲的，因此可以合併訪問。

對於計算能力3.x的設備，本地內存會被緩存在L1/L2 cahce；對於計算能力5.x和6.x設備，本地內存會被緩存到L2 cache。即便如此，其速度還是慢於寄存器。

5.3.2.4 共享內存(shared memory)

共享內存實際上是被分為多個存儲體(memory bank)。多個線程訪問同一個存儲體會造成串列化。

(存疑：存儲體其實是可以廣播的，因此多個線程讀同一個存儲體是不存在衝突的，只是寫會存在串列化問題)

因此，編寫內核時，需要認真設計，以避免存儲體訪問的衝突。

5.3.2.5 Texture and Surface Memory

略

5.4 最大化指令吞吐(Maximize Instruction Throughput)

可以使用如下方法來最大化指令吞吐：

盡量少使用吞吐率低的算數指令 盡量減少線程束內的分支 * 盡量減少指令數，如少用__syncthreads()，或者在合適的時候使用__restrict__

指令吞吐的定義：每個SP在每個時鐘周期內執行的操作數。如果一個線程束在一個時鐘周期內執行了N個操作，則指令吞吐為N/32。

5.4.1 算數指令(Arithmetic Instructions)

官方文檔這裡比較混亂，但主要有如下幾點： 不同架構的設備，不同指令有不同的指令吞吐，可以查表 有一些快速的內聯(inline)函數，如使用__fdividef()(快速浮點數除法)來代替普通的除法來加速 整形的除法和取余會比較慢，可能需要20個機器周期；因此對於n為2的冪次的情況，使用i>>log2(n)代替i/n，使用i&(n-1)來代替i%n 半精度(浮點數)運算(Half Precision Arithmetic)：可以使用half2數據類型，並使用對應的運算指令(如__hadd2, __hsub2, __hmul2, __hfma2等)，來讓一個周期內執行兩次運算，以節省指令帶寬。可以通過__halves2half2將兩個半精度浮點數合併為half2數據類型。 (半精度又是咋定義的？) * 數據類型轉換：當使用char或short，亦或是雙精度常量與單精度變數相互操作時，會觸發數據類型轉換，需要一定執行時間(實際上，char和short，不管是存儲在寄存器中，還是在運算時，都是以int型進行的)

5.4.2 控制流指令(Control Flow Instructions)

盡量避免向線程束中引入分支。

此外，可以使用#pragma unroll宏，來進行循環展開，減少控制指令。

5.4.3 同步指令(Synchronization Instruction)

下表為不同計算能力的設備，同步指令__syncthreads()需要消耗的指令周期為：

|計算能力|__syncthreads()消耗的指令周期| |---|---| |3.x|128| |5.x,6.1,6.2|64| |6.0|32| |7.x|16|

注意，__syncthreads()會造成線程塊中的線程等待，影響內核執行效率。


推薦閱讀：
相关文章