前言

GPU可以利用平行處理的方式，縮短執行時間，因此，這一次就來介紹多執行緒的程式設計方法及應用。

多執行緒的設定

上一篇 介紹GPU函數的設定是透過<<<...>>>指定三個參數：

1. 區塊(Block)數量。
2. 執行緒(Thread)數量：一個區塊要執行的執行緒數量。
3. 使用的共享記憶體(Shared memory)大小，此參數可不設定。

每一種GPU的數量均不相同，因此，我們先要偵測相關的數量限制，可以執行上一篇的deviceQuery.exe，預設建置的目錄在【v11.2\bin\win64\Debug】。

1. deviceProp.maxThreadsPerMultiProcessor：一個處理器(Multiprocessor)的最大執行緒數量，筆者的GPU卡是【GTX1050 Ti】，數量是2048。
2. deviceProp.maxThreadsPerBlock：筆者的GPU卡區塊的最大執行緒數量是1024。
3. deviceProp.sharedMemPerMultiprocessor：一個處理器(Multiprocessor)的最大共享記憶體大小，筆者的GPU卡是【98304 bytes】。
4. deviceProp.sharedMemPerBlock：一個區塊的最大共享記憶體大小，筆者的GPU卡是【49152 bytes】。

所以，如果要在不同的卡都能執行，必須在程式中計算可設定的執行緒數量，總數量超過程式就會發生異常。共享記憶體也是一樣的道理。【GTX1050 Ti】含6個處理器，我在以下的程式設定7*2048個執行緒也不會發生異常，我猜處理器有類似queue的功能，一旦執行緒超過容量限制，工作會排隊等待，直到處理器有空時才被處理。

程式撰寫

接下來，我們使用多執行緒進行向量的相加，一維陣列中的每個元素都獨立計算，每個元素計算交由一個執行緒處理。

1. 定義陣列含 1024 個元素。

#define N	1024

2. 使用GPU進行向量相加：threadIdx.x 表執行緒序號。

__global__ void vectorAdd(int* d_a, int* d_b, int* d_c) {
	// 使用第幾個執行緒計算 
	int tid = threadIdx.x;
	//printf("threadIdx = %d\n", tid);
	d_c[tid] = d_a[tid] + d_b[tid];
}

3. 設定CPU陣列(h_a、h_b)值。

    int h_a[N], h_b[N], h_c[N], h2_c[N];

	// 設定 a、b 的值
	for (int i = 0; i < N; i++) {
		h_a[i] = i; // 1,...,N
		h_b[i] = i * 2; // 2,...,2xN
	}

4. 主程式呼叫 vectorAdd 函數，採用 1 block x N threads，需先分配記憶體給GPU陣列(d_a、d_b)，再複製CPU陣列(h_a、h_b)值至GPU陣列(d_a、d_b)。

	// 分配記憶體
	cudaMalloc((void**)&d_a, N * sizeof(int));
	// 自 CPU變數 複製到 GPU變數
	cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice);
	cudaMalloc((void**)&d_b, N * sizeof(int));
	cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice);

	// 分配記憶體
	cudaMalloc((void**) &d_c, N * sizeof(int));

    // 使用GPU進行向量相加，1 block x N threads
    vectorAdd <<<1, N >>> (d_a, d_b, d_c);
    
    // 等待所有執行緒處理完畢
	cudaDeviceSynchronize();

• cudaMemcpy 參數：第一個為目標變數，第二個為來源變數，cudaMemcpyHostToDevice為CPU變數複製到GPU變數。CPU/GPU 變數不可混合運算，需先將變數全部轉移至GPU，才能使用GPU運算，CPU也是一樣。其他注意事項如下：
• 要列印變數，必須將GPU變數複製到CPU變數，才可正確顯示。
• vectorAdd 會被呼叫 1024 次，各由1個執行緒負責。
• 記得呼叫 cudaDeviceSynchronize()，等待所有執行緒處理完畢。
• GPU變數需在程式結束前釋放記憶體。

	cudaFree(d_a);
	cudaFree(d_b);
	cudaFree(d_c);

測試

1. 如果更改N為1025，則GPU計算結果均為0，因為，筆者的GPU卡是【GTX1050 Ti】，區塊的最大執行緒數量是1024。可以改成2個區塊(2x1024)：

vectorAdd <<<2, 1024 >>> (d_a, d_b, d_c);

vectorAdd 函數的陣列索引值要改為：

// blockIdx.x：第幾個block，blockDim.x：block內的執行緒數量
int tid = threadIdx.x + blockIdx.x * blockDim.x;

2. 使用clock()計算執行時間，CPU計算時間為0秒，GPU計算時間為0.188秒，含分配記憶體及CPU變數複製到GPU變數，若只考慮計算，則是0秒。

結語

在 GPU 撰寫多執行緒真是方便，只要呼叫時，指定個數即可，但是，要以程式檢查是否超出硬體限制條件，這部分可自deviceQuery截取相關程式碼，計算應使用的區塊數及執行緒數量。

平行處理要選擇多區塊或多執行緒呢? 使用多執行緒有兩個好處：

1. 資料共享：可透過共享記憶體，多執行緒可讀寫同一塊記憶體。
2. 同步：可以等非同步的多執行緒全部結束在進行下一段處理。
   我們會在下一篇實際運用這兩個特性，可提升運算效能。

完整程式放在 『GitHub』的VectorAdd目錄。使用多個區塊的改良程式放在VectorAdd_Improved目錄。