147
技術社區[雲棲]
CUDA從入門到精通（五）：線程並行
多線程我們應該都不陌生，在操作係統中，進程是資源分配的基本單元，而線程是CPU時間調度的基本單元（這裏假設隻有1個CPU）。
將線程的概念引申到CUDA程序設計中，我們可以認為線程就是執行CUDA程序的最小單元，前麵我們建立的工程代碼中，有個核函數概念不知各位童鞋還記得沒有，在GPU上每個線程都會運行一次該核函數。
但GPU上的線程調度方式與CPU有很大不同。CPU上會有優先級分配，從高到低，同樣優先級的可以采用時間片輪轉法實現線程調度。GPU上線程沒有優先級概念，所有線程機會均等，線程狀態隻有等待資源和執行兩種狀態，如果資源未就緒，那麼就等待；一旦就緒，立即執行。當GPU資源很充裕時，所有線程都是並發執行的，這樣加速效果很接近理論加速比；而GPU資源少於總線程個數時，有一部分線程就會等待前麵執行的線程釋放資源，從而變為串行化執行。
代碼還是用上一節的吧，改動很少，再貼一遍：
#include "cuda_runtime.h" //CUDA運行時API
#include "device_launch_parameters.h"
#include <stdio.h>
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t 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;
int num = 0;
cudaDeviceProp prop;
cudaStatus = cudaGetDeviceCount(&num);
for(int i = 0;i<num;i++)
{
cudaGetDeviceProperties(&prop,i);
}
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]);
// cudaThreadExit must be called before exiting in order for profiling and
// tracing tools such as Nsight and Visual Profiler to show complete traces.
cudaStatus = cudaThreadExit();
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaThreadExit failed!");
return 1;
}
return 0;
}
// 重點理解這個函數
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)
{
int *dev_a = 0; //GPU設備端數據指針
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;
}
// 分配GPU設備端內存
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;
}
// 拷貝數據到GPU
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;
}
// 運行核函數
addKernel<<<1, size>>>(dev_c, dev_a, dev_b);
// cudaThreadSynchronize waits for the kernel to finish, and returns
// any errors encountered during the launch.
cudaStatus = cudaThreadSynchronize(); //同步線程
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaThreadSynchronize 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); //釋放GPU設備端內存
cudaFree(dev_a);
cudaFree(dev_b);
return cudaStatus;
}
紅色部分即啟動核函數的調用過程，這裏看到調用方式和C不太一樣。<<<>>>表示運行時配置符號，裏麵1表示隻分配一個線程組（又稱線程塊、Block），size表示每個線程組有size個線程（Thread）。本程序中size根據前麵傳遞參數個數應該為5，所以運行的時候，核函數在5個GPU線程單元上分別運行了一次，總共運行了5次。這5個線程是如何知道自己“身份”的？是靠threadIdx這個內置變量，它是個dim3類型變量，接受<<<>>>中第二個參數，它包含x,y,z 3維坐標，而我們傳入的參數隻有一維，所以隻有x值是有效的。通過核函數中int i = threadIdx.x;這一句，每個線程可以獲得自身的id號，從而找到自己的任務去執行。
下節我們介紹塊並行。
最後更新：2017-04-03 16:48:42