cuda - 如何使用较少的线程在 CUDA 中调用 __device__ 函数

Question

我想从执行基数排序的内核内部调用独占扫描函数。但是独占扫描只需要一半的线程来完成它的工作。

独占扫描算法中需要几个 __syncthreads()。如果我在开始时有一个声明，例如

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

这些线程不会参与独占扫描同步线程，这是不允许的。
有没有办法解决这个问题。我确实调用了由 __syncthread()s 包围的独占扫描。

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 .您也可以尝试按照您指示的方式保留代码，没有工作的线程会提前返回/退出。它可能会起作用，可能会起作用。但不能保证它适用于 future 的架构或更新的工具链。