上一篇文章介紹了如何搭建CUDA編程環境，從這篇文章開始正式開始介紹如何使用CUDA進行編程。

異構計算架構

如下圖所示，一個典型的異構計算架構節點由一個多核CPU和一個或多個GPU組成，每個CPU和GPU都是獨立的設備，它們之間通過PCIe總線連接。GPU作為CPU的協處理器用于執行一些并行計算任務。CPU適合用于做一些邏輯控制任務，而GPU則適合作為CPU的協處理器用于做一些大計算量的數據并行計算任務。

一般我們稱CPU為host，稱GPU為device，相應地，一個異構計算的應用程序代碼也被分為兩部分：運行在CPU上的程序被稱為host代碼，運行在GPU上的程序被稱為device代碼。

一個Hello World

在我們學習一種新的編程語言的時候，一般都是先從打印一句"Hello World"開始，今天開始學習CUDA編程，按照國際慣例，也從打印"Hello World"開始。

首先新建一個CUDA C源文件hello_world.cu，然后輸入下面的內容：

#include <stdio.h>int main(void) {printf("Hello World from CPU\n");return 0; }

用nvcc進行編譯生成可執行文件：

nvcc hello_world.cu -o hello_world

運行可執行文件hello_world，沒有意外的話就能在終端看到打印出的"Hello World from CPU"了。這是在CPU上運行代碼打印了這句話，那么怎么在GPU上打印呢？

要在GPU上運行程序，我們需要寫一個能在GPU上執行的函數，然后在CPU上調用這個函數。來看一個例子：

#include <stdio.h>__global__ void HelloFromGPU(void) {printf("Hello from GPU\n"); }int main(void) {printf("Hello from CPU\n");HelloFromGPU<<<1, 5>>>();cudaDeviceReset();return 0; }

在上面的這段代碼里，我添加了一個用__global__函數類型限定符修飾的函數HelloFromGPU，然后在main函數中去調用它。在CUDA中，有以下3種函數類型限定符：

_global_

被__global__函數類型限定符修飾的函數被稱為內核函數，該函數在host上被調用，在device上執行，只能返回void類型，不能作為類的成員函數。調用__global__修飾的函數是異步的，也就是說它未執行完就會返回。

CUDA內核函數的限制：

• 只能訪問device內存
• 必須具有void返回類型
• 不支持可變數量的參數
• 不支持靜態變量
• 顯式異步行為

_device_

被__device__函數類型限定符修飾的函數只能在device上被調用，在device上執行，用于在device代碼中內部調用。

_host_

被__host__函數類型限定符修飾的函數只能在host上被調用，在host上執行，也就是host上的函數，__host__函數類型限定符可以省略。

與調用一般CPU函數不同的是，調用HelloFromGPU函數的時候會在后面寫上<<< >>>，意思這是從host端到device端的內核函數調用，里面的參數是執行配置，用來說明使用多少線程來執行內核函數。一個內核函數是通過一組線程來執行的，所有線程執行的同樣的代碼，我這里設置的是用了5個線程來執行。程序編譯成功后運行得到如下結果：

Hello from CPU Hello from GPU Hello from GPU Hello from GPU Hello from GPU Hello from GPU

可以看到，內核函數里的打印語句被執行了5次。

用CUDA實現數組相加

一個典型的CUDA程序結構一般由以下5個步驟組成：

分配GPU內存；

從CPU內存中拷貝數據到GPU內存中；

調用CUDA內核函數執行程序指定的計算任務；

從GPU內存中把數據拷貝回CPU內存中；

釋放GPU內存；

下面以一個數組相加的例子來展示以上過程，數組相加的計算過程比較簡單：數組a和數組b中對應下標的元素相加然后存入數組c中。

首先來看一下在CPU上實現數組相加的代碼：

