本文來源公眾號“DeepDriving”，僅用于學術分享，侵權刪，干貨滿滿。

原文鏈接：CUDA編程-02： 初識CUDA編程

上一篇文章DeepDriving | CUDA編程-01： 搭建CUDA編程環境-CSDN博客介紹了如何搭建CUDA編程環境，從這篇文章開始正式開始介紹如何使用CUDA進行編程。

1 異構計算架構

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

heterogeneous_architecture

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

heterogeneous_architecture

2 一個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__修飾的函數是異步的，也就是說它未執行完就會返回。

_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次。

3 用CUDA實現數組相加

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

1. 分配GPU內存；

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

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

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

5. 釋放GPU內存；

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

vector_add

首先來看一下在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?=?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;}//?分配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中線程和線程組的相關知識，我將在下一篇文章中進行闡述。

4 參考資料

• 《Professional CUDA C Programming》

• 《CUDA C Programming_Guide》

THE END !

文章結束，感謝閱讀。您的點贊，收藏，評論是我繼續更新的動力。大家有推薦的公眾號可以評論區留言，共同學習，一起進步。