版權(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 編程入門到精通（二）之 運(yùn)行第一個程序
• GPU 編程入門到精通（三）之 第一個 GPU 程序
• GPU 編程入門到精通（四）之 GPU 程序優(yōu)化
• GPU 編程入門到精通（五）之 GPU 程序優(yōu)化進(jìn)階

1. 數(shù)組平方和并行化

GPU 編程入門到精通（三）之 第一個 GPU 程序 中講到了如何利用 CUDA5.5 在 GPU 中運(yùn)行一個程序。通過程序的運(yùn)行，我們看到了 GPU 確實(shí)可以作為一個運(yùn)算器，但是，我們在前面的例子中并沒有正真的發(fā)揮 GPU 并行處理程序的能力，也就是說之前的例子只利用了 GPU 的一個線程，沒有發(fā)揮程序的并行性。

先來說說 CUDA5.5 中 GPU 的架構(gòu)。它是由 grid 組成，每個 grid 又可以由 block 組成，而每個 block 又可以細(xì)分為 thread。所以，線程是我們處理的最小的單元了。

接下來的例子通過修改前一個例子，把數(shù)組分割成若干個組（每個組由一個線程實(shí)現(xiàn)），每個組計算出一個和，然后在 CPU 中將分組的這幾個和加在一起，得到最終的結(jié)果。這種思想叫做歸約 。其實(shí)和分治思想差不多，就是先將大規(guī)模問題分解為小規(guī)模的問題，最后這些小規(guī)模問題整合得到最終解。

由于我的 GPU 支持的塊內(nèi)最大的線程數(shù)是 512 個，即 cudaGetDeviceProperties 中的 maxThreadsPerBlock 屬性。如何獲取這個屬性，請參看 GPU 編程入門到精通（二）之 運(yùn)行第一個程序 這一章節(jié)。 我們使用 512 個線程來實(shí)現(xiàn)并行加速。

好了，接下來就是寫程序的時候了。

1.1. 修改代碼

• 首先，在程序頭部增加一個關(guān)于線程數(shù)量的宏定義：

  // ======== define area ========#define DATA_SIZE 1048576 // 1M#define THREAD_NUM 512 // thread num

  其中，DATA_SIZE 表示處理數(shù)據(jù)個數(shù)， THREAD_NUM 表示我們將要使用 512 個線程。

• 其次，修改 GPU 部分的內(nèi)核函數(shù)

  const int size = DATA_SIZE / THREAD_NUM;const int tid = threadIdx.x;int tmp_sum = 0;for (int i = tid * size; i < (tid + 1) * size; i++) {tmp_sum += data[i] * data[i];}sum[tid] = tmp_sum;}

  此內(nèi)核程序的目的是把輸入的數(shù)據(jù)分?jǐn)偟?512 個線程上去計算部分和，并且 512 個部分和存放到 sum 數(shù)組中，最后在 CPU 中對 512 個部分和求和得到最終結(jié)果。

  此處對數(shù)據(jù)的遍歷方式請注意一下，我們是根據(jù)順序給每一個線程的，也就是如下表格所示：

  線程編號 數(shù)據(jù)下標(biāo)

  0	0 ~ 2047
  … …	… …
  511	1046528 ~ 1048575

• 然后，修改主函數(shù)部分
  主函數(shù)部分，只需要把 sum 改成數(shù)組就可以，并且設(shè)置一下調(diào)用 GPU 內(nèi)核函數(shù)的方式。

  // malloc space for datas in GPUcudaMalloc((void**) &sum, sizeof(int) * THREAD_NUM);// calculate the squares's sumsquaresSum<<<1, THREAD_NUM, 0>>>(gpuData, sum, time);

• 最后，在 CPU 內(nèi)增加部分和求和的代碼

  // print resultint tmp_result = 0;for (int i = 0; i < THREAD_NUM; ++i) {tmp_result += result[i];}printf("(GPU) sum:%d time:%ld\n", tmp_result, time_used);

1.2. 編譯運(yùn)行

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

2. 性能分析

經(jīng)過修改以后的程序，比之前的快了將近 36 倍（可以參考博文GPU 編程入門到精通（三）之 第一個 GPU 程序 進(jìn)行比較），可見并行化處理還是存在優(yōu)勢的。 不過仔細(xì)想一下，我們使用了 512 個線程， 可是性能怎么才提升了 36 倍，不應(yīng)該是 512 倍嗎？？？

這里就涉及到內(nèi)存的存取模式了，顯卡上面的內(nèi)存是 DRAM，是效率最高的存取方式，它是一種連續(xù)的存取方式。 前面我們的程序確實(shí)的連續(xù)讀取的呀，都挨個讀取了，怎么還是沒有達(dá)到預(yù)期的效果呢？？？

這里還需要考慮 thread 的執(zhí)行方式，GPU 編程入門到精通（三）之 第一個 GPU 程序 中說到，當(dāng)一個 thread 在等待內(nèi)存數(shù)據(jù)的時候， GPU 就會切換到下一個 thread。所以，實(shí)際執(zhí)行的順序類似于 thread0 —> thread1 —> … … —> thread511。

這就導(dǎo)致了同一個 thread 在讀取內(nèi)存是連續(xù)的， 但是對于整體而言，執(zhí)行的過程中讀取就不是連續(xù)的了（這里自己仔細(xì)想想，就明白了）。所以，正確的做法如下表格所示：

線程編號 數(shù)據(jù)下標(biāo)

0	0 ~ 512
… …	… …
511	511 ~ 1023

根據(jù)這個原理，修改內(nèi)核函數(shù)如下：

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

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

修改后程序，比之前的又快了 13 倍左右，可見，對內(nèi)存的讀取方式對于性能的影響很大。
至此，并行化后的程序較未并行化之前的程序，速度上快了 493 倍左右，可見，基本上發(fā)揮了 512 個線程的優(yōu)勢。

讓我們再來分析一下性能：

此 GPU 消耗的時鐘周期： 1595788 cycles GeForce G 103M 的 clockRate： 1.6 GHz 所以可以計算出 GPU 上運(yùn)行時間是： 時鐘周期 / clockRate = 997.3675 us 1 M 個 int 型數(shù)據(jù)有 4M Byte 的數(shù)據(jù)量，實(shí)際使用的 GPU 內(nèi)存帶寬是：數(shù)據(jù)量 / 運(yùn)行時間 = 4.01 GB/s

再來看看我的 GPU GeForce 103m 的內(nèi)存帶寬：運(yùn)行 SDK 目錄下面 /samples/1_Utilities/bandwidthTest

運(yùn)行后結(jié)果如下所示：

通過與系統(tǒng)參數(shù)的對比，可以知道，基本上達(dá)到了系統(tǒng)的極限性能。

總結(jié)

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

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