[llvm] [AMDGPU] Model structured control-flow lane masks as i1 (PR #209158)

[Pankaj Dwivedi via llvm-commits]

PankajDwivedi-25 wrote:

> Does this actually work? Have you run any executable tests?

I ran executable HIP tests on gfx90a (differential vs SelectionDAG and pre-patch GISel). A single divergent loop (for (i=0;i<=tid;i++) acc+=data[i]) passes on all three. 

// PASS on all paths: divergent loop, all lanes enter.
extern "C" __global__ void divergent_loop(int *out, const int *data) {
  int tid = threadIdx.x + blockIdx.x * blockDim.x;
  int acc = 0;
  for (int i = 0; i <= tid; i++) acc += data[i];
  out[tid] = acc;
}

But a divergent loop nested in a divergent if miscompiles only on this patch's GISel path: lanes with tid>=n (disabled by the outer if) get reactivated and overwrite their result; SDAG and pre-patch GISel are correct. 

// FAIL on this patch's GISel: post-loop store must run ONLY for tid<n.
// If exec is corrupted by the undef-seeded lane mask at end.cf, lanes
// tid>=n wrongly overwrite -1.
extern "C" __global__ void nested(int *out, const int *data, int n) {
  int tid = threadIdx.x + blockIdx.x * blockDim.x;
  out[tid] = -1;
  if (tid < n) {
    int acc = 0;
    for (int i = 0; i <= tid; i++) acc += data[i];
    out[tid] = acc;
  }
}

Root cause for this is the implicit-def concern raised earlier: since if.break is now a divergent i1, divergence lowering builds a lane-mask phi seeded with IMPLICIT_DEF, and end.cf's s_or exec, exec, <phi> ORs those garbage bits into exec.

https://github.com/llvm/llvm-project/pull/209158