cuda 程序的基本步驟如下：

• 在 cpu 中初始化數據
• 將輸入 transfer 到 GPU 中
• 利用分配好的 grid 和 block 啟動 kernel 函數
• 將計算結果 transfer 到 CPU 中
• 釋放申請的內存空間

從上面的步驟可以看出，一個 CUDA 程序主要包含兩部分，第一部分運行在 CPU 上，稱作 Host code，主要負責完成復雜的指令；第二部分運行在 GPU 上，稱作 Device code，主要負責并行地完成大量的簡單指令（如數值計算）；

2. 基本設施

運行在 GPU 中地函數稱作 kernel，該函數有這么幾個要求：

• 聲明時在返回類型前需要添加 "__globol__" 的標識
• 返回值只能是 void

__global__ void addKernel(int *c, const int *a, const int *b) {int i = threadIdx.x;c[i] = a[i] + b[i]; }

這就是一個合規的核函數。

除了聲明時的不同，和函數的調用也是不一樣的，需要以 “kernel_name <<<>>>();” 的形式調用。而在尖括號中間，則是定義了啟用了多少個 GPU 核，學習這一參數的使用，我們還需要知道下面幾個概念：

• dim3：一種數據類型，包含 x,y,z 三個 int 類型的成員，在初始化時一個 dim3 類型的變量時，成員值默認為 1
• grid : 一個 grid 中包含多個 block
• block: 一個 block 包含多個 thread

我們以一種更抽象的方式來理解 GPU 中程序的運行方式的話，可以這么看：

GPU 中的每個核可以獨立的運行一個線程，那我們就使用 thread 來代表 GPU 中的核，但一個 GPU 中的核數量很多，就需要有更高級的結構對全部用到的核進行約束、管理，這就是 block (塊)，一個塊中可以包含多個核，并且這些核在邏輯上的排布可以是三維的，在一個塊中我們可以使用一個 dim3 類型的量 threadIdx 來表示每個核所處的位置，threadIdx.x、threadIdx.y、threadIdx.z 分別表示在三個維度上的坐標；此外，每個塊還帶有一個 dim3 類型的屬性 blockDim，blockDim.x、blockDim.y、blockDim.z 分別表示該 block 三個維度上各有多少個核，這個 block 中的總核數為 blockDim.x * blockDim.y * blockDim.z；

我們一次使用的多個 block，最好能使用一個容器把他們都包起來，這就是 grid，類比于上文中 thread 和 block 的關系，block 和 grid 也有相似的關系。我們使用 blockIdx.x、blockIdx.y、blockIdx.z 表示每個 block 在 grid 中的位置；同樣，grid 也具有 gridDim.x、gridDim.y 和 gridDim.z 三個屬性以及三者相乘的總 block 數。

知道了上面這些知識后，我們可以對 “kernel_name <<<>>>();” 中尖括號中的參數做一個更具體的解釋，它應該被定義為在 GPU 中執行這一核函數的所有核的組織形式，以 "kernel_name <<< number_of_blocks, thread_per_block>>> (arguments)" 的形式使用，一個典型的示例如下：

int nx = 16; int ny = 4; dim3 block(8, 2); // z默認為1 dim3 grid(nx/8, ny/2); addKernel << <grid, block >> >(c, a, b);

這一示例中創建了一個有 (2*2) 個 block 的 grid，每個 block 中有 (8*2) 個 thread，下圖給出了更直觀的表述：

需要注意的是，對 block、grid 的尺寸定義并不是沒有限制的，一個 GPU 中的核的數量同樣是有限制的。對于一個 block 來說，總的核數不得超過 1024，x、y 維度都不得超過 1024，z 維度不得超過 64，如下圖

對于整個 grid 而言，x 維度上不得有超過?232?1232?1?個 thread，注意這里是 thread 而不是 block，在其 y 維度和 z 維度上 thread 數量不得超過 65536.

在 cuda 編程中我們經常會把數組的每一個元素分別放到單獨的一個核中處理，我們可以利用核的索引讀取數組中的數據進行操作，但由于 block、grid 的存在，索引的獲取需要一定的計算，在 exercise2 中給出了一個 3D 模型中取值的訓練，實現如下

__global__ void print_array(int *input) {int tid = (blockDim.x*blockDim.y)*threadIdx.z + blockDim.x*threadIdx.y + threadIdx.x;int xoffset = blockDim.x * blockDim.y * blockDim.z;int yoffset = blockDim.x * blockDim.y * blockDim.z * gridDim.x;int zoffset = blockDim.x * blockDim.y * blockDim.z * gridDim.x * gridDim.y;int gid = zoffset * blockIdx.z + yoffset * blockIdx.y + xoffset * blockIdx.x + tid;printf("blockIdx.x : %d, blockIdx.y : %d, blockIdx.z : %d,gid : %d, value: %d\n", blockIdx.x, blockIdx.y, blockIdx.z, gid, input[gid]); }

