本文介紹CUDA環境下兩個向量的加法運算。程式碼執行的系統環境為作業系統Ubuntu 16.04 LTS 64位，CUDA版本7.5，GCC版本5.4.0。專案Github下載地址為：CUDA向量加法Github專案

1. CUDA程式碼分析和實現

Step 1: 關於Host/Device 陣列指標

CUDA程式設計而言，我們習慣把CPU端稱為Host，GPU端稱為Device。基於Host端和Device端擁有各自不同的兩份記憶體地址空間的事實，程式碼層面上要求程式設計人員一方面維護兩份指標地址，即一份指向Host端記憶體的指標地址，一份指向Device端視訊記憶體空間的指標地址；另一方面必要時需要對Host端和Device端的記憶體資料進行相互訪存，拷貝資料。

Host端記憶體指標，

int *a_host;
int *b_host;
int *c_host;
int *devresult_host;

Device端記憶體指標，

a_host = (int *)malloc(arraysize*sizeof(int));
b_host = (int *)malloc(arraysize*sizeof(int));
c_host = (int *)malloc(arraysize*sizeof(int));

Host/Device端記憶體資料拷貝操作，

cudaMemcpy(a_dev, a_host, arraysize*
 
sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(b_dev, b_host, arraysize*sizeof(int), cudaMemcpyHostToDevice);

Step 2: 關於Device 並行執行緒配置

dim3 dimBlock()和dim3 dimGrid()是Device端關於執行緒配置的兩種資料結構，該資料結構具有兩個層面的含義：其一，初始化並行執行緒的數量；其二，初始化這些並行執行緒的空間組織方式，即並行執行緒的維度。blocks/threads的組織方式通常出於程式設計方便上的考慮。

例1：給定總共1024個threads和4個blocks，問：如果blocks以一維空間形式組織，threads以一維空間形式組織，列舉一個可行的並行執行緒排程模型？

int blocksize = 256;
int blocknum = 4;

dim3 dimBlock(blocksize, 1, 1);
dim3 dimGrid(blocknum, 1, 1);

例2：給定總共2048個threads和4個blocks，問：如果blocks以二維空間形式組織，threads以一維空間形式組織，列舉一個可行的並行執行緒排程模型？

int blocksize = 256;
int blocknum = 2;

dim3 dimBlock(blocksize, 1, 1);
dim3 dimGrid(blocknum, blocknum, 1);

例3：給定總共2048個threads和4個blocks，問：如果blocks以二維空間形式組織，threads以二維空間形式組織，列舉一個可行的並行執行緒排程模型？

int blocksize = 16;
int blocknum = 2;

dim3 dimBlock(blocksize, blocksize, 1);
dim3 dimGrid(blocknum, blocknum, 1);

本文中，blocks數量為blocknum，threads的數量 (每個block中包含的執行緒數量) 為512。blocks和threads均採用一維空間形式組織，即blocks的編號為0, 1, 2, 3, …, blocknum -1，每個block內部threads的編號為0, 1, 2, 3, …, 511。因此，我們可以用tid來表示這blocknum $\times$ 512個並行執行緒的唯一編號，如下，

int tid = blockIdx.x * blockDim.x + threadIdx.x

注：一維空間形式組織的blockDim.y = 1, blockDim.z=1, threadIdx.y=1 且 threadIdx.z = 1.

Step 3: 關於kernel的實現和呼叫

在資料從Host端拷貝到Device端之後，我們需要在Device端構造Kernel函式來計算這些資料。Kernel是CUDA並行結構的核心部分，一個kernel函式可以一次誘發多個執行緒併發執行，比方說，給定一個warp大小32的情況下，一個kernel函式可以誘發tid=0, 1, 2, …, 21的執行緒同時平行計算資料。如下，

__global__ void add_in_parallel(int *array_a, int *array_b, int *array_c)
{
   int tid = blockIdx.x * blockDim.x + threadIdx.x;
   array_c[tid] = array_a[tid] + array_b[tid];
}

__global__ 識別符號表示add_in_parallel()函式是一個Device端執行的kernel函式。

CUDA編譯器nvcc是一個修改版的C編譯器，語法上兩者的函式呼叫方式高度相似，如下，

add_in_parallel<<<dimGrid, dimBlock>>>(a_dev, b_dev, c_dev);

注：Kernel函式的呼叫需要事先定義blocks/threads的數量和空間組織形式，

Step 4: 關於Device 計算結果

將Device端的計算結果返回到Host端，如下，

cudaMemcpy(devresult_host, c_dev, arraysize*sizeof(int), cudaMemcpyDeviceToHost);

CPU端的計算結果，如下，

for (int i = 0; i < arraysize; i++)
{
   a_host[i] = i;
   b_host[i] = i;
   c_host[i] = a_host[i] + b_host[i];
}

校驗CPU/GPU兩者的計算結果的一致性，如下，

for (int i = 0; i < arraysize; i++)
{
   if (c_host[i]!=devresult_host[i])
   {
      status = 1;
   }
}

if (status)
{
   printf("Failed vervified.\n");
}
else
{
   printf("Sucessdully verified.\n");
}

Step 5: 關於釋放記憶體

釋放Host端的記憶體，如下，

free(a_host);
free(b_host);
free(c_host);

釋放Device端的記憶體，如下，

cudaFree(a_dev);
cudaFree(b_dev);
cudaFree(c_dev);

2. 編譯除錯

原始碼：vec_add.cu

編譯，如下，

$ nvcc vec_add.cu -o vec_add

執行，如下，

$ ./vec_add

3. 更多改進版本

Version 1: 增加計時器，包括資料拷貝耗時，kernel耗時

原始碼：vec_add_count.cu

Version 2: 增加GPU kernel warm up()函式

原始碼：vec_add_warmup.cu