版權聲明：本文為博主原創文章，未經博主允許不得轉載。

目錄(?)[+]

博主由于工作當中的需要，開始學習 GPU 上面的編程，主要涉及到的是基于 GPU 的深度學習方面的知識，鑒于之前沒有接觸過 GPU 編程，因此在這里特地學習一下 GPU 上面的編程。有志同道合的小伙伴，歡迎一起交流和學習，我的郵箱:caijinping220@gmail.com 。使用的是自己的老古董筆記本上面的 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 GPUcudaMalloc((void**) &sum, sizeof(int) * THREAD_NUM * BLOCK_NUM);// calculate the squares's sumsquaresSum<<<BLOCK_NUM, THREAD_NUM, 0>>>(gpuData, sum, time); 這里邊，sum 數組的長度計算方式變化了，但是大小沒有變化。另在在調用 GPU 內核函數的時候，參數發生了變化，需要告訴 GPU block 數 和 thread 數。不過這邊共享內存沒有使用。

• 最后，在 CPU 中計算部分和

  // print resultint 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 sumsquaresSum<<<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 deviceif (!InitCUDA()) {return 0;}printf("CUDA initialized.\n");// generate rand datasgenerateData(data, DATA_SIZE);// malloc space for datas in GPUint *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 sumsquaresSum<<<BLOCK_NUM, THREAD_NUM, THREAD_NUM * sizeof(int)>>>(gpuData, sum, time);// copy the result from GPU to HOSTint 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 spacescudaFree(gpuData);cudaFree(sum);cudaFree(time);// print resultint 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 calculatetmp_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 numbersint count; // get the cuda device countcudaGetDeviceCount(&count);if (count == 0) {fprintf(stderr, "There is no device.\n");return false;}// find the device >= 1.Xint 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 deviceif (i == count) {fprintf(stderr, "There is no device supporting CUDA 1.x.\n");return false;}// set cuda device cudaSetDevice(i);return true; }

---

總結

以上是生活随笔為你收集整理的GPU 编程入门到精通（五）之 GPU 程序优化进阶的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。