我看GTC的資料有一篇講inter block的同步
~www nvidia com/ content/GTC/documents/SC09_Feng.pdf(因為文章不能含有url,請自己加"點")
(詳細資料→論文標題:Inter-Block GPU Communication via Fast Barrier Synchronization)
大致上是說透過一個全局的變數來檢查是不是所有的block都直行到_gpu_sync_來達成同步
簡單來講就是避免掉透過返回CPU來達成inter block間的同步。
但是我看他的實驗都只有做到30個block而已,這是不是說他的卡剛好可以30個block同時運行才能這樣做
如果block的數量太大的話,變成block序列執行的那這個方法就不能用了。(這是我的理解)
只是如果block作的某一步後需要一直等待,那會不會先切換成其他的block去做?(我的理解是不行)
我一時想不到怎麼去實驗來驗證我的想法,所以說拋出來問看看大家的意見。
論文中gpu_sync同步的程式碼
//the mutex variable
__device__ volatile int g_mutex;
//GPU lock-based synchronization function
__device__ void __gpu_sync(int goalVal)
{
//thread ID in a block
int tid_in_block = threadIdx.x * blockDim.y + threadIdx.y;
// only thread 0 is used for synchronization
if (tid_in_block == 0)
{
atomicAdd((int *)&g_mutex, 1);
//only when all blocks add 1 to g_mutex will g_mutex equal to goalVal
while(g_mutex != goalVal) {//Do nothing here }
}
__syncthreads();
}