GPU 編程入門到精通（五）之 GPU 程序優化進階

本文轉載自查看原文 2015-12-26 19:37 5194

博主因為工作其中的須要，開始學習 GPU 上面的編程，主要涉及到的是基於 GPU 的深度學習方面的知識。鑒於之前沒有接觸過 GPU 編程。因此在這里特地學習一下 GPU 上面的編程。

有志同道合的小伙伴，歡迎一起交流和學習。我的郵箱: caijinping220@gmail.com 。使用的是自己的老古董筆記本上面的 Geforce 103m 顯卡，盡管顯卡相對於如今主流的系列已經很的弱，可是對於學習來說。還是能夠用的。本系列博文也遵從由簡單到復雜，記錄自己學習的過程。

0. 文件夾

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, 0>>>(gpuData, sum, time);

  這里邊。sum 數組的長度計算方式變化了，可是大小沒有變化。另在在調用 GPU 內核函數的時候，參數發生了變化。須要告訴 GPU block 數 和 thread 數。只是這邊共享內存沒有使用。

最后，在 CPU 中計算部分和

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

編譯執行以后。得到例如以下結果：

性能與直接使用 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 == 0) time[bid] = clock();

      shared[tid] = 0;
      // 把部分和結果放入共享內存中
      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 == 0) {
          for (int i = 1; i < THREAD_NUM; i++) {
              shared[0] += shared[i];
          }
          sum[bid] = shared[0];
      }

      if (tid == 0) 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 中額外執行求和的這部分代碼。導致了執行周期的變長，只是對應的，在 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 / 2;

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

    shared[tid] = 0;

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

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

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

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

主函數部分不許要改動，最后編譯執行結果例如以下:

性能有一部分的改善。

通過使用 GPU 的並行化編程。確實對性能會有非常大程度上的提升。

因為受限於 Geforce 103m 的內存帶寬，程序僅僅能優化到這一步，關於是否還有其它的方式優化，有待進一步學習。

4. 總結

通過這幾篇博文的討論，數組平方和的代碼優化到這一階段。

從但線程到多線程，再到共享內存，通過使用這幾種 GPU 上面的結構，做到了程序的優化。例如以下給出數組平方和的完整代碼：

/* ******************************************************************* ##### File Name: squareSum.cu ##### File Func: calculate the sum of inputs's square ##### Author: Caijinping ##### E-mail: caijinping220@gmail.com ##### 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 0;
    }
    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 * 2);
    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 * 2];
    cudaMemcpy(&result, sum, sizeof(int) * BLOCK_NUM, cudaMemcpyDeviceToHost);
    cudaMemcpy(&time_used, time, sizeof(clock_t) * BLOCK_NUM * 2, cudaMemcpyDeviceToHost);

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

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

    clock_t min_start, max_end;
    min_start = time_used[0];
    max_end = time_used[BLOCK_NUM];
    for (int i = 1; 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 = 0;

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

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

    return 0;
}

__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 / 2;

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

    shared[tid] = 0;

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

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

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

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


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[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
    printf("maxGridSize[0 - 2] : %d %d %d.\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
    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 == 0) {
        fprintf(stderr, "There is no device.\n");
        return false;
    }

    // find the device >= 1.X
    int i;
    for (i = 0; i < count; ++i) {
        cudaDeviceProp prop;
        if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
            if (prop.major >= 1) {
                //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 編程。
caijinping220@gmail.com
http://blog.csdn.net/xsc_c

免責聲明！

本站轉載的文章為個人學習借鑒使用，本站對版權不負任何法律責任。如果侵犯了您的隱私權益，請聯系本站郵箱yoyou2525@163.com刪除。