深刻理解CUDA點積運算

最近一直在學習CUDA並行計算的相關知識。在學習《GPU高性能編程CUDA實戰》（機械工業出版社）這本書時，遇到了一些問題，想了好長時間纔想明白，這裏我將本身的理解與你們分享一番，若是有錯誤的地方，歡迎請你們指點。

因爲在點積運算這個例子中，核函數是最關鍵也是最難懂的部分，所以在這裏我只詳細介紹一下核函數的部分。首先我闡釋一下大體的思路。按照書中的示例，進行點積運算的兩個向量長度爲33*1024，其中共使用了32個線程塊，每一個線程塊中使用了256個線程。咱們這裏就不作改變了。（詳情請參考本書第五章內容）

申請共享內存

首先咱們須要申請共享內存，在這個例子中聲明的是數組cache:

__shared__ float cache[threadsPerBlock];

這裏咱們須要明白的是，一旦這樣聲明數組，就會建立與線程塊的數量相同的數組cahce，即每一個線程塊都會對應一個這樣的數組cache。咱們都知道，共享內存是用於同一個線程塊內的線程之間交流的，不一樣線程塊之間是沒法經過共享內存進行交流的。另外，數組cache的大小是每一個線程塊中線程的個數，即線程塊的大小。

每一個線程單獨工做

如今讓咱們來看看每一個線程到底完成的是什麼工做！

若是你還記得前面計算任意長度的向量和的話，你就會很容易理解這個過程。若是向量長度不是特別長(假設大小等於總線程個數)的話，每一個線程只須要工做一次，即計算兩個元素的積並保存在中間變量temp裏。可是實際計算過程當中因爲向量長度過長，一次計算可能會計算不完，每一個線程須要屢次計算才能完成全部工做，所以temp保存的值可能爲多個元素乘積之和，以下圖所示

假設數組大小爲16，線程總數爲4。此時一次並行是沒法完成工做的，因此須要屢次並行，即每一個線程須要作四次工做纔可完成計算。

相應的代碼以下：

int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int cacheIndex = threadIdx.x;    

    float temp = 0;
    while (tid < N)
    {
        temp += a[tid] * b[tid];
        tid += blockDim.x * gridDim.x;
    }

若是你已經理解了上面這個過程，那麼你也應該會明白每一個線程塊移動的步長爲何是總線程的個數了，即tid += blockDim.x * gridDim.x這段代碼。

多個線程協同工做

這一章主要講的就是線程協做，因此咱們須要明白線程之間是如何協做的——經過共享內存。每一個線程將temp的值保存到每一個線程塊的共享內存(shared memory)中，即數組cache中,相應的代碼以下：

cache[cacheIndex] = temp;
            __syncthreads();

這樣每一個線程塊中對應的數組cache保存的就是每一個線程的計算結果。爲了節省帶寬，這裏又採用了並行計算中經常使用的歸約算法，來計算數組中全部值之和，並保存在第一個元素(cache[0])內。這樣每一個線程就經過共享內存(shared memory)進行數據交流了。具體代碼以下所示：

//歸約算法將每一個線程塊上的cache數組歸約爲一個值cache[0]，最終保存在數組c裏
    int i = blockDim.x /2;
    while (i != 0)
    {
        if (cacheIndex < i)
            cache[cacheIndex] += cache[cacheIndex + i];
        __syncthreads();        //確保每一個線程已經執行完前面的語句

        i /= 2;
    }

NOTE:不要遺漏__syncthreads()函數，另外關於歸約算法本書中有詳細的介紹，這裏就再也不贅述了。

保存歸約結果

如今每一個線程塊的計算結果已經保存到每一個共享數組cache的第一個元素cache[0]中，這樣能夠大大節省帶寬。下面就須要將這些歸約結果保存到全局內存(global memory)中。

觀察覈函數你會發現有一個傳入參數——數組c。這個數組是位於全局內存中，每次使用線程塊中線程ID爲0的線程來將每一個線程塊的歸約結果保存到該數組中，注意這裏每一個線程塊中的結果保存到數組c中與之相對應的位置，即c[blockIdx.x]。

//選擇每一個線程塊中線程索引爲0的線程將最終結果傳遞到全局內存中
    if (cacheIndex == 0)
        c[blockIdx.x] = cache[0];

到這裏核函數的工做已經結束，剩下的工做將交給主函數來完成，這裏就再也不贅述。

參考資料

1. GPU高性能編程CUDA實戰, Jason Sanders, Edward Kandrot, 機械工業出版社