505
CUDA從入門到精通（七）：流並行
前麵我們沒有講程序的結構，我想有些童鞋可能迫不及待想知道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}這個例子。代碼如下：
#include "cuda_runtime.h"
#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 = blockIdx.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;
}
// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)
{
int *dev_a = 0;
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;
}
// Allocate GPU buffers for three vectors (two input, one output) .
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;
}
// Copy input vectors from host memory to GPU buffers.
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;
}
cudaStream_t stream[5];
for(int i = 0;i<5;i++)
{
cudaStreamCreate(&stream[i]); //創建流
}
// Launch a kernel on the GPU with one thread for each element.
for(int i = 0;i<5;i++)
{
addKernel<<<1,1,0,stream[i]>>>(dev_c+i, dev_a+i, dev_b+i); //執行流
}
cudaDeviceSynchronize();
// 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:
for(int i = 0;i<5;i++)
{
cudaStreamDestroy(stream[i]); //銷毀流
}
cudaFree(dev_c);
cudaFree(dev_a);
cudaFree(dev_b);
return cudaStatus;
}
注意到，我們的核函數代碼仍然和塊並行的版本一樣，隻是在調用時做了改變，<<<>>>中的參數多了兩個，其中前兩個和塊並行、線程並行中的意義相同，仍然是線程塊數（這裏為1）、每個線程塊中線程數（這裏也是1）。第三個為0表示每個block用到的共享內存大小，這個我們後麵再講；第四個為流對象，表示當前核函數在哪個流上運行。我們創建了5個流，每個流上都裝載了一個核函數，同時傳遞參數有些不同，也就是每個核函數作用的對象也不同。這樣就實現了任務級別的並行，當我們有幾個互不相關的任務時，可以寫多個核函數，資源允許的情況下，我們將這些核函數裝載到不同流上，然後執行，這樣可以實現更粗粒度的並行。
好了，流並行就這麼簡單，我們處理任務時，可以根據需要，選擇最適合的並行方式。
最後更新：2017-04-03 16:48:42