精簡CUDA教程——CUDA Runtime API

tensorRT從零起步邁向高性能工業級部署（就業導向）課程筆記，講師講的不錯，可以去看原視頻支持下。

Runtime API 概述

環境

在這里插入圖片描述

圖中可以看到，Runtime API 是基于 Driver API 之上開發的一套 API。
之前提到過 Driver API 基本都是 cu 開頭的，而Runtime API 基本都是以 cuda 開頭的。

Runtime API 的特點

Runtime API 與 Driver API 最大的區別是懶加載，即在真正執行功能時才自動完成對應的動作，即：
- 第一個 Runtime API 調用時，會自動進行 cuInit 初始化，避免 Driver API 未初始化的錯誤；
- 第一個需要 context 的 API 調用時，會創建 context 并進行 context 關聯，和設置當前 context，調用 cuDevicePrimaryCtxRetain 實現；
- 絕大部分 api 都需要 context，例如查詢當前顯卡名稱、參數、內存分配釋放等
CUDA Runtime 是封裝了 CUDA Driver 的更高級別、更友好的 API
Runtime API 使用 cuDevicePrimaryCtxRetain 為每個設備設置 context，不再手動管理 context，并且不提供直接管理 context 的 API（可 Driver API 管理，通常不需要）
可以更友好地執行核函數，.cpp 可以與 .cu 文件無縫對接
Runtime API 對應 cuda_runtime.h 和 libcudart.so
Runtime API 隨 cudatoolkit 發布
主要知識點是核函數的使用、線程束布局、內存模型、流的使用
主要是為了實現歸約求和、放射變換、矩陣乘法、模型后處理，就可以解決絕大部分問題

錯誤處理

類似于在介紹 Driver API 時的情況，我們同樣提出 Runtime API 的錯誤處理方式：

#define checkRuntime(op)  __check_cuda_runtime((op), #op, __FILE__, __LINE__)bool __check_cuda_runtime(cudaError_t code, const char* op, const char* file, int line){if(code != cudaSuccess){const char* err_name = cudaGetErrorName(code);const char* err_message = cudaGetErrorString(code);printf("runtime error %s:%d  %s failed. \n  code = %s, message = %s\n", file, line, op, err_name, err_message);return false;}return true;
}

內存模型 pinned memory

內存模型是 CUDA 中很重要的知識點，主要理解 pinned_memory、global_memory、shared_memory 即可，其他的不太常用。
pinned_memory 屬于 host memory，而 global_memory、shared_memory 屬于 device memory。

下圖是的 Device Memory 的分類

在這里插入圖片描述

鎖定性和性能

對于主機內存，即整個 host memory 而言，操作系統在邏輯上將其區分為兩個大類：

pageable memory，可分頁內存
page lock memory (pinned memory)，頁鎖定內存/鎖頁內存

可以理解為 page lock memory 是酒店的 vip 房間，鎖定給你一個人使用。而 pageable memory 是普通房間，在酒店房間不夠的時候，選擇性地將你的房間騰出來（交換到硬盤上）給其他人使用，這樣就能容納更多人了。造成房間很多的假象，代價是性能很低。pageable memory 就是常見的虛擬內存的特性。

基于前面的理解，我們總結如下：

鎖定性
- pinned memory 具有鎖定特性，是穩定不會被交換的，這很重要，相當于每次去這個房間都一定能找到你
- pageable memory 沒有鎖定特性，對于第三方設備（如 GPU）去訪問時，因為無法感知內存是否被交換，可能得到不到正確的數據，相當于每次去房間找你，說不定你的房間正好被交換了
- 因此， GPU 可以直接訪問 pinned memory 而不能訪問 pageable memory
性能
- pageable memory 的性能比 pinned memory 差，因為我們的 pageable memory 很可能會被交換到硬盤上
- pageable memory 策略能使用內存假象，比如實際只有 8G 內存卻能使用 16G（借助 swap 交換），從而提高程序的運行數量
- pinned memory 也不能太多，會導致操作系統整體性能變差（可同時運行的程序變少），而且 8G 內存最多就 8G 鎖頁內存。

