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

目錄(?)[+]

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

---

0. 目錄

• GPU 編程入門到精通（一）之 CUDA 環(huán)境安裝
• GPU 編程入門到精通（二）之 運(yùn)行第一個(gè)程序
• GPU 編程入門到精通（三）之 第一個(gè) 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)系。我們?cè)谏厦娴某绦蛑型ㄟ^(guò)使用 512 個(gè)線程達(dá)到了 493 倍左右的性能提升，那么是不是可以繼續(xù)得到提升呢？？？

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

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

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

  // ======== define area ========#define DATA_SIZE 1048576 // 1M#define BLOCK_NUM 8 // block num#define THREAD_NUM 64 // thread num 通過(guò)在程序中添加 block 和 threads 的宏定義，這兩個(gè)定義是我們?cè)诤竺鏁?huì)用到的。他們決定了計(jì)算平方和使用的 CUDA 核心數(shù)。

• 接下來(lái)，修改內(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)存遍歷方式和前面講的是一致的，理解一下。同時(shí)記錄的時(shí)間是一個(gè)塊的開(kāi)始和結(jié)束時(shí)間，因?yàn)檫@里我們最后需要計(jì)算的是最早開(kāi)始和最晚結(jié)束的兩個(gè)時(shí)間差，即求出最糟糕的時(shí)間。

• 然后，就是主函數(shù)里面的具體實(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ù)組的長(zhǎng)度計(jì)算方式變化了，但是大小沒(méi)有變化。另在在調(diào)用 GPU 內(nèi)核函數(shù)的時(shí)候，參數(shù)發(fā)生了變化，需要告訴 GPU block 數(shù) 和 thread 數(shù)。不過(guò)這邊共享內(nèi)存沒(méi)有使用。

• 最后，在 CPU 中計(jì)算部分和

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

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

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

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

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

可以通過(guò)同步和共享內(nèi)存技術(shù)，實(shí)現(xiàn)在 GPU 上的 block 塊內(nèi)求取部分和，這樣最后只需要在 CPU 計(jì)算 16 個(gè)和就可以了。具體實(shí)現(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];}// 同步操作，必須等之前的線程都運(yùn)行結(jié)束，才能繼續(xù)后面的程序__syncthreads();// 同步完成之后，將部分和加到 shared[0] 上面，這里全都在一個(gè)線程內(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，每個(gè) block 中，各個(gè) thread 之間對(duì)于共享內(nèi)存是共享的，利用的是 GPU 上的內(nèi)存，所以速度很快，不必?fù)?dān)心 latency 的問(wèn)題。__syncthreads() 函數(shù)是 CUDA 的內(nèi)部函數(shù)，表示所有 threads 都必須同步到這個(gè)點(diǎn)，才會(huì)執(zhí)行接下來(lái)的代碼。我們要做的就是等待每個(gè) thread 計(jì)算結(jié)束以后，再來(lái)計(jì)算部分和，所以同步是必不可少的環(huán)節(jié)。把每個(gè) block 的部分和計(jì)算到 shared[0] 里面。

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

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

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

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

3. 加法樹(shù)

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

下面通過(guò)加法樹(shù)的方法，實(shí)現(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)的就是上訴描述的加法樹(shù)的結(jié)構(gòu)，注意這里第二個(gè) __syncthreads() 的使用，也就是說(shuō)，要進(jìn)行下一級(jí)流水線的計(jì)算，必須建立在前一級(jí)必須已經(jīng)計(jì)算完畢的情況下。

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

性能有一部分的改善。

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

4. 總結(jié)

通過(guò)這幾篇博文的討論，數(shù)組平方和的代碼優(yōu)化到這一階段。從但線程到多線程，再到共享內(nèi)存，通過(guò)使用這幾種 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 程序优化进阶的全部?jī)?nèi)容，希望文章能夠幫你解決所遇到的問(wèn)題。

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