GPU 程式設計入門到精通（五）之 GPU 程式優化進階

部落客由于工作當中的需要，開始學習 GPU 上面的程式設計，主要涉及到的是基于 GPU 的深度學習方面的知識，鑒于之前沒有接觸過 GPU 程式設計，是以在這裡特地學習一下 GPU 上面的程式設計。有志同道合的小夥伴，歡迎一起交流和學習，我的郵箱: [email protected] 。使用的是自己的老古董筆記本上面的 Geforce 103m 顯示卡，雖然顯示卡相對于現在主流的系列已經非常的弱，但是對于學習來說，還是可以用的。本系列博文也遵從由簡單到複雜，記錄自己學習的過程。

0. 目錄

GPU 程式設計入門到精通（一）之 CUDA 環境安裝
GPU 程式設計入門到精通（二）之運作第一個程式
GPU 程式設計入門到精通（三）之第一個 GPU 程式
GPU 程式設計入門到精通（四）之 GPU 程式優化
GPU 程式設計入門到精通（五）之 GPU 程式優化進階

1. 數組平方和并行化進階

GPU 程式設計入門到精通（四）之 GPU 程式優化這篇博文中提到了 grid、block、thread 三者之間的關系，知道了他們之間是逐漸包含的關系。我們在上面的程式中通過使用 512 個線程達到了 493 倍左右的性能提升，那麼是不是可以繼續得到提升呢？？？

答案是肯定的，這就要進一步考慮 GPU 的并行化處理了。前面的程式隻是使用了單個 block 下的 512 個線程，那麼，我們可不可以使用多個 block 來實作？？？

對，就是利用這個思想，達到進一步的并行化。這裡使用 8 個 block * 64 threads = 512 threads 實作。

首先，修改主函數宏定義，定義塊數量：

// ======== define area ========
  #define DATA_SIZE 1048576    // 1M
  #define BLOCK_NUM 8        // block num
  #define THREAD_NUM 64        // thread num

通過在程式中添加 block 和 threads 的宏定義，這兩個定義是我們在後面會用到的。他們決定了計算平方和使用的 CUDA 核心數。

接下來，修改核心函數：