3. 數據在 host 和 device 之間的遷移

我們前邊提到，cuda 的編程步驟是將數據移入 GPU，待計算完成后將其取出，官方對可能涉及到的內存操作類的操作都給出了接口。

首先是 cudaMemCpy 函數，其定義為

cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count, cudaMemcpyKind kind )

該函數是將數據從 CPU 移入到 GPU 或者從 GPU 移出到 CPU 中，參數 0 指向目標區域的地址，參數 1 指向數據的源地址，參數 2 表示要移動的數據的字節數，最后一個參數表示數據的移動方向（cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost 或 cudaMemcpyDeviceToDevice）

此外，對應 C 語言的內存空間操作，cuda 也推出了 CudaMalloc, CudaMemset, CudaFree 三個接口

cudaError_t cudaMalloc ( void** devPtr, size_t size ); cudaError_t cudaMemset ( void* devPtr, int value, size_t count ); cudaError_t cudaFree ( void* devPtr );

這里需要注意的一個點是 cudaMalloc 的第一參數的數據類型為 void**，這一點怎么理解呢？

這里我們結合一個示例進行解釋:

int *d_input; cudaMalloc((void **) &d_input, bytesize);

之所以使用 void，是因為這一步只管分配內存，不考慮如何解釋指針，所以只需要傳入待分配內存的地址，不需要傳入具體的類型，其他 API 中的 void* 也是同理。為什么是兩個 * 呢，這是因為我們在定義 d_input 時是定義了主存中的一個指針，它指向主存中的一個地址；而 & d_input 則是取得了存儲該指針值的地址，cudaMalloc 利用這一地址將在 GPU 中分配給該緩沖區的首地址賦值給 d_input。

利用上述的幾個接口函數，我們就可以實現一個基本的 cuda 程序的主函數：

int main() {const int arraySize = 64;const int byteSize = arraySize * sizeof(int);int *h_input,*d_input;h_input = (int*)malloc(byteSize);cudaMalloc((void **)&d_input,byteSize);srand((unsigned)time(NULL));for (int i = 0; i < 64; ++i){if(h_input[i] != NULL)h_input[i] = (int)rand()& 0xff;}cudaMemcpy(d_input, h_input, byteSize, cudaMemcpyHostToDevice);int nx = 4, ny = 4, nz = 4;dim3 block(2, 2, 2);dim3 grid(nx/2, ny/2, nz/2);print_array << < grid, block >> > (d_input);cudaDeviceSynchronize();cudaFree(d_input);free(h_input);return 0; }

其中 cudaDeviceSynchronize (); 的作用是在此處等待 GPU 中計算完成后再繼續執行后續的代碼。

4 錯誤處理

在 C++ 中，可以使用異常機制處理運行時錯誤，而 cuda 編程中由于 Host 和 Device 共同使用，難以利用異常機制，因此，cuda 提供了檢測運行時錯誤的機制。

看上面的 API 時會發現，每個函數的返回值類型都是 cudaError_t ，這正是 cuda 提供的錯誤檢測機制，如果返回值是 cudaSuccess 則說明執行正確，否則就是出現了錯誤。可以使用 cudaGetErrorString (?error?) 獲取返回值的代表的錯誤的文本。前面的代碼中沒有使用這一機制主要是為了便于閱讀，但實際的使用中這一機制是必不可少的，也會看到 VS 生成的 demo 代碼中就包含著大量的錯誤檢測代碼

cudaStatus = cudaSetDevice(0);if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?");goto Error;}cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMalloc failed!");goto Error;}if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMalloc failed!");goto Error;}cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMalloc failed!");goto Error;} ... ...

5 其他

不同的 block_size 計算耗時會不同，可以多嘗試后選擇計算的更快的參數（學 DL 的調參是吧，這也搞黑盒？）；考慮 GPU 的計算時間時要考慮數據移入移出 GPU 的時間。

不同的 GPU 有不同的性質，設備中也可能存在多個 GPU，在設計程序時需要考慮這些問題，cuda 也提供了訪問這些信息的接口

// 獲取設備數量 int deviceCount = 0; cudaGetDeviceCount(&deviceCount);//獲取第一個設備的各項性質 int devNo = 0; cudaDeviceProp iProp; cudaGetDeviceProperties(&iprop, devNo);

總結

以上是生活随笔為你收集整理的CUDA C编程入门的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。