cuda - 如何在没有分支发散的情况下检查 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 分歧会损害您的性能。这不是这里的情况。