160
CUDA從入門到精通（八）：線程通信
我們前麵幾節主要介紹了三種利用GPU實現並行處理的方式：線程並行，塊並行和流並行。在這些方法中，我們一再強調，各個線程所進行的處理是互不相關的，即兩個線程不回產生交集，每個線程都隻關注自己的一畝三分地，對其他線程毫無興趣，就當不存在。。。。
當然，實際應用中，這樣的例子太少了，也就是遇到向量相加、向量對應點乘這類才會有如此高的並行度，而其他一些應用，如一組數求和，求最大（小）值，各個線程不再是相互獨立的，而是產生一定關聯，線程2可能會用到線程1的結果，這時就需要利用本節的線程通信技術了。
線程通信在CUDA中有三種實現方式：
1. 共享存儲器；
2. 線程 同步；
3. 原子操作；
最常用的是前兩種方式，共享存儲器，術語Shared Memory，是位於SM中的特殊存儲器。還記得SM嗎，就是流多處理器，大核是也。一個SM中不僅包含若幹個SP（流處理器，小核），還包括一部分高速Cache，寄存器組，共享內存等，結構如圖所示：
從圖中可看出，一個SM內有M個SP，Shared Memory由這M個SP共同占有。另外指令單元也被這M個SP共享，即SIMT架構（單指令多線程架構），一個SM中所有SP在同一時間執行同一代碼。
為了實現線程通信，僅僅靠共享內存還不夠，需要有同步機製才能使線程之間實現有序處理。通常情況是這樣：當線程A需要線程B計算的結果作為輸入時，需要確保線程B已經將結果寫入共享內存中，然後線程A再從共享內存中讀出。同步必不可少，否則，線程A可能讀到的是無效的結果，造成計算錯誤。同步機製可以用CUDA內置函數：__syncthreads()；當某個線程執行到該函數時，進入等待狀態，直到同一線程塊（Block）中所有線程都執行到這個函數為止，即一個__syncthreads()相當於一個線程同步點，確保一個Block中所有線程都達到同步，然後線程進入運行狀態。
綜上兩點，我們可以寫一段線程通信的偽代碼如下：
//Begin
if this is thread B
write something to Shared Memory;
end if
__syncthreads();
if this is thread A
read something from Shared Memory;
end if
//End
上麵代碼在CUDA中實現時，由於SIMT特性，所有線程都執行同樣的代碼，所以在線程中需要判斷自己的身份，以免誤操作。
注意的是，位於同一個Block中的線程才能實現通信，不同Block中的線程不能通過共享內存、同步進行通信，而應采用原子操作或主機介入。
對於原子操作，如果感興趣可以翻閱《GPU高性能編程CUDA實戰》第九章“原子性”。
本節完。下節我們給出一個實例來看線程通信的代碼怎麼設計。
最後更新：2017-04-03 16:48:43