本帖經過多方整理，大多來自各路書籍《GPGPU編程技術》《cuda高性能》
1 grid 和 block都可以用三元向量來表示：
grid的數組元素是block
block的數組元素是grid
但是1.x計算能力的核心，grid的第三元必須為1.block的X和Y索引最大尺寸為512
2 通過__launch_bounds__(maxBlockSize,minBlocksPerMp)來限制每個block中最大的線程數，及每個多處理器上最少被激活的block數
3 SM streaming multiprocessor 多流處理器
SP scalar processor cores 標量處理核心
一個Block中的所有線程在一個多處理器上面並發執行。當這個Block的所有線程執行完后，再激活其他等待的Block.一個多處理器上也可以執行多個block。但是一個block卻不能拆分為多個處理器上面執行
對於同一個Block里面的線程：
1 同一個Block里的線程可以被同步
2 可以共同訪問多處理器里的共享存儲器
到2.x為止，多處理器 執行任務時，以32個並行線程為單位，稱為一個wrap。
當以個block到來的時候，會被分成線程號連續的多個wrap，然后多處理器上的SIMT控制器以wrap為單位控制調度線程。所以block中的線程數要是以32的整數倍來設計，就不會出現空閑的SP。組織WARP的時候，從線程號最小的開始
4 各個存儲器存儲位置及作用
5 寄存器放在SP中，如果溢出，會被放在設備處理器上面，發生嚴重滯后，影響性能。
1.0 4KB 2.0 16kb
共享存儲器位於SM中，大約兩個時鍾周期讀寫4B，靜態分配 __shared__ int shared[16];
1.0 16KB 2.0 48kb
6 共享存儲器，是以4個字節為單位的16個存儲器組
bank沖突：半個warp中的多線程訪問的數組元素處於同一個bank時，訪問串行化，發生沖突
避免沖突：最多的數據類型是int、float等占用4個字節的類型
7線程設計
float shared=data[base+tid]; base訪問的起始元素下標 tid線程號
如果要是char類型，每個元素占1個字節，就會沖突
float shared = data[base+4*tid];
8 共享存儲器廣播訪問：半個warp線程都訪問一個數據
9 補白策略
shared[tid]=global[tid];
int number = shared[tid*16];
int nRow = tid/16;
int nColumn = tid%16;
shared[nColumn*17+nRow] = global[tid];
int number = shared[17*tid];
10 一次性訪問全局存儲器：數據的起始地址應為每個線程訪問數據大小的16倍的整數倍
11 主機鎖頁存儲器：cudaHostMalloc()分配。
不參與操作系統分頁管理的存儲空間，訪問鎖頁文件不會耗費主機內存分頁管理方面的開銷。不會被操作系統放到硬盤的頁面文件中，因此比訪問普通的主機存儲器更快。
12 計算能力2.x的GPU上面，每個SM有獨立的一級緩存，有唯一的二級緩存
13 異步並發：
主機上的計算、
設備上的計算、
主機到設備上的傳輸、
設備到主機上的傳輸共同執行
14 設備存儲器 類型是DRAM，動態隨機存儲器。使用它最高效的方式就是順序讀取。為了保證順序：
__global__ static void sumof(int *pnNumber,int* pnResult,clock_t* pclock_tTime){
const int tid = threadIdx.x;
int nSum = 0;
int i;
clock_t clock_tStart;
if(tid == 0) clock_tStart = clock();
for(i = tid;i<DATA_SIZE;i+=THREAD_NUM){
nSum += pnNumber[i]*pnNumber[i];
}
pnResult[tid] = nSum;
if(tid == 0)
*pclock_tTime = clock()-clock_tStart;
}
每個block 在1.x的計算能力的GPU下，最多只有512的線程數
__global__ static void sumof(int *pnNumber,int* pnResult,clock_t* pclock_tTime){
const int tid = threadIdx.x;
const int bid = blockIdx.x;
int nSum = 0;
int i;
clock_t clock_tStart;
if(tid == 0) pclock_tTime[bid] = clock();
for(i = bid*THREAD_NUM+tid;i<DATA_SIZE;i+=BLOCK_NUM*THREAD_NUM){
nSum += pnNumber[i]*pnNumber[i];
}
pnResult[bid*THREAD_NUM+tid] = nSum;
if(tid == 0)
*pclock_tTime[bid+BLOCK_NUM] = clock();
}
15 用縮減樹避免bank沖突：
bank沖突指的是，一個warp內的線程同時訪問一個bank列，導致串行讀取數據
noffset = THREAD_NUM/2;
while(noffset > 0){
if(tid < offset)
nshared[tid] += nshared[tid+noffset];
}
noffset >>= 1;
__syncthreads();
16 CPU有強大的分支預測、程序堆棧、循環優化等針對控制采取的復雜邏輯。
GPU相對簡單，適合處理順序的，單一的，少循環，少跳轉的語句。
17 #progma unroll 5下面的程序循環5次
18 cuda中的同步
1》__syncthreads()同步
同一個warp內的線程總是被一同激活且一同被分配任務，因此不需要同步。因此最好把需要同步的線程放在同一個warp內，這樣就減少了__syncthreads()的指令
2》__threadfence() __threadfence_block()同步
前者針對grid的所有線程，后者針對block內的所有線程。告知線程，全局存儲器或共享存儲器已經被改變
3》cudaThreadSynchronize() 主機與設備間的同步
在主機程序里同步線程。該函數以上的設備線程完成后，控制權才交給cpu
4》volatile關鍵字
使用這個關鍵字定義數組，設備會知道這個數組隨時都會改變，就會自動重新讀取數組（但是不能保證線程間讀取的數據一致）