cuda控制分歧

Question

说我有 3 个共享内存阵列:a b c。我不确定以下线程排列是否会导致控制发散，

if (threadIdx < 64)
{
    if (threadIdx == 1)
        for (int i = 0; i < N; i++)
            c += a[threadIdx]*a[threadIdx];
    else
        for (int i = 0; i < N; i++)
            c += a[threadIdx]*b[threadIdx];
}

如果是这样，它会影响性能有多严重？有什么有效的方法来处理这个问题吗？谢谢

Answer 1

根据块的尺寸，第一个条件 threadIdx.x < 64 (注意 .x )可能根本不会引起任何分歧。例如，如果您有一个尺寸为 (128,1,1) 的块那么前两个经线(以锁步执行的 32 个线程组)将进入 if阻止而最后两个将绕过它。由于整个经线以一种方式或另一种方式进行，因此没有分歧。

条件像 threadIdx.x == 1会导致分歧，但它的成本非常适中。事实上，在许多情况下，CUDA 将能够用一条指令来实现条件表达式。例如，像 min 这样的操作, max , 和 abs通常会用一条指令来实现，绝对不会造成分歧。您可以在 PTX Manual 中阅读此类说明。 .

一般来说，您不应过分担心像上述那样适度的控制流分歧。在诸如此类的情况下，分歧会让您陷入困境

if (threadIdx.x % 4 == 0)
  // do expensive operation
else if (threadIdx.x % 4 == 1)
  // do expensive operation
else if (threadIdx.x % 4 == 2)
  // do expensive operation
else
  // do expensive operation

其中“昂贵的操作”将是需要 10 或 100 条指令的操作。在这种情况下，由 if 引起的背离语句会使效率降低 75%。

请记住，线程发散比 (1) 高级算法选择和 (2) 内存局部性/合并要少得多。很少有 CUDA 程序员应该关心您的示例中的那种分歧。