模型部署——cuda編程入門

CUDA中的線程與線程束

在這里插入圖片描述

kernel是在device上線程中并行執行的函數，核函數用__global__符號聲明，在調用時需要用<<<grid_size, block_size>>>來指定kernel要執行的線程數量。在CUDA中，每一個線程都要執行核函數，并且每個線程會分配一個唯一的線程號thread ID，這個ID值可以通過核函數的內置變量threadIdx來獲得。
synchronize是同步的意思，有幾種synchronize
cudaDeviceSynchronize: CPU與GPU端完成同步，CPU不執行之后的語句，直到這個語句以前的所有cuda操作結束
cudaStreamSynchronize: 跟cudaDeviceSynchronize很像，但是這個是針對某一個stream的。只同步指定的stream中的cpu/gpu操作，其他的不管
cudaThreadSynchronize: 現在已經不被推薦使用的方法
__syncthreads: 線程塊內同步
核函數編寫和調用舉例

#include <cuda_runtime.h>
#include <stdio.h>// 核函數
__global__ void print_idx_kernel(){printf("block idx: (%3d, %3d, %3d), thread idx: (%3d, %3d, %3d)\n",blockIdx.z, blockIdx.y, blockIdx.x,threadIdx.z, threadIdx.y, threadIdx.x);
}void print_one_dim(){int inputSize = 8;int blockDim = 4;int gridDim = inputSize / blockDim;dim3 block(blockDim);dim3 grid(gridDim);// 核函數調用print_idx_kernel<<<grid, block>>>();cudaDeviceSynchronize();
}

.cu與.cpp的相互引用及Makefile

編譯器：gcc g++ nvcc
舉個例子：

nvcc print_index.cu -o app -I /usr/local/cuda/include

獲取編譯器選項：

g++ --help
nvcc --help

Makefile編寫（是否可以使用CMakeLists.txt?）
.cpp中不能直接調用核函數，需要在.cu中提供調用接口

使用CUDA進行MATMUL計算

host端與device端數據傳輸

在這里插入圖片描述
host端與device端數據傳輸代碼實現：

void MatmulOnDevice(float *M_host, float *N_host, float* P_host, int width, int blockSize){/* 設置矩陣大小 */int size = width * width * sizeof(float);/* 分配M, N在GPU上的空間*/float *M_device;float *N_device;cudaMalloc(&M_device, size);cudaMalloc(&N_device, size);/* 分配M, N拷貝到GPU上*/cudaMemcpy(M_device, M_host, size, cudaMemcpyHostToDevice);cudaMemcpy(N_device, N_host, size, cudaMemcpyHostToDevice);/* 分配P在GPU上的空間*/float *P_device;cudaMalloc(&P_device, size);/* 調用kernel來進行matmul計算, 在這個例子中我們用的方案是：將一個矩陣切分成多個blockSize * blockSize的大小 */dim3 dimBlock(blockSize, blockSize);dim3 dimGrid(width / blockSize, width / blockSize);MatmulKernel <<<dimGrid, dimBlock>>> (M_device, N_device, P_device, width);/* 將結果從device拷貝回host*/cudaMemcpy(P_host, P_device, size, cudaMemcpyDeviceToHost);cudaDeviceSynchronize();/* Free */// free與malloc的順序是反著的cudaFree(P_device);cudaFree(N_device);cudaFree(M_device);
}

cuda core矩陣乘法核函數編寫

/* matmul的函數實現*/
__global__ void MatmulKernel(float *M_device, float *N_device, float *P_device, int width){/* 我們設定每一個thread負責P中的一個坐標的matmul所以一共有width * width個thread并行處理P的計算*/// 確定負責計算的結果元素的索引int y = blockIdx.y * blockDim.y + threadIdx.y;int x = blockIdx.x * blockDim.x + threadIdx.x;float P_element = 0;/* 對于每一個P的元素，我們只需要循環遍歷width次M和N中的元素就可以了*/for (int k = 0; k < width; k ++){float M_element = M_device[y * width + k];float N_element = N_device[k * width + x];P_element += M_element * N_element;}P_device[y * width + x] = P_element;
}

