如何在线程数较少的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); 

免责声明：写在浏览器中，完全未经测试！

不管是否值得这个麻烦，都是另一个问题。