GPU之CUDA 第一彈

1.GUP與CUDA

GPU：Graphics Processing Units

CUDA：Compute Unified Device Architecture

2.主要概念

2.1 主機：CPU+記憶體

2.2 裝置：GPU+顯存

2.3 Grid（線程格）：一個Kernel對應一個Grid，一個Grid包含若幹Block

2.4 Block（線程塊）：Kernel以Block為機關執行，各Block并行運作，無法通信，沒有執行順序

2.5 thread（線程）：

可以認為線程是執行CUDA程式的最小單元。

GPU資源充裕時，所有線程都并發執行；GPU資源少于總線程個數時，部分線程需等待前面執行的線程釋放資源，即變為串行化執行。

2.6 warp（線程束）：32個并行線程為warp塊

2.7 Kernel（核函數）：

通過__global__定義，在主機端調用（需先配置設定空間再調用）。

調用時聲明執行參數<<< m,n,a,b >>>，m個線程塊，每個線程塊有n個線程，每個線程塊使用了a位元組的動态配置設定的共享存儲器，核心在流b中執行。

Kernel調用是異步的，即主機僅把要執行的Kernel順序送出給GPU，并不等待其執行完成。這種異步調用導緻Kernel無法傳回函數值。

2.8 dim3結構類型：

3個由unsigned int型組成的結構體，unsigned int x;unsigned int y;unsigned z

gridDim為grid尺寸，blockDim為block尺寸，blockIdx為block在grid中的索引值，threadIdx為線程索引号。

CUDA定位公式：tid=blockIdx.x*blockDim.x+threadIdx.x

3.常用記憶體函數

3.1 cudaMalloc((void**)&d_a,sizeof(float)*n)：在GPU上配置設定記憶體

3.2 cudaMemcpy(d_a,a,sizeof(float)*n,cudaMemcpyHostToDevice)：在主機記憶體和GPU顯存間互傳資料（上述為主機到裝置）

3.3 cudaFree(d_a)：釋放記憶體

4.GPU記憶體分類

4.1 全局記憶體(global memory)：

可讀寫

通俗意義上的裝置記憶體

通路全局記憶體造成的延時很高，高達通路寄存器變量延時的600倍。GPU的資料重用是達到高性能的關鍵。

高性能CUDA程式應充分利用GPU的内部資源（寄存器、共享記憶體等）來避開全局記憶體瓶頸。

4.2 共享記憶體(shared memory）：

可讀寫

位于裝置記憶體，添加關鍵字__shared__聲明

4.3 寄存器（registers）：

可讀寫

4.4 Local memory：

可讀寫

4.5 常量記憶體(constant memory）：

隻讀

4.6 紋理記憶體(texture memory）：

隻讀

總結：每個線程都有自己的registers和local memory

    同一個block中的每個線程共享一份共享記憶體

    所有線程（包括不同block的thread）都共享一份全局記憶體、常量記憶體、紋理記憶體。

    不同grid有各自的全局記憶體、常量記憶體、紋理記憶體。

5.SM與SP

5.1 SM（流多處理器）：

同block的線程被配置設定在一個SM中

目前CUDA定義中，一個SM中至多配置設定8個blocks 

5.2 SP（标量流處理器）：

GPU中最基本的處理單元即為SP

在 G80/G92 的架構下，有 128 個 SP，以 8 個 SP 為一組，組成 16 個 SM，再以兩個 SM 為一個 TPC，共分成 8 個 TPC 來運作。

在新一代的 GT200 裡，SP 則是增加到 240 個，還是以 8 個 SP 組成一個 SM，但是改成以 3 個 SM 組成一個 TPC，共 10 組 TPC。

6.錯誤處理

cudaErrot_t：CUDA的錯誤類型。

cudaError_t cudaStatus;  
cudaStatus = cudaSetDevice(0);  
if (cudaStatus != cudaSuccess) {  
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
        goto Error;  
}  
           

7.函數限定符

7.1 __device__：在GPU上調用，在GPU上執行。

7.2 __global__：在CPU上調用，在GPU上執行。

7.3 __host__：在CPU上調用，在CPU上執行。

8.變量限定符

8.1 __device__：聲明的資料存放在顯存中，所有線程均可通路，且主機也可以通過運作時庫通路。

8.2 __shared__：資料存放在共享存儲器上，隻有所在塊内的線程可以通路，其他塊内的線程不能通路。

8.3 __constant__：資料存放在常量存儲器上，可以被所有線程通路，主機也可以通過運作時庫通路。