NO IMAGE

1. 使用 Thrust

Thrust 是一個開源的 C 庫,用於開發高效能並行應用程式,以 C 標準模板庫為藍本實現。

官方文件見這裡:CUDA Thrust

/* … */
float *fMatrix_Device; // 指向裝置視訊記憶體
int iMatrixSize = iRow * iCol; // 矩陣元素個數

cudaMalloc((void**)&fMatrix_Device, iMatrixSize * sizeof(float)); // 在視訊記憶體中為矩陣開闢空間
cudaMemcpy(fMatrix_Device, fMatrix_Host, iMatrixSize * sizeof(float), cudaMemcpyHostToDevice); // 將資料拷貝到視訊記憶體

thrust::device_ptr<float> dev_ptr(fMatrix_Device);
float thrustResult = thrust::reduce(dev_ptr, dev_ptr size_t(iMatrixSize), (float)0, thrust::plus<float>());

其中,fMatrix_Host 為指向主機記憶體的矩陣的頭指標。

2. 我的 Reduction

/**
* 每個 warp 自動同步,不用 __syncthreads();
* volatile : 加上關鍵字volatile的變數將被定義為敏感變數,意思是加了volatile
*            的變數在記憶體中的值可能會隨時發生變化,當程式要去讀取這個變數時,
             必須要從記憶體中讀取,而不是從快取中讀取
* sdata  陣列頭指標,陣列位於共享記憶體
* tid    執行緒索引
*/
__device__ void warpReduce(volatile float *sdata, int tid)
{
    sdata[tid] = sdata[tid 32];
    sdata[tid] = sdata[tid 16];
    sdata[tid] = sdata[tid 8];
    sdata[tid] = sdata[tid 4];
    sdata[tid] = sdata[tid 2];
    sdata[tid] = sdata[tid 1];
}

/**
* 優化:解決了 reduce3 中存在的多餘同步操作(每個warp預設自動同步)。
* globalInputData  輸入資料,位於全域性記憶體
* globalOutputData 輸出資料,位於全域性記憶體
*/
__global__ void reduce4(float *globalInputData, float *globalOutputData, unsigned int n)
{
    __shared__ float sdata[BLOCK_SIZE];

    // 座標索引
    unsigned int tid = threadIdx.x;
    unsigned int index = blockIdx.x*(blockDim.x * 2) threadIdx.x;
    unsigned int indexWithOffset = index blockDim.x;

    if (index >= n) sdata[tid] = 0;
    else if (indexWithOffset >= n) sdata[tid] = globalInputData[index];
    else sdata[tid] = globalInputData[index] globalInputData[indexWithOffset];

    __syncthreads();

    // 在共享記憶體中對每一個塊進行規約計算
    for (unsigned int s = blockDim.x / 2; s>32; s >>= 1)
    {
        if (tid < s) sdata[tid] = sdata[tid s];

        __syncthreads();
    }
    if (tid < 32) warpReduce(sdata, tid);

    // 把計算結果從共享記憶體寫回全域性記憶體
    if (tid == 0) globalOutputData[blockIdx.x] = sdata[0];
}

/**
* 計算 reduce4 函式的時間
* fMatrix_Host  矩陣頭指標
* iRow          矩陣行數
* iCol          矩陣列數
* @return       和
*/
float RuntimeOfReduce4(float *fMatrix_Host, const int iRow, const int iCol)
{
    float *fReuslt = (float*)malloc(sizeof(float));;
    float *fMatrix_Device; // 指向裝置視訊記憶體
    int iMatrixSize = iRow * iCol; // 矩陣元素個數

    cudaMalloc((void**)&fMatrix_Device, iMatrixSize * sizeof(float)); // 在視訊記憶體中為矩陣開闢空間
    cudaMemcpy(fMatrix_Device, fMatrix_Host, iMatrixSize * sizeof(float), cudaMemcpyHostToDevice); // 將資料拷貝到視訊記憶體

    /* … */
    for (int i = 1, int iNum = iMatrixSize; i < iMatrixSize; i = 2 * i * BLOCK_SIZE)
    {
        int iBlockNum = (iNum (2 * BLOCK_SIZE) – 1) / (2 * BLOCK_SIZE);
        reduce4<<<iBlockNum, BLOCK_SIZE>>>(fMatrix_Device, fMatrix_Device, iNum);
        iNum = iBlockNum;
    }

    cudaMemcpy(fReuslt, fMatrix_Device, sizeof(float), cudaMemcpyDeviceToHost); // 將資料拷貝到記憶體

    /* … */
    cudaFree(fMatrix_Device);// 釋放視訊記憶體空間

    return fReuslt[0];
}

上述程式是優化的最終版本,優化的主要內容包括: 

1. 避免每個 Warp 中出現分支導致效率低下。 
2. 減少取餘操作。 
3. 減小不必要的同步操作,每個warp都是預設同步的,不用額外的同步操作。 
4. 減小執行緒的閒置,提高並行度

3. 時間對比

資料的大小為:

iRow = 1000; 
iCol = 1000;

時間為:

ReduceThrust 的執行時間為:0.179968ms.
494497
Reduce0 的執行時間為:0.229152ms.
494497
Reduce1 的執行時間為:0.134816ms.
494497
Reduce2 的執行時間為:0.117504ms.
494497
Reduce3 的執行時間為:0.086016ms.
494497
Reduce4 的執行時間為:0.07424ms.
494497
CPU的執行時間為:1 ms.
494497

資料的大小為:

iRow = 2000; 
iCol = 2000;

時間為:

ReduceThrust 的執行時間為:0.282944ms.
1.97828e 006
Reduce0 的執行時間為:0.779776ms.
1.97828e 006
Reduce1 的執行時間為:0.42624ms.
1.97828e 006
Reduce2 的執行時間為:0.343744ms.
1.97828e 006
Reduce3 的執行時間為:0.217248ms.
1.97828e 006
Reduce4 的執行時間為:0.160416ms.
1.97828e 006
CPU的執行時間為:3 ms.
1.97828e 006

資料的大小為:

iRow = 4000; 
iCol = 4000;

時間為:

ReduceThrust 的執行時間為:0.536832ms.
7.91319e 006
Reduce0 的執行時間為:2.9919ms.
7.91319e 006
Reduce1 的執行時間為:1.56054ms.
7.91319e 006
Reduce2 的執行時間為:1.26618ms.
7.91319e 006
Reduce3 的執行時間為:0.726016ms.
7.91319e 006
Reduce4 的執行時間為:0.531712ms.
7.91319e 006
CPU的執行時間為:11 ms.
7.91319e 006

資料的大小為:

iRow = 6000; 
iCol = 6000;

時間為:

ReduceThrust 的執行時間為:0.988992ms.
1.7807e 007
Reduce4 的執行時間為:1.09286ms.
1.7807e 007
CPU的執行時間為:25 ms.
1.7807e 007

資料的大小為:

iRow = 11000; 
iCol = 11000;

時間為:

ReduceThrust 的執行時間為:2.9208ms.
5.98583e 007
Reduce4 的執行時間為:3.36998ms.
5.98583e 007
CPU的執行時間為:85 ms.
5.98583e 007

從上可以看出,2 中介紹的幾種優化方式取得了良好的效果;另外,當資料量較少時,我自己優化的規約函式比 Thrust 中的規約更高效,但是當資料量大於 4000 * 4000 時,Thrust 更高效,因此還有優化的空間。

4. 完整程式碼

GitHub