執行緒的索引計算
只需要知並行執行緒的初始索引，以及如何確定遞增的量值，我們希望每個並行執行緒從不同的索引開始，因此就需要對執行緒索引和執行緒塊索引進行線性化，每個執行緒的其實索引按照以下公式來計算：
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》