如何在没有分支分歧的情况下检查 CUDA 内核中数组的边界答案

【问题标题】：How to check boundary of array in CUDA Kernel without branch divergence如何在没有分支分歧的情况下检查 CUDA 内核中数组的边界
【发布时间】：2014-10-28 06:20:10
【问题描述】：

在下面的内核中，我使用if 语句来避免超出范围的计算。但如果我理解正确，“if”语句会导致分支发散，从而减慢计算速度——如果我在这里错了，请纠正我。

我的问题：如何在内核中处理超出范围的计算时避免使用 if 语句？

__global__ void vector_add(float *a, float *b, float *c)
{
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    if(index < N)
          c[index] = a[index]*a[index] + b[index]*b[index];
}
//kernel call here
vector_add<<< (N + (THREADS_PER_BLOCK+1)) / THREADS_PER_BLOCK, THREADS_PER_BLOCK >>>( d_a, d_b, d_c );

【问题讨论】：

你为什么声称在你的执行路径中有任何分歧分支？带有index >= N 的线程将只是等待或根本不启动，并且不会序列化任何内容，因为您的代码 sn-p 中没有 else 语句。

标签： cuda

【解决方案1】：

虽然它在技术上被称为“分歧”（因为并非所有线程都以相同的方式评估条件），但它完全无害。

不评估true 的谓词的线程将被禁用：这不是性能问题，因为这些线程无论如何都不会参与计算。您不会丢失任何实际的工作线程。在N 与1 mod 32 （或任何经线大小）一致的病态情况下，只有一条经线几乎完全“浪费”了，但同样，这不是性能问题。

当 warp 中的线程采用需要串行执行的不同路径时，Warp 发散会损害您的性能。这不是这里的情况。

【讨论】：

@Cicada 虽然您的回答对于上述情况似乎是正确的，但我不想将其概括为不同的场景。 Warp 的禁用通道是宝贵的计算资源，在到达重新收敛点之前一直不可用。例如，我宁愿让一个经线的 32 个线程来做类似的工作，而不是每个通道的 32 个经线来做这项工作。