CPU和GPU架構
處理器結構主要要考慮的兩個指標：延遲和吞吐量。
- 延遲：一條指令從發布到返回結果所經歷的時間。
- 吞吐量：單位時間內處理的指令的條數。
CPU： 延遲到向內核
GPU： 吞吐導向內核
CPUs
內存大
- 多級緩存結構提高訪問速度
有復雜的控制單元
- 分支預測機制（判斷運算分支）
- 流水線數據前送機制（判斷經常用到的數據，並將其前送）
運算單元強大
- 整型浮點型運算速度快
GPUs
緩存小
- 提高內存吞吐
控制簡單
- 沒有分支預測
- 沒有數據轉發
精簡運算單元
- 多長延時流水線以實現高吞吐量
- 需要大量線程來容忍延遲
GPU&CPU特點
CPUs：連續計算部分，延遲優先。CPU比GPU單條復雜指令延遲快10倍以上。
GPUs：並行計算部分，吞吐優先。GPU比CPU單位時間內執行指令數量10倍以上。
什么樣的問題適合GPU
計算密集：數值計算的比例要遠大於內存操作，因此內存訪問的延時可以被計算掩蓋。
數據並行：大任務可以被拆解為執行相同指令的小任務，因此對復雜流程控制的需求較低。
- 內存訪問次數少
- 控制簡單
- 計算簡單
- 並行度高
GPU編程與CUDA
CUDA (Compute Unified Device Architecture)，初衷是為GPU增加一個易用的編程接口，讓開發者無需學習復雜的着色語言或者圖形處理原語。
OpenCL(Open Computing Languge)是2008年發布的異構平台並行編程的開發標准，也是一個編程框架。OpenCL相比於CUDA，支持的平台更多，除了GPU還支持CPU、DSP、FPGA等設備。
GPU編程並行計算整體流程
void GPUkernel(float* A, float* B, float* C, int n){
1. // Allocate device memory for A, B and C
// copy A and B from host memory (CPU) to device memory (GPU)
2. // Kernel launch code - to have the device
// to perform the actual vector addition
3. // copy C from the device memory to host memory
// free device vectors
}
CUDA編程術語
硬件側
-
Device=GPU
-
Host=CPU
-
Kernel=GPU上運行的函數
-
CUDA中的內存模型分為以下幾個層次：
- 每個線程處理器（SP）都用自己的registers（寄存器）
- 每個SP都有自己的local memory（局部內存）， register和local memory只能被線程自己訪問
- 每個多核處理器（SM）內都有自己的shared memeory（共享內存），shared memory可以被線程塊內所有線程訪問，不過shared memory的大小較小
- 一個GPU的所有SM共有一塊global memory（全局內存）， 不用線程塊的線程都可使用
軟件側
- CUDA中的內存模型分為以下幾個層次：
- 線程處理器（SP）對應線程（thread）
- 多核處理器（SM）對應線程塊（thread block）
- 設備端（device）對應線程塊組合體（grid）
- 一個kernel其實由一個grid來執行
- 一個kernel一次只能在一個GPU上執行
線程塊：可擴展的集合體
線程是內存模型在軟件側最基本的執行單位，線程塊就是線程的組合體。
將線程數組分成多個塊
- 塊內的線程通過共享內存、原子操作和屏障同步進行協作（shared memory， atomic operations and barrier synvhronization）
- 不同塊中的線程不能協作
特點：
- 線程塊內的所有線程的計算和內存都是互相獨立的
- 線程塊內的共享內存都可訪問
- 可以用公有時鍾來同步
網格（grid）：並行線程塊組合體
CUDA核函數由線程網格（數組）執行
- 每個線程都有一個索引，用於計算內存地址和做出控制決策
特點：
- 網格內的線程塊是互相獨立的
- 網格內的全局內存可由網格內的所有線程塊訪問
- 可用公有時鍾同步網格內所有線程塊
線程塊id&線程id：定位獨立線程
也是通過索引來確定線程的地址， 但GPU不止需要線程的索引，還需要線程塊的索引（可能有多維）。
dim3 dimGrid(M, N);
dim3 dimBlock(P, Q, S);
threadId.x = blockIdx.x * blockDim.x + threadIdx.x;
threadId.y = blockIdx.y * blockDim.y + threadIdx.y;
線程束（warp）
- SM采用的SMIT（Single-Instruction， Multiple-Thread， 單指令多線程）架構，warp（線約束）是最基本的執行單元，一個warp包含32個並行thread，這些thread以不同數據資源執行相同的指令。warp本質上是線程在GPU上運行的最小單元。
- 當一個kernel被執行時，grid中的線程塊被分配到SM上，一個線程塊的thread只能在一個SM上調度，SM一般可以調度多個線程塊，大量的thread可能被分到不同的SM上。每個threade擁有它自己的程序計數器和狀態寄存器，並且用該線程自己的數據執行指令，這就是所謂的Single Instruction Multiple Thread（SIMT）。
- 由於warp的大小為32， 所以block包含的thread的大小一般設置為32的倍數。
並行計算實例：向量相加
CPU
// Compute vector sum C=A+B
void vecAdd(float* A, float* B, float *c, int n){
for (i = 0, i < n, i++)
c[i] = A[i] + B[i];
}
int main(){
//Memory alocation for A_h, B_h and C_h
// I/O to read A_h and B_h, N elements
vecAdd(A_h, B_h, C_h, N);
}
GPU
void vecAdd(float* A, float* B, float* C, int n){
int size = n * sizeof(float);
float* A_d, *B_d, *C_d;
// Transfer A and B to device memory
cudaMalloc((void **) &A_d, size);
CUDAMemcpy(A_d, A, size, cudaMemcpyHostToDevice);
cudaMalloc((void **) &B_d, size);
cudaMemcpy(B_d, B, size, cudaMemcpyHostToDevice);
// Allocate device memory for
cudaMalloc((void **) &C_d, size);
// Kernel invocation code - to be shown later
// Transfer C from device to host
cudaMemcpy(C, C_d, size, cudaMemcpyDeviceToHost);
// Free device memory for A, B, C
cudaFree(A_d);
cudaFree(B_d);
cudaFree(C_d);
}