在一個kernel內block的執行是調度的?

我看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();
 }

楼主您好,这个其实很容易想明白。

该文是通过原子+1, 然后通过直接读取来判断(一般认为不可以,实际也是可以的,例如这种情况,立刻原子返回后,再次读取的可能是n次原子+1之前的值,但最终还是在一定延迟后得到这个值的)。

但此文有个假设,就是期待的数目个blocks(goablVal), 需要能同时存在于GPU上。例如如果需要100个blocks, 卡只能上40个。那么这40个block长期循环运行,而目前是无法preemption式的调度它们下去的,从而导致其他60个blocks无法上来,从而死锁。

从CC3.5起,提供了有限能力的调度能力,但目前(资料)看起来是仅限于一个block和它所启动的子kernel之前的,而不能在同个kernel的多个block之间抢占式的调度。(也许以后会软件升级,例如提供yield()之类的kernel内部能调用的api, 以便能主动切换下去?或者直接实现抢占式调度,让长期运行的block下去让给其他block运行机会?这个要看将来是否能实现。)

所以这个目前只能允许<=最大驻留数目的blocks进行同步,否则会卡死。

感謝版主的回應!今天仔細的把論文看完,確實有提到這種方法只能在少量的block上運作。

文章中也提到這樣的方法同步法不見得會比用CPU調用kernel還快,

不過好處是shared memory可以不用釋放掉,因此需要連計算方法都要改寫才會比較好。

平常再作大量計算時需要調用很大量的block,這種方法就不可行了!

心得很好。欢迎常来!这里是你永远的家(虽然你还没有设置论坛头像)。

我剛剛翻了手冊想了一下,似乎這樣的方法不適用kepler架構!!

我的看法是kepler架構讓一個SMX有192個核心,

與fermi架構比同樣的核心下SM數量減少,而SM對應到的是block的執行

所以跟fermi比起來能溝通的block是更少的。

尤其在計算版本3.0的卡最為尷尬,能溝通的block數量少也不能像CC3.5一樣動態配置

如果我這樣的理解是對的話!看來要存錢買新卡了!

折算到SP, 的确少。但是这个方法本身就不是一个安全的方法。所以总是不建议您使用的。

以及,3.5的Dynamic Parallism, 无法满足您的这个(全局同步)的要求的。我建议最方便的方式还是改写算法,去除对全局同步的要求。