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

目錄(?)[+]

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

---

0. 目錄

• GPU 編程入門到精通（一）之 CUDA 環境安裝
• GPU 編程入門到精通（二）之 運行第一個程序
• GPU 編程入門到精通（三）之 第一個 GPU 程序
• GPU 編程入門到精通（四）之 GPU 程序優化
• GPU 編程入門到精通（五）之 GPU 程序優化進階

1. 數組平方和并行化

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

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

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

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

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

1.1. 修改代碼

• 首先，在程序頭部增加一個關于線程數量的宏定義：

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

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

• 其次，修改 GPU 部分的內核函數

  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;}

  此內核程序的目的是把輸入的數據分攤到 512 個線程上去計算部分和，并且 512 個部分和存放到 sum 數組中，最后在 CPU 中對 512 個部分和求和得到最終結果。

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

  線程編號 數據下標

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

• 然后，修改主函數部分
  主函數部分，只需要把 sum 改成數組就可以，并且設置一下調用 GPU 內核函數的方式。

  // 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 內增加部分和求和的代碼

  // 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. 編譯運行

編譯后，運行結果如下所示：

2. 性能分析

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

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

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

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

線程編號 數據下標

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

根據這個原理，修改內核函數如下：

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

編譯運行后結果如下所示：

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

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

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

再來看看我的 GPU GeForce 103m 的內存帶寬：運行 SDK 目錄下面 /samples/1_Utilities/bandwidthTest

運行后結果如下所示：

通過與系統參數的對比，可以知道，基本上達到了系統的極限性能。

總結

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

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