cuda編程知識普及

本帖通過多方整理，大多來自各路書籍《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關鍵字

 

　　使用這個關鍵字定義數組，設備會知道這個數組隨時都會改變，就會自動從新讀取數組（可是不能保證線程間讀取的數據一致）