cuda core 每個線程執行核函數計算一個結果元素
GPU剛開始執行核函數的時候，會存在一個warmup階段，耗時會比較長
CPU與GPU的浮點運算會存在誤差，誤差控制在e-4以內是ok的
CUDA中規定，一個block中可以分配的thread的數量最大是1024個線程，如果大于1024會顯示配置錯誤
為什么block size = 1的時候比等于16的時候慢很多？

cuda中的error handler

在這里插入圖片描述

獲取GPU的硬件信息

利用cuda runtime api打印硬件信息 & LOG

#include <stdio.h>
#include <cuda_runtime.h>
#include <string>#include "utils.hpp"int main(){int count;int index = 0;cudaGetDeviceCount(&count);while (index < count) {cudaSetDevice(index);cudaDeviceProp prop;cudaGetDeviceProperties(&prop, index);LOG("%-40s",             "*********************Architecture related**********************");LOG("%-40s%d%s",         "Device id: ",                   index, "");LOG("%-40s%s%s",         "Device name: ",                 prop.name, "");LOG("%-40s%.1f%s",       "Device compute capability: ",   prop.major + (float)prop.minor / 10, "");LOG("%-40s%.2f%s",       "GPU global meory size: ",       (float)prop.totalGlobalMem / (1<<30), "GB");LOG("%-40s%.2f%s",       "L2 cache size: ",               (float)prop.l2CacheSize / (1<<20), "MB");LOG("%-40s%.2f%s",       "Shared memory per block: ",     (float)prop.sharedMemPerBlock / (1<<10), "KB");LOG("%-40s%.2f%s",       "Shared memory per SM: ",        (float)prop.sharedMemPerMultiprocessor / (1<<10), "KB");LOG("%-40s%.2f%s",       "Device clock rate: ",           prop.clockRate*1E-6, "GHz");LOG("%-40s%.2f%s",       "Device memory clock rate: ",    prop.memoryClockRate*1E-6, "Ghz");LOG("%-40s%d%s",         "Number of SM: ",                prop.multiProcessorCount, "");LOG("%-40s%d%s",         "Warp size: ",                   prop.warpSize, "");LOG("%-40s",             "*********************Parameter related************************");LOG("%-40s%d%s",         "Max block numbers: ",           prop.maxBlocksPerMultiProcessor, "");LOG("%-40s%d%s",         "Max threads per block: ",       prop.maxThreadsPerBlock, "");LOG("%-40s%d:%d:%d%s",   "Max block dimension size:",     prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2], "");LOG("%-40s%d:%d:%d%s",   "Max grid dimension size: ",     prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2], "");index ++;printf("\n");}return 0;
}

Roofline model（待補充）

Nsight system and Nsight compute

谷歌搜索下載:官網鏈接

Nsight system

參考鏈接

