版權(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ù)組平方和并行化

GPU 編程入門到精通（三）之 第一個(gè) GPU 程序 中講到了如何利用 CUDA5.5 在 GPU 中運(yùn)行一個(gè)程序。通過(guò)程序的運(yùn)行，我們看到了 GPU 確實(shí)可以作為一個(gè)運(yùn)算器，但是，我們?cè)谇懊娴睦又胁](méi)有正真的發(fā)揮 GPU 并行處理程序的能力，也就是說(shuō)之前的例子只利用了 GPU 的一個(gè)線程，沒(méi)有發(fā)揮程序的并行性。

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

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

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

好了，接下來(lái)就是寫程序的時(shí)候了。

1.1. 修改代碼

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

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

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

• 其次，修改 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 個(gè)線程上去計(jì)算部分和，并且 512 個(gè)部分和存放到 sum 數(shù)組中，最后在 CPU 中對(duì) 512 個(gè)部分和求和得到最終結(jié)果。

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

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

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

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

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

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

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

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

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

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

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

讓我們?cè)賮?lái)分析一下性能：

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

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

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

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

總結(jié)

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

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