前面我們沒有講程式的結構，我想有些童鞋可能迫不及待想知道CUDA程式到底是怎麼一個執行過程。好的，這一節在介紹流之前，先把CUDA程式結構簡要說一下。

CUDA程式檔案字尾為.cu，有些編譯器可能不認識這個字尾的檔案，我們可以在VS2008的Tools->Options->Text Editor->File Extension裡新增cu字尾到VC++中，如下圖：

一個.cu檔案內既包含CPU程式（稱為主機程式），也包含GPU程式（稱為裝置程式）。如何區分主機程式和裝置程式？根據宣告，凡是掛有“__global__”或者“__device__”字首的函式，都是在GPU上執行的裝置程式，不同的是__global__裝置程式可被主機程式呼叫，而__device__裝置程式則只能被裝置程式呼叫。

沒有掛任何字首的函式，都是主機程式。主機程式顯示宣告可以用__host__字首。裝置程式需要由NVCC進行編譯，而主機程式只需要由主機編譯器（如VS2008中的cl.exe，Linux上的GCC）。主機程式主要完成裝置環境初始化，資料傳輸等必備過程，裝置程式只負責計算。

主機程式中，有一些“cuda”打頭的函式，這些都是CUDA Runtime API，即執行時函式，主要負責完成裝置的初始化、記憶體分配、記憶體拷貝等任務。我們前面第三節用到的函式cudaGetDeviceCount()，cudaGetDeviceProperties()，cudaSetDevice()都是執行時API。這些函式的具體引數宣告我們不必一一記下來，拿出第三節的官方利器就可以輕鬆查詢，讓我們開啟這個檔案：

開啟後，在pdf搜尋欄中輸入一個執行時函式，例如cudaMemcpy，查到的結果如下：

可以看到，該API函式的引數形式為，第一個表示目的地，第二個表示來源地，第三個引數表示位元組數，第四個表示型別。如果對型別不瞭解，直接點選超連結，得到詳細解釋如下：

可見，該API可以實現從主機到主機、主機到裝置、裝置到主機、裝置到裝置的記憶體拷貝過程。同時可以發現，利用該API手冊可以很方便地查詢我們需要用的這些API函式，所以以後編CUDA程式一定要把它開啟，隨時準備查詢，這樣可以大大提高程式設計效率。

好了，進入今天的主題：流並行。

前面已經介紹了執行緒並行和塊並行，知道了執行緒並行為細粒度的並行，而塊並行為粗粒度的並行，同時也知道了CUDA的執行緒組織情況，即Grid-Block-Thread結構。一組執行緒並行處理可以組織為一個block，而一組block並行處理可以組織為一個Grid，很自然地想到，Grid只是一個網格，我們是否可以利用多個網格來完成並行處理呢？答案就是利用流。

流可以實現在一個裝置上執行多個核函式。前面的塊並行也好，執行緒並行也好，執行的核函式都是相同的（程式碼一樣，傳遞引數也一樣）。而流並行，可以執行不同的核函式，也可以實現對同一個核函式傳遞不同的引數，實現任務級別的並行。

CUDA中的流用cudaStream_t型別實現，用到的API有以下幾個：cudaStreamCreate(cudaStream_t * s)用於建立流，cudaStreamDestroy(cudaStream_t s)用於銷燬流，cudaStreamSynchronize()用於單個流同步，cudaDeviceSynchronize()用於整個裝置上的所有流同步，cudaStreamQuery()用於查詢一個流的任務是否已經完成。具體的含義可以查詢API手冊。

下面我們將前面的兩個例子中的任務改用流實現，仍然是{1,2,3,4,5}+{10,20,30,40,50} = {11,22,33,44,55}這個例子。程式碼如下：

1. #include "cuda_runtime.h"
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 = blockIdx.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. // Helper function for using CUDA to add vectors in parallel.
43. cudaError_t addWithCuda(int *c, constint *a, constint *b, size_t size)  
44. {  
45.     int *dev_a = 0;  
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.     // Allocate GPU buffers for three vectors (two input, one output)    .
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.     // Copy input vectors from host memory to GPU buffers.
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. <span style="BACKGROUND-COLOR: #ff6666">    cudaStream_t stream[5];  
89.     for(int i = 0;i<5;i++)  
90.     {  
91.         cudaStreamCreate(&stream[i]);   //建立流
92.     }  
93. </span>    // Launch a kernel on the GPU with one thread for each element.
94. <span style="BACKGROUND-COLOR: #ff6666">    for(int i = 0;i<5;i++)  
95.     {  
96.         addKernel<<<1,1,0,stream[i]>>>(dev_c+i, dev_a+i, dev_b+i);    //執行流
97.     }  
98.     cudaDeviceSynchronize();  
99. </span>    // cudaThreadSynchronize waits for the kernel to finish, and returns
100.     // any errors encountered during the launch.
101.     cudaStatus = cudaThreadSynchronize();  
102.     if (cudaStatus != cudaSuccess)   
103.     {  
104.         fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);  
105.         goto Error;  
106.     }  
107.     // Copy output vector from GPU buffer to host memory.
108.     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);  
109.     if (cudaStatus != cudaSuccess)   
110.     {  
111.         fprintf(stderr, "cudaMemcpy failed!");  
112.         goto Error;  
113.     }  
114. Error:  
115. <span style="BACKGROUND-COLOR: #ff6666">    for(int i = 0;i<5;i++)  
116.     {  
117.         cudaStreamDestroy(stream[i]);   //銷燬流
118.     }  
119. </span>    cudaFree(dev_c);  
120.     cudaFree(dev_a);  
121.     cudaFree(dev_b);      
122.     return cudaStatus;  
123. }  

注意到，我們的核函式程式碼仍然和塊並行的版本一樣，只是在呼叫時做了改變，<<<>>>中的引數多了兩個，其中前兩個和塊並行、執行緒並行中的意義相同，仍然是執行緒塊數（這裡為1）、每個執行緒塊中執行緒數（這裡也是1）。第三個為0表示每個block用到的共享記憶體大小，這個我們後面再講；第四個為流物件，表示當前核函式在哪個流上執行。我們建立了5個流，每個流上都裝載了一個核函式，同時傳遞引數有些不同，也就是每個核函式作用的物件也不同。這樣就實現了任務級別的並行，當我們有幾個互不相關的任務時，可以寫多個核函式，資源允許的情況下，我們將這些核函式裝載到不同流上，然後執行，這樣可以實現更粗粒度的並行。

好了，流並行就這麼簡單，我們處理任務時，可以根據需要，選擇最適合的並行方式。