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

0. 目錄

1. 陣列平方和並行化

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

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

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

由於我的 GPU 支援的塊內最大的執行緒數是 512 個，即 cudaGetDeviceProperties 中的 maxThreadsPerBlock

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

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 GPU
    cudaMalloc((void**) &sum, sizeof(int) * THREAD_NUM);
  
    // calculate the squares's sum
    squaresSum<<<1, THREAD_NUM, 0>>>(gpuData, sum, time);
  

• 最後，在 CPU 內增加部分和求和的程式碼

    // print result
    int 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 在讀取記憶體是連續的， 但是對於整體而言，執行的過程中讀取就不是連續的了（這裡自己仔細想想，就明白了）。所以，正確的做法如下表格所示：

根據這個原理，修改核心函式如下：

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

執行後結果如下所示：

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