詳解第一個CUDA程式kernel.cu
CUDA是一個基於NVIDIA GPU的平行計算平臺和程式設計模型，通過呼叫CUDA提供的API，可以開發高效能的並行程式。CUDA安裝好之後，會自動配置好VS編譯環境，按照UCDA模板新建一個工程“Hello CUDA”：
建好之後，發現該工程下已經存在一個專案 kernel.cu。這個是CUDA程式設計的入門示例，實現的功能是兩個整型陣列相加，程式碼如下：
#include "cuda_runtime.h" #include "device_launch_parameters.h" #include <stdio.h> cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size); __global__ void addKernel(int *c, const int *a, const int *b) { int i = threadIdx.x; c[i] = a[i] + b[i]; } int main() { const int arraySize = 5; const int a[arraySize] = { 1, 2, 3, 4, 5 }; const int b[arraySize] = { 10, 20, 30, 40, 50 }; int c[arraySize] = { 0 }; // Add vectors in parallel. cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize); if (cudaStatus != cudaSuccess) { fprintf(stderr, "addWithCuda failed!"); return 1; } printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n", c[0], c[1], c[2], c[3], c[4]); // cudaDeviceReset must be called before exiting in order for profiling and // tracing tools such as Nsight and Visual Profiler to show complete traces. cudaStatus = cudaDeviceReset(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaDeviceReset failed!"); return 1; } return 0; } // Helper function for using CUDA to add vectors in parallel. cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size) { int *dev_a = 0; int *dev_b = 0; int *dev_c = 0; cudaError_t cudaStatus; // Choose which GPU to run on, change this on a multi-GPU system. cudaStatus = cudaSetDevice(0); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?"); goto Error; } // Allocate GPU buffers for three vectors (two input, one output) . cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int)); 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; } cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int)); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc failed!"); goto Error; } // Copy input vectors from host memory to GPU buffers. cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } // Launch a kernel on the GPU with one thread for each element. addKernel << <1, size >> > (dev_c, dev_a, dev_b); // Check for any errors launching the kernel cudaStatus = cudaGetLastError(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus)); goto Error; } // cudaDeviceSynchronize waits for the kernel to finish, and returns // any errors encountered during the launch. cudaStatus = cudaDeviceSynchronize(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus); goto Error; } // Copy output vector from GPU buffer to host memory. cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } Error: cudaFree(dev_c); cudaFree(dev_a); cudaFree(dev_b); return cudaStatus; }
程式首先定義了一個函式addWithCuda，它是呼叫GPU運算的入口函式，返回型別是cudaError_t。
cudaError_t是一個列舉型別，可以作為幾乎所有CUDA函式的返回型別，用來檢測函式執行期間發生的不同型別的錯誤，一共有80多個錯誤型別，可以在driver_types.h標頭檔案中檢視每一個整型對應的錯誤型別，如果返回0，代表執行成功。
第二個函式addKernel在最前有一個修飾符“__global__”，這個修飾符告訴編譯器，被修飾的函式應該編譯為在GPU而不是在CPU上執行，所以這個函式將被交給編譯裝置程式碼的編譯器——NVCC編譯器來處理，其他普通的函式或語句將交給主機編譯器處理。
這裡“裝置”的概念可以理解為GPU和其視訊記憶體組成的運算單元，“主機”可以理解為CPU和系統記憶體組成的運算單元。在GPU上執行的函式稱為核函式。
addKernel函式定義：
__global__ void addKernel(int *c, const int *a, const int *b)
{
int i = threadIdx.x;
c[i] = a[i] + b[i];
}這個核函式裡有一個陌生的threadIdx.x，表示的是thread在x方向上的索引號，理解這個之前得先了解一下GPU執行緒的層次結構：
CUDA中的執行緒（thread）是裝置中並行運算結構中的最小單位，類似於主機中的執行緒的概念，thread可以以一維、二維、三維的形式組織在一起，threadIdx.x表示的是thread在x方向的索引號，還可能存在thread在y和z方向的索引號threadIdx.y和threadIdx.z。
一維、二維或三維的thread組成一個執行緒塊（Block），一維、二維或三維的執行緒塊（Block）組合成一個執行緒塊網格（Grid），執行緒塊網格（Grid）可以是一維或二維的。通過網格塊（Grid）->執行緒塊（Block）->執行緒（thread）的 順序可以定位到每一個並且唯一的執行緒。
回到程式中的addKernel函式上來，這個函式會被GPU上的多個執行緒同時執行一次，執行緒間彼此沒有通訊，相互獨立。到底會有多少個執行緒來分別執行核函式，是在“<<<>>>”符號裡定義的。“<<<>>>”表示執行時配置符號，在本程式中的定義是<<<1，size>>>，表示分配了一個執行緒塊（Block），每個執行緒塊有分配了size個執行緒，“<<<>>>”中的 引數並不是傳遞給裝置程式碼的引數，而是定義主機程式碼執行時如何啟動裝置程式碼。以上定義的這些執行緒都是一個維度上的，可以通過thredaIdx.x來獲取執行當前計算任務的執行緒的ID號。
cudaSetDevice函式用來設定要在哪個GPU上執行，如果只有一個GPU，設定為cudaSetDevice（0）；
cudaMalloc函式用來為參與運算的資料分配視訊記憶體空間，函式原型：cudaError_t cudaMalloc(void **p, size_t s)；
cudaMemcpy函式用於主機記憶體和裝置視訊記憶體以及主機與主機之間，裝置與裝置之間相互拷貝資料，函式原型：
cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind);
第一個引數dst是目標資料地址，第二個引數src是源資料地址，第三個引數count是資料大小，第四個引數kind定義資料拷貝的型別，有如下幾類列舉型別：
/**
* CUDA memory copy types
*/
enum __device_builtin__ cudaMemcpyKind
{
cudaMemcpyHostToHost = 0, /**< Host -> Host */
cudaMemcpyHostToDevice = 1, /**< Host -> Device */
cudaMemcpyDeviceToHost = 2, /**< Device -> Host */
cudaMemcpyDeviceToDevice = 3, /**< Device -> Device */
cudaMemcpyDefault = 4 /**< Direction of the transfer is inferred from the pointer values. Requires unified virtual addressing */
};接下來在呼叫核函式時候添加了執行時配置符號“<<<>>>”，定義執行緒塊和執行緒的數量，如<<<1,5>>>表示定義了一個執行緒塊，每個執行緒塊包含了5個執行緒。
cudaGetLastError函式用於返回最新的一個執行時呼叫錯誤，對於任何CUDA錯誤，都可以通過函式cudaGetErrorString函式來獲取錯誤的詳細資訊。
cudaDeviceSynchronize函式提供了一個阻塞，用於等待所有的執行緒都執行完各自的計算任務，然後繼續往下執行。
cudaFree函式用於釋放申請的視訊記憶體空間。
cudaDeviceReset函式用於釋放所有申請的視訊記憶體空間和重置裝置狀態；
第一個CUDA程式kernel.cu涉及的內容主要就是這些。CUDA的使用步驟如下：
- 主機程式碼執行
- 傳輸資料給GPU
- 確定Grid、Block大小
- 呼叫核心函式，GPU多執行緒執行程式
- 傳輸運算結果給CPU
- 繼續主機程式碼執行
期間涉及到在裝置上的一些視訊記憶體空間申請、銷燬等操作，從記憶體到視訊記憶體上資料的相互拷貝是一個比較耗時的過程，應該儘量減少這種操作。