programming model

kernels

類似於c函式，函式定義使用global宣告，使用<<<…>>>形式的execution configuration決定kernal執行的執行緒數，使用threadIdx變數可以獲得每一個執行緒對應的id。

// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}
int main()
{
    ...
    // Kernel invocation with N threads
 

    VecAdd<<<1, N>>>(A, B, C);
    ...
}

execution configuration定義形式為

<<< Dg, Db, Ns, S >>>
Dg: 表示grid的維度和大小，對應的block個數為Dg.x * Dg.y * Dg.z
Db: 表示block的維度和大小，對用的thread個數為Db.x * Db.y * Db.z
Ns: 表示shared memory的大小，可選引數
S : 表示cudaStream_t的型別，可選引數

thread hierarchy

thread：
block: 多個thread，個數型別可以是int/dim3，一般最多1024個thread
grid: 多個block，個數型別可以是int/dim3
block之間獨立執行，block內的thread可以通過shared memory共享資料，通過呼叫__syncthreads()函式實現同步。
- threadIdx: block內thread的id
- blockIdx: grid內的block的id
- blockDim: 一個block的thread的個數

memory hierarchy

每個thread有local memory
每個block有shared memory，block內的thread都可以訪問
所有的thread都可以訪問global memory/const memory/texture memory，其中const memory和texture memory只讀。

heterogeneous programming

host和device維護DRAM中不同的記憶體空間，分別稱為host memory和device memory。

compute capability

表示GPU硬體的特性，也稱為“SM version”，由兩部分組成”X.Y”，同一個X的版本號對應的核心架構一致，其中5表示基於Maxwell架構，3表示Kepler架構，2表示Fermi架構，1表示Tesla架構。從CUDA7.0開始不再支援Tesla架構。

programming interface

NVCC編譯

編譯分為offline compilation或者just-in-time compilation
1. offline compilation
cuda的程式碼可以是PTX或者c，都需要通過nvcc進行編譯，編譯流程
- 分離host code和device code
- 編譯device code為PTX彙編格式或者cubin二進位制格式
- 使用編譯好的kernel替換host code中的kernel函式（以<<<…>>>語法標記的）
- 修改後的host code使用其他工具編譯
- nvcc編譯後的程式可以連結到編譯好的host code或者直接通過CUDA driver API進行載入執行

1. just-in-time compilation
   執行時載入的PTX程式，使用device driver進一步編譯成二進位制程式，稱為just-in-time compilation.
   這種編譯方式增加了載入時間，但是可以使用新的device driver帶來的特性。

-code指定二進位制程式執行的目標裝置的架構，比如-code=sm_35產生的二進位制程式執行在compute capability 3.5
-arch指定的架構針對c程式碼到PTX彙編的編譯過程
支援部分c/c++的語法格式
64-bit模式編譯的device code只能被64-bit模式編譯的host code支援；32-bit的同理。
32-bit的nvcc可以使用-m64選項編譯64-bit的device code；64-bit的nvcc可以使用-m32選項編譯32-bit的device code。

cuda c runtime

通過cudart庫檔案實現: libcudart.a或者libcudart.so。

initialization

第一次呼叫runtime函式的時候執行初始化，為device建立cuda context，所有的host執行緒共享。
cudaDeviceReset()函式銷燬context。

device memory

分為linear memory和CUDA arrays兩種。
linear memory是device上的40-bit的地址空間。
- cudaMalloc/cudaFree/cudaMemcpy
- cudaMallocPitch/cudaMalloc3D/cudaMemcpy2D/cudaMemcpy3D: 分配或者複製2D/3D陣列，滿足對齊要求/高效，儘量使用

shared memory

使用__shared__標示，shared memory比global memory快，儘量使用shared memory。

page-locked host memory

優點：
- page-locked host memory和device memory之間記憶體拷貝可以和kernel執行同時進行
- 可以map到device的地址空間
- 具有front-side bus的系統，page-locked host memory和device memory有更高的頻寬

非同步並行

以下操作相互獨立，可以同步執行：
- host計算
- device計算
- host和device回見的記憶體轉移
- device內部和device之間的記憶體轉移

多裝置系統

cudaGetDeviceCount: 獲取裝置數
cudaSetDevice: 設定使用裝置

統一虛擬地址空間

對於64-bit的程序，host和compute capability 2.0及更高版本的device使用一個地址空間

程序間通訊

使用Inter Process Communication API完成

version

兩個版本需要關注：compute capability有關硬體裝置的特性；CUDA driver api有關driver API和runtime。
driver API的版本使用CUDA_VERSION定義，支援向後相容，也就是低版本的程式可以在高版本的driver上面執行。

compute modes

可以使用nvidia-smi設定為一下幾種模式：
- default: 多個host執行緒可以同時使用device
- exclusive-process: 只能在device建立一個CUDA context，建立context的程序中的多個執行緒可以使用
- exclusive-process-and-thread: 只允許建立一個context，而且context一次只能被一個執行緒使用
- prohibit: 禁止建立CUDA context

c語言擴充套件

函式型別：

• __device__: 在device上執行，只能被device呼叫
• __global__: 標示kernel函式，在device上執行，可以被host/device呼叫
• __host__: 在host執行，只能被host呼叫
• __noinline__: 函式不inline
• __forceinline__: 函式inline

變數型別：

• __device__: device上的變數
• __constant__: constant memory space上的變數
• __shared__: shared memory，所有以這種形式宣告的指標具有相同的起點
• __managed__: host和device都可以讀寫的地址
• __restrict__: 避免aliasing問題

內建vector

make_<type name>

vector是結構體，可以通過x/y/z/w來獲取第1/2/3/4個元素。
dim3用來表示維度，等同於uint3。

內建變數

gridDim: grid維度
blockIdx: grid中的block index
blockDim: block維度
threadIdx: block中thread的index
warpSize: 執行緒中的warp size

效能優化

並行

並行庫，比如cuBLAS/cuFFT
並行編譯器，比如使用progma
並行程式碼

debug

timing

1. CPU Timer
   注意CUDA API函式很多是非同步的，在使用計時函式前需要呼叫cudaDeviceSynchronize()函式以同步CPU執行緒和GPU。

2. GPU Timer
   使用event計時

cudaEvent_t start, stop;
float time;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord( start, 0 );
kernel<<<grid,threads>>> ( d_odata, d_idata, size_x, size_y,
NUM_REPS);
cudaEventRecord( stop, 0 );
cudaEventSynchronize( stop );
cudaEventElapsedTime( &time, start, stop );
cudaEventDestroy( start );
cudaEventDestroy( stop );

參考