多執行緒我們應該都不陌生，在作業系統中，程序是資源分配的基本單元，而執行緒是CPU時間排程的基本單元（這裡假設只有1個CPU）。

將執行緒的概念引申到CUDA程式設計中，我們可以認為執行緒就是執行CUDA程式的最小單元，前面我們建立的工程程式碼中，有個核函式概念不知各位童鞋還記得沒有，在GPU上每個執行緒都會執行一次該核函式。

但GPU上的執行緒排程方式與CPU有很大不同。CPU上會有優先順序分配，從高到低，同樣優先順序的可以採用時間片輪轉法實現執行緒排程。GPU上執行緒沒有優先順序概念，所有執行緒機會均等，執行緒狀態只有等待資源和執行兩種狀態，如果資源未就緒，那麼就等待；一旦就緒，立即執行。當GPU資源很充裕時，所有執行緒都是併發執行的，這樣加速效果很接近理論加速比；而GPU資源少於匯流排程個數時，有一部分執行緒就會等待前面執行的執行緒釋放資源，從而變為序列化執行。

程式碼還是用上一節的吧，改動很少，再貼一遍：

1. #include "cuda_runtime.h"           //CUDA執行時API
2. #include "device_launch_parameters.h"   
3. #include <stdio.h>
4. cudaError_t addWithCuda(int *c, constint *a, constint *b, size_t size);  
5. __global__ void addKernel(int *c, constint *a, constint *b)  
6. {  
7.     int i = threadIdx.x;  
8.     c[i] = a[i] + b[i];  
9. }  
10. int main()  
11. {  
12.     constint arraySize = 5;  
13.     constint a[arraySize] = { 1, 2, 3, 4, 5 };  
14.     constint b[arraySize] = { 10, 20, 30, 40, 50 };  
15.     int c[arraySize] = { 0 };  
16.     // Add vectors in parallel.
17.     cudaError_t cudaStatus;  
18.     int num = 0;  
19.     cudaDeviceProp prop;  
20.     cudaStatus = cudaGetDeviceCount(&num);  
21.     for(int i = 0;i<num;i++)  
22.     {  
23.         cudaGetDeviceProperties(&prop,i);  
24.     }  
25.     cudaStatus = addWithCuda(c, a, b, arraySize);  
26.     if (cudaStatus != cudaSuccess)   
27.     {  
28.         fprintf(stderr, "addWithCuda failed!");  
29.         return 1;  
30.     }  
31.     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]);  
32.     // cudaThreadExit must be called before exiting in order for profiling and
33.     // tracing tools such as Nsight and Visual Profiler to show complete traces.
34.     cudaStatus = cudaThreadExit();  
35.     if (cudaStatus != cudaSuccess)   
36.     {  
37.         fprintf(stderr, "cudaThreadExit failed!");  
38.         return 1;  
39.     }  
40.     return 0;  
41. }  
42. // 重點理解這個函式
43. cudaError_t addWithCuda(int *c, constint *a, constint *b, size_t size)  
44. {  
45.     int *dev_a = 0; //GPU裝置端資料指標
46.     int *dev_b = 0;  
47.     int *dev_c = 0;  
48.     cudaError_t cudaStatus;     //狀態指示
49.     // Choose which GPU to run on, change this on a multi-GPU system.
50.     cudaStatus = cudaSetDevice(0);  //選擇執行平臺
51.     if (cudaStatus != cudaSuccess)   
52.     {  
53.         fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
54.         goto Error;  
55.     }  
56.     // 分配GPU裝置端記憶體
57.     cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));  
58.     if (cudaStatus != cudaSuccess)   
59.     {  
60.         fprintf(stderr, "cudaMalloc failed!");  
61.         goto Error;  
62.     }  
63.     cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));  
64.     if (cudaStatus != cudaSuccess)   
65.     {  
66.         fprintf(stderr, "cudaMalloc failed!");  
67.         goto Error;  
68.     }  
69.     cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));  
70.     if (cudaStatus != cudaSuccess)   
71.     {  
72.         fprintf(stderr, "cudaMalloc failed!");  
73.         goto Error;  
74.     }  
75.     // 拷貝資料到GPU
76.     cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);  
77.     if (cudaStatus != cudaSuccess)   
78.     {  
79.         fprintf(stderr, "cudaMemcpy failed!");  
80.         goto Error;  
81.     }  
82.     cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);  
83.     if (cudaStatus != cudaSuccess)   
84.     {  
85.         fprintf(stderr, "cudaMemcpy failed!");  
86.         goto Error;  
87.     }  
88.     // 執行核函式
89. <span style="BACKGROUND-COLOR: #ff6666"><strong>    addKernel<<<1, size>>>(dev_c, dev_a, dev_b);</strong>  
90. </span>    // cudaThreadSynchronize waits for the kernel to finish, and returns
91.     // any errors encountered during the launch.
92.     cudaStatus = cudaThreadSynchronize();   //同步執行緒
93.     if (cudaStatus != cudaSuccess)   
94.     {  
95.         fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);  
96.         goto Error;  
97.     }  
98.     // Copy output vector from GPU buffer to host memory.
99.     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);      //拷貝結果回主機
100.     if (cudaStatus != cudaSuccess)   
101.     {  
102.         fprintf(stderr, "cudaMemcpy failed!");  
103.         goto Error;  
104.     }  
105. Error:  
106.     cudaFree(dev_c);    //釋放GPU裝置端記憶體
107.     cudaFree(dev_a);  
108.     cudaFree(dev_b);      
109.     return cudaStatus;  
110. }  

紅色部分即啟動核函式的呼叫過程，這裡看到呼叫方式和C不太一樣。<<<>>>表示執行時配置符號，裡面1表示只分配一個執行緒組（又稱執行緒塊、Block），size表示每個執行緒組有size個執行緒（Thread）。本程式中size根據前面傳遞引數個數應該為5，所以執行的時候，核函式在5個GPU執行緒單元上分別運行了一次，總共運行了5次。這5個執行緒是如何知道自己“身份”的？是靠threadIdx這個內建變數，它是個dim3型別變數，接受<<<>>>中第二個引數，它包含x,y,z 3維座標，而我們傳入的引數只有一維，所以只有x值是有效的。通過核函式中int i = threadIdx.x;這一句，每個執行緒可以獲得自身的id號，從而找到自己的任務去執行。