深度學習的興起，使得多執行緒以及GPU程式設計逐漸成為演算法工程師無法規避的問題。這裡主要記錄自己的GPU自學歷程。

目錄

十、 流並行

我們前面學習的CUDA並行程式設計，基本上都是在一批資料上利用大量執行緒實現並行。 除此之外， NVIDIA系列GPU還支援另外一種型別的並行性 —— 流。

GPU中的流並行類似於CPU上的任務並行，即每個流都可以看作是一個獨立的任務，每個流中的程式碼操作順序執行。

下面從流並行的基礎到使用來說明。

10.1 頁鎖定記憶體

流並行的使用需要有硬體支援：即必須是支援裝置重疊功能的GPU。

通過下面的程式碼查詢裝置是否支援裝置重疊功能：

cudaDeviceProp mprop;
cudaGetDeviceProperties(&mprop,0);
if (!mprop.deviceOverlap)
{
    cout << "Device not support overlaps, so stream is invalid!" << endl;
}

只有支援裝置重疊，GPU在執行一個核函式的同時，才可以同時在裝置與主機之間執行復制操作。 當然，這種複製操作需要在一種特殊的記憶體上才可以進行 —— 頁鎖定記憶體。

• 頁鎖定記憶體： 需要由cudaHostAlloc()分配，又稱為固定記憶體（Pinned Memory）或者不可分頁記憶體。 作業系統將不會對這塊記憶體分頁並交換到磁碟上，從而確保了該記憶體始終駐留在實體記憶體中，因為這塊記憶體將不會被破壞或者重新定位。 由於gpu知道記憶體的實體地址，因此可以通過“直接記憶體訪問（Direct Memory Access，DMA）” 直接在gpu和主機之間複製資料。
• 可分頁記憶體： malloc()分配的記憶體是標準的、可分頁的（Pagable）主機記憶體。 可分頁記憶體面臨著重定位的問題，因此使用可分頁記憶體進行復制時，複製可能執行兩次操作：從可分頁記憶體複製到一塊“臨時”頁鎖定記憶體，然後從頁鎖定記憶體複製到GPU。

雖然在頁鎖定記憶體上執行復制操作效率比較高，但消耗實體記憶體更多。因此，通常對cudaMemcpy()呼叫的源記憶體或者目標記憶體才使用，而且使用完畢立即釋放。

10.2 流並行機制

流並行是指我們可以建立多個流來執行多個任務， 但每個流都是一個需要按照順序執行的操作佇列。 那麼我們如何實現程式加速？ 其核心就在於，在頁鎖定記憶體上的資料複製是獨立於核函式執行的，即我們可以在執行核函式的同時進行資料複製。

這裡的複製需要使用cudaMemcpyAsync()，一個以非同步執行的函式。呼叫cudaMemcpyAsync()時，只是放置一個請求，表示在流中執行一次記憶體複製操作。當函式返回時，我們無法確保複製操作已經結束。我們能夠得到的保證是，複製操作肯定會當下一個被放入流中的操作之前執行。(相比之下，cudaMemcpy()是一個同步執行函式。當函式返回時，複製操作已完成。)

以計算 a + b = c為例，假如我們建立了兩個流，每個流都是按順序執行：

複製a(主機到GPU) -> 複製b(主機到GPU) -> 核函式計算 -> 複製c(GPU到主機)

如上圖，複製操作和核函式執行是分開的，但由於每個流內部需要按順序執行，因此複製c的操作需要等待核函式執行完畢。 於是，整個程式執行的時間線如下圖：(箭頭表示需要等待)

從上面的時間線我們可以啟發式的思考下：如何調整每個流當中的操作順序來獲得最大的收益？ 提高重疊率。

如下圖所示，假如複製一份資料的時間和執行一次核函式的時間差不多，那麼我們可以採用交叉執行的策略：

由於流0的a和b已經準備完成，因此當複製流1的b時，可以同步執行流0的核函式。 這樣整個時間線，相較於之前的操作很明顯少掉了兩塊操作。

10.3 流並行示例

與流相關的常用函式如下：

// 建立與銷燬
cudaStream_t stream//定義流 
cudaStreamCreate(cudaStream_t * s)//建立流 
cudaStreamDestroy(cudaStream_t s)//銷燬流 

//同步 
cudaStreamSynchronize()//同步單個流：等待該流上的命令都完成 
cudaDeviceSynchronize()//同步所有流：等待整個裝置上流都完成 
cudaStreamWaitEvent()//等待某個事件結束後執行該流上的命令 
cudaStreamQuery()//查詢一個流任務是否完成 

//回撥 
cudaStreamAddCallback()//在任何點插入回撥函式 

//優先順序 
cudaStreamCreateWithPriority() 
cudaDeviceGetStreamPriorityRange()