_global__ static void squaresSum(int *data, int *sum, clock_t *time)
  {
      const int tid = threadIdx.x;
      const int bid = blockIdx.x;

      for (int i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {
          tmp_sum += data[i] * data[i];
      }

      sum[bid * THREAD_NUM + tid] = tmp_sum;
  }

注意：這裡的記憶體周遊方式和前面講的是一緻的，了解一下。

  同時記錄的時間是一個塊的開始和結束時間，因為這裡我們最後需要計算的是最早開始和最晚結束的兩個時間差，即求出最糟糕的時間。

然後，就是主函數裡面的具體實作了：

// malloc space for datas in GPU
  cudaMalloc((void**) &sum, sizeof(int) * THREAD_NUM * BLOCK_NUM);

  // calculate the squares's sum
  squaresSum<<<BLOCK_NUM, THREAD_NUM, >>>(gpuData, sum, time);

這裡邊，sum 數組的長度計算方式變化了，但是大小沒有變化。另在在調用 GPU 核心函數的時候，參數發生了變化，需要告訴 GPU block 數 和 thread 數。不過這邊共享記憶體沒有使用。

最後，在 CPU 中計算部分和

// print result
  int tmp_result = ;
  for (int i = ; i < THREAD_NUM * BLOCK_NUM; ++i) {
      tmp_result += result[i];
  }

編譯運作以後，得到如下結果：

GPU 程式設計入門到精通（五）之 GPU 程式優化進階
性能與直接使用 512 個線程基本一緻。因為受到 GPU 記憶體帶寬的限制，GPU 程式設計入門到精通（四）之 GPU 程式優化中的優化，已經接近極限，是以通過 block 方式，效果不明顯。

2. 線程同步和共享記憶體

前面的程式，計算求和的工作在 CPU 中完成，總共需要在 CPU 中做 512 次加法運算，那麼有沒有辦法減少 CPU 中執行加法的次數呢？？？

可以通過同步和共享記憶體技術，實作在 GPU 上的 block 塊内求取部分和，這樣最後隻需要在 CPU 計算 16 個和就可以了。具體實作方法如下：

首先，在修改核心函數，定義一塊共享記憶體，用

 __shared__

訓示：

__global__ static void squaresSum(int *data, int *sum, clock_t *time)
  {
      // define of shared memory
      __shared__ int shared[BLOCK_NUM];

      const int tid = threadIdx.x;
      const int bid = blockIdx.x;

      if (tid == ) time[bid] = clock();

      shared[tid] = ;
      // 把部分和結果放入共享記憶體中
      for (int i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {
          shared[tid] += data[i] * data[i];
      }
      // 同步操作，必須等之前的線程都運作結束，才能繼續後面的程式
      __syncthreads();
      // 同步完成之後，将部分和加到 shared[0] 上面，這裡全都在一個線程内完成
      if (tid == ) {
          for (int i = ; i < THREAD_NUM; i++) {
              shared[] += shared[i];
          }
          sum[bid] = shared[];
      }

      if (tid == ) time[bid + BLOCK_NUM] = clock();
  }

利用 __shared__ 聲明的變量是 shared memory，每個 block 中，各個 thread 之間對于共享記憶體是共享的，利用的是 GPU 上的記憶體，是以速度很快，不必擔心 latency 的問題。

  __syncthreads() 函數是 CUDA 的内部函數，表示所有 threads 都必須同步到這個點，才會執行接下來的代碼。我們要做的就是等待每個 thread 計算結束以後，再來計算部分和，是以同步是必不可少的環節。把每個 block 的部分和計算到 shared[0] 裡面。

接下來，修改 main 函數:
```
// calculate the squares's sum
  squaresSum<<<BLOCK_NUM, THREAD_NUM, THREAD_NUM * sizeof(int)>>>(gpuData, sum, time);
           
```
編譯運作後結果如下：

GPU 程式設計入門到精通（五）之 GPU 程式優化進階
其實和前一版程式相比，時間上沒有什麼優勢，原因在于，我們需要在 GPU 中額外運作求和的這部分代碼，導緻了運作周期的變長，不過相應的，在 CPU 中的運作時間會減少。

3. 加法樹

我們在這個程式中，隻當每個 block 的 thread0 的時候，計算求和的工作，這樣做影響了執行的效率，其實求和可以并行化處理的，也就是通過加法樹來實作并行化。舉個例子，要計算 8 個數的和，我們沒必要用一個 for 循環，逐個相加，而是可以通過第一級流水線實作兩兩相加，變成 4 個數，第二級流水實作兩兩相加，變成 2 個數，第三級流水實作兩兩相加，求得最後的和。

下面通過加法樹的方法，實作最後的求和，修改核心函數如下：

__global__ static void squaresSum(int *data, int *sum, clock_t *time)
{
    __shared__ int shared[BLOCK_NUM];
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;

    int offset = THREAD_NUM / ;

    if (tid == ) time[bid] = clock();

    shared[tid] = ;

    for (int i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {
        shared[tid] += data[i] * data[i];
    }

    __syncthreads();
    while (offset > ) {
        if (tid < offset) {
            shared[tid] += shared[tid + offset];
        }
        offset >>= ;
        __syncthreads();
    }

    if (tid == ) {
        sum[bid] = shared[];
        time[bid + BLOCK_NUM] = clock();
    }
}

此程式實作的就是上訴描述的加法樹的結構，注意這裡第二個 __syncthreads() 的使用，也就是說，要進行下一級流水線的計算，必須建立在前一級必須已經計算完畢的情況下。

主函數部分不許要修改，最後編譯運作結果如下:

GPU 程式設計入門到精通（五）之 GPU 程式優化進階
性能有一部分的改善。

通過使用 GPU 的并行化程式設計，确實對性能會有很大程度上的提升。由于受限于 Geforce 103m 的記憶體帶寬，程式隻能優化到這一步，關于是否還有其他的方式優化，有待進一步學習。

4. 總結

通過這幾篇博文的讨論，數組平方和的代碼優化到這一階段。從但線程到多線程，再到共享記憶體，通過使用這幾種 GPU 上面的結構，做到了程式的優化。如下給出數組平方和的完整代碼：

/* *******************************************************************
##### File Name: squareSum.cu
##### File Func: calculate the sum of inputs's square
##### Author: Caijinping
##### E-mail: [email protected]
##### Create Time: 2014-5-7
* ********************************************************************/

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>

// ======== define area ========
#define DATA_SIZE 1048576    // 1M
#define BLOCK_NUM 8            // block num
#define THREAD_NUM 64        // thread num

// ======== global area ========
int data[DATA_SIZE];

void printDeviceProp(const cudaDeviceProp &prop);
bool InitCUDA();
void generateData(int *data, int size);
__global__ static void squaresSum(int *data, int *sum, clock_t *time);

int main(int argc, char const *argv[])
{
    // init CUDA device
    if (!InitCUDA()) {
        return ;
    }
    printf("CUDA initialized.\n");

    // generate rand datas
    generateData(data, DATA_SIZE);

    // malloc space for datas in GPU
    int *gpuData, *sum;
    clock_t *time;
    cudaMalloc((void**) &gpuData, sizeof(int) * DATA_SIZE);
    cudaMalloc((void**) &sum, sizeof(int) * BLOCK_NUM);
    cudaMalloc((void**) &time, sizeof(clock_t) * BLOCK_NUM * );
    cudaMemcpy(gpuData, data, sizeof(int) * DATA_SIZE, cudaMemcpyHostToDevice);

    // calculate the squares's sum
    squaresSum<<<BLOCK_NUM, THREAD_NUM, THREAD_NUM * sizeof(int)>>>(gpuData, sum, time);

    // copy the result from GPU to HOST
    int result[BLOCK_NUM];
    clock_t time_used[BLOCK_NUM * ];
    cudaMemcpy(&result, sum, sizeof(int) * BLOCK_NUM, cudaMemcpyDeviceToHost);
    cudaMemcpy(&time_used, time, sizeof(clock_t) * BLOCK_NUM * , cudaMemcpyDeviceToHost);

    // free GPU spaces
    cudaFree(gpuData);
    cudaFree(sum);
    cudaFree(time);

    // print result
    int tmp_result = ;
    for (int i = ; i < BLOCK_NUM; ++i) {
        tmp_result += result[i];
    }

    clock_t min_start, max_end;
    min_start = time_used[];
    max_end = time_used[BLOCK_NUM];
    for (int i = ; i < BLOCK_NUM; ++i)    {
        if (min_start > time_used[i]) min_start = time_used[i];
        if (max_end < time_used[i + BLOCK_NUM]) max_end = time_used[i + BLOCK_NUM];
    }
    printf("(GPU) sum:%d time:%ld\n", tmp_result, max_end - min_start);

    // CPU calculate
    tmp_result = ;

    for (int i = ; i < DATA_SIZE; ++i)    {
        tmp_result += data[i] * data[i];
    }

    printf("(CPU) sum:%d\n", tmp_result);

    return ;
}

__global__ static void squaresSum(int *data, int *sum, clock_t *time)
{
    __shared__ int shared[BLOCK_NUM];
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;

    int offset = THREAD_NUM / ;

    if (tid == ) time[bid] = clock();

    shared[tid] = ;

    for (int i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {
        shared[tid] += data[i] * data[i];
    }

    __syncthreads();
    while (offset > ) {
        if (tid < offset) {
            shared[tid] += shared[tid + offset];
        }
        offset >>= ;
        __syncthreads();
    }

    if (tid == ) {
        sum[bid] = shared[];
        time[bid + BLOCK_NUM] = clock();
    }
}

// ======== used to generate rand datas ========
void generateData(int *data, int size)
{
    for (int i = ; i < size; ++i) {
        data[i] = rand() % ;
    }
}


void printDeviceProp(const cudaDeviceProp &prop)
{
    printf("Device Name : %s.\n", prop.name);
    printf("totalGlobalMem : %d.\n", prop.totalGlobalMem);
    printf("sharedMemPerBlock : %d.\n", prop.sharedMemPerBlock);
    printf("regsPerBlock : %d.\n", prop.regsPerBlock);
    printf("warpSize : %d.\n", prop.warpSize);
    printf("memPitch : %d.\n", prop.memPitch);
    printf("maxThreadsPerBlock : %d.\n", prop.maxThreadsPerBlock);
    printf("maxThreadsDim[0 - 2] : %d %d %d.\n", prop.maxThreadsDim[], prop.maxThreadsDim[], prop.maxThreadsDim[]);
    printf("maxGridSize[0 - 2] : %d %d %d.\n", prop.maxGridSize[], prop.maxGridSize[], prop.maxGridSize[]);
    printf("totalConstMem : %d.\n", prop.totalConstMem);
    printf("major.minor : %d.%d.\n", prop.major, prop.minor);
    printf("clockRate : %d.\n", prop.clockRate);
    printf("textureAlignment : %d.\n", prop.textureAlignment);
    printf("deviceOverlap : %d.\n", prop.deviceOverlap);
    printf("multiProcessorCount : %d.\n", prop.multiProcessorCount);
}

bool InitCUDA()
{
    //used to count the device numbers
    int count;    

    // get the cuda device count
    cudaGetDeviceCount(&count);
    if (count == ) {
        fprintf(stderr, "There is no device.\n");
        return false;
    }

    // find the device >= 1.X
    int i;
    for (i = ; i < count; ++i) {
        cudaDeviceProp prop;
        if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
            if (prop.major >= ) {
                //printDeviceProp(prop);
                break;
            }
        }
    }

    // if can't find the device
    if (i == count) {
        fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
        return false;
    }

    // set cuda device 
    cudaSetDevice(i);

    return true;
}

歡迎大家和我一起讨論和學習 GPU 程式設計。

[email protected]

http://blog.csdn.net/xsc_c

GPU 程式設計入門到精通（五）之 GPU 程式優化進階

0. 目錄

1. 數組平方和并行化進階

2. 線程同步和共享記憶體

3. 加法樹

4. 總結

繼續閱讀

【CUDA-C/C++】任意次元矩陣乘

WSL2 的docker裡使用顯示卡的安裝

關于安裝pytorch的一些問題總結

jetson nano ubuntu 安裝opencv4 cuda10 pytorch

Jetson Nano Ubuntu編譯OpenCV4.4.0+opencv_contrib(帶CUDA)

TensorFlow錯題集

Android 過度繪制簡單的分析概念對比分析參考文章

pyTorch 使用多GPU訓練

【cpu or gpu】【tensorflow】怎麼檢視用的是CPU還是GPU

大學、碩士、博士的差別是什麼？

yolov7 tensorrt模型加速部署【實戰】

linxu下CUDA靜态庫-上

《cuda并行程式設計》勘誤（3）

《cuda并行程式設計》勘誤（2）

ImportError: libcublas.so.9.0: cannot open shared object file: No such file or directory（完美解決）

一種解決思路： ImportError: libcublas.so.10.0: cannot open shared object file: No such file