GPU Glossary
/
Performance·1 min read

Branch Efficiency

How often all threads in a warp take the same path at a conditional; high branch efficiency means little warp divergence.

Branch efficiency measures how often all threads in a warp take the same execution path when they hit a conditional statement.

It is the ratio of uniform control flow decisions to total branch instructions executed. Uniformity is measured per warp, so high branch efficiency means an absence of warp divergence.

Not every conditional hurts#

The bounds check that appears in most kernels usually has very high branch efficiency:

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

Almost every warp has threads that all agree on idx < n. Only a single warp, whose thread indices straddle n, disagrees.

Uniformity in time vs in space#

CPUs also care about uniform branching, but mainly uniformity over time. Branch prediction and speculative execution hardware gathers data about a branch each time it is encountered, and performance improves as it learns.

GPUs care about uniformity in space. The threads of a warp run concurrently but on different data, and performance improves when they branch the same way.

Key idea A CPU wants a branch to behave the same each time it runs; a GPU wants it to behave the same across all 32 threads of a warp at once.

  • Warp divergence — what low branch efficiency indicates.
  • Warp — the group whose threads must agree.