CUDA通訊機制

分類：

gpu/cuda/opencl2009-08-25 21:57

1479人閱讀

收藏舉報

cuda

compiler

thread

function任務

cuda體系架構支援怎麼樣的通訊方法呢？往下看嘍@_@

1. __syncthreads()

block內的執行緒同步。block內所有執行緒都執行到這一位置(bar指令)，先到的要等後來的，到齊了後再繼續後面的任務。執行結果對block內所有執行緒可見

2. memory fence

l __threadfence()

grid內的執行緒同步。保證該語句前的，grid中所有執行緒發出的訪存指令（對global memory / shared memory）都已結束。執行結果對grid內所有執行緒可見

l __threadfence_block()

block內的執行緒同步。保證該語句前的，block中所有執行緒發出的訪存指令（對global memory / shared memory）都已結束。執行結果對block內所有執行緒可見

4. cudathreadsynchronize()

gpu與cpu同步。kernel啟動後控制權非同步返回，該函式用以確定所有裝置端線程均已執行結束。

5. volatile

__global__ void mykernel(int* result,int* myarray)

{int tid = threadidx.x;

int ref1 = myarray[tid] * 1;

myarray[tid + 1] = 2;

int ref2 = myarray[tid] * 1;

result[tid] = ref1 * ref2;

上面這段程式，myarray初值為1。這段程式跑出來的結果result全是1，說明計算ref2時myarray[tid]並如有如約變成2（由tid-1修改），程式設計師是想用改變後的值2來著。為什麼會出現這種情況呢？

是這樣，在計算ref1時從視訊記憶體取出myarray[tid]，然後在計算ref2時compiler不會再去顯訪問一次而是去復用剛剛取到的結果，myarray[tid+1]=2這句操作呢是要寫回視訊記憶體，所以嘍^o^ 可以通過加上volatile關鍵字來解決。

__global__ void mykernel(int* result,volatile

int* myarray)

volatile限定符將變數宣告為敏感變數，compiler認為其它執行緒可能隨時會修改變數的值，因此每次對該變數的引用都會被編譯成一次真實的訪存指令。注意，在上面的**中即使將myarray宣告為valatile仍不能確保ref的值是2，這是因為執行緒tid可能在myarray[tid]被賦為2之前就已經進行了讀操作，怎麼辦呢？做乙個同步就好嘍。

6. atomic function

當多個執行緒同時訪問global memory / shared memory同一位置時，原子函式來確保對這個32-/64-bit字的read-modify-write原子操作。保證每個執行緒能夠實現對共享可寫資料的互斥操作：在乙個操作完成前，其它任何執行緒都無法訪問此位址。

7. vote

cuda2.0引入的新特性，1.2計算能力的硬體支援。vote指令作用域是乙個warp。

int __all(int predicate); //warp

中所有執行緒的判斷表示式結果為真，則返回1，否則0

int __any(int predicate); //warp

中只要乙個thread的判斷表示式結果為真，則返回1，否則0

CUDA通訊機制

CUDA通訊機制

Socket通訊機制

socket 通訊機制

CUDA通訊機制

CUDA通訊機制

Socket通訊機制

socket 通訊機制

相關推薦