版權(quán)聲明：本文為博主原創(chuàng)文章，未經(jīng)博主允許不得轉(zhuǎn)載。

目錄(?)[+]

博主由于工作當(dāng)中的需要，開始學(xué)習(xí) GPU 上面的編程，主要涉及到的是基于 GPU 的深度學(xué)習(xí)方面的知識，鑒于之前沒有接觸過 GPU 編程，因此在這里特地學(xué)習(xí)一下 GPU 上面的編程。有志同道合的小伙伴，歡迎一起交流和學(xué)習(xí)，我的郵箱:caijinping220@gmail.com 。使用的是自己的老古董筆記本上面的 Geforce 103m 顯卡，雖然顯卡相對于現(xiàn)在主流的系列已經(jīng)非常的弱，但是對于學(xué)習(xí)來說，還是可以用的。本系列博文也遵從由簡單到復(fù)雜，記錄自己學(xué)習(xí)的過程。

---

0. 目錄

• GPU 編程入門到精通（一）之 CUDA 環(huán)境安裝
• GPU 編程入門到精通（二）之 運行第一個程序
• GPU 編程入門到精通（三）之 第一個 GPU 程序
• GPU 編程入門到精通（四）之 GPU 程序優(yōu)化
• GPU 編程入門到精通（五）之 GPU 程序優(yōu)化進(jìn)階

1. 數(shù)組平方和并行化進(jìn)階

GPU 編程入門到精通（四）之 GPU 程序優(yōu)化 這篇博文中提到了 grid、block、thread 三者之間的關(guān)系，知道了他們之間是逐漸包含的關(guān)系。我們在上面的程序中通過使用 512 個線程達(dá)到了 493 倍左右的性能提升，那么是不是可以繼續(xù)得到提升呢？？？

答案是肯定的，這就要進(jìn)一步考慮 GPU 的并行化處理了。前面的程序只是使用了單個 block 下的 512 個線程，那么，我們可不可以使用多個 block 來實現(xiàn)？？？

對，就是利用這個思想，達(dá)到進(jìn)一步的并行化。這里使用 8 個 block * 64 threads = 512 threads 實現(xiàn)。

• 首先，修改主函數(shù)宏定義，定義塊數(shù)量：

  // ======== define area ========#define DATA_SIZE 1048576 // 1M#define BLOCK_NUM 8 // block num#define THREAD_NUM 64 // thread num 通過在程序中添加 block 和 threads 的宏定義，這兩個定義是我們在后面會用到的。他們決定了計算平方和使用的 CUDA 核心數(shù)。

• 接下來，修改內(nèi)核函數(shù)：

  _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;} 注意：這里的內(nèi)存遍歷方式和前面講的是一致的，理解一下。同時記錄的時間是一個塊的開始和結(jié)束時間，因為這里我們最后需要計算的是最早開始和最晚結(jié)束的兩個時間差，即求出最糟糕的時間。

• 然后，就是主函數(shù)里面的具體實現(xiàn)了：

  // 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 數(shù)組的長度計算方式變化了，但是大小沒有變化。另在在調(diào)用 GPU 內(nèi)核函數(shù)的時候，參數(shù)發(fā)生了變化，需要告訴 GPU block 數(shù) 和 thread 數(shù)。不過這邊共享內(nèi)存沒有使用。

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

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

編譯運行以后，得到如下結(jié)果：

性能與直接使用 512 個線程基本一致。因為受到 GPU 內(nèi)存帶寬的限制，GPU 編程入門到精通（四）之 GPU 程序優(yōu)化 中的優(yōu)化，已經(jīng)接近極限，所以通過 block 方式，效果不明顯。

2. 線程同步和共享內(nèi)存

前面的程序，計算求和的工作在 CPU 中完成，總共需要在 CPU 中做 512 次加法運算，那么有沒有辦法減少 CPU 中執(zhí)行加法的次數(shù)呢？？？

可以通過同步和共享內(nèi)存技術(shù)，實現(xiàn)在 GPU 上的 block 塊內(nèi)求取部分和，這樣最后只需要在 CPU 計算 16 個和就可以了。具體實現(xiàn)方法如下：

• 首先，在修改內(nèi)核函數(shù)，定義一塊共享內(nèi)存，用 __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;// 把部分和結(jié)果放入共享內(nèi)存中for (int i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {shared[tid] += data[i] * data[i];}// 同步操作，必須等之前的線程都運行結(jié)束，才能繼續(xù)后面的程序__syncthreads();// 同步完成之后，將部分和加到 shared[0] 上面，這里全都在一個線程內(nèi)完成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 之間對于共享內(nèi)存是共享的，利用的是 GPU 上的內(nèi)存，所以速度很快，不必?fù)?dān)心 latency 的問題。__syncthreads() 函數(shù)是 CUDA 的內(nèi)部函數(shù)，表示所有 threads 都必須同步到這個點，才會執(zhí)行接下來的代碼。我們要做的就是等待每個 thread 計算結(jié)束以后，再來計算部分和，所以同步是必不可少的環(huán)節(jié)。把每個 block 的部分和計算到 shared[0] 里面。

• 接下來，修改 main 函數(shù):

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

  編譯運行后結(jié)果如下：

  其實和前一版程序相比，時間上沒有什么優(yōu)勢，原因在于，我們需要在 GPU 中額外運行求和的這部分代碼，導(dǎo)致了運行周期的變長，不過相應(yīng)的，在 CPU 中的運行時間會減少。

3. 加法樹

我們在這個程序中，只當(dāng)每個 block 的 thread0 的時候，計算求和的工作，這樣做影響了執(zhí)行的效率，其實求和可以并行化處理的，也就是通過加法樹來實現(xiàn)并行化。舉個例子，要計算 8 個數(shù)的和，我們沒必要用一個 for 循環(huán)，逐個相加，而是可以通過第一級流水線實現(xiàn)兩兩相加，變成 4 個數(shù)，第二級流水實現(xiàn)兩兩相加，變成 2 個數(shù)，第三級流水實現(xiàn)兩兩相加，求得最后的和。

下面通過加法樹的方法，實現(xiàn)最后的求和，修改內(nèi)核函數(shù)如下：

__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();} } 此程序?qū)崿F(xiàn)的就是上訴描述的加法樹的結(jié)構(gòu)，注意這里第二個 __syncthreads() 的使用，也就是說，要進(jìn)行下一級流水線的計算，必須建立在前一級必須已經(jīng)計算完畢的情況下。

主函數(shù)部分不許要修改，最后編譯運行結(jié)果如下:

性能有一部分的改善。

通過使用 GPU 的并行化編程，確實對性能會有很大程度上的提升。由于受限于 Geforce 103m 的內(nèi)存帶寬，程序只能優(yōu)化到這一步，關(guān)于是否還有其他的方式優(yōu)化，有待進(jìn)一步學(xué)習(xí)。

4. 總結(jié)

通過這幾篇博文的討論，數(shù)組平方和的代碼優(yōu)化到這一階段。從但線程到多線程，再到共享內(nèi)存，通過使用這幾種 GPU 上面的結(jié)構(gòu)，做到了程序的優(yōu)化。如下給出數(shù)組平方和的完整代碼：

/* ******************************************************************* ##### 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; }

---

總結(jié)

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

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯，歡迎將生活随笔推薦給好友。