數據傳輸到GPU

在這里插入圖片描述

pinned memory 可以直接傳送數據到 GPU
而 pageable memory ，由于并不鎖定，需要先傳到 pinned memory。

關于內存其他幾個點

GPU 可以直接訪問 pinned memory，稱為 DMA (Direct Memort Access)
對于 GPU 訪問而言，距離計算單元越近，效率越高，所以：

SharedMemory > GlobalMemory > PinnedMemory
代碼中，
- 由 new/malloc 分配的是 pageable memory
- 由 cudaMallocHost 分配的是 PinnedMemory
- 由 cudaMalloc 分配的是 GlobalMemory
盡量多用 PinnedMemory 儲存 host 數據，或者顯式處理 Host 到 Device 時，用 PinnedMemory 做緩存，都是提高性能的關鍵

流 stream

流是一種基于 context 之上的任務管道（任務隊列）抽象，一個 context 可以創建 n 個流
流是異步控制的主要方式
nullptr 表示默認流，每個線程都有自己的默認流。

生活中的例子

同步（串行）	異步

在這個例子中，男朋友的微信消息，就是任務隊列，流的一種抽象
女朋友發出指令之后，她可以做任何事情，無需等待指令執行完畢。即異步操作中，執行的代碼加入流的隊列之后，立即返回，不耽誤時間。
女朋友發的指令被送到流中排隊，男朋友根據流的隊列，順序執行。
女朋友選擇性，在需要的時候等待所有的執行結果
新建一個流，就是新建一個男朋友，給他發指令就是發微信，可以新建很多個男朋友
通過 cudaEvent 可以選擇性等待任務隊列中的部分任務是否就緒

注意

要十分注意，指令發出后，流隊列中儲存的是指令參數，不能在任務加入隊列后立即釋放參數指針，這會導致流隊列執行該指令時指針失效而出錯。應當在十分肯定流已經不需要這個指針之后，才進行修改或釋放，否則會有非預期行為出現。

就比如，女朋友讓男朋友去賣西瓜并轉給了他錢，但是卻在男朋友買瓜成功前將轉賬撤了回去，這時就無法知道男朋友在水果店會發生什么，比如會不會跟老板打起來之類的。因此，要保證買瓜行為順利完成（行為符合預期），在買瓜成功前就不能動買瓜的錢。

核函數

簡介

核函數是 cuda 編程的關鍵
通過 xxx.cu 創建一個 cudac 程序文件，并把 cu 文件交給 nvcc 編譯，才能識別 cuda 語法；
__xxx__ 修飾
- __global__ 表示為核函數，由 host 調用；
- __device__ 表示設備函數，由 device 調用；
- __host__ 表示主機函數，由 host 調用；
- __shared__ 表示變量為共享變量。
- 可能存在上述多個關鍵字修飾同一個函數，如 __device__ 和 __host__ 修飾的函數，既可以設備上調用，也可以在主機上調用
host 調用核函數：
```
function<<<gridDim, blockDim, sharedMemorySize, stream>>>(args, ...)
```
gridDim 和 blockDim 的變量類型為 dim3，是一個三維的值；

function 函數總共啟動的線程數目可以這樣計算：n_threads = gridDim.x * gridDim.y * gridDim.z * blockDim.x * blockDim.y * blockDim.z

詳細請參考線程束的相關知識
只有 __global__ 修飾的函數才可以用 <<< >>> 的方式調用s
調用核函數是傳值的，不能傳引用，可以傳遞類，結構體等，核函數可以使模板
核函數的返回值必須是 void
核函數的執行是異步的，也就是立即返回的
線程 layout 主要用到 blockDim、gridDim
和函數內訪問線程索引主要用到 threadIdx、blockIdx、blockDim、gridDim 這些內置變量

