如何在CUDA内核中检查数组边界而无分支分歧

Question

在以下内核中，我使用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 );

Answer 1

从技术上讲，它被称为“发散”（因为并非经纱中的所有线程都相同地评估条件），但这是完全无害的。

不评估谓词的线程true将被禁用：这不是性能问题，因为这些线程无论如何也不会参与计算。您不会丢失任何实际的工作线程。在N符合1 mod 32（或任何经线大小）的病理情况下，仅存在几乎完全“浪费”的经线，但这又不是性能问题。

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