CUDA精講（1）

背景

隨著深度學習的發展，平行計算的需求也越來越多，不論是演算法工程師還是搞效能優化的，知道點cuda程式設計也是當前的必備技能之一。但是目前關於gpu資源很雜，重點不突出，我們需要在極短的時間內get到cuda的精華，本系列就由此而誕生。

文章框架

cuda精講系列文章主要由三部分構成：

1. cuda基礎
2. GPU架構
3. cuda程式設計優化
   每一部分都是通過知識點的方式將重點提煉出來，方便快速檢視。

cuda基礎

1. cuda的基本概念
   host: CPU
   device: gpu
   從軟體層面講：
   （1）kernel：就是開發者開發的一個在gpu上要執行的函式程式碼
   （2）thread：用來執行一個kernel的執行緒
   （3）block：多個thread的組合
   （4）grid：多個block的組合
   （5）warp：block以一個warp為排程單位進行執行緒排程。因此block中的執行緒數儘量是32的整數倍。每32個連續執行緒號的執行緒被安排在同一個wrap裡
   從硬體層面講：
   （1）SP（core）：一個物理上的SP對應一個邏輯上的thread
   （2）SM（multicore processor）：多個sp的組成，處於SM內的sp可以共享shared_mem，並且同一個block裡的執行緒可以被同步。一個物理上的SM可以對應一個或多個邏輯上的block
   （3）Device：就是一個GPU裝置，一個GPU包括了多個SM。
   
2. cuda程式設計框架

__global__ void my_kernel(...){
    ...
}

int main() {
	...
	cudaMalloc(...)
	// host to device
	cudaMemcpy(...)
	...
	my_kernel<<<nblock, blocksize>>>(...)
	...
	// device to host
	cudaMemcpy(...)
	...
}

4. cuda程式設計常用語法：

• kernel launch：就是gpu上執行的kernel函式，因為執行kernel通常要從cpu登入到gpu因此稱為launch，kernel函式語法為：

kernelFunc<<<nB, nT>>>

nB:block的數量
nT：thread的數量

• 內建變數
  threadIdx;blockIdx;blockDim;gridDim
• 同步語句
  __syncthreads()
• 生存週期生命：
  global void kernelFunc(…)，執行在device中，在cpu中呼叫
  device void GlobalVar; 裝置變數
  shared void SharedVar; 定義每個block的shared memory中定義的變數

5. cuda函式的非同步性
   cuda中的大部分函式其實都是非同步的，比如：

• kernel launches
• 記憶體拷貝中可以指定option：cudaMemcpyAsync，cudaMemsetAsync
• cudaEvent function

6. CPU與GPU的同步語句:

• Device based: cudaDeviceSynchronize(),在某一個裝置GPUkernel執行結束後在執行cpu語句。
• Context（Thread） based: cudaThreadSynchronize()，所有裝置的所有kernel都執行完畢後在執行。
• Stream based: cudaStreamSynchronize(stream-id)
• Event based：某個裝置的kernel事件執行完畢後cpu語句再執行。

7. 多GPU程式設計，後續補上

• 單機多卡：利用pthread、openMP
• 多機多卡：MPI，計算交給GPU，具體的通訊細節交給MPI

8. GPU之間的資料共享
   （1）直接拷貝 host to device
   

（2）零拷貝：將device的記憶體直接對映到host的記憶體當中，但是在host中被共享的資料需要被page-locked（pinned）：由此擴充套件的hostAPI有,

• cudaMallocHost()，分配pinned的host mem

（3）p2p記憶體拷貝：
直接拷貝資料從GPUA到GPUB當中，API為：cudaMemcpyPeer(void dst, int
dstDevice, const void src, int srcDevice,size_t count)

9. 動態並行（Dynamic parallelism, dp）
   避免從cpu到gpu的多次kernel登入，可以實現從gpu到gpu之間進行kernel登入
   
   通過dp可實現動態的blocksize與gridsize分配：