CUDA 中的分歧 - 从内核中的线程退出答案

【问题标题】：Divergence in CUDA - exit from a thread in kernelCUDA 中的分歧 - 从内核中的线程退出
【发布时间】：2013-01-29 22:35:24
【问题描述】：

我想知道如何退出线程索引太大的线程。我看到了两种可能性：

int i = threadIdx.x;
if(i >= count)
    return;
// do logic

或

int i = threadIdx.x;
if(i < count) {
    // do logic
}

我知道，两者都是正确的，但哪一个对性能的影响更大？

【问题讨论】：

【解决方案1】：

虽然两者在性能上是一样的，但是你应该考虑到第一个是不推荐的。

在内核中返回线程可能会导致其余代码出现意外行为。

我所说的意外行为是指与分组在一个扭曲中的最小线程单元相关的任何问题。例如，如果您的内核中有一个 if / else 块，这种情况称为线程发散，在正常情况下，它会导致线程保持空闲，而其他线程执行一些指令。

CUDA by Example Book，第 5 章，线程协作：

但是在 __syncthreads() 的情况下，结果有点悲惨。 CUDA 架构保证在块中的每个线程都执行完 __syncthreads() 之前，没有线程会前进到超出 __syncthreads() 的指令

所以，它主要与内核中的线程同步有关。你可以在这里找到关于这个主题的一个很好的问题/答案：Can I use __syncthreads() after having dropped threads?

正如我最后指出的，我也使用了那个不好的做法，没有出现任何问题，但不能保证将来会出现问题。这是我不推荐的东西

【讨论】：