CUDA經常使用概念及注意點

線程的索引計算

只須要知並行線程的初始索引，以及如何肯定遞增的量值，咱們但願每一個並行線程從不一樣的索引開始，所以就須要對線程索引和線程塊索引進行線性化，每一個線程的其實索引按照如下公式來計算：

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

線程塊數量限制：65536

線程塊中線程數量限制：512

共享內存和同步

共享內存

__share__添加到變量聲明,使得聲明的變量駐留在共享內存中。cuda c編譯器對共享內存中的變量和普通變量採用不一樣的處理策略，對於GPU啓動的每一個線程塊，cuda c都將建立該變量的一個副本。線程塊中的全部線程都會共享這塊內存。但線程卻沒法看到也不能修改其餘線程塊的變量副本。這是同一個線程塊中的不一樣線程進行通訊和協做的基礎。

共享內存緩衝區駐留在物理gpu上，訪問時延極低。

同步

__syncthreads();對線程塊中的線程進行同步。線程發散的容易出現，使得部分場景下的線程同步頗有必要。

線程發散：當某些線程須要執行一些指令，而其餘線程不須要執行時，這種狀況叫作線程發散。在正常環境中，發散的分支會使得某些線程處於空閒狀態，而其餘線程將執行線程中的代碼。在__syncthreads()狀況中，線程發散形成的結果有些糟糕，cuda架構將確保，除非線程塊中全部的線程都執行了同步操做，不然沒有任何線程能夠執行同步操做以後的指令。

常量內存與事件

常量內存：NVIDIA提供64k的常量內存，有效減小內存寬帶。__constant__ 將變量的訪問限制爲只讀。

從主機內存複製到GPU上的常量內存，使用方法cudaMemcpyToSymbol()進行復制。

性能提高緣由

（1）對常量內存的單次操做能夠廣播到其餘鄰近線程，節約15次的讀寫操做；

當處理常量內存時，NVIDIA硬件將單次內存讀取操做廣播到每一個半線程束。

（2）常量內存的數據將緩存起來，所以對相同地址的連續訪問不會產生額外的內存通訊量。

線程束：warp

在cuda架構中，線程束指的是一個包含32個線程的集合，這些個線程被編制在一塊兒，而且以步調一致（LockStep）的形式執行，在程序中的每一行，線程束中的每一個線程都將在不一樣的數據上執行相同的命令。

事件API

cuda的事件本質上其實就是一個時間戳，這個時間戳就是在用戶指定的時間上記錄的。得到一個時間戳只有兩個步驟：建立一個事件，記錄一個事件。

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
// gpu執行操做
...
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);

cudaEventDestroy(start);
cudaEventDestroy(stop);

文理內存（Texture Memory）

簡介：

​ 與常量內存相似，只讀，緩存在芯片上，減小對內存的請求，並提供更高效的內存帶寬。專門爲在內存訪問模式中存在大量空間局部性（special locality）的圖形應用程序而設計。在某個計算應用程序中，這意味着一個線程讀取的位置可能與鄰近線程讀取的位置很是接近。紋理內存專門爲加速這種內存訪問模式。

紋理存儲器中的數據是以一維、二維或者三維數組的形式存儲在顯存中，能夠經過緩存加速訪問，而且能夠聲明大小比常量存儲器要大得多。在kernel中訪問紋理存儲器的操做稱爲紋理拾取。將顯存中的數據和紋理參考系關聯的操做，稱爲將數據和紋理綁定。顯存中能夠綁定到紋理的數據有兩種，分別是普通的線性存儲器和cuda數組。

使用步驟：

（1）須要將輸入的數據聲明爲texture類型的引用；聲明變量在gpu上；

（2） gpu中分配內存，經過cudaBindTexture()將變量綁定到內存緩衝區。告訴cuda運行時兩件事情：

• 咱們但願將指定的緩衝區做爲紋理來使用；
• 咱們但願將紋理引用做爲紋理的「名字」。

（3）啓動核函數，讀取核函數中的紋理時，須要經過特殊的函數來告訴GPU將讀取請求轉發到紋理內存而不是標準的全局內存，使用編譯器內置函數：tex1Dfetch();

（4）釋放緩衝區，清除與紋理的綁定，cudaUnbindTexture();

流

頁鎖定內存

頁鎖定主機內存，固定內存，不可分頁內存，OS將不會對這塊內存分頁而且交換到磁盤上。從而確保該內存始終駐留在物理內存中。OS能夠安全的使某個應用程序訪問該內存的物理地址，這塊內存將不會被破壞或者從新定位。

• malloc分配的是標準的、可分頁的主機內存；
• cudaHostAlloc將分配頁鎖定的主機內存。

建議：僅對cudaMemcpy（）調用的源內存或者目標內存，才能使用頁鎖定內存，而且在不須要使用他們時，當即釋放。

流

支持設備重疊功能的設備，支持設備重疊功能的GPU可以在執行一個CUDA C核函數的同時，還能在設備和主機之間進行復制操做。

一些新的GPU設備同時支持核函數和兩次的複製操做，一次是從主機到設備，一次是從設備到主機。在任何支持內存複製和核函數的執行相互重疊的設備上，當使用多個流時，應用程序的總體性能都能獲得提高。

判斷設備是否支持計算與內存複製操做的重疊：

int main( void ) {
	cudaDeviceProp prop;
	int whichDevice;
	HANDLE_ERROR( cudaGetDevice(&whichDevice) );
	HANDLE_ERROR( cudaGetDeviceProperties(&prop, whichDevice) );
	if(!prop.deviceOverlap) {
      	printf("Device will not handle overlaps");
      	return 0;
	}
}

多GPU系統上的CUDA C

零拷貝內存：能夠在cuda C核函數中，直接訪問這種類型的主機內存，因爲這種內存不須要複製到GPU，所以稱爲零拷貝內存。經過cudaHostAlloc進行分配，最後一個參數採用：cudaHostAllocMapped.

判斷設備是否支持映射主機內存：

int main( void ) {
	cudaDeviceProp prop;
	int whichDevice;
	HANDLE_ERROR( cudaGetDevice(&whichDevice) );
	HANDLE_ERROR( cudaGetDeviceProperties(&prop, whichDevice) );
	if(prop.canMapHostMemory != 1) {
      	printf("Device can not map memory");
      	return 0;
	}
}

當輸入內存和輸出內存都只是用一次時，那麼在獨立GPU上使用零拷貝內存將帶來性能提高。

判斷某個GPU時集成的仍是獨立：

cudaGetDeviceProperties()獲取屬性結構體，該結構中的域：integrated，若是是設備是集成GPU，該值爲true，不然爲false。

注意：多GPU場景下，每一個gpu若是都要運行gpu程序的話，都須要主機cpu啓動單獨的線程進行資源控制，都有對應本身的線程。

《Programming Massively Parallel Processors: a Hands-On Approach》