如何在線程數較少的CUDA中調用__device__函數

Question

我想從內核中調用獨佔掃描函數來進行基數排序。但排他掃描只需要一半的線程來完成它的工作。

獨家掃描算法需要幾個__syncthreads（）。如果我在開始時聲明像

if（threadIdx.x> NTHREADS/2）return;

這些線程不會參與獨佔掃描syncthreads，這是不允許的。 有沒有辦法解決這個問題。我確實打電話給由__syncthread（）包圍的獨家掃描。

Answer 1

像這樣的東西應該工作（不使用提前返回）：

__syncthreads(); // at entry to exclusive scan region 
// begin exclusive scan function 
if (threadIdx.x < NTHREADS/2) { 
    // do first phase of exclusive scan up to first syncthreads 
    } 
__syncthreads(); // first syncthreads in exclusive scan function 
if (threadIdx.x < NTHREADS/2) { 
    // do second phase of exclusive scan up to second syncthreads 
    } 
__syncthreads(); // second syncthreads in exclusive scan function 
(... etc.) 
__syncthreads(); // at exit from exclusive scan region 

這有點乏味，但它是我知道的堅持法律條文上__syncthreads()usage的唯一途徑。您也可以嘗試按照您指示的方式離開代碼，而不做任何工作的線程會盡早返回/退出。它可能正常工作，可能會工作。但是不能保證它能用於未來的架構或更新的工具鏈。

Answer 2

只是想指出的替代：
您還可以使用內聯彙編相當於__syncthreads()，它允許使用可選參數參與的線程數，可從計算能力2.0起。像這樣的東西應該工作：

#define __syncthreads_active(active_threads) asm volatile("bar.sync 0, %0;" :: "r"(active_threads)); 

if(threadIdx.x >= NTHREADS/2) return; 

int active_warps = (NTHREADS/2 + warpSize)/warpSize; 
int active_threads = active_warps * warpSize; // hopefully the compiler will optimize this to a simple active_threads = (NTHREADS/2 + warpSize) & ~32 

__syncthreads_active(active_threads); 
// do some work... 
__syncthreads_active(active_threads); 
// do some more work... 
__syncthreads_active(active_threads); 

免責聲明：寫在瀏覽器中，完全未經測試！

不管是否值得這個麻煩，都是另一個問題。