#include <iostream>void VectorAddCPU(const float *const a, const float *const b, float *const c,const int n) {for (int i = 0; i < n; ++i) {c[i] = a[i] + b[i];} }int main(void) {// alloc memory for hostconst size_t size = 1024;float *ha = new float[size]();float *hb = new float[size]();float *hc = new float[size]();for (int i = 0; i < size; ++i) {ha[i] = i;hb[i] = size - i;}VectorAddCPU(ha, hb, hc, size);delete[] ha;delete[] hb;delete[] hc;return 0; }

函數VectorAddCPU用了一個for循環在CPU上實現數組相加的過程。如果要用GPU來實現該過程，則調用CUDA的API按照前面說的5個步驟編寫代碼：

#include <cuda_runtime.h> #include <iostream>__global__ void VectorAddGPU(const float *const a, const float *const b,float *const c, const int n) {int i = blockDim.x * blockIdx.x + threadIdx.x;if (i < n) {c[i] = a[i] + b[i];} }int main(void) {// 分配CPU內存const size_t size = 10240;const size_t n = size / sizeof(float);float *ha = new float[n]();float *hb = new float[n]();float *hc = new float[n]();for (int i = 0; i < n; ++i) {ha[i] = i;hb[i] = size - i;}// 分配GPU內存float *da = nullptr;float *db = nullptr;float *dc = nullptr;cudaMalloc((void **)&da, size);cudaMalloc((void **)&db, size);cudaMalloc((void **)&dc, size);// 把數據從CPU拷貝到GPUcudaMemcpy(da, ha, size, cudaMemcpyHostToDevice);cudaMemcpy(db, hb, size, cudaMemcpyHostToDevice);cudaMemcpy(dc, hc, size, cudaMemcpyHostToDevice);const int thread_per_block = 256;const int block_per_grid = (size + thread_per_block - 1) / thread_per_block;// 調用核函數VectorAddGPU<<<block_per_grid, thread_per_block>>>(da, db, dc, size);// 把數據從GPU拷貝回CPUcudaMemcpy(hc, dc, size, cudaMemcpyDeviceToHost);// 釋放GPU內存cudaFree(da);cudaFree(db);cudaFree(dc);// 釋放CPU內存delete[] ha;delete[] hb;delete[] hc;return 0; }

在這段代碼里，用到了幾個CUDA的內存管理函數：

cudaMalloc

函數原型：

cudaError_t cudaMalloc(void** devPtr, size_t size)

該函數用于在GPU上分配指定大小的內存空間，類似于標準C語言中的malloc函數。

cudaMemcpy

函數原型：

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

該函數用于內存拷貝，類似于標準C語言中的memcpy函數。數據拷貝的流向由參數kind指定，分為以下4種方式：

• cudaMemcpyHostToHost
• cudaMemcpyHostToDevice
• cudaMemcpyDeviceToHost
• cudaMemcpyDeviceToDevice

從它們的字面意思就能知道，如果是從host拷貝數據到device，那么kind參數應該設置為cudaMemcpyHostToDevice；如果是從device拷貝數據到host，那么kind參數應該設置為cudaMemcpyDeviceToHost。上面的代碼體現了這一點。

cudaFree

udaError_t cudaFree(void* devPtr)

該函數用于釋放已分配的GPU內存空間，類似于標準C語言中的free函數。

在上面的代碼中，內核函數VectorAddGPU用于實現數組相加，與前面的代碼中VectorAddCPU函數不同的是，VectorAddGPU函數里面并沒有用for循環，因為在GPU中是將數據進行并行化劃分，然后通過線程組去實現計算過程的，線程組中的每個線程都是執行c[i] = a[i] + b[i]這個計算過程。在這個程序里，數組的長度為1024，我設置了每個線程組中的線程數量為256，總共用了4個線程組去進行計算。

關于GPU中線程和線程組的相關知識，我將在下一篇文章中進行闡述。

參考資料

• 《Professional CUDA C Programming》
• 《CUDA C Programming_Guide》

總結

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

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