GPU程式設計(四):並行規約優化

目錄

• 前言
• cuda-gdb
• 未優化並行規約
• 優化後並行規約
• 結果分析
• 最後

---

前言

• 之前第三篇也看到了, 並行方面GPU真的是無往不利, 現在再看下第二個例子, 並行規約. 通過這次的例子會發現, 需要了解GPU架構, 然後寫出與之對應的演算法的, 兩者結合才能得到令人驚歎的結果.
• 這次也會簡要介紹下cuda-gdb的用法, 其實和gdb用法幾乎一樣, 也就是多了個cuda命令.

---

cuda-gdb

如果之前沒有用過gdb, 可以速學一下, 就幾個指令.
想要用cuda-gdb對程式進行除錯, 首先你要確保你的gpu沒有在執行作業系統介面, 比方說, 我用的是ubuntu, 我就需要用sudo service lightdm stop關閉圖形介面, 進入tty1這種字元介面.
當然用ssh遠端訪問也是可以的.
接下來, 使用第二篇中矩陣加法的例子. 但是注意, 編譯的使用需要改變一下, 加入-g -G引數, 其實和gdb是相似的.

nvcc -g -G CUDAAdd.cu -o CUDAAdd.o

然後使用cuda-gdb CUDAAdd.o即可對程式進行除錯.

在除錯之前, 我把程式碼貼出來:

#include <stdio.h>

__global__ void add(float * x, float *y, float * z, int n){
        int index = threadIdx.x + blockIdx.x * blockDim.x;
        int stride = blockDim.x * gridDim.x;

        for (int i = index; i < n; i += stride){
                z[i] = x[i] + y[i];
        }
}

int main()
{
    int N = 1 << 20;
    int nBytes = N * sizeof(float);

    float *x, *y, *z;
    cudaMallocManaged((void**)&x, nBytes);
    cudaMallocManaged((void**)&y, nBytes);
    cudaMallocManaged((void**)&z, nBytes);

    for (int i = 0; i < N; ++i)
    {
        x[i] = 10.0;
        y[i] = 20.0;
    }

    dim3 blockSize(256);
    // 4096
    dim3 gridSize((N + blockSize.x - 1) / blockSize.x);

    add << < gridSize, blockSize >> >(x, y, z, N);

    cudaDeviceSynchronize();

    float maxError = 0.0;
    for (int i = 0; i < N; i++){
                maxError = fmax(maxError, (float)(fabs(z[i] - 30.0)));
    }
    printf ("max default: %.4f
", maxError);

    cudaFree(x);
    cudaFree(y);
    cudaFree(z);

    return 0;
}

之後就是常規操作了, 新增斷點, 執行, 下一步, 檢視想看的資料. 不同點是cuda的指令, 例如cuda block(1,0,0)可以從一開始block(0,0,0)切換到block(1,0,0).

---

未優化並行規約

如果按照常規的思路, 兩兩進行進行加法運算. 每次步長翻倍即可, 從演算法的角度來說, 這是沒啥問題的. 但是沒有依照GPU架構進行設計.

#include <stdio.h>

const int    threadsPerBlock = 512;
const int    N        = 2048;
const int    blocksPerGrid    = (N + threadsPerBlock - 1) / threadsPerBlock; /* 4 */

__global__ void ReductionSum( float * d_a, float * d_partial_sum )
{
    /* 申請共享記憶體, 存在於每個block中 */
    __shared__ float partialSum[threadsPerBlock];

    /* 確定索引 */
    int    i    = threadIdx.x + blockIdx.x * blockDim.x;
    int    tid    = threadIdx.x;

    /* 傳global memory資料到shared memory */
    partialSum[tid] = d_a[i];

    /* 傳輸同步 */
    __syncthreads();

    /* 在共享儲存器中進行規約 */
    for ( int stride = 1; stride < blockDim.x; stride *= 2 )
    {
        if ( tid % (2 * stride) == 0 )
            partialSum[tid] += partialSum[tid + stride];
        __syncthreads();
    }

    /* 將當前block的計算結果寫回輸出陣列 */
    if ( tid == 0 )
        d_partial_sum[blockIdx.x] = partialSum[0];
}


int main()
{
    int size = sizeof(float);

    /* 分配視訊記憶體空間 */
    float    * d_a;
    float    * d_partial_sum;

    cudaMallocManaged( (void * *) &d_a, N * size );
    cudaMallocManaged( (void * *) &d_partial_sum, blocksPerGrid * size );

    for ( int i = 0; i < N; ++i )
        d_a[i] = i;

    /* 呼叫核心函式 */
    ReductionSum << < blocksPerGrid, threadsPerBlock >> > (d_a, d_partial_sum);

    cudaDeviceSynchronize();

    /* 將部分和求和 */
    int sum = 0;
    for ( int i = 0; i < blocksPerGrid; ++i )
        sum += d_partial_sum[i];

    printf( "sum = %d
", sum );

    /* 釋放視訊記憶體空間 */
    cudaFree( d_a );
    cudaFree( d_partial_sum );

    return(0);
}

---

優化後並行規約

其實需要改動的地方非常小, 改變步長即可.

__global__ void ReductionSum( float * d_a, float * d_partial_sum )
{
    // 相同, 略去
    /* 在共享儲存器中進行規約 */
    for ( int stride = blockDim.x / 2; stride > 0; stride /= 2 )
    {
        if ( tid < stride )
            partialSum[tid] += partialSum[tid + stride];
        __syncthreads();
    }
    // 相同, 略去
}

---

結果分析

之前的文章裡面也說過warp.
warp: GPU執行程式時的排程單位, 目前cuda的warp的大小為32, 同在一個warp的執行緒, 以不同資料資源執行相同的指令, 這就是所謂SIMT.
說人話就是, 這32個執行緒必須要幹相同的事情, 如果有執行緒動作不一致, 就需要等待一波執行緒完成自己的工作, 然後再去做另外一件事情.
所以, 用圖說話就是, 第二種方案可以更快將warp閒置, 交給GPU排程, 所以, 肯定是第二種更快.

圖一在運算依次之後, 沒有warp可以空閒, 而圖二直接空閒2個warp. 圖一到了第二次可以空閒2個warp, 而圖二已經空閒3個warp. 我這副圖只是示意圖, 如果是實際的, 差距會更大.

所以來看下執行耗時, 會發現差距還是很大的, 幾乎是差了一倍. 不過GPU確實算力太猛, 這樣看還不太明顯, 有意放大資料量會更加明顯.

最後

所以GPU又一次展示了強大的算力, 而且, 這次也看到了只是小小變動, 讓演算法更貼合架構, 就讓運算耗時減半, 所以在優化方面可以做的工作真的是太多了, 之後還有更多優化相關的文章, 有意見或者建議, 評論區見哦~

---