深入理解CUDA點積運算

本部落格已遷移至segmentfault-cuda程式設計學習專欄共享記憶體之bank衝突。本博主未來關於cuda平行計算的文章都將發布在segmentfault-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)

如果你已經理解了上面這個過程，那麼你也應該會明白每個執行緒塊移動的步長為什麼是匯流排程的個數了，即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)

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

現在每個執行緒塊的計算結果已經儲存到每個共享陣列cache的第乙個元素cache[0]中，這樣可以大大節省頻寬。下面就需要將這些歸約結果儲存到全域性記憶體(global memory)中。

觀察核函式你會發現有乙個傳入引數——陣列c。這個陣列是位於全域性記憶體中，每次使用執行緒塊中線程id為0的執行緒來將每個執行緒塊的歸約結果儲存到該陣列中，注意這裡每個執行緒塊中的結果儲存到陣列c中與之相對應的位置，即c[blockidx.x]。

//選擇每個執行緒塊中線程索引為0的執行緒將最終結果傳遞到全域性記憶體中
if (cacheindex == 0)
c[blockidx.x] = cache[0];

到這裡核函式的工作已經結束，剩下的工作將交給主函式來完成，這裡就不再贅述。

gpu高效能程式設計cuda實戰, jason sanders, edward kandrot, 機械工業出版社

深入理解CUDA點積運算

Cuda計算點積

點積運算叉積運算

深入理解C語言深入理解指標

深入理解CUDA點積運算

Cuda計算點積

點積運算叉積運算

深入理解C語言 深入理解指標

相關推薦

深入理解C語言深入理解指標