如何用更少的线程在 CUDA 中调用 __device__ 函数答案

【问题标题】：How to call __device__ function in CUDA with fewer threads如何用更少的线程在 CUDA 中调用 __device__ 函数
【发布时间】：2013-03-19 11:17:39
【问题描述】：

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

独占扫描算法需要几个__syncthreads()。如果我一开始就有这样的声明

if(threadIdx.x > NTHREADS/2) 返回；

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

【问题讨论】：

标签： cuda

【解决方案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 上的法律条文的唯一方法。您也可以尝试按照您指示的方式保留代码，使不工作的线程提前返回/退出。它可能会起作用，也可能会起作用。但不能保证它适用于未来的架构或更新的工具链。

【讨论】：

【解决方案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);

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

这是否值得麻烦是另一个问题。

【讨论】：

我学到了一些关于内联 asm 的知识，所以 +1，但是 __syncthreads 已经并且总是增加一个 warp 中的线程数——不多也不少——即使一个线程中只有一个线程分歧分支命中__syncthreads。无论如何，这实际上使它成为per-warp barrier 指令。