下面給出一個2個流執行a + b = c的示例， 我們假設資料量非常大，需要將資料拆分，每次計算一部分。

#include <iostream>
#include "cuda_runtime.h"


using namespace std;

#define N   (1024*256)  // 每次處理的資料量
#define SIZE   (N*20)  //資料總量

// 核函式，a + b = c
__global__ void add(int* a, int* b, int* c)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N)
        c[i] = a[i] + b[i];
}

int main()
{
    // 獲取0號GPU的屬性並判斷是否支援裝置重疊功能
    cudaDeviceProp mprop;
    cudaGetDeviceProperties(&mprop,0);
    if (!mprop.deviceOverlap)
    {
        cout << "Device not support overlaps, so stream is invalid!" << endl;
        return 0;
    }

    // 建立計時事件
    cudaEvent_t     start, stop;
    cudaEventCreate(&start); cudaEventCreate(&stop);
    float           elapsedTime;

    // 建立流
    cudaStream_t    stream0, stream1;
    cudaStreamCreate(&stream0);
    cudaStreamCreate(&stream1);

    // 開闢主機頁鎖定記憶體，並隨機初始化資料
    int *host_a, *host_b, *host_c;
    cudaHostAlloc((void**)&host_a, SIZE*sizeof(int), cudaHostAllocDefault);
    cudaHostAlloc((void**)&host_b, SIZE*sizeof(int), cudaHostAllocDefault);
    cudaHostAlloc((void**)&host_c, SIZE*sizeof(int), cudaHostAllocDefault);
    for (size_t i = 0; i < SIZE; i++)
    {
        host_a[i] = rand();
        host_b[i] = rand();
    }

    // 宣告並開闢相關變數記憶體
    int *dev_a0, *dev_b0, *dev_c0;   //用於流0的資料
    int *dev_a1, *dev_b1, *dev_c1;   //用於流1的資料
    cudaMalloc((void**)&dev_a0,N*sizeof(int));
    cudaMalloc((void**)&dev_b0, N*sizeof(int));
    cudaMalloc((void**)&dev_c0, N*sizeof(int));
    cudaMalloc((void**)&dev_a1, N*sizeof(int));
    cudaMalloc((void**)&dev_b1, N*sizeof(int));
    cudaMalloc((void**)&dev_c1, N*sizeof(int));

    /************************  核心計算部分    ***************************/
    cudaEventRecord(start, 0);
    for (size_t i = 0; i < SIZE; i += 2*N)
    {

        // 複製流0資料a
        cudaMemcpyAsync(dev_a0, host_a + i,   N*sizeof(int), cudaMemcpyHostToDevice, stream0);
        // 複製流1資料a
        cudaMemcpyAsync(dev_a1, host_a + i+N, N*sizeof(int), cudaMemcpyHostToDevice, stream1);
        // 複製流0資料b
        cudaMemcpyAsync(dev_b0, host_b + i,   N*sizeof(int), cudaMemcpyHostToDevice, stream0);
        // 複製流1資料b
        cudaMemcpyAsync(dev_b1, host_b + i+N, N*sizeof(int), cudaMemcpyHostToDevice, stream1);
        // 執行流0核函式
        add << <N / 256, 256, 0, stream0 >> >(dev_a0, dev_b0, dev_c0);
        // 執行流1核函式
        add << <N / 256, 256, 0, stream1 >> >(dev_a1, dev_b1, dev_c1);
        // 複製流0資料c
        cudaMemcpyAsync(host_c + i*N,   dev_c0, N*sizeof(int), cudaMemcpyDeviceToHost, stream0);
        // 複製流1資料c
        cudaMemcpyAsync(host_c + i*N+N, dev_c1, N*sizeof(int), cudaMemcpyDeviceToHost, stream1);
    }

    // 流同步
    cudaStreamSynchronize(stream0);
    cudaStreamSynchronize(stream1);

    // 處理計時
    cudaEventSynchronize(stop);
    cudaEventRecord(stop, 0);
    cudaEventElapsedTime(&elapsedTime, start, stop);
    cout << "GPU time: " << elapsedTime << "ms" << endl;

    // 銷燬所有開闢的記憶體
    cudaFreeHost(host_a); cudaFreeHost(host_b); cudaFreeHost(host_c);
    cudaFree(dev_a0); cudaFree(dev_b0); cudaFree(dev_c0);
    cudaFree(dev_a1); cudaFree(dev_b1); cudaFree(dev_c1);

    // 銷燬流以及計時事件
    cudaStreamDestroy(stream0); cudaStreamDestroy(stream1);
    cudaEventDestroy(start);    cudaEventDestroy(stop);

    return 0;
}

參考資料

• 《CUDA by Example: An Introduction to General-Purpose GPU Programming》 中文名《GPU高效能程式設計CUDA實戰》