CUDA编程：概述
CUDA編程
- GPU的架構
- 每個GPU由N個SM組成，1個SM分為2個SMP，1個SMP有16個DPUnit和32個CudaCore以及一些特殊函數處理模塊
- 比如說， RTX 2080Ti，具有68個SM，總共有 68 x 64 個CudaCore
- SM全稱（Streaming Multiprocessor）
- WarpSize一般為32（這個大部分都這個數）
- 最小調度單位，也就是最小調度單位是一個SMP
- 如果啟動線程數小于32，他依舊啟動32線程，僅僅是其他線程處于非激活狀態deactive
- 內存，分為GlobalMemory、SharedMemory、緩存、寄存器等
- 理解內存的排布，是優化編程的基礎
- 距離Core越近的內存，速度越快
- 架構名稱
- 不同的架構，具有不同的顯卡硬件設置，比如說游戲級更注重渲染，工業級更注重科學計算或者是長期運行（所以散熱等更注重）。
- 比如說，工業級，會增加例如TensorCore、DLA、FP16、INT8、INT4等NN的支持。而這些，游戲級別大部分沒有
- 架構一般有：圖靈(Turing、游戲級)、帕斯卡(PASCAL，游戲級)、開普勒(Kepler，游戲級)、麥克斯維爾(Maxwell，游戲級)、VOLTA(工業級)
- 編程上
- 線程的抽象形式，是使用gridDim和blockDim指定你需要啟動的線程數
- gridDim是用x、y、z三個維度表示，其取值具有最大上限，例如RTX2080Ti上是（deviceQuery可以查到）：2147483647, 65535, 65535
- blockDim是x、y、z三個維度表示，例如RTX2080Ti上是：1024, 1024, 64
- 因此對于線程數量的計算，是可以認為是6個維度的張量來表示。線程數，則是6個維度乘積
- 對于核函數中，position的計算，使用blockIdx和threadIdx做索引（共6個索引），得到當前線程在線程數中的絕對id
- 計算形式是
- 假設shape是 A x B x C x D x E x F
- 假設index是 a x b x c x d x e x f
- position = ((((a x B + b) x C + c) x D + d) x E + e) x F + f
- 對于grid和block的例子，其shape為：
- gridDim.z
- gridDim.y
- gridDim.x
- blockDim.z
- blockDim.y
- blockDim.x
- 對于grid和block的例子，其index為：
- blockIdx.z
- blockIdx.y
- blockIdx.x
- threadIdx.z
- threadIdx.y
- threadIdx.x
- 對于常規使用時，通常只用到2個維度
- gridDim = dim3(blocks)
- dim3的構造函數是dim3(int x, int y=1, int z=1)
- 所以gridDim.x = blocks，gridDim.y = 1, gridDim.z = 1
- blockDim = dim3(threads)
- dim3的構造函數是dim3(int x, int y=1, int z=1)
- 所以blockDim.x = threads，blockDim.y = 1, blockDim.z = 1
- 計算position時：
- 其shape為：
- gridDim.z = 1
- gridDim.y = 1
- gridDim.x = blocks
- blockDim.z = 1
- blockDim.y = 1
- blockDim.x = threads
- 其index為：
- blockIdx.z = 0
- blockIdx.y = 0
- blockIdx.x = u
- threadIdx.z = 0
- threadIdx.y = 0
- threadIdx.x = v
- position = u * threads + v = blockIdx.x * blockDim.x + threadIdx.x
- 其shape為：
- gridDim = dim3(blocks)
- threads的數量，要求會給warp_size的倍數
- 通常二維時，可以固定為512，如果顯卡比較差，可以給小。視情況可以給大
- blocks = ceil(jobs / (float)threads)
- 計算形式是
- 修飾符
- __global__ 修飾，指host可以調用的函數，通常咱們認為是cuda核函數，通過kernel<<<gridDim, blockDim, memory, stream>>>(args)啟動
- 傳值和傳地址
- 傳值不需要考慮復制到設備，即時你的值是一坨結構體（struct）
- 傳地址，如果需要在核函數中訪問地址數據，則必須把地址指向的數據復制到顯卡上。否則操作異常
- __device__ 修飾，指device調用的函數
- __host__ 修飾，主機函數，這個默認寫的函數就是這個
- __global__ 修飾，指host可以調用的函數，通常咱們認為是cuda核函數，通過kernel<<<gridDim, blockDim, memory, stream>>>(args)啟動
- static，是C語言語法，編譯的符號，只在當前cpp中有效。不會參與全局的鏈接。如果不加static，默認是全局都會進行鏈接
- 內存布局
- opencv加載的圖像，顏色空間格式是RGB的，內存排布是BGRBGRBGR。如果按照維度看，則是HWC
- 平時訓練，可以使用BGR或者RGB，這個沒有要求
- 顏色空間格式，有很多種：RGB、YUV、HSV、HSL
- cuda流，cudaStream_t
- 實現cuda核以及cuda操作pipline的一個任務隊列，可以理解為一個線程+任務隊列
- 如果使用nullptr，則是默認流
- cudaMemcpy，是同步復制，他是由cudaMemcpyAsync(nullptr)和cudaDeviceSynchronize()實現
- cudaMemcpyAsync，是異步復制
- 實際時候用時，避免采用同步，盡可能全部異步。這樣能夠極大的提高顯卡使用率
- cudaError_t cudaMemcpy(dst, src, size, kind)
- dst，是目標指針
- src，是來源指針
- size，是復制的大小
- kind，是復制的方式，cudaMemcpyHostToDevice、DeviceToHost、HostToHost、DeviceToDevice
- cudaError_t cudaMalloc(void** ptr, size_t size)
- ptr，是主機的二級指針，分配后的 地址，通過二級指針修改
- 你想要的是，分配后的地址
- 你還想要的是，如果出錯了，錯誤代碼給我
- cuda規定了，函數的定義形式，返回值是錯誤
- void* ptr = nullptr;
- 在棧上開辟了8個字節空間，儲存了nullptr
- &ptr = 取ptr在棧的位置，設為P
- cudaMalloc(&ptr, 100); 傳給 cudaMalloc的是，ptr在棧上的位置P
- cudaMalloc(void** ptr2, size_t size)
- ptr2 = ptr在棧上的位置，P
- ptr2 = 引用棧的P位置上，8個字節。解釋為void 。此時*ptr2 值是 nullpter
- G = malloc(size) 分配size個大小空間，并得到地址，這個地址假設是GPU的地址G
- ptr2 = G 把G復制到ptr2中。*ptr2又是個引用 P位置的8個字節
- G也是8個字節，所以直接G復制到P上8個字節位置做填充
- ptr就被修改了
- cudaMalloc(void** ptr2, size_t size)
- size，是分配的大小
- ptr，是主機的二級指針，分配后的 地址，通過二級指針修改
- cudaEvent_t
- 加入事件到流中，根據需要進行等待，或者統計時間
- 原子操作，atomicAdd
- 實現globalMemory或者sharedMemory上的內存加指定值，并返回舊值
- 用來解決多線程異步操作同一個變量的計數和索引問題
- yolov5的后處理用到
- 內存釋放，cudaFree
- 線程的抽象形式，是使用gridDim和blockDim指定你需要啟動的線程數
總結
- 上一篇： ABIs [armeabi] are n
- 下一篇： 一个使用指针的简单程序