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

Question

我想在进行基数排序的内核中调用一个独占扫描函数。但是，独占扫描只需要一半的线程来完成其工作。

独占扫描算法中需要几个__syncthreads()。如果我在开头有一个语句：

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

这些线程将不参与独占扫描的同步线程，这是不允许的。有没有什么方法解决这个问题？我确实将对独占扫描的调用包围在__syncthread()中。

- user1428701

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()用法法律条文的唯一方法。您也可以尝试按您指示的方式保留代码，使不执行任何工作的线程提前返回/退出。它可能会奏效，可能会奏效。但不能保证它将在未来的架构或更新的工具链中起作用。

- Robert Crovella

---

无法保证它适用于未来的架构。我可以证明这一点。我需要修复的遗留库在 Kepler、Maxwell、Pascal、Turing 上运行良好多年，但在 Ampere 上突然出现死锁问题。原因是代码中在不执行任务的线程中调用 return，但稍后调用了 __syncthreads()。我们仍然不知道是由于架构还是GPU更快的原因。 - Huy Le

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);

免责声明：本文在浏览器中编写，且未经任何测试！

但是否值得尝试另外是一个问题。

- tera

---

我学到了一些关于内联汇编的知识，所以加1。但是__syncthreads已经总是按照warp中线程数增加，即使只有一个分支中的线程命中__syncthreads，也不会更多或更少。这实际上使它成为每个warp屏障指令。 - chappjc