安裝目錄:
ls /usr/local/bin |grep nsys
nsys
nsys-ui啟動GUI界面
sudo ./nsys-ui(不加sudo會存在權限問題）

舉個例子：
配置可執行文件以及感興趣內容：
在這里插入圖片描述可視化分析：
詳細使用手冊：官網文檔

Nsight compute

查看可安裝版本：
sudo apt policy nsight-compute-2022.2.1
安裝：
sudo apt install nsight-compute-2022.2.1
查看安裝位置：
dpkg -L nsight-compute-2022.2.1
路徑：/opt/nvidia/nsight-compute/2022.2.1/
文件：ncu ncu-ui等啟動：
sudo ./ncu-ui

舉個例子：
基本配置：replay mode: application
在這里插入圖片描述選擇感興趣內容：
launch即可，第一次運行會比較慢，會重復運行很多次。
結果：
不知道為什么roofline model沒有正常顯示出來，需要查一查？

擴展知識

共享內存以及BANK CONFLICT

shared memory

硬件結構

SM（Streaming Multiprocessor）
在CUDA編程模型中，線程被組織成線程塊（block），多個線程塊組成一個網格（grid）。每個線程塊被分配到一個SM中執行，而SM內部的warp調度器會將線程塊中的線程分成多個warp進行執行。
當一個warp中的線程需要等待某些操作（例如內存訪問）完成時，SM可以切換到另一個warp繼續執行，從而提高計算效率。
在這里插入圖片描述

核函數編寫

#include "cuda_runtime_api.h"
#include "utils.hpp"#define BLOCKSIZE 16/* 使用shared memory把計算一個tile所需要的數據分塊存儲到訪問速度快的memory中
*/
__global__ void MatmulSharedStaticKernel(float *M_device, float *N_device, float *P_device, int width){__shared__ float M_deviceShared[BLOCKSIZE][BLOCKSIZE];__shared__ float N_deviceShared[BLOCKSIZE][BLOCKSIZE];/* 對于x和y, 根據blockID, tile大小和threadID進行索引*/int x = blockIdx.x * blockDim.x + threadIdx.x;int y = blockIdx.y * blockDim.y + threadIdx.y;float P_element = 0.0;int ty = threadIdx.y;int tx = threadIdx.x;/* 對于每一個P的元素，我們只需要循環遍歷width / tile_width 次就okay了，這里有點繞，畫圖理解一下*/for (int m = 0; m < width / BLOCKSIZE; m ++) {M_deviceShared[ty][tx] = M_device[y * width + (m * BLOCKSIZE + tx)];N_deviceShared[ty][tx] = N_device[(m * BLOCKSIZE + ty)* width + x];__syncthreads(); // 上述兩句所有thread都會執行，等待所有thread執行完成for (int k = 0; k < BLOCKSIZE; k ++) {P_element += M_deviceShared[ty][k] * N_deviceShared[k][tx];}__syncthreads();}P_device[y * width + x] = P_element;
}__global__ void MatmulSharedDynamicKernel(float *M_device, float *N_device, float *P_device, int width, int blockSize){/* 聲明動態共享變量的時候需要加extern，同時需要是一維的 注意這里有個坑, 不能夠像這樣定義： __shared__ float M_deviceShared[];__shared__ float N_deviceShared[];因為在cuda中定義動態共享變量的話，無論定義多少個他們的地址都是一樣的。所以如果想要像上面這樣使用的話，需要用兩個指針分別指向shared memory的不同位置才行*/extern __shared__ float deviceShared[];int stride = blockSize * blockSize;/* 對于x和y, 根據blockID, tile大小和threadID進行索引*/int x = blockIdx.x * blockSize + threadIdx.x;int y = blockIdx.y * blockSize + threadIdx.y;float P_element = 0.0;int ty = threadIdx.y;int tx = threadIdx.x;/* 對于每一個P的元素，我們只需要循環遍歷width / tile_width 次就okay了 */for (int m = 0; m < width / blockSize; m ++) {deviceShared[ty * blockSize + tx] = M_device[y * width + (m * blockSize + tx)];deviceShared[stride + (ty * blockSize + tx)] = N_device[(m * blockSize + ty)* width + x];__syncthreads();for (int k = 0; k < blockSize; k ++) {P_element += deviceShared[ty * blockSize + k] * deviceShared[stride + (k * blockSize + tx)];}__syncthreads();}if (y < width && x < width) {P_device[y * width + x] = P_element;}
}

動態共享內存比靜態共享內存速度慢，沒有特殊情況下，使用靜態共享內存。

cuda event進行時間測算

在這里插入圖片描述

BANK CONFLICT（存儲體沖突）

在shared memory中什么是bank?

在這里插入圖片描述

什么時候會發生bank conflict

在這里插入圖片描述

按行存儲，按列訪問的時候，會發生bank conflict:
在這里插入圖片描述

如何減緩bank conflict

在這里插入圖片描述

代碼實現

#include "cuda_runtime_api.h"
#include "utils.hpp"#define BLOCKSIZE 16/* 使用shared memory把計算一個tile所需要的數據分塊存儲到訪問速度快的memory中
*/
__global__ void MatmulSharedStaticConflictPadKernel(float *M_device, float *N_device, float *P_device, int width){/* 添加一個padding，可以防止bank conflict發生，結合圖理解一下*/__shared__ float M_deviceShared[BLOCKSIZE][BLOCKSIZE + 1];__shared__ float N_deviceShared[BLOCKSIZE][BLOCKSIZE + 1];/* 對于x和y, 根據blockID, tile大小和threadID進行索引*/int x = blockIdx.x * BLOCKSIZE + threadIdx.x;int y = blockIdx.y * BLOCKSIZE + threadIdx.y;float P_element = 0.0;int ty = threadIdx.y;int tx = threadIdx.x;/* 對于每一個P的元素，我們只需要循環遍歷width / tile_width 次就okay了，這里有點繞，畫圖理解一下*/for (int m = 0; m < width / BLOCKSIZE; m ++) {/* 這里為了實現bank conflict, 把tx與tx的順序顛倒，同時索引也改變了*/M_deviceShared[tx][ty] = M_device[x * width + (m * BLOCKSIZE + ty)];N_deviceShared[tx][ty] = M_device[(m * BLOCKSIZE + tx)* width + y];__syncthreads();for (int k = 0; k < BLOCKSIZE; k ++) {P_element += M_deviceShared[tx][k] * N_deviceShared[k][ty];}__syncthreads();}/* 列優先 */P_device[x * width + y] = P_element;
}__global__ void MatmulSharedDynamicConflictPadKernel(float *M_device, float *N_device, float *P_device, int width, int blockSize){/* 聲明動態共享變量的時候需要加extern，同時需要是一維的 注意這里有個坑, 不能夠像這樣定義： __shared__ float M_deviceShared[];__shared__ float N_deviceShared[];因為在cuda中定義動態共享變量的話，無論定義多少個他們的地址都是一樣的。所以如果想要像上面這樣使用的話，需要用兩個指針分別指向shared memory的不同位置才行*/extern __shared__ float deviceShared[];int stride = (blockSize + 1) * blockSize;/* 對于x和y, 根據blockID, tile大小和threadID進行索引*/int x = blockIdx.x * blockSize + threadIdx.x;int y = blockIdx.y * blockSize + threadIdx.y;float P_element = 0.0;int ty = threadIdx.y;int tx = threadIdx.x;/* 對于每一個P的元素，我們只需要循環遍歷width / tile_width 次就okay了 */for (int m = 0; m < width / blockSize; m ++) {/* 這里為了實現bank conflict, 把tx與tx的順序顛倒，同時索引也改變了*/deviceShared[tx * (blockSize + 1) + ty]             = M_device[x * width + (m * blockSize + ty)];deviceShared[stride + (tx * (blockSize + 1) + ty)]  = N_device[(m * blockSize + tx) * width + y];__syncthreads();for (int k = 0; k < blockSize; k ++) {P_element += deviceShared[tx * (blockSize + 1) + k] * deviceShared[stride + (k * (blockSize + 1 ) + ty)];}__syncthreads();}/* 列優先 */P_device[x * width + y] = P_element;
}

STREAM和EVENT

什么是stream

在這里插入圖片描述

參考下述鏈接，理解cuda編程的一些基礎概念：
理解CUDA中的thread,block,grid和warp
cuda stream的使用

多流編程實現

單流：
在這里插入圖片描述
多流：

利用nsight systems進行分析：

如何利用多流進行隱藏訪存和核函數執行延遲的調度

舉一個栗子：
在這里插入圖片描述

使用CUDA進行預處理/后處理

雙線性插值

在這里插入圖片描述

雙線性插值的cuda實現

在這里插入圖片描述

查看圖片大小：
identity xx.png
可視化圖片：
feh xx.png