性能

分支效率branch efficiency

分支效率用于衡量的是当遇到条件语句时，线程束 (warp) 中的所有线程 (thread) 选择相同执行路径的频率。

分支效率的计算方式为：统一控制流决策数量与执行的分支指总令数之比。控制流统一性是在线程束 (warp) 级别进行测量的，因此分支效率反映了线程束分歧 (warp divergence) 的缺失程度。

并非所有条件语句都会降低分支效率。大多数内核 (kernel) 中常见的 "边界检查" 代码片段

int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n)

通常具有非常高的分支效率，因为大多数线程束 (warp) 中的线程 (thread) 往往具有相同的条件判断结果，仅在少数线程束 (warp) 中的线程 (thread) 索引会同时分布在 n 的上下两侧。

虽然 CPU 也关心分支行为的统一性，但其核心关注点是时间维度的分支行为一致性（作为硬件控制的分支预测和推测执行机制的一部分）。具体而言，随着CPU内部电路在程序执行过程中不断积累某条分支的历史数据，性能会逐步优化。

而 GPU 更关注空间维度的统一性。也就是说，线程束 (warp) 内部的一致性，线程束 (warp) 中的线程 (thread) 在时间上并发执行，但映射到不同的数据，如果这些线程 (thread) 能够统一分支，性能就会提高。

英文原文(Modal GPU Glossary)

Branch efficiency measures how often all threads in a warp take the same execution path when encountering conditional statements.

Branch efficiency is calculated as the ratio of uniform control flow decisions to total branch instructions executed. Control flow uniformity is measured at the level of warps, and so branch efficiency indicates the absence of warp divergence.

Not all conditionals reduce branch efficiency. The common "bounds-check" fragment that appears in most kernels

int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n)

will generally have very high branch efficiency, since most warps will be composed of threads that all have the same value for the conditional, save for a single warp whose threads' indices are above and below n.

While CPUs also care about the uniformity of branching behavior, they tend to care primarily about uniformity of branch behavior over time, as part of hardware-controlled branch prediction and speculative execution. That is, as circuits within the CPU accumulate data about a branch as it is encountered multiple times during program execution, the performance should improve.

GPUs instead care about uniformity in space. That is, uniformity is measured within warps, whose threads execute concurrently in time but are mapped onto distinct data, and performance improves if those threads branch uniformly.

相关词条