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
- 下一篇: 一个使用指针的简单程序