線程索引計算

共涉及四個變量：blockDim、gridDim、threadIdx、blockIdx ，其中前兩者可以認為是形狀，后兩者可以認為是對應的索引。就像我們 PyTorch 中如果一個張量的形狀為 $(2, 3)$ ，那么對應的，其兩個維度上索引的取值范圍就是： $0 ? 1, 0 ? 2$ 。

在這里插入圖片描述

線程索引 id 計算方法：左乘右加，如上圖所示。

共享內存

由 __shared__ 關鍵字修飾
共享內存因為更靠近計算單元，所以訪問速度更快
共享內存通常可以作為訪問全局內存的緩存使用
可以利用共享內存實現線程間的通信
通常與 __syncthreads 同時出現，這個函數是同步 block 內的所有線程，全部執行到這一行才往下繼續執行

如：
```
__shared__ int shared_value1;
__shared__ int shared_value2;if (threadIdx.x == 0) {if (blockIdx.x == 0) {shared_value1 = 123;shared_value2 = 55;}else {shared_value1 = 331;shared_value2 = 8;}__syncthreads();printf("...")
}
```
其他 threadIdx.x 不為 0 的線程不會進到判斷語句，但是會卡在 __syncthreads() ，等待 threadIdx.x 為 0 的線程設置好共享內存，再一起繼續向下執行。
共享內存使用方式：通常是在線程 id 為 0 的時候從 global memory 取值，然后 __syncthreads ，然后再使用
動態共享內存與靜態共享內存
- 動態共享內存的聲明需要加 extern 關鍵字，不需要指定數組大小，如：
```
extern __shared__ char dynamic_shared_memory[];
```
- 靜態共享內存的聲明需要指定數組大小，如：
```
const size_t static_shared_memory_size = 6 * 1024; // 6KB
__shared__ char static_shared_memory[static_shared_memory_size];
```

warp affine 實戰

chapter: 1.6, caption: vector-add, description: 使用cuda核函數實現向量加法
chapter: 1.7, caption: shared-memory, description: 共享內存的操作
chapter: 1.8, caption: reduce-sum, description: 規約求和的實現，利用共享內存，高性能
chapter: 1.9, caption: atomic, description: 原子操作，實現動態數組的操作
chapter: 1.10, caption: warpaffine, description: 仿射變換雙線性插值的實現，yolov5的預處理
chapter: 1.11, caption: cublas-gemm, description: 通用矩陣乘法的cuda核函數實現，以及cublasSgemm的調用
chapter: 1.12, caption: yolov5-postprocess, description: 使用cuda核函數實現yolov5的后處理案例

TODO

thrust

相當于 cuda 的 stl，但并不常用

錯誤處理

若核函數出錯，由于它是異步的，立即執行 cudaPeekAtLastError 只會拿到對輸入參數校驗是否正確的狀態，而不會拿到核函數是否正確執行的狀態。

需要等待核函數真正執行完畢之后才知道當前核函數是否出錯，一般通過設備同步或者流同步進行等待

錯誤分為可恢復和不可恢復兩種

可恢復
- 參數配置錯誤，例如 block 越界（一般最大值是 1024），shared memory 超出大小范圍（一般是 64KB）等
- 通過 cudaGetlastError 可以獲取錯誤代碼，同時把當前狀態恢復為success
- 該種錯誤可以在調用核函數之后立即通過 cudaGetLastError / cudaPeekAtLastError 拿到
- 該種錯誤在下一個函數調用時會覆蓋
不可恢復
- 核函數執行錯誤，例如訪問越界等
- 該錯誤會傳遞到之后所有的 cuda 操作上
- 錯誤狀態通常需要等到核函數執行完畢才能夠拿到，也就是有可能在后續的任何流程中突然異常（因為是異步的）