[RFC] A New Divergence Analysis for LLVM

TL;DR This RFC is a joint effort by Intel and Saarland University to bring the divergence analysis of the Region Vectorizer [1,2,3,4,5] (dubbed the vectorization analysis of RV) to LLVM. The implementation is available on github for feedback [0]. The existing divergence analysis infrastructure in LLVM has conceptual limitations (structured control, SCEV based). The new analysis resolves bugs for the GPU backends from the first patch on and will be used by VPlan [6] in the mid term. The implementation is based on RV, the Region Vectorizer.

– RV –

The Region Vectorizer (github [1]) is an analysis and transformation framework for outer-loop and whole-function vectorization. RV vectorizes arbitrary, reducible control flow including nested divergent loops through partial control-flow linearization [2]. RV is being used by the Impala [7] and the PACXX [3,8] high performance programming frameworks and implements OpenMP #pragma omp simd and #pragma omp declare simd.

– The Divergence Analysis –

The Divergence Analysis determines how instructions will behave if executed in lockstep for multiple threads or vector lanes. The loop vectorizer (VPlan) uses divergence information to generate correct and efficient vector code. The GPU backends rely on it to re-structure the CFG.
An instruction is uniform if it evaluates to the same result for all (active) threads. Otherwise, it is is called divergent. Divergent branch conditions cause control to diverge at branches and can turn loops into divergent loops. Divergent branches induce divergence in phi nodes, if the branch divergence implies that a phi node could be reached by different threads from two distinct predecessors at once or from two distinct loop iterations at once (temporal divergence).

Example A: Divergence analysis result with regard to the outer loop.
for (int thread_id = 0; thread_id < n; ++thread_id) {`` // thread_id divergent
V = 0; // V uniform
if (A[thread_id] > 0) { // divergent branch
V = 1; // V uniform
// divergent phi for V (control induced)
int j = 0; // uniform
for (; j < m; ++j) { // uniform loop exit
if (B[j] > 50) { // uniform branch
C[j][thread_id] = 42;
if (D[j] < V) break; // divergent loop exit
x = j; // divergent (temporal; due to divergent loop)

– Motivation –

The LoopVectorizer and VPlan currently use SCEV to detect uniform values. The StructurizeCFG pass and the AMDGPU backends rely on LLVM’s existing DivergenceAnalysis pass for SPMD kernels. Both approaches have shortcomings that are resolved by the new analysis.

  • Why the current DivergenceAnalysis is broken -

LLVM already includes a DivergenceAnalysis for GPU kernels. It is used in the GPU backends and to re-structure CFGs (StructurizeCFG pass). However, the existing DivergenceAnalysis of LLVM is broken for unstructured CFGs (bug report [9]):

Example B: undetected divergence in unstructured control.
divergent branch to B and C
uniform branch to C and D
%phi.var.hidden = phi .. <- undetected DIVERGENT phi
br D
%phi.var = phi .. <- detected DIVERGENT phi in IPostDom of entry

The existing DivergenceAnalysis expects control to re-converge only at the immediate post dominator of a branch, which is not the case for unstructured control as shown in the example. The new analysis supports unstructured control flow.

– Why SCEV is not enough –

SCEV is restricted to induction variables and (piecewise) polynomial functions. The divergence analysis on the other hand is able to detect uniform instructions and branches in arbitrary programs. The key takeaway here is that “there are uniform values that are not invariant.” (http://lists.llvm.org/pipermail/llvm-dev/2017-December/119522.html).

Example A: non-affine loop (Test NonAffineUniLoop.ll)
for (int i = 0; i < n; ++i) {
for (int len = 1; len * 2 < n; len *= 2) {
for (int k = 0; k < n k += 2 * len) {
A[k * n + i] += A[(k + len) * n + i];

The variable “len” is uniform with regards to the outer i-loop. SCEV does not detect this:
%mul44 U: [0,-1) S: [-2147483648,2147483647)
Exits: <<Unknown>>
LoopDispositions: { %for.body8.lr.ph: Variant, %for.body: Variant, %for.body8: Invariant }

– The new DivergenceAnalysis –

The proposed DivergenceAnalysis detects divergence in unstructured, reducible control flow and nested divergent loops. The implementation provides light-weight frontends for use in VPlan (LoopDivergenceAnalysis) and the analysis of GPU kernels (GPUDivergenceAnalysis). Both frontends build on a single, shared implementation of the actual divergence analysis algorithm.

Apart from solving current issues with divergence detection in LLVM, the new analysis will enable the adoption of advanced vectorization techniques in VPlan that are already implemented in RV, for example partial linearization [2], automatic conversion of divergent loops, etc).

— Planned patches —

– Patch set 1 (reference impl. [0]) –
o New DivergenceAnalysis class, which implements a basic divergence analysis (uniform/divergent) with support for unstructured, reducible control and nested divergent loops. The analysis operates on LLVM IR.GPUDivergenceAnalysis: analysis frontend for SPMD kernels (fixes the unstructured control bug [9]). Acts as a drop-in replacement for the existing DivergenceAnalysis in LLVM (opt-in flag -use-rv-da).
o LoopDivergenceAnalysis: frontend for loop vectorization in preparation of VPlan [6].
o The existing DivergenceAnalysis gets renamed to KernelDivergenceAnalysis.
o LIT tests for both frontends.

– Patch set 2 –
o Integration with VPlan to detect uniform values in outer loop vectorization (opt-in flag -vectorizer-use-da).
o More expressive analysis lattice (uniform/consecutive/strided… + alignment). This is the full sa lattice (stride+alignment lattice) used in RV [3].

– Patch set 3 –
o Templatize the divergence analysis to operate directly on VPlans also (block/instruction type parametric).

We provide the implementation of this RFC as a fork of LLVM on github [0]. We welcome your feedback. This implementation will be the source of upcoming patches.



[0] VPlan+RV: [1] RV on github: [2] Partial Control-Flow Linearization, Moll, Hack, PLDI ‘18 (to appear): [3] PACXXv2 + RV: An LLVM-based Portable High-Performance Programming Model: [4] VPlan+RV SRC talk at US DevMtg ‘17: [5] VPlan+RV Lightning talk at EuroLLVM ‘18: [6] VPlan proposal: [7] Impala: [8] PACXX: [9] DivergenceAnalysis bug