[test] Remove occurences of br undef in CodeGen/AMDGPU tests

[kritgpt]

Summary

This patch removes br undef as suggested by https://discourse.llvm.org/t/please-dont-use-br-undef-in-tests-aka-please-avoid-test-cases-with-ub/63115. Script used: https://gist.github.com/kritgpt/69bfc7ef78e1784716c693d531a58adf.

[nlopes]

tests failing

[kritgpt]

Split naive cases from files that need manual checking.

Patch now contains only naive cases.

[nlopes]

LGTM

[nlopes]

This broke the build. I had to revert the patch.
You need to rebase your patch to latest git main & update the wave32.ll test.

llvm/test/CodeGen/AMDGPU/GlobalISel/br-constant-invalid-sgpr-copy.ll

 
 bb0:
   br i1 true, label %.exit5, label %bb0
 
 .exit5:
   ret void
 }
 
-define void @br_undef() {
-; WAVE64-LABEL: br_undef:
+define void @br_undef(i1 %c0) {
+; WAVE64-LABEL: br_[[%c0:%.*]]:
 ; WAVE64:       ; %bb.0: ; %.exit
 ; WAVE64-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; WAVE64-NEXT:  .LBB2_1: ; %bb0
 ; WAVE64-NEXT:    ; =>This Inner Loop Header: Depth=1
 ; WAVE64-NEXT:    ; implicit-def: $sgpr4
 ; WAVE64-NEXT:    s_and_b32 s4, s4, 1
 ; WAVE64-NEXT:    s_cmp_lg_u32 s4, 0
 ; WAVE64-NEXT:    s_cbranch_scc1 .LBB2_1
 ; WAVE64-NEXT:  ; %bb.2: ; %.exit5
 ; WAVE64-NEXT:    s_setpc_b64 s[30:31]
 ;
-; WAVE32-LABEL: br_undef:
+; WAVE32-LABEL: br_[[%c1:%.*]]:
 ; WAVE32:       ; %bb.0: ; %.exit
 ; WAVE32-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; WAVE32-NEXT:    s_waitcnt_vscnt null, 0x0
 ; WAVE32-NEXT:  .LBB2_1: ; %bb0
 ; WAVE32-NEXT:    ; =>This Inner Loop Header: Depth=1
 ; WAVE32-NEXT:    ; implicit-def: $sgpr4
 ; WAVE32-NEXT:    s_and_b32 s4, s4, 1
 ; WAVE32-NEXT:    s_cmp_lg_u32 s4, 0
 ; WAVE32-NEXT:    s_cbranch_scc1 .LBB2_1
 ; WAVE32-NEXT:  ; %bb.2: ; %.exit5
 ; WAVE32-NEXT:    s_setpc_b64 s[30:31]
 .exit:
   br label %bb0
 
 bb0:
-  br i1 undef, label %.exit5, label %bb0
+  br i1 %c0, label %.exit5, label %bb0
 
 .exit5:
   ret void
 }
 
 define void @br_poison() {
 ; WAVE64-LABEL: br_poison:
 ; WAVE64:       ; %bb.0: ; %.exit

llvm/test/CodeGen/AMDGPU/agpr-copy-no-free-registers.ll

; GFX90A-NEXT:    s_setpc_b64 s[30:31]
   %v32 = call i32 asm sideeffect "; def $0","=${v32}"()
   %agpr1 = call i32 asm sideeffect "; copy ", "={a1},a,~{a[0:15]}"(<16 x float> %agpr0)
   %agpr2 = call i32 asm sideeffect "; copy ", "={a2},a,{a[0:15]}"(i32 %agpr1, <16 x float> %mfma)
   call void asm sideeffect "; use $0 $1","{a3},{v[0:31]}"(i32 %agpr2, <32 x i32> %vgpr0)
   call void asm sideeffect "; use $0","${v32}"(i32 %v32)
   ret void
 }
 
-define amdgpu_kernel void @introduced_copy_to_sgpr(i64 %arg, i32 %arg1, i32 %arg2, i64 %arg3, <2 x half> %arg4, <2 x half> %arg5) #3 {
+define amdgpu_kernel void @introduced_copy_to_sgpr(i64 %arg, i32 %arg1, i32 %arg2, i64 %arg3, <2 x half> %arg4, <2 x half> %arg5, i1 %c0) #3 {
 ; GFX908-LABEL: introduced_copy_to_sgpr:
 ; GFX908:       ; %bb.0: ; %bb
 ; GFX908-NEXT:    global_load_ushort v16, v[0:1], off glc
 ; GFX908-NEXT:    s_load_dwordx4 s[0:3], s[4:5], 0x0
 ; GFX908-NEXT:    s_load_dwordx2 s[6:7], s[4:5], 0x10
 ; GFX908-NEXT:    s_load_dword s9, s[4:5], 0x18
 ; GFX908-NEXT:    s_mov_b32 s8, 0
 ; GFX908-NEXT:    s_mov_b32 s5, s8
bb:
   %i = load volatile i16, ptr addrspace(4) undef, align 2
   %i6 = zext i16 %i to i64
   %i7 = udiv i32 %arg1, %arg2
   %i8 = zext i32 %i7 to i64
   br label %bb9
 
 bb9:                                              ; preds = %bb12, %bb
   %i10 = phi i64 [ %arg3, %bb ], [ %i13, %bb12 ]
-  br i1 undef, label %bb14, label %bb12
+  br i1 %c0, label %bb14, label %bb12
 
 bb12:                                             ; preds = %bb58, %bb9
   %i13 = add nuw nsw i64 %i10, %i8
   br label %bb9
 
 bb14:                                             ; preds = %bb9
   %i11 = icmp slt i64 %i10, 0
   %i15 = load i64, ptr addrspace(1) null, align 8

llvm/test/CodeGen/AMDGPU/bug-vopc-commute.ll

 ; RUN: llc -march=amdgcn -mcpu=verde -verify-machineinstrs < %s | FileCheck %s
 ; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck %s
 
 ; CHECK-LABEL: {{^}}main:
 ;
 ; Test for compilation only. This generated an invalid machine instruction
 ; by trying to commute the operands of a V_CMP_EQ_i32_e32 instruction, both
 ; of which were in SGPRs.
-define amdgpu_vs float @main(i32 %v) {
+define amdgpu_vs float @main(i32 %v, i1 %c0) {
 main_body:
   %d1 = call float @llvm.amdgcn.s.buffer.load.f32(<4 x i32> undef, i32 960, i32 0)
   %d2 = call float @llvm.amdgcn.s.buffer.load.f32(<4 x i32> undef, i32 976, i32 0)
-  br i1 undef, label %ENDIF56, label %IF57
+  br i1 %c0, label %ENDIF56, label %IF57
 
 IF57:                                             ; preds = %ENDIF
   %v.1 = mul i32 %v, 2
   br label %ENDIF56
 
 ENDIF56:                                          ; preds = %IF57, %ENDIF
   %v.2 = phi i32 [ %v, %main_body ], [ %v.1, %IF57 ]
   %d1.i = bitcast float %d1 to i32

llvm/test/CodeGen/AMDGPU/cf-loop-on-constant.ll

for.body:
   %arrayidx = getelementptr float, ptr addrspace(3) %ptr, i32 %tmp
   %vecload = load float, ptr addrspace(3) %arrayidx, align 4
   %add = fadd float %vecload, 1.0
   store float %add, ptr addrspace(3) %arrayidx, align 8
   %inc = add i32 %indvar, 1
   br i1 false, label %for.body, label %for.exit
 }
 
-define amdgpu_kernel void @loop_const_undef(ptr addrspace(3) %ptr, i32 %n) nounwind {
+define amdgpu_kernel void @loop_const_undef(ptr addrspace(3) %ptr, i32 %n, i1 %c0) nounwind {
 ; GCN-LABEL: loop_const_undef:
 ; GCN:       ; %bb.0: ; %entry
 ; GCN-NEXT:    s_load_dword s0, s[0:1], 0x9
 ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 ; GCN-NEXT:    v_mov_b32_e32 v0, s0
 ; GCN-NEXT:    s_mov_b32 m0, -1
 ; GCN-NEXT:    ds_read_b32 v1, v0 offset:128
 ; GCN-NEXT:    s_waitcnt lgkmcnt(0)

 for.body:
   %indvar = phi i32 [ %inc, %for.body ], [ 0, %entry ]
   %tmp = add i32 %indvar, 32
   %arrayidx = getelementptr float, ptr addrspace(3) %ptr, i32 %tmp
   %vecload = load float, ptr addrspace(3) %arrayidx, align 4
   %add = fadd float %vecload, 1.0
   store float %add, ptr addrspace(3) %arrayidx, align 8
   %inc = add i32 %indvar, 1
-  br i1 undef, label %for.body, label %for.exit
+  br i1 %c0, label %for.body, label %for.exit
 }
 
 define amdgpu_kernel void @loop_arg_0(ptr addrspace(3) %ptr, i32 %n) nounwind {
 ; GCN-LABEL: loop_arg_0:
 ; GCN:       ; %bb.0: ; %entry
 ; GCN-NEXT:    v_mov_b32_e32 v0, 0
 ; GCN-NEXT:    s_mov_b32 m0, -1
 ; GCN-NEXT:    ds_read_u8 v0, v0

llvm/test/CodeGen/AMDGPU/cgp-bitfield-extract.ll

 ; GCN: s_cbranch_scc{{[0-1]}}
 
 ; GCN: s_bfe_u32 s{{[0-9]+}}, s{{[0-9]+}}, 0x70008
 ; GCN: .LBB0_3:
 ; GCN: s_bfe_u32 s{{[0-9]+}}, s{{[0-9]+}}, 0x80008
 
 ; GCN: buffer_store_dword
 ; GCN: s_endpgm
-define amdgpu_kernel void @sink_ubfe_i32(ptr addrspace(1) %out, i32 %arg1) #0 {
+define amdgpu_kernel void @sink_ubfe_i32(ptr addrspace(1) %out, i32 %arg1, i1 %c0) #0 {
 entry:
   %shr = lshr i32 %arg1, 8
-  br i1 undef, label %bb0, label %bb1
+  br i1 %c0, label %bb0, label %bb1
 
 bb0:
   %val0 = and i32 %shr, 255
   store volatile i32 0, ptr addrspace(1) undef
   br label %ret
 
 bb1:
   %val1 = and i32 %shr, 127

 ; OPT-NEXT: %val1 = and i32 %1, 127
 ; OPT: br label
 
 ; OPT: ret:
 ; OPT: store
 ; OPT: ret
 
 ; GCN-LABEL: {{^}}sink_sbfe_i32:
-define amdgpu_kernel void @sink_sbfe_i32(ptr addrspace(1) %out, i32 %arg1) #0 {
+define amdgpu_kernel void @sink_sbfe_i32(ptr addrspace(1) %out, i32 %arg1, i1 %c0) #0 {
 entry:
   %shr = ashr i32 %arg1, 8
-  br i1 undef, label %bb0, label %bb1
+  br i1 %c0, label %bb0, label %bb1
 
 bb0:
   %val0 = and i32 %shr, 255
   store volatile i32 0, ptr addrspace(1) undef
   br label %ret
 
 bb1:
   %val1 = and i32 %shr, 127

 ; VI: v_mov_b32_e32 v{{[0-9]+}}, 0x7f
 
 ; GCN: .LBB2_3:
 ; SI: s_bfe_u32 s{{[0-9]+}}, s{{[0-9]+}}, 0x80004
 ; VI: v_mov_b32_e32 v{{[0-9]+}}, 0xff
 
 ; GCN: buffer_store_short
 ; GCN: s_endpgm
-define amdgpu_kernel void @sink_ubfe_i16(ptr addrspace(1) %out, i16 %arg1) #0 {
+define amdgpu_kernel void @sink_ubfe_i16(ptr addrspace(1) %out, i16 %arg1, i1 %c0) #0 {
 entry:
   %shr = lshr i16 %arg1, 4
-  br i1 undef, label %bb0, label %bb1
+  br i1 %c0, label %bb0, label %bb1
 
 bb0:
   %val0 = and i16 %shr, 255
   store volatile i16 0, ptr addrspace(1) undef
   br label %ret
 
 bb1:
   %val1 = and i16 %shr, 127

 ; GCN: s_cbranch_scc{{[0-1]}} .LBB3_2
 ; GCN: v_alignbit_b32 v[[LO:[0-9]+]], s{{[0-9]+}}, v{{[0-9]+}}, 30
 ; GCN: v_and_b32_e32 v{{[0-9]+}}, 0x7f, v[[LO]]
 
 ; GCN: .LBB3_3:
 ; GCN: v_and_b32_e32 v{{[0-9]+}}, 0xff, v[[LO]]
 
 ; GCN: buffer_store_dwordx2
-define amdgpu_kernel void @sink_ubfe_i64_span_midpoint(ptr addrspace(1) %out, i64 %arg1) #0 {
+define amdgpu_kernel void @sink_ubfe_i64_span_midpoint(ptr addrspace(1) %out, i64 %arg1, i1 %c0) #0 {
 entry:
   %shr = lshr i64 %arg1, 30
-  br i1 undef, label %bb0, label %bb1
+  br i1 %c0, label %bb0, label %bb1
 
 bb0:
   %val0 = and i64 %shr, 255
   store volatile i32 0, ptr addrspace(1) undef
   br label %ret
 
 bb1:
   %val1 = and i64 %shr, 127

 ; GCN: s_cbranch_scc{{[0-1]}} .LBB4_2
 
 ; GCN: s_bfe_u32 s{{[0-9]+}}, s{{[0-9]+}}, 0x7000f
 
 ; GCN: .LBB4_3:
 ; GCN: s_bfe_u32 s{{[0-9]+}}, s{{[0-9]+}}, 0x8000f
 
 ; GCN: buffer_store_dwordx2
-define amdgpu_kernel void @sink_ubfe_i64_low32(ptr addrspace(1) %out, i64 %arg1) #0 {
+define amdgpu_kernel void @sink_ubfe_i64_low32(ptr addrspace(1) %out, i64 %arg1, i1 %c0) #0 {
 entry:
   %shr = lshr i64 %arg1, 15
-  br i1 undef, label %bb0, label %bb1
+  br i1 %c0, label %bb0, label %bb1
 
 bb0:
   %val0 = and i64 %shr, 255
   store volatile i32 0, ptr addrspace(1) undef
   br label %ret
 
 bb1:
   %val1 = and i64 %shr, 127

 ; GCN-LABEL: {{^}}sink_ubfe_i64_high32:
 ; GCN: s_cbranch_scc{{[0-1]}} .LBB5_2
 ; GCN: s_bfe_u32 s{{[0-9]+}}, s{{[0-9]+}}, 0x70003
 
 ; GCN: .LBB5_3:
 ; GCN: s_bfe_u32 s{{[0-9]+}}, s{{[0-9]+}}, 0x80003
 
 ; GCN: buffer_store_dwordx2
-define amdgpu_kernel void @sink_ubfe_i64_high32(ptr addrspace(1) %out, i64 %arg1) #0 {
+define amdgpu_kernel void @sink_ubfe_i64_high32(ptr addrspace(1) %out, i64 %arg1, i1 %c0) #0 {
 entry:
   %shr = lshr i64 %arg1, 35
-  br i1 undef, label %bb0, label %bb1
+  br i1 %c0, label %bb0, label %bb1
 
 bb0:
   %val0 = and i64 %shr, 255
   store volatile i32 0, ptr addrspace(1) undef
   br label %ret
 
 bb1:
   %val1 = and i64 %shr, 127

llvm/test/CodeGen/AMDGPU/coalescer_distribute.ll

 ; RUN: llc -mtriple=amdgcn-- -verify-machineinstrs -o /dev/null %s
 ; This testcase produces a situation with unused value numbers in subregister
 ; liveranges that get distributed by ConnectedVNInfoEqClasses.
 
-define amdgpu_kernel void @hoge() {
+define amdgpu_kernel void @hoge(i1 %c0, i1 %c1, i1 %c2, i1 %c3, i1 %c4) {
 bb:
   %tmp = tail call i32 @llvm.amdgcn.workitem.id.x()
-  br i1 undef, label %bb2, label %bb23
+  br i1 %c0, label %bb2, label %bb23
 
 bb2:
-  br i1 undef, label %bb6, label %bb8
+  br i1 %c1, label %bb6, label %bb8
 
 bb6:
   %tmp7 = or i64 undef, undef
   br label %bb8
 
 bb8:
   %tmp9 = phi i64 [ %tmp7, %bb6 ], [ undef, %bb2 ]
   %tmp10 = icmp eq i32 %tmp, 0
   br i1 %tmp10, label %bb11, label %bb23
 
 bb11:
-  br i1 undef, label %bb20, label %bb17
+  br i1 %c2, label %bb20, label %bb17
 
 bb17:
   br label %bb20
 
 bb20:
   %tmp21 = phi i64 [ undef, %bb17 ], [ %tmp9, %bb11 ]
   %tmp22 = trunc i64 %tmp21 to i32
   br label %bb23
 
 bb23:
   %tmp24 = phi i32 [ %tmp22, %bb20 ], [ undef, %bb8 ], [ undef, %bb ]
   br label %bb25
 
 bb25:
   %tmp26 = phi i32 [ %tmp24, %bb23 ], [ undef, %bb25 ]
-  br i1 undef, label %bb25, label %bb30
+  br i1 %c3, label %bb25, label %bb30
 
 bb30:
-  br i1 undef, label %bb32, label %bb34
+  br i1 %c4, label %bb32, label %bb34
 
 bb32:
   %tmp33 = zext i32 %tmp26 to i64
   br label %bb34
 
 bb34:
   ret void
 }
 
 declare i32 @llvm.amdgcn.workitem.id.x()

llvm/test/CodeGen/AMDGPU/combine-add-zext-xor.ll

 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
 ; RUN: llc -march=amdgcn -mcpu=gfx1010 -verify-machineinstrs < %s | FileCheck --check-prefix=GFX1010 %s
 ; RUN: llc -march=amdgcn -mcpu=gfx1100 -verify-machineinstrs < %s | FileCheck --check-prefix=GFX1100 %s
 
 ; Test that unused lanes in the s_xor result are masked out with v_cndmask.
 
-define i32 @combine_add_zext_xor() {
+define i32 @combine_add_zext_xor(i1 %c0) {
 ; GFX1010-LABEL: combine_add_zext_xor:
 ; GFX1010:       ; %bb.0: ; %.entry
 ; GFX1010-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; GFX1010-NEXT:    s_waitcnt_vscnt null, 0x0
 ; GFX1010-NEXT:    v_mov_b32_e32 v1, 0
 ; GFX1010-NEXT:    s_branch .LBB0_2
 ; GFX1010-NEXT:  .LBB0_1: ; %bb9
 ; GFX1010-NEXT:    ; in Loop: Header=BB0_2 Depth=1

 ; GFX1100-NEXT:    s_branch .LBB0_1
 ; GFX1100-NEXT:  .LBB0_4: ; %.exit
 ; GFX1100-NEXT:    s_setpc_b64 s[30:31]
 .entry:
   br label %.a
 
 .a:                                               ; preds = %bb9, %.entry
   %.2 = phi i32 [ 0, %.entry ], [ %i11, %bb9 ]
-  br i1 undef, label %bb9, label %bb
+  br i1 %c0, label %bb9, label %bb
 
 bb:                                               ; preds = %.a
   %.i3 = call i32 @llvm.amdgcn.raw.buffer.load.i32(<4 x i32> undef, i32 %.2, i32 64, i32 1)
   %i5 = icmp eq i32 %.i3, 0
   br label %bb9
 
 bb9:                                              ; preds = %bb, %.a
   %.2.0.in.in = phi i1 [ %i5, %bb ], [ undef, %.a ]
   %.2.0.in = xor i1 %.2.0.in.in, true
   %.2.0 = zext i1 %.2.0.in to i32
   %i11 = add i32 %.2, %.2.0
   %i12 = icmp sgt i32 %.2, -1050
   br i1 %i12, label %.a, label %.exit
 
 .exit:                                            ; preds = %bb9
   ret i32 %.2.0
 }
 
 ; Test that unused lanes in the s_xor result are masked out with v_cndmask.
 
-define i32 @combine_sub_zext_xor() {
+define i32 @combine_sub_zext_xor(i1 %c0) {
 ; GFX1010-LABEL: combine_sub_zext_xor:
 ; GFX1010:       ; %bb.0: ; %.entry
 ; GFX1010-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; GFX1010-NEXT:    s_waitcnt_vscnt null, 0x0
 ; GFX1010-NEXT:    v_mov_b32_e32 v1, 0
 ; GFX1010-NEXT:    s_branch .LBB1_2
 ; GFX1010-NEXT:  .LBB1_1: ; %bb9
 ; GFX1010-NEXT:    ; in Loop: Header=BB1_2 Depth=1

 ; GFX1100-NEXT:    s_branch .LBB1_1
 ; GFX1100-NEXT:  .LBB1_4: ; %.exit
 ; GFX1100-NEXT:    s_setpc_b64 s[30:31]
 .entry:
   br label %.a
 
 .a:                                               ; preds = %bb9, %.entry
   %.2 = phi i32 [ 0, %.entry ], [ %i11, %bb9 ]
-  br i1 undef, label %bb9, label %bb
+  br i1 %c0, label %bb9, label %bb
 
 bb:                                               ; preds = %.a
   %.i3 = call i32 @llvm.amdgcn.raw.buffer.load.i32(<4 x i32> undef, i32 %.2, i32 64, i32 1)
   %i5 = icmp eq i32 %.i3, 0
   br label %bb9
 
 bb9:                                              ; preds = %bb, %.a
   %.2.0.in.in = phi i1 [ %i5, %bb ], [ undef, %.a ]
   %.2.0.in = xor i1 %.2.0.in.in, true
   %.2.0 = zext i1 %.2.0.in to i32
   %i11 = sub i32 %.2, %.2.0
   %i12 = icmp sgt i32 %.2, -1050
   br i1 %i12, label %.a, label %.exit
 
 .exit:                                            ; preds = %bb9
   ret i32 %.2.0
 }
 
 ; Test that unused lanes in the s_or result are masked out with v_cndmask.
 
-define i32 @combine_add_zext_or() {
+define i32 @combine_add_zext_or(i1 %c0) {
 ; GFX1010-LABEL: combine_add_zext_or:
 ; GFX1010:       ; %bb.0: ; %.entry
 ; GFX1010-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; GFX1010-NEXT:    s_waitcnt_vscnt null, 0x0
 ; GFX1010-NEXT:    v_mov_b32_e32 v1, 0
 ; GFX1010-NEXT:    s_branch .LBB2_2
 ; GFX1010-NEXT:  .LBB2_1: ; %bb9
 ; GFX1010-NEXT:    ; in Loop: Header=BB2_2 Depth=1

 ; GFX1100-NEXT:    s_branch .LBB2_1
 ; GFX1100-NEXT:  .LBB2_4: ; %.exit
 ; GFX1100-NEXT:    s_setpc_b64 s[30:31]
 .entry:
   br label %.a
 
 .a:                                               ; preds = %bb9, %.entry
   %.2 = phi i32 [ 0, %.entry ], [ %i11, %bb9 ]
-  br i1 undef, label %bb9, label %bb
+  br i1 %c0, label %bb9, label %bb
 
 bb:                                               ; preds = %.a
   %.i3 = call i32 @llvm.amdgcn.raw.buffer.load.i32(<4 x i32> undef, i32 %.2, i32 64, i32 1)
   %i5 = icmp eq i32 %.i3, 0
   br label %bb9
 
 bb9:                                              ; preds = %bb, %.a
   %.2.0.in.in = phi i1 [ %i5, %bb ], [ undef, %.a ]
   %t = icmp sgt i32 %.2, -1050
   %.2.0.in = or i1 %.2.0.in.in, %t
   %.2.0 = zext i1 %.2.0.in to i32
   %i11 = add i32 %.2, %.2.0
   %i12 = icmp sgt i32 %.2, -1050
   br i1 %i12, label %.a, label %.exit
 
 .exit:                                            ; preds = %bb9
   ret i32 %.2.0
 }
 
 ; Test that unused lanes in the s_or result are masked out with v_cndmask.
 
-define i32 @combine_sub_zext_or() {
+define i32 @combine_sub_zext_or(i1 %c0) {
 ; GFX1010-LABEL: combine_sub_zext_or:
 ; GFX1010:       ; %bb.0: ; %.entry
 ; GFX1010-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; GFX1010-NEXT:    s_waitcnt_vscnt null, 0x0
 ; GFX1010-NEXT:    v_mov_b32_e32 v1, 0
 ; GFX1010-NEXT:    s_branch .LBB3_2
 ; GFX1010-NEXT:  .LBB3_1: ; %bb9
 ; GFX1010-NEXT:    ; in Loop: Header=BB3_2 Depth=1

 ; GFX1100-NEXT:    s_branch .LBB3_1
 ; GFX1100-NEXT:  .LBB3_4: ; %.exit
 ; GFX1100-NEXT:    s_setpc_b64 s[30:31]
 .entry:
   br label %.a
 
 .a:                                               ; preds = %bb9, %.entry
   %.2 = phi i32 [ 0, %.entry ], [ %i11, %bb9 ]
-  br i1 undef, label %bb9, label %bb
+  br i1 %c0, label %bb9, label %bb
 
 bb:                                               ; preds = %.a
   %.i3 = call i32 @llvm.amdgcn.raw.buffer.load.i32(<4 x i32> undef, i32 %.2, i32 64, i32 1)
   %i5 = icmp eq i32 %.i3, 0
   br label %bb9
 
 bb9:                                              ; preds = %bb, %.a
   %.2.0.in.in = phi i1 [ %i5, %bb ], [ undef, %.a ]
   %t = icmp sgt i32 %.2, -1050
   %.2.0.in = or i1 %.2.0.in.in, %t
   %.2.0 = zext i1 %.2.0.in to i32
   %i11 = sub i32 %.2, %.2.0
   %i12 = icmp sgt i32 %.2, -1050
   br i1 %i12, label %.a, label %.exit
 
 .exit:                                            ; preds = %bb9
   ret i32 %.2.0
 }
 
 ; Test that unused lanes in the s_and result are masked out with v_cndmask.
 
-define i32 @combine_add_zext_and() {
+define i32 @combine_add_zext_and(i1 %c0) {
 ; GFX1010-LABEL: combine_add_zext_and:
 ; GFX1010:       ; %bb.0: ; %.entry
 ; GFX1010-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; GFX1010-NEXT:    s_waitcnt_vscnt null, 0x0
 ; GFX1010-NEXT:    v_mov_b32_e32 v1, 0
 ; GFX1010-NEXT:    s_branch .LBB4_2
 ; GFX1010-NEXT:  .LBB4_1: ; %bb9
 ; GFX1010-NEXT:    ; in Loop: Header=BB4_2 Depth=1

 ; GFX1100-NEXT:    s_branch .LBB4_1
 ; GFX1100-NEXT:  .LBB4_4: ; %.exit
 ; GFX1100-NEXT:    s_setpc_b64 s[30:31]
 .entry:
   br label %.a
 
 .a:                                               ; preds = %bb9, %.entry
   %.2 = phi i32 [ 0, %.entry ], [ %i11, %bb9 ]
-  br i1 undef, label %bb9, label %bb
+  br i1 %c0, label %bb9, label %bb
 
 bb:                                               ; preds = %.a
   %.i3 = call i32 @llvm.amdgcn.raw.buffer.load.i32(<4 x i32> undef, i32 %.2, i32 64, i32 1)
   %i5 = icmp eq i32 %.i3, 0
   br label %bb9
 
 bb9:                                              ; preds = %bb, %.a
   %.2.0.in.in = phi i1 [ %i5, %bb ], [ undef, %.a ]
   %t = icmp sgt i32 %.2, -1050
   %.2.0.in = and i1 %.2.0.in.in, %t
   %.2.0 = zext i1 %.2.0.in to i32
   %i11 = add i32 %.2, %.2.0
   %i12 = icmp sgt i32 %.2, -1050
   br i1 %i12, label %.a, label %.exit
 
 .exit:                                            ; preds = %bb9
   ret i32 %.2.0
 }
 
 ; Test that unused lanes in the s_and result are masked out with v_cndmask.
 
-define i32 @combine_sub_zext_and() {
+define i32 @combine_sub_zext_and(i1 %c0) {
 ; GFX1010-LABEL: combine_sub_zext_and:
 ; GFX1010:       ; %bb.0: ; %.entry
 ; GFX1010-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; GFX1010-NEXT:    s_waitcnt_vscnt null, 0x0
 ; GFX1010-NEXT:    v_mov_b32_e32 v1, 0
 ; GFX1010-NEXT:    s_branch .LBB5_2
 ; GFX1010-NEXT:  .LBB5_1: ; %bb9
 ; GFX1010-NEXT:    ; in Loop: Header=BB5_2 Depth=1

 ; GFX1100-NEXT:    s_branch .LBB5_1
 ; GFX1100-NEXT:  .LBB5_4: ; %.exit
 ; GFX1100-NEXT:    s_setpc_b64 s[30:31]
 .entry:
   br label %.a
 
 .a:                                               ; preds = %bb9, %.entry
   %.2 = phi i32 [ 0, %.entry ], [ %i11, %bb9 ]
-  br i1 undef, label %bb9, label %bb
+  br i1 %c0, label %bb9, label %bb
 
 bb:                                               ; preds = %.a
   %.i3 = call i32 @llvm.amdgcn.raw.buffer.load.i32(<4 x i32> undef, i32 %.2, i32 64, i32 1)
   %i5 = icmp eq i32 %.i3, 0
   br label %bb9
 
 bb9:                                              ; preds = %bb, %.a
   %.2.0.in.in = phi i1 [ %i5, %bb ], [ undef, %.a ]

llvm/test/CodeGen/AMDGPU/dagcomb-shuffle-vecextend-non2.ll

 ; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
 
 ; We are only checking that instruction selection can succeed in this case. This
 ; cut down test results in no instructions, but that's fine.
 ;
 ; See https://llvm.org/PR33743 for details of the bug being addressed
 ;
 ; Checking that shufflevector with 3-vec mask is handled in
 ; combineShuffleToVectorExtend
 ;
 ; GCN: s_endpgm
 
-define amdgpu_ps void @main(i32 %in1) local_unnamed_addr {
+define amdgpu_ps void @main(i32 %in1, i1 %c0) local_unnamed_addr {
 .entry:
-  br i1 undef, label %bb12, label %bb
+  br i1 %c0, label %bb12, label %bb
 
 bb:
   %__llpc_global_proxy_r5.12.vec.insert = insertelement <4 x i32> undef, i32 %in1, i32 3
   %tmp3 = shufflevector <4 x i32> %__llpc_global_proxy_r5.12.vec.insert, <4 x i32> undef, <3 x i32> <i32 undef, i32 undef, i32 1>
   %tmp4 = bitcast <3 x i32> %tmp3 to <3 x float>
   %a2.i123 = extractelement <3 x float> %tmp4, i32 2
   %tmp5 = bitcast float %a2.i123 to i32
   %__llpc_global_proxy_r2.0.vec.insert196 = insertelement <4 x i32> undef, i32 %tmp5, i32 0

llvm/test/CodeGen/AMDGPU/early-if-convert.ll

done:
   store i32 %value, ptr addrspace(1) %out
   ret void
 }
 
 ; GCN-LABEL: {{^}}ifcvt_undef_scc:
 ; GCN: {{^}}; %bb.0:
 ; GCN-NEXT: s_load_dwordx2
 ; GCN-NEXT: s_cselect_b32 s{{[0-9]+}}, 0, 1{{$}}
-define amdgpu_kernel void @ifcvt_undef_scc(i32 %cond, ptr addrspace(1) %out) {
+define amdgpu_kernel void @ifcvt_undef_scc(i32 %cond, ptr addrspace(1) %out, i1 %c0) {
 entry:
-  br i1 undef, label %else, label %if
+  br i1 %c0, label %else, label %if
 
 if:
   br label %done
 
 else:
   br label %done
 
 done:

llvm/test/CodeGen/AMDGPU/extract-subvector-16bit.ll

 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
 ; RUN: llc -march=amdgcn -mtriple=amdgcn-- -verify-machineinstrs -o - %s | FileCheck -check-prefix=SI %s
 ; RUN: llc -march=amdgcn -mtriple=amdgcn-- -mcpu=gfx900 -verify-machineinstrs -o - %s | FileCheck -check-prefix=GFX9 %s
 
-define <4 x i16> @vec_8xi16_extract_4xi16(ptr addrspace(1) %p0, ptr addrspace(1) %p1) {
+define <4 x i16> @vec_8xi16_extract_4xi16(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) {
 ; SI-LABEL: vec_8xi16_extract_4xi16:
 ; SI:       ; %bb.0:
 ; SI-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; SI-NEXT:    s_cbranch_scc0 .LBB0_2
 ; SI-NEXT:  ; %bb.1: ; %F
 ; SI-NEXT:    s_mov_b32 s6, 0
 ; SI-NEXT:    s_mov_b32 s7, 0xf000
 ; SI-NEXT:    s_mov_b32 s4, s6

 ; GFX9-NEXT:    v_or_b32_sdwa v3, v0, s4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
 ; GFX9-NEXT:    v_pk_ashrrev_i16 v0, 15, v2 op_sel_hi:[0,1]
 ; GFX9-NEXT:    v_or_b32_e32 v2, 0xffff8000, v0
 ; GFX9-NEXT:    v_or_b32_sdwa v0, v0, s4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
 ; GFX9-NEXT:    s_mov_b32 s4, 0x5040100
 ; GFX9-NEXT:    v_perm_b32 v0, v0, v2, s4
 ; GFX9-NEXT:    v_perm_b32 v1, v3, v1, s4
 ; GFX9-NEXT:    s_setpc_b64 s[30:31]
-  br i1 undef, label %T, label %F
+  br i1 %c0, label %T, label %F
 
 T:
   %t = load volatile <8 x i16>, ptr addrspace(1) %p0
   br label %exit
 
 F:
   %f = load volatile <8 x i16>, ptr addrspace(1) %p1
   br label %exit
 
 exit:
   %m = phi <8 x i16> [ %t, %T ], [ %f, %F ]
   %v2 = shufflevector <8 x i16> %m, <8 x i16> undef, <4 x i32> <i32 0, i32 1, i32 2, i32 2>
   %b2 = icmp sgt <4 x i16> %v2, <i16 -1, i16 -1, i16 -1, i16 -1>
   %r2 = select <4 x i1> %b2, <4 x i16> <i16 -32768, i16 -32768, i16 -32768, i16 -32768>, <4 x i16> <i16 -1, i16 -1, i16 -1, i16 -1>
   ret <4 x i16> %r2
 }
 
-define <4 x i16> @vec_8xi16_extract_4xi16_2(ptr addrspace(1) %p0, ptr addrspace(1) %p1) {
+define <4 x i16> @vec_8xi16_extract_4xi16_2(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) {
 ; SI-LABEL: vec_8xi16_extract_4xi16_2:
 ; SI:       ; %bb.0:
 ; SI-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; SI-NEXT:    s_cbranch_scc0 .LBB1_2
 ; SI-NEXT:  ; %bb.1: ; %F
 ; SI-NEXT:    s_mov_b32 s6, 0
 ; SI-NEXT:    s_mov_b32 s7, 0xf000
 ; SI-NEXT:    s_mov_b32 s4, s6

 ; GFX9-NEXT:    v_or_b32_sdwa v2, v0, s4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
 ; GFX9-NEXT:    v_pk_ashrrev_i16 v0, 15, v4 op_sel_hi:[0,1]
 ; GFX9-NEXT:    v_or_b32_e32 v3, 0xffff8000, v0
 ; GFX9-NEXT:    v_or_b32_sdwa v0, v0, s4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
 ; GFX9-NEXT:    s_mov_b32 s4, 0x5040100
 ; GFX9-NEXT:    v_perm_b32 v0, v0, v3, s4
 ; GFX9-NEXT:    v_perm_b32 v1, v2, v1, s4
 ; GFX9-NEXT:    s_setpc_b64 s[30:31]
-  br i1 undef, label %T, label %F
+  br i1 %c0, label %T, label %F
 
 T:
   %t = load volatile <8 x i16>, ptr addrspace(1) %p0
   br label %exit
 
 F:
   %f = load volatile <8 x i16>, ptr addrspace(1) %p1
   br label %exit
 
 exit:
   %m = phi <8 x i16> [ %t, %T ], [ %f, %F ]
   %v2 = shufflevector <8 x i16> %m, <8 x i16> undef, <4 x i32> <i32 4, i32 5, i32 6, i32 7>
   %b2 = icmp sgt <4 x i16> %v2, <i16 -1, i16 -1, i16 -1, i16 -1>
   %r2 = select <4 x i1> %b2, <4 x i16> <i16 -32768, i16 -32768, i16 -32768, i16 -32768>, <4 x i16> <i16 -1, i16 -1, i16 -1, i16 -1>
   ret <4 x i16> %r2
 }
 
-define <4 x half> @vec_8xf16_extract_4xf16(ptr addrspace(1) %p0, ptr addrspace(1) %p1) {
+define <4 x half> @vec_8xf16_extract_4xf16(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) {
 ; SI-LABEL: vec_8xf16_extract_4xf16:
 ; SI:       ; %bb.0:
 ; SI-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; SI-NEXT:    s_cbranch_scc0 .LBB2_2
 ; SI-NEXT:  ; %bb.1: ; %F
 ; SI-NEXT:    s_mov_b32 s6, 0
 ; SI-NEXT:    s_mov_b32 s7, 0xf000
 ; SI-NEXT:    s_mov_b32 s4, s6

 ; GFX9-NEXT:    v_cndmask_b32_e32 v6, v4, v3, vcc
 ; GFX9-NEXT:    v_cmp_ge_f16_e32 vcc, 0.5, v2
 ; GFX9-NEXT:    v_cndmask_b32_e32 v0, v3, v4, vcc
 ; GFX9-NEXT:    v_cmp_le_f16_sdwa vcc, v2, v1 src0_sel:WORD_1 src1_sel:DWORD
 ; GFX9-NEXT:    v_cndmask_b32_e32 v1, v3, v4, vcc
 ; GFX9-NEXT:    v_pack_b32_f16 v0, v0, v1
 ; GFX9-NEXT:    v_pack_b32_f16 v1, v5, v6
 ; GFX9-NEXT:    s_setpc_b64 s[30:31]
-  br i1 undef, label %T, label %F
+  br i1 %c0, label %T, label %F
 
 T:
   %t = load volatile <8 x half>, ptr addrspace(1) %p0
   br label %exit
 
 F:
   %f = load volatile <8 x half>, ptr addrspace(1) %p1
   br label %exit
 
 exit:
   %m = phi <8 x half> [ %t, %T ], [ %f, %F ]
   %v2 = shufflevector <8 x half> %m, <8 x half> undef, <4 x i32> <i32 0, i32 1, i32 2, i32 2>
   %b2 = fcmp ugt <4 x half> %v2, <half 0xH3800, half 0xH3800, half 0xH3800, half 0xH3800>
   %r2 = select <4 x i1> %b2, <4 x half> <half 0xH3900, half 0xH3900, half 0xH3900, half 0xH3900>, <4 x half> <half 0xH3D00, half 0xH3D00, half 0xH3D00, half 0xH3D00>
   ret <4 x half> %r2
 }
 
-define <4 x i16> @vec_16xi16_extract_4xi16(ptr addrspace(1) %p0, ptr addrspace(1) %p1) {
+define <4 x i16> @vec_16xi16_extract_4xi16(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) {
 ;
 ; SI-LABEL: vec_16xi16_extract_4xi16:
 ; SI:       ; %bb.0:
 ; SI-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; SI-NEXT:    s_cbranch_scc0 .LBB3_2
 ; SI-NEXT:  ; %bb.1: ; %F
 ; SI-NEXT:    s_mov_b32 s6, 0
 ; SI-NEXT:    s_mov_b32 s7, 0xf000

 ; GFX9-NEXT:    v_or_b32_sdwa v2, v0, s4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
 ; GFX9-NEXT:    v_pk_ashrrev_i16 v0, 15, v4 op_sel_hi:[0,1]
 ; GFX9-NEXT:    v_or_b32_e32 v3, 0xffff8000, v0
 ; GFX9-NEXT:    v_or_b32_sdwa v0, v0, s4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
 ; GFX9-NEXT:    s_mov_b32 s4, 0x5040100
 ; GFX9-NEXT:    v_perm_b32 v0, v0, v3, s4
 ; GFX9-NEXT:    v_perm_b32 v1, v2, v1, s4
 ; GFX9-NEXT:    s_setpc_b64 s[30:31]
-  br i1 undef, label %T, label %F
+  br i1 %c0, label %T, label %F
 
 T:
   %t = load volatile <16 x i16>, ptr addrspace(1) %p0
   br label %exit
 
 F:
   %f = load volatile <16 x i16>, ptr addrspace(1) %p1
   br label %exit
 
 exit:
   %m = phi <16 x i16> [ %t, %T ], [ %f, %F ]
   %v2 = shufflevector <16 x i16> %m, <16 x i16> undef, <4 x i32> <i32 0, i32 1, i32 2, i32 2>
   %b2 = icmp sgt <4 x i16> %v2, <i16 -1, i16 -1, i16 -1, i16 -1>
   %r2 = select <4 x i1> %b2, <4 x i16> <i16 -32768, i16 -32768, i16 -32768, i16 -32768>, <4 x i16> <i16 -1, i16 -1, i16 -1, i16 -1>
   ret <4 x i16> %r2
 }
 
-define <4 x i16> @vec_16xi16_extract_4xi16_2(ptr addrspace(1) %p0, ptr addrspace(1) %p1) {
+define <4 x i16> @vec_16xi16_extract_4xi16_2(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) {
 ;
 ; SI-LABEL: vec_16xi16_extract_4xi16_2:
 ; SI:       ; %bb.0:
 ; SI-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; SI-NEXT:    s_cbranch_scc0 .LBB4_2
 ; SI-NEXT:  ; %bb.1: ; %F
 ; SI-NEXT:    s_mov_b32 s6, 0
 ; SI-NEXT:    s_mov_b32 s7, 0xf000

 ; GFX9-NEXT:    v_or_b32_sdwa v2, v0, s4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
 ; GFX9-NEXT:    v_pk_ashrrev_i16 v0, 15, v6 op_sel_hi:[0,1]
 ; GFX9-NEXT:    v_or_b32_e32 v3, 0xffff8000, v0
 ; GFX9-NEXT:    v_or_b32_sdwa v0, v0, s4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:DWORD
 ; GFX9-NEXT:    s_mov_b32 s4, 0x5040100
 ; GFX9-NEXT:    v_perm_b32 v0, v0, v3, s4
 ; GFX9-NEXT:    v_perm_b32 v1, v2, v1, s4
 ; GFX9-NEXT:    s_setpc_b64 s[30:31]
-  br i1 undef, label %T, label %F
+  br i1 %c0, label %T, label %F
 
 T:
   %t = load volatile <16 x i16>, ptr addrspace(1) %p0
   br label %exit
 
 F:
   %f = load volatile <16 x i16>, ptr addrspace(1) %p1
   br label %exit
 
 exit:
   %m = phi <16 x i16> [ %t, %T ], [ %f, %F ]
   %v2 = shufflevector <16 x i16> %m, <16 x i16> undef, <4 x i32> <i32 4, i32 5, i32 6, i32 7>
   %b2 = icmp sgt <4 x i16> %v2, <i16 -1, i16 -1, i16 -1, i16 -1>
   %r2 = select <4 x i1> %b2, <4 x i16> <i16 -32768, i16 -32768, i16 -32768, i16 -32768>, <4 x i16> <i16 -1, i16 -1, i16 -1, i16 -1>
   ret <4 x i16> %r2
 }
 
-define <4 x half> @vec_16xf16_extract_4xf16(ptr addrspace(1) %p0, ptr addrspace(1) %p1) {
+define <4 x half> @vec_16xf16_extract_4xf16(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) {
 ;
 ; SI-LABEL: vec_16xf16_extract_4xf16:
 ; SI:       ; %bb.0:
 ; SI-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; SI-NEXT:    s_cbranch_scc0 .LBB5_2
 ; SI-NEXT:  ; %bb.1: ; %F
 ; SI-NEXT:    s_mov_b32 s6, 0
 ; SI-NEXT:    s_mov_b32 s7, 0xf000

 ; GFX9-NEXT:    v_cndmask_b32_e32 v6, v3, v2, vcc
 ; GFX9-NEXT:    v_cmp_ge_f16_e32 vcc, 0.5, v4
 ; GFX9-NEXT:    v_cndmask_b32_e32 v0, v2, v3, vcc
 ; GFX9-NEXT:    v_cmp_le_f16_sdwa vcc, v4, v1 src0_sel:WORD_1 src1_sel:DWORD
 ; GFX9-NEXT:    v_cndmask_b32_e32 v1, v2, v3, vcc
 ; GFX9-NEXT:    v_pack_b32_f16 v0, v0, v1
 ; GFX9-NEXT:    v_pack_b32_f16 v1, v5, v6
 ; GFX9-NEXT:    s_setpc_b64 s[30:31]
-  br i1 undef, label %T, label %F
+  br i1 %c0, label %T, label %F
 
 T:
   %t = load volatile <16 x half>, ptr addrspace(1) %p0
   br label %exit
 
 F:
   %f = load volatile <16 x half>, ptr addrspace(1) %p1
   br label %exit
 
 exit:
   %m = phi <16 x half> [ %t, %T ], [ %f, %F ]
   %v2 = shufflevector <16 x half> %m, <16 x half> undef, <4 x i32> <i32 0, i32 1, i32 2, i32 2>
   %b2 = fcmp ugt <4 x half> %v2, <half 0xH3800, half 0xH3800, half 0xH3800, half 0xH3800>
   %r2 = select <4 x i1> %b2, <4 x half> <half 0xH3900, half 0xH3900, half 0xH3900, half 0xH3900>, <4 x half> <half 0xH3D00, half 0xH3D00, half 0xH3D00, half 0xH3D00>
   ret <4 x half> %r2
 }

llvm/test/CodeGen/AMDGPU/extract-subvector.ll

 ; GCN: buffer_load_ushort
 ; GCN: buffer_load_ushort
 ; GCN: buffer_load_ushort
 ; GCN: buffer_load_ushort
 ; GCN: buffer_load_ushort
 ; GCN: v_bfe_i32
 ; GCN: v_bfe_i32
 
-define <2 x i16> @extract_2xi16(ptr addrspace(1) %p0, ptr addrspace(1) %p1) {
-  br i1 undef, label %T, label %F
+define <2 x i16> @extract_2xi16(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) {
+  br i1 %c0, label %T, label %F
 
 T:
   %t = load volatile <8 x i16>, ptr addrspace(1) %p0
   br label %exit
 
 F:
   %f = load volatile <8 x i16>, ptr addrspace(1) %p1
   br label %exit
 
 exit:
   %m = phi <8 x i16> [ %t, %T ], [ %f, %F ]
   %v2 = shufflevector <8 x i16> %m, <8 x i16> undef, <2 x i32> <i32 0, i32 1>
   %b2 = icmp sgt <2 x i16> %v2, <i16 -1, i16 -1>
   %r2 = select <2 x i1> %b2, <2 x i16> <i16 -32768, i16 -32768>, <2 x i16> <i16 -1, i16 -1>
   ret <2 x i16> %r2
 }
 
 ; GCN-LABEL: extract_2xi64
 ; GCN-COUNT-2: v_cndmask_b32
-define <2 x i64> @extract_2xi64(ptr addrspace(1) %p0, ptr addrspace(1) %p1) {
-  br i1 undef, label %T, label %F
+define <2 x i64> @extract_2xi64(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) {
+  br i1 %c0, label %T, label %F
 
 T:
   %t = load volatile <8 x i64>, ptr addrspace(1) %p0
   br label %exit
 
 F:
   %f = load volatile <8 x i64>, ptr addrspace(1) %p1
   br label %exit
 
 exit:
   %m = phi <8 x i64> [ %t, %T ], [ %f, %F ]
   %v2 = shufflevector <8 x i64> %m, <8 x i64> undef, <2 x i32> <i32 0, i32 1>
   %b2 = icmp sgt <2 x i64> %v2, <i64 -1, i64 -1>
   %r2 = select <2 x i1> %b2, <2 x i64> <i64 -32768, i64 -32768>, <2 x i64> <i64 -1, i64 -1>
   ret <2 x i64> %r2
 }
 
 ; GCN-LABEL: extract_4xi64
 ; GCN-COUNT-4: v_cndmask_b32
-define <4 x i64> @extract_4xi64(ptr addrspace(1) %p0, ptr addrspace(1) %p1) {
-  br i1 undef, label %T, label %F
+define <4 x i64> @extract_4xi64(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) {
+  br i1 %c0, label %T, label %F
 
 T:
   %t = load volatile <8 x i64>, ptr addrspace(1) %p0
   br label %exit
 
 F:
   %f = load volatile <8 x i64>, ptr addrspace(1) %p1
   br label %exit
 
 exit:
   %m = phi <8 x i64> [ %t, %T ], [ %f, %F ]
   %v2 = shufflevector <8 x i64> %m, <8 x i64> undef, <4 x i32> <i32 0, i32 1, i32 2, i32 3>
   %b2 = icmp sgt <4 x i64> %v2, <i64 -1, i64 -1, i64 -1, i64 -1>
   %r2 = select <4 x i1> %b2, <4 x i64> <i64 -32768, i64 -32768, i64 -32768, i64 -32768>, <4 x i64> <i64 -1, i64 -1, i64 -1, i64 -1>
   ret <4 x i64> %r2
 }
 
 ; GCN-LABEL: extract_8xi64
 ; GCN-COUNT-8: v_cndmask_b32
-define <8 x i64> @extract_8xi64(ptr addrspace(1) %p0, ptr addrspace(1) %p1) {
-  br i1 undef, label %T, label %F
+define <8 x i64> @extract_8xi64(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) {
+  br i1 %c0, label %T, label %F
 
 T:
   %t = load volatile <16 x i64>, ptr addrspace(1) %p0
   br label %exit
 
 F:
   %f = load volatile <16 x i64>, ptr addrspace(1) %p1
   br label %exit
 
 exit:
   %m = phi <16 x i64> [ %t, %T ], [ %f, %F ]
   %v2 = shufflevector <16 x i64> %m, <16 x i64> undef, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7>
   %b2 = icmp sgt <8 x i64> %v2, <i64 -1, i64 -1, i64 -1, i64 -1, i64 -1, i64 -1, i64 -1, i64 -1>
   %r2 = select <8 x i1> %b2, <8 x i64> <i64 -32768, i64 -32768, i64 -32768, i64 -32768, i64 -32768, i64 -32768, i64 -32768, i64 -32768>, <8 x i64> <i64 -1, i64 -1, i64 -1, i64 -1, i64 -1, i64 -1, i64 -1, i64 -1>
   ret <8 x i64> %r2
 }
 
 ; GCN-LABEL: extract_2xf64
 ; GCN-COUNT-2: v_cndmask_b32
-define <2 x double> @extract_2xf64(ptr addrspace(1) %p0, ptr addrspace(1) %p1) {
-  br i1 undef, label %T, label %F
+define <2 x double> @extract_2xf64(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) {
+  br i1 %c0, label %T, label %F
 
 T:
   %t = load volatile <8 x double>, ptr addrspace(1) %p0
   br label %exit
 
 F:
   %f = load volatile <8 x double>, ptr addrspace(1) %p1
   br label %exit
 
 exit:
   %m = phi <8 x double> [ %t, %T ], [ %f, %F ]
   %v2 = shufflevector <8 x double> %m, <8 x double> undef, <2 x i32> <i32 0, i32 1>
   %b2 = fcmp ogt <2 x double> %v2, <double -1.0, double -1.0>
   %r2 = select <2 x i1> %b2, <2 x double> <double -2.0, double -2.0>, <2 x double> <double -1.0, double -1.0>
   ret <2 x double> %r2
 }
 
 ; GCN-LABEL: extract_4xf64
 ; GCN-COUNT-4: v_cndmask_b32
-define <4 x double> @extract_4xf64(ptr addrspace(1) %p0, ptr addrspace(1) %p1) {
-  br i1 undef, label %T, label %F
+define <4 x double> @extract_4xf64(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) {
+  br i1 %c0, label %T, label %F
 
 T:
   %t = load volatile <8 x double>, ptr addrspace(1) %p0
   br label %exit
 
 F:
   %f = load volatile <8 x double>, ptr addrspace(1) %p1
   br label %exit
 
 exit:
   %m = phi <8 x double> [ %t, %T ], [ %f, %F ]
   %v2 = shufflevector <8 x double> %m, <8 x double> undef, <4 x i32> <i32 0, i32 1, i32 2, i32 3>
   %b2 = fcmp ogt <4 x double> %v2, <double -1.0, double -1.0, double -1.0, double -1.0>
   %r2 = select <4 x i1> %b2, <4 x double> <double -2.0, double -2.0, double -2.0, double -2.0>, <4 x double> <double -1.0, double -1.0, double -1.0, double -1.0>
   ret <4 x double> %r2
 }
 
 ; GCN-LABEL: extract_8xf64
 ; GCN-COUNT-8: v_cndmask_b32
-define <8 x double> @extract_8xf64(ptr addrspace(1) %p0, ptr addrspace(1) %p1) {
-  br i1 undef, label %T, label %F
+define <8 x double> @extract_8xf64(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) {
+  br i1 %c0, label %T, label %F
 
 T:
   %t = load volatile <16 x double>, ptr addrspace(1) %p0
   br label %exit
 
 F:
   %f = load volatile <16 x double>, ptr addrspace(1) %p1
   br label %exit
 
 exit:
   %m = phi <16 x double> [ %t, %T ], [ %f, %F ]
   %v2 = shufflevector <16 x double> %m, <16 x double> undef, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7>
   %b2 = fcmp ogt <8 x double> %v2, <double -1.0, double -1.0, double -1.0, double -1.0, double -1.0, double -1.0, double -1.0, double -1.0>
   %r2 = select <8 x i1> %b2, <8 x double> <double -2.0, double -2.0, double -2.0, double -2.0, double -2.0, double -2.0, double -2.0, double -2.0>, <8 x double> <double -1.0, double -1.0, double -1.0, double -1.0, double -1.0, double -1.0, double -1.0, double -1.0>
   ret <8 x double> %r2
 }

llvm/test/CodeGen/AMDGPU/i1-copy-implicit-def.ll

 ; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=SI %s
 ; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck -check-prefix=SI %s
 
 ; SILowerI1Copies was not handling IMPLICIT_DEF
 ; SI-LABEL: {{^}}br_implicit_def:
 ; SI: %bb.0:
 ; SI-NEXT: s_cbranch_scc1
-define amdgpu_kernel void @br_implicit_def(ptr addrspace(1) %out, i32 %arg) #0 {
+define amdgpu_kernel void @br_implicit_def(ptr addrspace(1) %out, i32 %arg, i1 %c0) #0 {
 bb:
-  br i1 undef, label %bb1, label %bb2
+  br i1 %c0, label %bb1, label %bb2
 
 bb1:
   store volatile i32 123, ptr addrspace(1) %out
   ret void
 
 bb2:
   ret void
 }
 
 attributes #0 = { nounwind }

llvm/test/CodeGen/AMDGPU/i1-copy-phi.ll

 declare i32 @llvm.amdgcn.workitem.id.x() #0
 
 attributes #0 = { nounwind readnone }
 
 ; Make sure this won't crash.
 ; SI-LABEL: {{^}}vcopy_i1_undef
 ; SI: v_cndmask_b32_e64
 ; SI: v_cndmask_b32_e64
-define <2 x float> @vcopy_i1_undef(ptr addrspace(1) %p) {
+define <2 x float> @vcopy_i1_undef(ptr addrspace(1) %p, i1 %c0) {
 entry:
-  br i1 undef, label %exit, label %false
+  br i1 %c0, label %exit, label %false
 
 false:
   %x = load <2 x float>, ptr addrspace(1) %p
   %cmp = fcmp one <2 x float> %x, zeroinitializer
   br label %exit
 
 exit:
   %c = phi <2 x i1> [ undef, %entry ], [ %cmp, %false ]
   %ret = select <2 x i1> %c, <2 x float> <float 2.0, float 2.0>, <2 x float> <float 4.0, float 4.0>
   ret <2 x float> %ret
 }

llvm/test/CodeGen/AMDGPU/implicit-def-muse.ll

 ; RUN: llc -march=amdgcn -stop-after=amdgpu-isel -verify-machineinstrs -o - %s | FileCheck %s
 
 ; CHECK-LABEL: vcopy_i1_undef
 ; CHECK: [[IMPDEF0:%[0-9]+]]:sreg_64 = IMPLICIT_DEF
 ; CHECK-NOT: COPY
 ; CHECK: [[IMPDEF1:%[0-9]+]]:sreg_64 = IMPLICIT_DEF
 ; CHECK-NOT: COPY [[IMPDEF0]]
 ; CHECK-NOT: COPY [[IMPDEF1]]
 ; CHECK: .false:
-define <2 x float> @vcopy_i1_undef(ptr addrspace(1) %p) {
+define <2 x float> @vcopy_i1_undef(ptr addrspace(1) %p, i1 %c0) {
 entry:
-  br i1 undef, label %exit, label %false
+  br i1 %c0, label %exit, label %false
 
 false:
   %x = load <2 x float>, ptr addrspace(1) %p
   %cmp = fcmp one <2 x float> %x, zeroinitializer
   br label %exit
 
 exit:
   %c = phi <2 x i1> [ undef, %entry ], [ %cmp, %false ]
   %ret = select <2 x i1> %c, <2 x float> <float 2.0, float 2.0>, <2 x float> <float 4.0, float 4.0>
   ret <2 x float> %ret
 }

llvm/test/CodeGen/AMDGPU/infinite-loop.ll

 loop:
   store volatile i32 999, ptr addrspace(1) %out, align 4
   br label %loop
 
 return:
   ret void
 }
 
-define amdgpu_kernel void @infinite_loops(ptr addrspace(1) %out) {
+define amdgpu_kernel void @infinite_loops(ptr addrspace(1) %out, i1 %c0) {
 ; SI-LABEL: infinite_loops:
 ; SI:       ; %bb.0: ; %entry
 ; SI-NEXT:    s_load_dwordx2 s[0:1], s[0:1], 0x9
 ; SI-NEXT:    s_mov_b64 s[2:3], -1
 ; SI-NEXT:    s_cbranch_scc1 .LBB2_4
 ; SI-NEXT:  ; %bb.1:
 ; SI-NEXT:    s_mov_b32 s3, 0xf000
 ; SI-NEXT:    s_mov_b32 s2, -1

 ; SI-NEXT:    buffer_store_dword v0, off, s[0:3], 0
 ; SI-NEXT:    s_waitcnt vmcnt(0)
 ; SI-NEXT:    s_mov_b64 vcc, vcc
 ; SI-NEXT:    s_cbranch_vccz .LBB2_6
 ; SI-NEXT:  .LBB2_7: ; %DummyReturnBlock
 ; SI-NEXT:    s_endpgm
 ; IR-LABEL: @infinite_loops(
 ; IR-NEXT:  entry:
-; IR-NEXT:    br i1 undef, label [[LOOP1:%.*]], label [[LOOP2:%.*]]
+; IR-NEXT:    br i1 [[%c0:%.*]], label [[LOOP1:%.*]], label [[LOOP2:%.*]]
 ; IR:       loop1:
 ; IR-NEXT:    store volatile i32 999, ptr addrspace(1) [[OUT:%.*]], align 4
 ; IR-NEXT:    br i1 true, label [[LOOP1]], label [[DUMMYRETURNBLOCK:%.*]]
 ; IR:       loop2:
 ; IR-NEXT:    store volatile i32 888, ptr addrspace(1) [[OUT]], align 4
 ; IR-NEXT:    br i1 true, label [[LOOP2]], label [[DUMMYRETURNBLOCK]]
 ; IR:       DummyReturnBlock:
 ; IR-NEXT:    ret void
 ;
 entry:
-  br i1 undef, label %loop1, label %loop2
+  br i1 %c0, label %loop1, label %loop2
 
 loop1:
   store volatile i32 999, ptr addrspace(1) %out, align 4
   br label %loop1
 
 loop2:
   store volatile i32 888, ptr addrspace(1) %out, align 4
   br label %loop2

llvm/test/CodeGen/AMDGPU/inline-asm.ll

 entry:
   tail call void asm sideeffect "; sgpr96 $0", "s"(<3 x i32> <i32 10, i32 11, i32 12>) #1
   ret void
 }
 
 ; Check aggregate types are handled properly.
 ; CHECK-LABEL: mad_u64
 ; CHECK: v_mad_u64_u32
-define void @mad_u64(i32 %x) {
+define void @mad_u64(i32 %x, i1 %c0) {
 entry:
-  br i1 undef, label %exit, label %false
+  br i1 %c0, label %exit, label %false
 
 false:
   %s0 = tail call { i64, i64 } asm sideeffect "v_mad_u64_u32 $0, $1, $2, $3, $4", "=v,=s,v,v,v"(i32 -766435501, i32 %x, i64 0)
   br label %exit
 
 exit:
   %s1 = phi { i64, i64} [ undef, %entry ], [ %s0, %false]
   %v0 = extractvalue { i64, i64 } %s1, 0
   %v1 = extractvalue { i64, i64 } %s1, 1
   tail call void asm sideeffect "; use $0", "v"(i64 %v0)
   tail call void asm sideeffect "; use $0", "v"(i64 %v1)
   ret void
 }

llvm/test/CodeGen/AMDGPU/ipra-return-address-save-restore.ll

 
 ; Function Attrs: nofree nosync nounwind readnone speculatable willreturn
 declare <4 x float> @llvm.fmuladd.v4f32(<4 x float>, <4 x float>, <4 x float>) #0
 
 ; Function Attrs: argmemonly nofree nosync nounwind willreturn
 declare void @llvm.lifetime.end.p5(i64 immarg, ptr addrspace(5) nocapture) #1
 
 ; Function Attrs: norecurse
-define internal fastcc void @svm_node_closure_bsdf(ptr addrspace(1) %sd, ptr %stack, <4 x i32> %node, ptr %offset, i32 %0, i8 %trunc, float %1, float %2, float %mul80, i1 %cmp412.old, <4 x i32> %3, float %4, i32 %5, i1 %cmp440, i1 %cmp442, i1 %or.cond1306, float %.op, ptr addrspace(1) %arrayidx.i.i2202, ptr addrspace(1) %retval.0.i.i22089, ptr addrspace(1) %retval.1.i221310, i1 %cmp575, ptr addrspace(1) %num_closure_left.i2215, i32 %6, i1 %cmp.i2216, i32 %7, i64 %idx.ext.i2223, i32 %sub5.i2221) #2 {
+define internal fastcc void @svm_node_closure_bsdf(ptr addrspace(1) %sd, ptr %stack, <4 x i32> %node, ptr %offset, i32 %0, i8 %trunc, float %1, float %2, float %mul80, i1 %cmp412.old, <4 x i32> %3, float %4, i32 %5, i1 %cmp440, i1 %cmp442, i1 %or.cond1306, float %.op, ptr addrspace(1) %arrayidx.i.i2202, ptr addrspace(1) %retval.0.i.i22089, ptr addrspace(1) %retval.1.i221310, i1 %cmp575, ptr addrspace(1) %num_closure_left.i2215, i32 %6, i1 %cmp.i2216, i32 %7, i64 %idx.ext.i2223, i32 %sub5.i2221, i1 %c0, i1 %c1, i1 %c2) #2 {
 ; GCN-LABEL: {{^}}svm_node_closure_bsdf:
 ; GCN-DAG: v_writelane_b32 [[CSR_VGPR:v[0-9]+]], s30,
 ; GCN-DAG: v_writelane_b32 [[CSR_VGPR]], s31,
 ; GCN: s_movk_i32 s30, 0x60
 ; GCN-NOT: s31
 ; GCN-DAG: v_readlane_b32 s31, [[CSR_VGPR]],
 ; GCN-DAG: v_readlane_b32 s30, [[CSR_VGPR]],
 ; GCN: s_waitcnt vmcnt(0)
 ; GCN: s_setpc_b64 s[30:31]
 entry:
   %8 = extractelement <4 x i32> %node, i64 0
   %cmp.i.not = icmp eq i32 undef, 0
-  br i1 undef, label %common.ret.critedge, label %cond.true
+  br i1 %c0, label %common.ret.critedge, label %cond.true
 
 cond.true:                                        ; preds = %entry
   %9 = load float, ptr null, align 4
   %phi.cmp = fcmp oeq float %9, 0.000000e+00
   br i1 %phi.cmp, label %common.ret, label %cond.true20
 
 cond.true20:                                      ; preds = %cond.true
   %v = zext i8 %trunc to i32

 if.then534:                                       ; preds = %bsdf_alloc.exit2188
   %.op7 = fmul nsz float undef, 0.000000e+00
   %mul558 = select i1 %cmp440, float 0.000000e+00, float %1
   %15 = tail call float @llvm.amdgcn.fmed3.f32(float 0.000000e+00, float 0.000000e+00, float 0.000000e+00)
   store float %mul558, ptr addrspace(1) null, align 4
   br label %if.end627.sink.split
 
 if.else568:                                       ; preds = %if.then413
-  br i1 undef, label %bsdf_alloc.exit2214, label %if.then.i2198
+  br i1 %c1, label %bsdf_alloc.exit2214, label %if.then.i2198
 
 if.then.i2198:                                    ; preds = %if.else568
-  br i1 undef, label %closure_alloc.exit.i2210, label %if.end.i.i2207
+  br i1 %c2, label %closure_alloc.exit.i2210, label %if.end.i.i2207
 
 if.end.i.i2207:                                   ; preds = %if.then.i2198
   %arrayidx.i.i22028 = getelementptr inbounds %struct.ShaderData, ptr addrspace(1) %sd, i64 0, i32 30, i64 undef
   br label %closure_alloc.exit.i2210
 
 closure_alloc.exit.i2210:                         ; preds = %if.end.i.i2207, %if.then.i2198
   %retval.0.i.i220899 = phi ptr addrspace(1) [ %arrayidx.i.i2202, %if.end.i.i2207 ], [ null, %if.then.i2198 ]
   br i1 false, label %bsdf_alloc.exit2214, label %if.end.i2212

llvm/test/CodeGen/AMDGPU/large-constant-initializer.ll

 ; RUN: llc -march=amdgcn -mcpu=tahiti < %s
 ; RUN: llc -march=amdgcn -mcpu=tonga < %s
 ; CHECK: s_endpgm
 
 @gv = external unnamed_addr addrspace(4) constant [239 x i32], align 4
 
-define amdgpu_kernel void @opencv_cvtfloat_crash(ptr addrspace(1) %out, i32 %x) nounwind {
+define amdgpu_kernel void @opencv_cvtfloat_crash(ptr addrspace(1) %out, i32 %x, i1 %c0) nounwind {
   %val = load i32, ptr addrspace(4) getelementptr ([239 x i32], ptr addrspace(4) @gv, i64 0, i64 239), align 4
   %mul12 = mul nsw i32 %val, 7
-  br i1 undef, label %exit, label %bb
+  br i1 %c0, label %exit, label %bb
 
 bb:
   %cmp = icmp slt i32 %x, 0
   br label %exit
 
 exit:
   ret void
 }
 

llvm/test/CodeGen/AMDGPU/mdt-preserving-crash.ll

 ; RUN: llc < %s
 target datalayout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7"
 target triple = "amdgcn-amd-amdhsa"
 
 @_RSENC_gDcd_______________________________ = external protected addrspace(1) externally_initialized global [4096 x i8], align 16
 
-define protected amdgpu_kernel void @_RSENC_PRInit__________________________________() local_unnamed_addr #0 {
+define protected amdgpu_kernel void @_RSENC_PRInit__________________________________(i1 %c0) local_unnamed_addr #0 {
 entry:
   %runtimeVersionCopy = alloca [128 x i8], align 16, addrspace(5)
   %licenseVersionCopy = alloca [128 x i8], align 16, addrspace(5)
   %pD10 = alloca [128 x i8], align 16, addrspace(5)
   br label %if.end
 
 if.end:                                           ; preds = %entry
   %0 = load i32, ptr undef, align 4
   %mul = mul i32 %0, 3
   %cmp13 = icmp eq i32 %mul, 989619
   br i1 %cmp13, label %cleanup.cont, label %if.end15
 
 if.end15:                                         ; preds = %if.end
-  br i1 undef, label %cleanup.cont, label %lor.lhs.false17
+  br i1 %c0, label %cleanup.cont, label %lor.lhs.false17
 
 lor.lhs.false17:                                  ; preds = %if.end15
   br label %while.cond.i
 
 while.cond.i:                                     ; preds = %while.cond.i, %lor.lhs.false17
   switch i32 undef, label %if.end60 [
     i32 0, label %while.cond.i
     i32 3, label %if.end60.loopexit857

llvm/test/CodeGen/AMDGPU/mfma-cd-select.ll

 }
 
 ; We could avoid scan to find calls since we see these during lowering before selection.
 ; However, in SDag lowering and selection is done block by block, so it would only work
 ; in Global ISel.
 
 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32_call_multi_bb:
 ; GCN: v_mfma_f32_32x32x1{{.*}} a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9:]+}}, a[{{[0-9:]+}}]
-define amdgpu_kernel void @test_mfma_f32_32x32x1f32_call_multi_bb(ptr addrspace(1) %arg) #0 {
+define amdgpu_kernel void @test_mfma_f32_32x32x1f32_call_multi_bb(ptr addrspace(1) %arg, i1 %c0) #0 {
 bb1:
   %in.1 = load <32 x float>, ptr addrspace(1) %arg
   %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %in.1, i32 1, i32 2, i32 3)
   store <32 x float> %mai.1, ptr addrspace(1) %arg
-  br i1 undef, label %bb2, label %bb3
+  br i1 %c0, label %bb2, label %bb3
   br label %bb2
 
 bb2:
   call void @foo()
   br label %bb3
 
 bb3:
   ret void

llvm/test/CodeGen/AMDGPU/noclobber-barrier.ll

 ; GCN: s_load_dword s
 ; GCN: s_waitcnt lgkmcnt(0)
 ; GCN: s_waitcnt lgkmcnt(0)
 ; GCN: s_barrier
 ; GCN-NOT: global_load_dword
 ; GCN: s_load_dword s
 ; GCN-NOT: global_load_dword
 ; GCN: global_store_dword
-define amdgpu_kernel void @memory_phi_no_clobber(ptr addrspace(1) %arg) {
+define amdgpu_kernel void @memory_phi_no_clobber(ptr addrspace(1) %arg, i1 %c0) {
 ; CHECK-LABEL: @memory_phi_no_clobber(
 ; CHECK-NEXT:  bb:
 ; CHECK-NEXT:    [[I:%.*]] = load i32, ptr addrspace(1) [[ARG:%.*]], align 4, !amdgpu.noclobber !0
-; CHECK-NEXT:    br i1 undef, label [[IF_THEN:%.*]], label [[IF_ELSE:%.*]], !amdgpu.uniform !0
+; CHECK-NEXT:    br i1 [[%c0:%.*]], label [[IF_THEN:%.*]], label [[IF_ELSE:%.*]], !amdgpu.uniform !0
 ; CHECK:       if.then:
 ; CHECK-NEXT:    tail call void @llvm.amdgcn.s.barrier()
 ; CHECK-NEXT:    br label [[IF_END:%.*]], !amdgpu.uniform !0
 ; CHECK:       if.else:
 ; CHECK-NEXT:    fence syncscope("workgroup") release
 ; CHECK-NEXT:    br label [[IF_END]], !amdgpu.uniform !0
 ; CHECK:       if.end:
 ; CHECK-NEXT:    [[I1:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[ARG]], i64 1, !amdgpu.uniform !0
 ; CHECK-NEXT:    [[I2:%.*]] = load i32, ptr addrspace(1) [[I1]], align 4, !amdgpu.noclobber !0
 ; CHECK-NEXT:    [[I3:%.*]] = add i32 [[I2]], [[I]]
 ; CHECK-NEXT:    [[I4:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[ARG]], i64 2
 ; CHECK-NEXT:    store i32 [[I3]], ptr addrspace(1) [[I4]], align 4
 ; CHECK-NEXT:    ret void
 ;
 bb:
   %i = load i32, ptr addrspace(1) %arg, align 4
-  br i1 undef, label %if.then, label %if.else
+  br i1 %c0, label %if.then, label %if.else
 
 if.then:
   tail call void @llvm.amdgcn.s.barrier()
   br label %if.end
 
 if.else:
   fence syncscope("workgroup") release
   br label %if.end
 
 if.end:
   %i1 = getelementptr inbounds i32, ptr addrspace(1) %arg, i64 1
   %i2 = load i32, ptr addrspace(1) %i1, align 4
   %i3 = add i32 %i2, %i
   %i4 = getelementptr inbounds i32, ptr addrspace(1) %arg, i64 2
   store i32 %i3, ptr addrspace(1) %i4, align 4
   ret void
 }
 
 ; GCN-LABEL: {{^}}memory_phi_clobber1:
 ; GCN: s_load_dword s
 ; GCN: s_barrier
 ; GCN: global_store_dword
 ; GCN: global_load_dword
 ; GCN: global_store_dword
-define amdgpu_kernel void @memory_phi_clobber1(ptr addrspace(1) %arg) {
+define amdgpu_kernel void @memory_phi_clobber1(ptr addrspace(1) %arg, i1 %c0) {
 ; CHECK-LABEL: @memory_phi_clobber1(
 ; CHECK-NEXT:  bb:
 ; CHECK-NEXT:    [[I:%.*]] = load i32, ptr addrspace(1) [[ARG:%.*]], align 4, !amdgpu.noclobber !0
-; CHECK-NEXT:    br i1 undef, label [[IF_THEN:%.*]], label [[IF_ELSE:%.*]], !amdgpu.uniform !0
+; CHECK-NEXT:    br i1 [[%c1:%.*]], label [[IF_THEN:%.*]], label [[IF_ELSE:%.*]], !amdgpu.uniform !0
 ; CHECK:       if.then:
 ; CHECK-NEXT:    [[GEP:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[ARG]], i64 3
 ; CHECK-NEXT:    store i32 1, ptr addrspace(1) [[GEP]], align 4
 ; CHECK-NEXT:    br label [[IF_END:%.*]], !amdgpu.uniform !0
 ; CHECK:       if.else:
 ; CHECK-NEXT:    tail call void @llvm.amdgcn.s.barrier()
 ; CHECK-NEXT:    br label [[IF_END]], !amdgpu.uniform !0
 ; CHECK:       if.end:
 ; CHECK-NEXT:    [[I1:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[ARG]], i64 1, !amdgpu.uniform !0
 ; CHECK-NEXT:    [[I2:%.*]] = load i32, ptr addrspace(1) [[I1]], align 4
 ; CHECK-NEXT:    [[I3:%.*]] = add i32 [[I2]], [[I]]
 ; CHECK-NEXT:    [[I4:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[ARG]], i64 2
 ; CHECK-NEXT:    store i32 [[I3]], ptr addrspace(1) [[I4]], align 4
 ; CHECK-NEXT:    ret void
 ;
 bb:
   %i = load i32, ptr addrspace(1) %arg, align 4
-  br i1 undef, label %if.then, label %if.else
+  br i1 %c0, label %if.then, label %if.else
 
 if.then:
   %gep = getelementptr inbounds i32, ptr addrspace(1) %arg, i64 3
   store i32 1, ptr addrspace(1) %gep, align 4
   br label %if.end
 
 if.else:
   tail call void @llvm.amdgcn.s.barrier()

 }
 
 ; GCN-LABEL: {{^}}memory_phi_clobber2:
 ; GCN-DAG: s_load_dword s
 ; GCN-DAG: global_store_dword
 ; GCN: s_barrier
 ; GCN: global_load_dword
 ; GCN: global_store_dword
-define amdgpu_kernel void @memory_phi_clobber2(ptr addrspace(1) %arg) {
+define amdgpu_kernel void @memory_phi_clobber2(ptr addrspace(1) %arg, i1 %c0) {
 ; CHECK-LABEL: @memory_phi_clobber2(
 ; CHECK-NEXT:  bb:
 ; CHECK-NEXT:    [[I:%.*]] = load i32, ptr addrspace(1) [[ARG:%.*]], align 4, !amdgpu.noclobber !0
-; CHECK-NEXT:    br i1 undef, label [[IF_THEN:%.*]], label [[IF_ELSE:%.*]], !amdgpu.uniform !0
+; CHECK-NEXT:    br i1 [[%c2:%.*]], label [[IF_THEN:%.*]], label [[IF_ELSE:%.*]], !amdgpu.uniform !0
 ; CHECK:       if.then:
 ; CHECK-NEXT:    tail call void @llvm.amdgcn.s.barrier()
 ; CHECK-NEXT:    br label [[IF_END:%.*]], !amdgpu.uniform !0
 ; CHECK:       if.else:
 ; CHECK-NEXT:    [[GEP:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[ARG]], i64 3
 ; CHECK-NEXT:    store i32 1, ptr addrspace(1) [[GEP]], align 4
 ; CHECK-NEXT:    br label [[IF_END]], !amdgpu.uniform !0
 ; CHECK:       if.end:
 ; CHECK-NEXT:    [[I1:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[ARG]], i64 1, !amdgpu.uniform !0
 ; CHECK-NEXT:    [[I2:%.*]] = load i32, ptr addrspace(1) [[I1]], align 4
 ; CHECK-NEXT:    [[I3:%.*]] = add i32 [[I2]], [[I]]
 ; CHECK-NEXT:    [[I4:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[ARG]], i64 2
 ; CHECK-NEXT:    store i32 [[I3]], ptr addrspace(1) [[I4]], align 4
 ; CHECK-NEXT:    ret void
 ;
 bb:
   %i = load i32, ptr addrspace(1) %arg, align 4
-  br i1 undef, label %if.then, label %if.else
+  br i1 %c0, label %if.then, label %if.else
 
 if.then:
   tail call void @llvm.amdgcn.s.barrier()
   br label %if.end
 
 if.else:
   %gep = getelementptr inbounds i32, ptr addrspace(1) %arg, i64 3
   store i32 1, ptr addrspace(1) %gep, align 4

llvm/test/CodeGen/AMDGPU/promote-alloca-to-lds-phi.ll

 ; RUN: opt -S -mtriple=amdgcn-unknown-amdhsa -mcpu=kaveri -passes=amdgpu-promote-alloca < %s | FileCheck %s
 
 
 ; CHECK-LABEL: @branch_ptr_var_same_alloca(
 ; CHECK: [[GEP:%[0-9]+]] = getelementptr inbounds [256 x [64 x i32]], ptr addrspace(3) @branch_ptr_var_same_alloca.alloca, i32 0, i32 %{{[0-9]+}}
 
 ; CHECK: if:
 ; CHECK: %arrayidx0 = getelementptr inbounds [64 x i32], ptr addrspace(3) [[GEP]], i32 0, i32 %a
 
 ; CHECK: else:
 ; CHECK: %arrayidx1 = getelementptr inbounds [64 x i32], ptr addrspace(3) [[GEP]], i32 0, i32 %b
 
 ; CHECK: endif:
 ; CHECK: %phi.ptr = phi ptr addrspace(3) [ %arrayidx0, %if ], [ %arrayidx1, %else ]
 ; CHECK: store i32 0, ptr addrspace(3) %phi.ptr, align 4
-define amdgpu_kernel void @branch_ptr_var_same_alloca(i32 %a, i32 %b) #0 {
+define amdgpu_kernel void @branch_ptr_var_same_alloca(i32 %a, i32 %b, i1 %c0) #0 {
 entry:
   %alloca = alloca [64 x i32], align 4, addrspace(5)
-  br i1 undef, label %if, label %else
+  br i1 %c0, label %if, label %else
 
 if:
   %arrayidx0 = getelementptr inbounds [64 x i32], ptr addrspace(5) %alloca, i32 0, i32 %a
   br label %endif
 
 else:
   %arrayidx1 = getelementptr inbounds [64 x i32], ptr addrspace(5) %alloca, i32 0, i32 %b
   br label %endif
 
 endif:
   %phi.ptr = phi ptr addrspace(5) [ %arrayidx0, %if ], [ %arrayidx1, %else ]
   store i32 0, ptr addrspace(5) %phi.ptr, align 4
   ret void
 }
 
 ; CHECK-LABEL: @branch_ptr_phi_alloca_null_0(
 ; CHECK: %phi.ptr = phi ptr addrspace(3) [ %arrayidx0, %if ], [ null, %entry ]
-define amdgpu_kernel void @branch_ptr_phi_alloca_null_0(i32 %a, i32 %b) #0 {
+define amdgpu_kernel void @branch_ptr_phi_alloca_null_0(i32 %a, i32 %b, i1 %c0) #0 {
 entry:
   %alloca = alloca [64 x i32], align 4, addrspace(5)
-  br i1 undef, label %if, label %endif
+  br i1 %c0, label %if, label %endif
 
 if:
   %arrayidx0 = getelementptr inbounds [64 x i32], ptr addrspace(5) %alloca, i32 0, i32 %a
   br label %endif
 
 endif:
   %phi.ptr = phi ptr addrspace(5) [ %arrayidx0, %if ], [ null, %entry ]
   store i32 0, ptr addrspace(5) %phi.ptr, align 4
   ret void
 }
 
 ; CHECK-LABEL: @branch_ptr_phi_alloca_null_1(
 ; CHECK: %phi.ptr = phi ptr addrspace(3)  [ null, %entry ], [ %arrayidx0, %if ]
-define amdgpu_kernel void @branch_ptr_phi_alloca_null_1(i32 %a, i32 %b) #0 {
+define amdgpu_kernel void @branch_ptr_phi_alloca_null_1(i32 %a, i32 %b, i1 %c0) #0 {
 entry:
   %alloca = alloca [64 x i32], align 4, addrspace(5)
-  br i1 undef, label %if, label %endif
+  br i1 %c0, label %if, label %endif
 
 if:
   %arrayidx0 = getelementptr inbounds [64 x i32], ptr addrspace(5) %alloca, i32 0, i32 %a
   br label %endif
 
 endif:
   %phi.ptr = phi ptr addrspace(5) [ null, %entry ], [ %arrayidx0, %if ]
   store i32 0, ptr addrspace(5) %phi.ptr, align 4

 ; CHECK: %arrayidx0 = getelementptr inbounds [64 x i32], ptr addrspace(5) %alloca, i32 0, i32 %a
 
 ; CHECK: else:
 ; CHECK: %arrayidx1 = call ptr addrspace(5) @get_unknown_pointer()
 
 ; CHECK: endif:
 ; CHECK: %phi.ptr = phi ptr addrspace(5) [ %arrayidx0, %if ], [ %arrayidx1, %else ]
 ; CHECK: store i32 0, ptr addrspace(5) %phi.ptr, align 4
-define amdgpu_kernel void @branch_ptr_alloca_unknown_obj(i32 %a, i32 %b) #0 {
+define amdgpu_kernel void @branch_ptr_alloca_unknown_obj(i32 %a, i32 %b, i1 %c0) #0 {
 entry:
   %alloca = alloca [64 x i32], align 4, addrspace(5)
-  br i1 undef, label %if, label %else
+  br i1 %c0, label %if, label %else
 
 if:
   %arrayidx0 = getelementptr inbounds [64 x i32], ptr addrspace(5) %alloca, i32 0, i32 %a
   br label %endif
 
 else:
   %arrayidx1 = call ptr addrspace(5) @get_unknown_pointer()
   br label %endif

llvm/test/CodeGen/AMDGPU/promote-alloca-to-lds-select.ll

define amdgpu_kernel void @lds_promoted_alloca_select_input_select(i32 %a, i32 %b, i32 %c, i1 %c1, i1 %c2) #0 {
   %ptr1 = getelementptr inbounds [16 x i32], ptr addrspace(5) %alloca, i32 0, i32 %b
   %ptr2 = getelementptr inbounds [16 x i32], ptr addrspace(5) %alloca, i32 0, i32 %c
   %select0 = select i1 %c1, ptr addrspace(5) %ptr0, ptr addrspace(5) %ptr1
   %select1 = select i1 %c2, ptr addrspace(5) %select0, ptr addrspace(5) %ptr2
   store i32 0, ptr addrspace(5) %select1, align 4
   ret void
 }
 
-define amdgpu_kernel void @lds_promoted_alloca_select_input_phi(i32 %a, i32 %b, i32 %c) #0 {
+define amdgpu_kernel void @lds_promoted_alloca_select_input_phi(i32 %a, i32 %b, i32 %c, i1 %c0) #0 {
 entry:
   %alloca = alloca [16 x i32], align 4, addrspace(5)
   %ptr0 = getelementptr inbounds [16 x i32], ptr addrspace(5) %alloca, i32 0, i32 %a
   %ptr1 = getelementptr inbounds [16 x i32], ptr addrspace(5) %alloca, i32 0, i32 %b
   store i32 0, ptr addrspace(5) %ptr0
-  br i1 undef, label %bb1, label %bb2
+  br i1 %c0, label %bb1, label %bb2
 
 bb1:
   %ptr2 = getelementptr inbounds [16 x i32], ptr addrspace(5) %alloca, i32 0, i32 %c
   %select0 = select i1 undef, ptr addrspace(5) undef, ptr addrspace(5) %ptr2
   store i32 0, ptr addrspace(5) %ptr1
   br label %bb2
 
 bb2:

llvm/test/CodeGen/AMDGPU/reg-coalescer-sched-crash.ll

 ; RUN: llc -march=amdgcn -verify-machineinstrs -o /dev/null < %s
 ; RUN: llc -march=amdgcn -mcpu=tonga -mattr=-flat-for-global -verify-machineinstrs -o /dev/null < %s
 
 ; The register coalescer introduces a verifier error which later
 ; results in a crash during scheduling.
 
 declare i32 @llvm.amdgcn.workitem.id.x() #0
 
-define amdgpu_kernel void @reg_coalescer_breaks_dead(ptr addrspace(1) nocapture readonly %arg, i32 %arg1, i32 %arg2, i32 %arg3) #1 {
+define amdgpu_kernel void @reg_coalescer_breaks_dead(ptr addrspace(1) nocapture readonly %arg, i32 %arg1, i32 %arg2, i32 %arg3, i1 %c0) #1 {
 bb:
   %id.x = call i32 @llvm.amdgcn.workitem.id.x()
   %cmp0 = icmp eq i32 %id.x, 0
   br i1 %cmp0, label %bb3, label %bb4
 
 bb3:                                              ; preds = %bb
   %tmp = ashr exact i32 undef, 8
   br label %bb6
 
 bb4:                                              ; preds = %bb6, %bb
   %tmp5 = phi <2 x i32> [ zeroinitializer, %bb ], [ %tmp13, %bb6 ]
-  br i1 undef, label %bb15, label %bb16
+  br i1 %c0, label %bb15, label %bb16
 
 bb6:                                              ; preds = %bb6, %bb3
   %tmp7 = phi <2 x i32> [ zeroinitializer, %bb3 ], [ %tmp13, %bb6 ]
   %tmp8 = add nsw i32 0, %arg1
   %tmp9 = add nsw i32 %tmp8, 0
   %tmp10 = sext i32 %tmp9 to i64
   %tmp11 = getelementptr inbounds <2 x i32>, ptr addrspace(1) %arg, i64 %tmp10
   %tmp12 = load <2 x i32>, ptr addrspace(1) %tmp11, align 8

llvm/test/CodeGen/AMDGPU/rename-disconnected-bug.ll

 ; RUN: llc -verify-machineinstrs -o /dev/null %s
 ; Check that renameDisconnectedComponents() does not create vregs without a
 ; definition on every path (there should at least be IMPLICIT_DEF instructions).
 target triple = "amdgcn--"
 
-define amdgpu_kernel void @func() {
+define amdgpu_kernel void @func(i1 %c0, i1 %c1, i1 %c2) {
 B0:
-  br i1 undef, label %B1, label %B2
+  br i1 %c0, label %B1, label %B2
 
 B1:
   br label %B2
 
 B2:
   %v0 = phi <4 x float> [ zeroinitializer, %B1 ], [ <float 0.0, float 0.0, float 0.0, float undef>, %B0 ]
-  br i1 undef, label %B20.1, label %B20.2
+  br i1 %c1, label %B20.1, label %B20.2
 
 B20.1:
   br label %B20.2
 
 B20.2:
   %v2 = phi <4 x float> [ zeroinitializer, %B20.1 ], [ %v0, %B2 ]
-  br i1 undef, label %B30.1, label %B30.2
+  br i1 %c2, label %B30.1, label %B30.2
 
 B30.1:
   %sub = fsub <4 x float> %v2, undef
   br label %B30.2
 
 B30.2:
   %v3 = phi <4 x float> [ %sub, %B30.1 ], [ %v2, %B20.2 ]
   %ve0 = extractelement <4 x float> %v3, i32 0
   store float %ve0, ptr addrspace(3) undef, align 4
   ret void
 }

llvm/test/CodeGen/AMDGPU/si-annotate-cf-noloop.ll

 ; RUN: opt -mtriple=amdgcn-- -S -structurizecfg -si-annotate-control-flow -simplifycfg-require-and-preserve-domtree=1 %s | FileCheck -check-prefix=OPT %s
 ; RUN: llc -march=amdgcn -verify-machineinstrs -simplifycfg-require-and-preserve-domtree=1 < %s | FileCheck -check-prefix=GCN %s
 
 
 ; OPT-LABEL: @annotate_unreachable_noloop(
 ; OPT-NOT: call i1 @llvm.amdgcn.loop
 
 ; GCN-LABEL: {{^}}annotate_unreachable_noloop:
 ; GCN: s_cbranch_scc1
 ; GCN-NOT: s_endpgm
 ; GCN: .Lfunc_end0
-define amdgpu_kernel void @annotate_unreachable_noloop(ptr addrspace(1) noalias nocapture readonly %arg) #0 {
+define amdgpu_kernel void @annotate_unreachable_noloop(ptr addrspace(1) noalias nocapture readonly %arg, i1 %c0) #0 {
 bb:
   %tmp = tail call i32 @llvm.amdgcn.workitem.id.x()
   br label %bb1
 
 bb1:                                              ; preds = %bb
   %tmp2 = sext i32 %tmp to i64
   %tmp3 = getelementptr inbounds <4 x float>, ptr addrspace(1) %arg, i64 %tmp2
   %tmp4 = load <4 x float>, ptr addrspace(1) %tmp3, align 16
-  br i1 undef, label %bb5, label %bb3
+  br i1 %c0, label %bb5, label %bb3
 
 bb3:                                              ; preds = %bb1
   %tmp6 = extractelement <4 x float> %tmp4, i32 2
   %tmp7 = fcmp olt float %tmp6, 0.000000e+00
   br i1 %tmp7, label %bb4, label %bb5 ; crash goes away if these are swapped
 
 bb4:                                              ; preds = %bb3
   unreachable

 
 ; OPT-LABEL: @uniform_annotate_ret_noloop(
 ; OPT-NOT: call i1 @llvm.amdgcn.loop
 
 ; GCN-LABEL: {{^}}uniform_annotate_ret_noloop:
 ; GCN: s_cbranch_scc1
 ; GCN: s_endpgm
 ; GCN: .Lfunc_end
-define amdgpu_kernel void @uniform_annotate_ret_noloop(ptr addrspace(1) noalias nocapture readonly %arg, i32 %tmp) #0 {
+define amdgpu_kernel void @uniform_annotate_ret_noloop(ptr addrspace(1) noalias nocapture readonly %arg, i32 %tmp, i1 %c0) #0 {
 bb:
   br label %bb1
 
 bb1:                                              ; preds = %bb
   %tmp2 = sext i32 %tmp to i64
   %tmp3 = getelementptr inbounds <4 x float>, ptr addrspace(1) %arg, i64 %tmp2
   %tmp4 = load <4 x float>, ptr addrspace(1) %tmp3, align 16
-  br i1 undef, label %bb5, label %bb3
+  br i1 %c0, label %bb5, label %bb3
 
 bb3:                                              ; preds = %bb1
   %tmp6 = extractelement <4 x float> %tmp4, i32 2
   %tmp7 = fcmp olt float %tmp6, 0.000000e+00
   br i1 %tmp7, label %bb4, label %bb5 ; crash goes away if these are swapped
 
 bb4:                                              ; preds = %bb3
   ret void

llvm/test/CodeGen/AMDGPU/si-annotate-cf-unreachable.ll

 ; RUN: opt -mtriple=amdgcn-- -S -structurizecfg -si-annotate-control-flow %s | FileCheck -check-prefix=OPT %s
 ; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
 
 
 ; OPT-LABEL: @annotate_unreachable(
 ; OPT: call { i1, i64 } @llvm.amdgcn.if.i64(
 ; OPT-NOT: call void @llvm.amdgcn.end.cf
 
 
 ; GCN-LABEL: {{^}}annotate_unreachable:
 ; GCN: s_and_saveexec_b64
 ; GCN-NOT: s_endpgm
 ; GCN: .Lfunc_end0
-define amdgpu_kernel void @annotate_unreachable(ptr addrspace(1) noalias nocapture readonly %arg) #0 {
+define amdgpu_kernel void @annotate_unreachable(ptr addrspace(1) noalias nocapture readonly %arg, i1 %c0) #0 {
 bb:
   %tmp = tail call i32 @llvm.amdgcn.workitem.id.x()
   br label %bb1
 
 bb1:                                              ; preds = %bb
   %tmp2 = sext i32 %tmp to i64
   %tmp3 = getelementptr inbounds <4 x float>, ptr addrspace(1) %arg, i64 %tmp2
   %tmp4 = load <4 x float>, ptr addrspace(1) %tmp3, align 16
-  br i1 undef, label %bb3, label %bb5  ; label order reversed
+  br i1 %c0, label %bb3, label %bb5  ; label order reversed
 
 bb3:                                              ; preds = %bb1
   %tmp6 = extractelement <4 x float> %tmp4, i32 2
   %tmp7 = fcmp olt float %tmp6, 0.000000e+00
   br i1 %tmp7, label %bb4, label %bb5
 
 bb4:                                              ; preds = %bb3
   unreachable

llvm/test/CodeGen/AMDGPU/si-spill-cf.ll

 ; RUN: llc -march=amdgcn < %s -verify-machineinstrs | FileCheck -check-prefix=SI %s
 ; RUN: llc -march=amdgcn -mcpu=tonga < %s -verify-machineinstrs | FileCheck -check-prefix=SI %s
 
 ; If this occurs it is likely due to reordering and the restore was
 ; originally supposed to happen before SI_END_CF.
 
 ; SI: s_or_b64 exec, exec, [[SAVED:s\[[0-9]+:[0-9]+\]|[a-z]+]]
 ; SI-NOT: v_readlane_b32 [[SAVED]]
 
-define amdgpu_ps void @main() #0 {
+define amdgpu_ps void @main(i1 %c0, i1 %c1, i1 %c2, i1 %c3, i1 %c4, i1 %c5, i1 %c6, i1 %c7, i1 %c8, i1 %c9, i1 %c10, i1 %c11) #0 {
 main_body:
   %tmp = call float @llvm.amdgcn.s.buffer.load.f32(<4 x i32> undef, i32 16, i32 0)
   %tmp1 = call float @llvm.amdgcn.s.buffer.load.f32(<4 x i32> undef, i32 32, i32 0)
   %tmp2 = call float @llvm.amdgcn.s.buffer.load.f32(<4 x i32> undef, i32 80, i32 0)
   %tmp3 = call float @llvm.amdgcn.s.buffer.load.f32(<4 x i32> undef, i32 84, i32 0)
   %tmp4 = call float @llvm.amdgcn.s.buffer.load.f32(<4 x i32> undef, i32 88, i32 0)
   %tmp5 = call float @llvm.amdgcn.s.buffer.load.f32(<4 x i32> undef, i32 96, i32 0)
   %tmp6 = call float @llvm.amdgcn.s.buffer.load.f32(<4 x i32> undef, i32 100, i32 0)
ELSE2596:                                         ; preds = %ELSE2593
   br i1 %tmp224, label %ENDIF2594, label %ELSE2632
 
 ENDIF2594:                                        ; preds = %ELSE2788, %ELSE2785, %ELSE2782, %ELSE2779, %IF2775, %ELSE2761, %ELSE2758, %IF2757, %ELSE2704, %ELSE2686, %ELSE2671, %ELSE2668, %IF2667, %ELSE2632, %ELSE2596, %ELSE2593
   %temp894.2 = phi float [ 0.000000e+00, %IF2667 ], [ 0.000000e+00, %ELSE2671 ], [ 0.000000e+00, %IF2757 ], [ 0.000000e+00, %ELSE2761 ], [ %temp894.0, %ELSE2758 ], [ 0.000000e+00, %IF2775 ], [ 0.000000e+00, %ELSE2779 ], [ 0.000000e+00, %ELSE2782 ], [ %.2848, %ELSE2788 ], [ 0.000000e+00, %ELSE2785 ], [ 0.000000e+00, %ELSE2593 ], [ 0.000000e+00, %ELSE2632 ], [ 0.000000e+00, %ELSE2704 ], [ 0.000000e+00, %ELSE2686 ], [ 0.000000e+00, %ELSE2668 ], [ 0.000000e+00, %ELSE2596 ]
   %tmp225 = fmul float %temp894.2, undef
   br label %ENDIF2564
 
 ELSE2632:                                         ; preds = %ELSE2596
-  br i1 undef, label %ENDIF2594, label %ELSE2650
+  br i1 %c0, label %ENDIF2594, label %ELSE2650
 
 ELSE2650:                                         ; preds = %ELSE2632
   %tmp226 = fcmp oeq float %temp292.11, %tmp110
   %tmp227 = fcmp olt float %tmp110, %tmp111
   %tmp228 = and i1 %tmp226, %tmp227
   br i1 %tmp228, label %IF2667, label %ELSE2668
 
 IF2667:                                           ; preds = %ELSE2650
-  br i1 undef, label %ENDIF2594, label %ELSE2671
+  br i1 %c1, label %ENDIF2594, label %ELSE2671
 
 ELSE2668:                                         ; preds = %ELSE2650
   %tmp229 = fcmp oeq float %temp292.11, %tmp128
   %tmp230 = fcmp olt float %tmp128, undef
   %tmp231 = and i1 %tmp229, %tmp230
   br i1 %tmp231, label %ENDIF2594, label %ELSE2686
 
 ELSE2671:                                         ; preds = %IF2667
   br label %ENDIF2594
 
 ELSE2686:                                         ; preds = %ELSE2668
   %tmp232 = fcmp oeq float %temp292.11, %tmp145
   %tmp233 = fcmp olt float %tmp145, undef
   %tmp234 = and i1 %tmp232, %tmp233
   br i1 %tmp234, label %ENDIF2594, label %ELSE2704
 
 ELSE2704:                                         ; preds = %ELSE2686
   %tmp235 = fcmp oeq float %temp292.11, %tmp180
   %tmp236 = fcmp olt float %tmp180, undef
   %tmp237 = and i1 %tmp235, %tmp236
   br i1 %tmp237, label %ENDIF2594, label %ELSE2740
 
 ELSE2740:                                         ; preds = %ELSE2704
-  br i1 undef, label %IF2757, label %ELSE2758
+  br i1 %c2, label %IF2757, label %ELSE2758
 
 IF2757:                                           ; preds = %ELSE2740
-  br i1 undef, label %ENDIF2594, label %ELSE2761
+  br i1 %c3, label %ENDIF2594, label %ELSE2761
 
 ELSE2758:                                         ; preds = %ELSE2740
-  br i1 undef, label %IF2775, label %ENDIF2594
+  br i1 %c4, label %IF2775, label %ENDIF2594
 
 ELSE2761:                                         ; preds = %IF2757
   br label %ENDIF2594
 
 IF2775:                                           ; preds = %ELSE2758
   %tmp238 = fcmp olt float undef, undef
   br i1 %tmp238, label %ENDIF2594, label %ELSE2779
 
 ELSE2779:                                         ; preds = %IF2775
-  br i1 undef, label %ENDIF2594, label %ELSE2782
+  br i1 %c5, label %ENDIF2594, label %ELSE2782
 
 ELSE2782:                                         ; preds = %ELSE2779
-  br i1 undef, label %ENDIF2594, label %ELSE2785
+  br i1 %c6, label %ENDIF2594, label %ELSE2785
 
 ELSE2785:                                         ; preds = %ELSE2782
   %tmp239 = fcmp olt float undef, 0.000000e+00
   br i1 %tmp239, label %ENDIF2594, label %ELSE2788
 
 ELSE2788:                                         ; preds = %ELSE2785
   %tmp240 = fcmp olt float 0.000000e+00, undef
   %.2848 = select i1 %tmp240, float -1.000000e+00, float 1.000000e+00
ELSE2797:                                         ; preds = %ENDIF2564
   %tmp262 = fcmp olt float %tmp257, %tmp258
   %tmp263 = and i1 %tmp261, %tmp262
   br i1 %tmp263, label %ENDIF2795, label %ELSE2800
 
 ENDIF2795:                                        ; preds = %ELSE2824, %ELSE2821, %ELSE2818, %ELSE2815, %ELSE2812, %ELSE2809, %ELSE2806, %ELSE2803, %ELSE2800, %ELSE2797, %ENDIF2564
   br label %LOOP
 
 ELSE2800:                                         ; preds = %ELSE2797
-  br i1 undef, label %ENDIF2795, label %ELSE2803
+  br i1 %c7, label %ENDIF2795, label %ELSE2803
 
 ELSE2803:                                         ; preds = %ELSE2800
   %tmp264 = fsub float %tmp20, undef
   %tmp265 = fsub float %tmp21, undef
   %tmp266 = fsub float %tmp22, undef
   %tmp267 = fmul float %tmp264, undef
   %tmp268 = fmul float %tmp265, undef
   %tmp269 = fmul float %tmp266, 0.000000e+00
ELSE2806:                                         ; preds = %ELSE2803
   %tmp302 = fcmp ogt float %tmp300, 0.000000e+00
   %tmp303 = fcmp olt float %tmp300, 1.000000e+00
   %tmp304 = and i1 %tmp302, %tmp303
   %tmp305 = fcmp olt float %tmp300, %tmp301
   %tmp306 = and i1 %tmp304, %tmp305
   br i1 %tmp306, label %ENDIF2795, label %ELSE2809
 
 ELSE2809:                                         ; preds = %ELSE2806
-  br i1 undef, label %ENDIF2795, label %ELSE2812
+  br i1 %c8, label %ENDIF2795, label %ELSE2812
 
 ELSE2812:                                         ; preds = %ELSE2809
-  br i1 undef, label %ENDIF2795, label %ELSE2815
+  br i1 %c9, label %ENDIF2795, label %ELSE2815
 
 ELSE2815:                                         ; preds = %ELSE2812
-  br i1 undef, label %ENDIF2795, label %ELSE2818
+  br i1 %c10, label %ENDIF2795, label %ELSE2818
 
 ELSE2818:                                         ; preds = %ELSE2815
-  br i1 undef, label %ENDIF2795, label %ELSE2821
+  br i1 %c11, label %ENDIF2795, label %ELSE2821
 
 ELSE2821:                                         ; preds = %ELSE2818
   %tmp307 = fsub float %tmp56, undef
   %tmp308 = fsub float %tmp57, undef
   %tmp309 = fsub float %tmp58, undef
   %tmp310 = fmul float %tmp307, undef
   %tmp311 = fmul float %tmp308, 0.000000e+00
   %tmp312 = fmul float %tmp309, undef

llvm/test/CodeGen/AMDGPU/simplifydemandedbits-recursion.ll

 
 ; Function Attrs: nounwind readnone speculatable
 declare i32 @llvm.amdgcn.workitem.id.x() #0
 
 ; Function Attrs: nounwind readnone speculatable
 declare float @llvm.fmuladd.f32(float, float, float) #0
 
 ; CHECK: s_endpgm
-define amdgpu_kernel void @foo(ptr addrspace(1) noalias nocapture readonly %arg, ptr addrspace(1) noalias nocapture readonly %arg1, ptr addrspace(1) noalias nocapture %arg2, float %arg3) local_unnamed_addr !reqd_work_group_size !0 {
+define amdgpu_kernel void @foo(ptr addrspace(1) noalias nocapture readonly %arg, ptr addrspace(1) noalias nocapture readonly %arg1, ptr addrspace(1) noalias nocapture %arg2, float %arg3, i1 %c0, i1 %c1, i1 %c2, i1 %c3, i1 %c4, i1 %c5) local_unnamed_addr !reqd_work_group_size !0 {
 bb:
   %tmp = tail call i32 @llvm.amdgcn.workitem.id.y()
   %tmp4 = tail call i32 @llvm.amdgcn.workitem.id.x()
   %tmp5 = and i32 %tmp, 15
   %tmp6 = mul nuw nsw i32 %tmp5, 21
   %tmp7 = sub i32 %tmp6, 0
   %tmp8 = add i32 %tmp7, 0
   %tmp9 = add i32 %tmp8, 0
   br label %bb12
 
 bb11:                                             ; preds = %bb30
-  br i1 undef, label %bb37, label %bb38
+  br i1 %c0, label %bb37, label %bb38
 
 bb12:                                             ; preds = %bb30, %bb
   br i1 false, label %.preheader, label %.loopexit145
 
 .loopexit145:                                     ; preds = %.preheader, %bb12
   br label %bb13
 
 bb13:                                             ; preds = %.loopexit, %.loopexit145
   %tmp14 = phi i32 [ %tmp5, %.loopexit145 ], [ %tmp20, %.loopexit ]
   %tmp15 = add nsw i32 %tmp14, -3
   %tmp16 = mul i32 %tmp14, 21
-  br i1 undef, label %bb17, label %.loopexit
+  br i1 %c1, label %bb17, label %.loopexit
 
 bb17:                                             ; preds = %bb13
   %tmp18 = mul i32 %tmp15, 224
   %tmp19 = add i32 undef, %tmp18
   br label %bb21
 
 .loopexit:                                        ; preds = %bb21, %bb13
   %tmp20 = add nuw nsw i32 %tmp14, 16
-  br i1 undef, label %bb13, label %bb26
+  br i1 %c2, label %bb13, label %bb26
 
 bb21:                                             ; preds = %bb21, %bb17
   %tmp22 = phi i32 [ %tmp4, %bb17 ], [ %tmp25, %bb21 ]
   %tmp23 = add i32 %tmp22, %tmp16
   %tmp24 = getelementptr inbounds float, ptr addrspace(3) @0, i32 %tmp23
   store float undef, ptr addrspace(3) %tmp24, align 4
   %tmp25 = add nuw i32 %tmp22, 8
-  br i1 undef, label %bb21, label %.loopexit
+  br i1 %c3, label %bb21, label %.loopexit
 
 bb26:                                             ; preds = %.loopexit
   br label %bb31
 
 .preheader:                                       ; preds = %.preheader, %bb12
   %tmp27 = phi i32 [ %tmp28, %.preheader ], [ undef, %bb12 ]
   %tmp28 = add nuw i32 %tmp27, 128
   %tmp29 = icmp ult i32 %tmp28, 1568
   br i1 %tmp29, label %.preheader, label %.loopexit145
 
 bb30:                                             ; preds = %bb31
-  br i1 undef, label %bb11, label %bb12
+  br i1 %c4, label %bb11, label %bb12
 
 bb31:                                             ; preds = %bb31, %bb26
   %tmp32 = phi i32 [ %tmp9, %bb26 ], [ undef, %bb31 ]
   %tmp33 = getelementptr inbounds [462 x float], ptr addrspace(3) @0, i32 0, i32 %tmp32
   %tmp34 = load float, ptr addrspace(3) %tmp33, align 4
   %tmp35 = tail call float @llvm.fmuladd.f32(float %tmp34, float undef, float undef)
   %tmp36 = tail call float @llvm.fmuladd.f32(float undef, float undef, float %tmp35)
-  br i1 undef, label %bb30, label %bb31
+  br i1 %c5, label %bb30, label %bb31
 
 bb37:                                             ; preds = %bb11
   br label %bb38
 
 bb38:                                             ; preds = %bb37, %bb11
   ret void
 }
 
 attributes #0 = { nounwind readnone speculatable }
 
 !0 = !{i32 8, i32 16, i32 1}

llvm/test/CodeGen/AMDGPU/sink-image-sample.ll

 ; RUN: llc -mtriple=amdgcn-mesa-mesa3d -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN %s
 ; RUN: llc -mtriple=amdgcn-mesa-mesa3d -mcpu=gfx1010 -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN %s
 
 ; Test that image.sample instruction is sunk across the branch and not left in the first block. Since the kill may terminate the shader there might be no need for sampling the image.
 
 ; GCN-LABEL: {{^}}sinking_img_sample:
 ; GCN-NOT: image_sample
 ; GCN: branch
 ; GCN: image_sample
 ; GCN: exp null
 
-define amdgpu_ps float @sinking_img_sample() {
+define amdgpu_ps float @sinking_img_sample(i1 %c0) {
 main_body:
   %i = call <3 x float> @llvm.amdgcn.image.sample.2d.v3f32.f32(i32 7, float undef, float undef, <8 x i32> undef, <4 x i32> undef, i1 false, i32 0, i32 0)
-  br i1 undef, label %endif1, label %if1
+  br i1 %c0, label %endif1, label %if1
 
 if1:                                              ; preds = %main_body
   call void @llvm.amdgcn.kill(i1 false) #4
   br label %exit
 
 endif1:                                           ; preds = %main_body
   %i22 = extractelement <3 x float> %i, i32 2
   %i23 = call nsz arcp contract float @llvm.fma.f32(float %i22, float 0.000000e+00, float 0.000000e+00) #1

llvm/test/CodeGen/AMDGPU/skip-if-dead.ll

bb:
   br i1 %loop.cond, label %bb, label %exit
 
 exit:
   store volatile i32 8, ptr addrspace(1) undef
   ret void
 }
 
 ; bug 28550
-define amdgpu_ps void @phi_use_def_before_kill(float inreg %x) #0 {
+define amdgpu_ps void @phi_use_def_before_kill(float inreg %x, i1 %c0) #0 {
 ; SI-LABEL: phi_use_def_before_kill:
 ; SI:       ; %bb.0: ; %bb
 ; SI-NEXT:    v_add_f32_e64 v1, s0, 1.0
 ; SI-NEXT:    v_cmp_lt_f32_e32 vcc, 0, v1
 ; SI-NEXT:    v_cndmask_b32_e64 v0, 0, -1.0, vcc
 ; SI-NEXT:    v_cmp_nlt_f32_e32 vcc, 0, v1
 ; SI-NEXT:    s_andn2_b64 exec, exec, vcc
 ; SI-NEXT:    s_cbranch_scc0 .LBB11_6

 ; GFX11-NEXT:    exp mrt0 off, off, off, off done
 ; GFX11-NEXT:    s_endpgm
 bb:
   %tmp = fadd float %x, 1.000000e+00
   %tmp1 = fcmp olt float 0.000000e+00, %tmp
   %tmp2 = select i1 %tmp1, float -1.000000e+00, float 0.000000e+00
   %cmp.tmp2 = fcmp olt float %tmp2, 0.0
   call void @llvm.amdgcn.kill(i1 %cmp.tmp2)
-  br i1 undef, label %phibb, label %bb8
+  br i1 %c0, label %phibb, label %bb8
 
 phibb:
   %tmp5 = phi float [ %tmp2, %bb ], [ 4.0, %bb8 ]
   %tmp6 = fcmp oeq float %tmp5, 0.000000e+00
   br i1 %tmp6, label %bb10, label %end
 
 bb8:
   store volatile i32 8, ptr addrspace(1) undef

llvm/test/CodeGen/AMDGPU/smrd.ll

 ; on a divergent shader input is correctly recognized as divergent. This was
 ; reduced from an actual regression. Yes, the %unused argument matters, as
 ; well as the fact that %arg4 is a vector.
 ;
 ; GCN-LABEL: {{^}}arg_divergence:
 ; GCN: buffer_load_dword v0, v0,
 ; GCN-NEXT: s_waitcnt
 ; GCN-NEXT: ; return to shader part epilog
-define amdgpu_cs float @arg_divergence(i32 inreg %unused, <3 x i32> %arg4) #0 {
+define amdgpu_cs float @arg_divergence(i32 inreg %unused, <3 x i32> %arg4, i1 %c0) #0 {
 main_body:
-  br i1 undef, label %if1, label %endif1
+  br i1 %c0, label %if1, label %endif1
 
 if1:                                              ; preds = %main_body
   store i32 0, ptr addrspace(3) undef, align 4
   br label %endif1
 
 endif1:                                           ; preds = %if1, %main_body
   %tmp13 = extractelement <3 x i32> %arg4, i32 0
   %tmp97 = call float @llvm.amdgcn.s.buffer.load.f32(<4 x i32> undef, i32 %tmp13, i32 0)

llvm/test/CodeGen/AMDGPU/split-smrd.ll

 ; RUN: llc -march=amdgcn -mcpu=tahiti -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
 
 ; FIXME: Move this to sgpr-copy.ll when this is fixed on VI.
 ; Make sure that when we split an smrd instruction in order to move it to
 ; the VALU, we are also moving its users to the VALU.
 
 ; GCN-LABEL: {{^}}split_smrd_add_worklist:
 ; GCN: image_sample v{{[0-9]+}}, v[{{[0-9]+:[0-9]+}}], s[{{[0-9]+:[0-9]+}}], s[{{[0-9]+:[0-9]+}}] dmask:0x1
-define amdgpu_ps void @split_smrd_add_worklist(ptr addrspace(4) inreg %arg) #0 {
+define amdgpu_ps void @split_smrd_add_worklist(ptr addrspace(4) inreg %arg, i1 %c0) #0 {
 bb:
   %tmp = call float @llvm.amdgcn.s.buffer.load.f32(<4 x i32> undef, i32 96, i32 0)
   %tmp1 = bitcast float %tmp to i32
-  br i1 undef, label %bb2, label %bb3
+  br i1 %c0, label %bb2, label %bb3
 
 bb2:                                              ; preds = %bb
   unreachable
 
 bb3:                                              ; preds = %bb
   %tmp4 = bitcast float %tmp to i32
   %tmp5 = add i32 %tmp4, 4
   %tmp6 = sext i32 %tmp5 to i64

llvm/test/CodeGen/AMDGPU/subreg-coalescer-crash.ll

 ; RUN: llc -march=amdgcn -mcpu=tahiti -verify-machineinstrs -o - %s | FileCheck -check-prefix=GCN %s
 ; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs -o - %s | FileCheck -check-prefix=GCN %s
 
 ; GCN-LABEL:{{^}}row_filter_C1_D0:
-define amdgpu_kernel void @row_filter_C1_D0() #0 {
+define amdgpu_kernel void @row_filter_C1_D0(i1 %c0, i1 %c1, i1 %c2) #0 {
 entry:
-  br i1 undef, label %for.inc.1, label %do.body.preheader
+  br i1 %c0, label %for.inc.1, label %do.body.preheader
 
 do.body.preheader:                                ; preds = %entry
   %tmp = insertelement <4 x i32> zeroinitializer, i32 undef, i32 1
-  br i1 undef, label %do.body56.1, label %do.body90
+  br i1 %c1, label %do.body56.1, label %do.body90
 
 do.body90:                                        ; preds = %do.body56.2, %do.body56.1, %do.body.preheader
   %tmp1 = phi <4 x i32> [ %tmp6, %do.body56.2 ], [ %tmp5, %do.body56.1 ], [ %tmp, %do.body.preheader ]
   %tmp2 = insertelement <4 x i32> %tmp1, i32 undef, i32 2
   %tmp3 = insertelement <4 x i32> %tmp2, i32 undef, i32 3
-  br i1 undef, label %do.body124.1, label %do.body.1562.preheader
+  br i1 %c2, label %do.body124.1, label %do.body.1562.preheader
 
 do.body.1562.preheader:                           ; preds = %do.body124.1, %do.body90
   %storemerge = phi <4 x i32> [ %tmp3, %do.body90 ], [ %tmp7, %do.body124.1 ]
   %tmp4 = insertelement <4 x i32> undef, i32 undef, i32 1
   br label %for.inc.1
 
 do.body56.1:                                      ; preds = %do.body.preheader
   %tmp5 = insertelement <4 x i32> %tmp, i32 undef, i32 1

 for.inc.1:                                        ; preds = %do.body.1562.preheader, %entry
   %storemerge591 = phi <4 x i32> [ zeroinitializer, %entry ], [ %storemerge, %do.body.1562.preheader ]
   %add.i495 = add <4 x i32> %storemerge591, undef
   unreachable
 }
 
 ; GCN-LABEL: {{^}}foo:
 ; GCN: s_endpgm
-define amdgpu_ps void @foo() #0 {
+define amdgpu_ps void @foo(i1 %c0, i1 %c1, i1 %c2, i1 %c3, i1 %c4, i1 %c5, i1 %c6) #0 {
 bb:
-  br i1 undef, label %bb2, label %bb1
+  br i1 %c0, label %bb2, label %bb1
 
 bb1:                                              ; preds = %bb
-  br i1 undef, label %bb4, label %bb6
+  br i1 %c1, label %bb4, label %bb6
 
 bb2:                                              ; preds = %bb4, %bb
   %tmp = phi float [ %tmp5, %bb4 ], [ 0.000000e+00, %bb ]
-  br i1 undef, label %bb9, label %bb13
+  br i1 %c2, label %bb9, label %bb13
 
 bb4:                                              ; preds = %bb7, %bb6, %bb1
   %tmp5 = phi float [ undef, %bb1 ], [ undef, %bb6 ], [ %tmp8, %bb7 ]
   br label %bb2
 
 bb6:                                              ; preds = %bb1
-  br i1 undef, label %bb7, label %bb4
+  br i1 %c3, label %bb7, label %bb4
 
 bb7:                                              ; preds = %bb6
   %tmp8 = fmul float undef, undef
   br label %bb4
 
 bb9:                                              ; preds = %bb2
   %tmp10 = call <4 x float> @llvm.amdgcn.image.sample.1d.v4f32.f32(i32 15, float undef, <8 x i32> undef, <4 x i32> undef, i1 0, i32 0, i32 0)
   %tmp11 = extractelement <4 x float> %tmp10, i32 1
   %tmp12 = extractelement <4 x float> %tmp10, i32 3
   br label %bb14
 
 bb13:                                             ; preds = %bb2
-  br i1 undef, label %bb23, label %bb24
+  br i1 %c4, label %bb23, label %bb24
 
 bb14:                                             ; preds = %bb27, %bb24, %bb9
   %tmp15 = phi float [ %tmp12, %bb9 ], [ undef, %bb27 ], [ 0.000000e+00, %bb24 ]
   %tmp16 = phi float [ %tmp11, %bb9 ], [ undef, %bb27 ], [ %tmp25, %bb24 ]
   %tmp17 = fmul float 1.050000e+01, %tmp16
   %tmp18 = fmul float 1.150000e+01, %tmp15
   call void @llvm.amdgcn.exp.f32(i32 0, i32 15, float %tmp18, float %tmp17, float %tmp17, float %tmp17, i1 true, i1 true) #0
   ret void
 
 bb23:                                             ; preds = %bb13
-  br i1 undef, label %bb24, label %bb26
+  br i1 %c5, label %bb24, label %bb26
 
 bb24:                                             ; preds = %bb26, %bb23, %bb13
   %tmp25 = phi float [ %tmp, %bb13 ], [ %tmp, %bb26 ], [ 0.000000e+00, %bb23 ]
-  br i1 undef, label %bb27, label %bb14
+  br i1 %c6, label %bb27, label %bb14
 
 bb26:                                             ; preds = %bb23
   br label %bb24
 
 bb27:                                             ; preds = %bb24
   br label %bb14
 }
 
 
 declare void @llvm.amdgcn.exp.f32(i32, i32, float, float, float, float, i1, i1) #0
 declare <4 x float> @llvm.amdgcn.image.sample.1d.v4f32.f32(i32, float, <8 x i32>, <4 x i32>, i1, i32, i32) #1
 
 attributes #0 = { nounwind }
 attributes #1 = { nounwind readonly }

llvm/test/CodeGen/AMDGPU/switch-default-block-unreachable.ll

 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -verify-machineinstrs -stop-after=amdgpu-isel -o - %s | FileCheck -check-prefix=GCN %s
-define void @test() #1 {
+define void @test(i1 %c0) #1 {
   ; Clean up the unreachable blocks introduced with LowerSwitch pass.
   ; This test ensures that, in the pass flow, UnreachableBlockElim pass
   ; follows the LowerSwitch. Otherwise, this testcase will crash
   ; immediately after the instruction selection due to the incomplete
   ; PHI node in an MBB whose incoming values were never codegenerated.
   ;
   ; GCN-LABEL: name: test
   ; GCN: bb.{{[0-9]+}}.entry:
   ; GCN: bb.{{[0-9]+}}.Flow1:
   ; GCN: bb.{{[0-9]+}}.entry.true.blk:
   ; GCN: bb.{{[0-9]+}}.entry.false.blk:
   ; GCN: bb.{{[0-9]+}}.switch.blk:
 
   ; GCN-NOT: bb.{{[0-9]+}}.preheader.blk
   ; GCN-NOT: bb.{{[0-9]+}}.pre.false.blk:
   ; GCN-NOT: bb.{{[0-9]+}}.unreach.blk:
   ; GCN-NOT: PHI
 
   ; GCN: bb.{{[0-9]+}}.Flow:
   ; GCN: bb.{{[0-9]+}}.UnifiedReturnBlock:
   entry:
     %idx = tail call i32 @llvm.amdgcn.workitem.id.x() #0
-    br i1 undef, label %entry.true.blk, label %entry.false.blk
+    br i1 %c0, label %entry.true.blk, label %entry.false.blk
 
   entry.true.blk:                                   ; preds = %entry
     %exit.cmp = icmp ult i32 %idx, 3
     br i1 %exit.cmp, label %switch.blk, label %exit
 
   entry.false.blk:                                  ; preds = %entry
     unreachable
 

llvm/test/CodeGen/AMDGPU/undefined-subreg-liverange.ll

 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
 ; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck %s
 ; RUN: llc -march=amdgcn -verify-machineinstrs -early-live-intervals < %s | FileCheck %s
 
 ; We may have subregister live ranges that are undefined on some paths. The
 ; verifier should not complain about this.
 
-define amdgpu_kernel void @func() #0 {
+define amdgpu_kernel void @func(i1 %c0, i1 %c1) #0 {
 ; CHECK-LABEL: func:
 ; CHECK:       ; %bb.0: ; %B0
 ; CHECK-NEXT:    s_mov_b32 s0, 0
 ; CHECK-NEXT:    s_cbranch_scc1 .LBB0_2
 ; CHECK-NEXT:  ; %bb.1: ; %B30.1
 ; CHECK-NEXT:    s_mov_b32 s0, 0x7fc00000
 ; CHECK-NEXT:  .LBB0_2: ; %B30.2
 ; CHECK-NEXT:    v_mov_b32_e32 v0, s0
 ; CHECK-NEXT:    s_mov_b32 m0, -1
 ; CHECK-NEXT:    ds_write_b32 v0, v0
 ; CHECK-NEXT:    s_endpgm
 B0:
-  br i1 undef, label %B1, label %B2
+  br i1 %c0, label %B1, label %B2
 
 B1:
   br label %B2
 
 B2:
   %v0 = phi <4 x float> [ zeroinitializer, %B1 ], [ <float 0.0, float 0.0, float 0.0, float undef>, %B0 ]
-  br i1 undef, label %B30.1, label %B30.2
+  br i1 %c1, label %B30.1, label %B30.2
 
 B30.1:
   %sub = fsub <4 x float> %v0, undef
   br label %B30.2
 
 B30.2:
   %v3 = phi <4 x float> [ %sub, %B30.1 ], [ %v0, %B2 ]
   %ve0 = extractelement <4 x float> %v3, i32 0

llvm/test/CodeGen/AMDGPU/unhandled-loop-condition-assertion.ll

 ; RUN: llc -O0 -verify-machineinstrs -asm-verbose=0 -march=amdgcn < %s | FileCheck -check-prefix=SI -check-prefix=COMMON %s
 ; RUN: llc -O0 -verify-machineinstrs -asm-verbose=0 -march=amdgcn -mcpu=tonga < %s | FileCheck -check-prefix=SI -check-prefix=COMMON %s
 ; XUN: llc -O0 -verify-machineinstrs -asm-verbose=0 -march=r600 -mcpu=redwood < %s | FileCheck -check-prefix=EG -check-prefix=COMMON %s
 
 ; SI hits an assertion at -O0, evergreen hits a not implemented unreachable.
 
 ; COMMON-LABEL: {{^}}branch_true:
-define amdgpu_kernel void @branch_true(ptr addrspace(1) nocapture %main, i32 %main_stride) #0 {
+define amdgpu_kernel void @branch_true(ptr addrspace(1) nocapture %main, i32 %main_stride, i1 %c0) #0 {
 entry:
   br i1 true, label %for.end, label %for.body.lr.ph
 
 for.body.lr.ph:                                   ; preds = %entry
   %add.ptr.sum = shl i32 %main_stride, 1
   %add.ptr1.sum = add i32 %add.ptr.sum, %main_stride
   %add.ptr4.sum = shl i32 %main_stride, 2
   br label %for.body
 
 for.body:                                         ; preds = %for.body, %for.body.lr.ph
   %main.addr.011 = phi ptr addrspace(1) [ %main, %for.body.lr.ph ], [ %add.ptr6, %for.body ]
   %0 = load i32, ptr addrspace(1) %main.addr.011, align 4
   %add.ptr = getelementptr inbounds i8, ptr addrspace(1) %main.addr.011, i32 %main_stride
   %1 = load i32, ptr addrspace(1) %add.ptr, align 4
   %add.ptr1 = getelementptr inbounds i8, ptr addrspace(1) %main.addr.011, i32 %add.ptr.sum
   %2 = load i32, ptr addrspace(1) %add.ptr1, align 4
   %add.ptr2 = getelementptr inbounds i8, ptr addrspace(1) %main.addr.011, i32 %add.ptr1.sum
   %3 = load i32, ptr addrspace(1) %add.ptr2, align 4
   %add.ptr3 = getelementptr inbounds i8, ptr addrspace(1) %main.addr.011, i32 %add.ptr4.sum
   %4 = load i32, ptr addrspace(1) %add.ptr3, align 4
   %add.ptr6 = getelementptr inbounds i8, ptr addrspace(1) %main.addr.011, i32 undef
-  br i1 undef, label %for.end, label %for.body
+  br i1 %c0, label %for.end, label %for.body
 
 for.end:                                          ; preds = %for.body, %entry
   ret void
 }
 
 ; COMMON-LABEL: {{^}}branch_false:
 ; SI: s_cbranch_scc1
 ; SI: s_endpgm
-define amdgpu_kernel void @branch_false(ptr addrspace(1) nocapture %main, i32 %main_stride) #0 {
+define amdgpu_kernel void @branch_false(ptr addrspace(1) nocapture %main, i32 %main_stride, i1 %c0) #0 {
 entry:
   br i1 false, label %for.end, label %for.body.lr.ph
 
 for.body.lr.ph:                                   ; preds = %entry
   %add.ptr.sum = shl i32 %main_stride, 1
   %add.ptr1.sum = add i32 %add.ptr.sum, %main_stride
   %add.ptr4.sum = shl i32 %main_stride, 2
   br label %for.body
 
 for.body:                                         ; preds = %for.body, %for.body.lr.ph
   %main.addr.011 = phi ptr addrspace(1) [ %main, %for.body.lr.ph ], [ %add.ptr6, %for.body ]
   %0 = load i32, ptr addrspace(1) %main.addr.011, align 4
   %add.ptr = getelementptr inbounds i8, ptr addrspace(1) %main.addr.011, i32 %main_stride
   %1 = load i32, ptr addrspace(1) %add.ptr, align 4
   %add.ptr1 = getelementptr inbounds i8, ptr addrspace(1) %main.addr.011, i32 %add.ptr.sum
   %2 = load i32, ptr addrspace(1) %add.ptr1, align 4
   %add.ptr2 = getelementptr inbounds i8, ptr addrspace(1) %main.addr.011, i32 %add.ptr1.sum
   %3 = load i32, ptr addrspace(1) %add.ptr2, align 4
   %add.ptr3 = getelementptr inbounds i8, ptr addrspace(1) %main.addr.011, i32 %add.ptr4.sum
   %4 = load i32, ptr addrspace(1) %add.ptr3, align 4
   %add.ptr6 = getelementptr inbounds i8, ptr addrspace(1) %main.addr.011, i32 undef
-  br i1 undef, label %for.end, label %for.body
+  br i1 %c0, label %for.end, label %for.body
 
 for.end:                                          ; preds = %for.body, %entry
   ret void
 }
 
-; COMMON-LABEL: {{^}}branch_undef:
+; COMMON-LABEL: {{^}}branch_[[%c0:%.*]]:
 ; SI: s_cbranch_scc1
 ; SI: s_cbranch_scc1
 ; SI: s_endpgm
-define amdgpu_kernel void @branch_undef(ptr addrspace(1) nocapture %main, i32 %main_stride) #0 {
+define amdgpu_kernel void @branch_undef(ptr addrspace(1) nocapture %main, i32 %main_stride, i1 %c0, i1 %c1) #0 {
 entry:
-  br i1 undef, label %for.end, label %for.body.lr.ph
+  br i1 %c0, label %for.end, label %for.body.lr.ph
 
 for.body.lr.ph:                                   ; preds = %entry
   %add.ptr.sum = shl i32 %main_stride, 1
   %add.ptr1.sum = add i32 %add.ptr.sum, %main_stride
   %add.ptr4.sum = shl i32 %main_stride, 2
   br label %for.body
 
 for.body:                                         ; preds = %for.body, %for.body.lr.ph
   %main.addr.011 = phi ptr addrspace(1) [ %main, %for.body.lr.ph ], [ %add.ptr6, %for.body ]
   %0 = load i32, ptr addrspace(1) %main.addr.011, align 4
   %add.ptr = getelementptr inbounds i8, ptr addrspace(1) %main.addr.011, i32 %main_stride
   %1 = load i32, ptr addrspace(1) %add.ptr, align 4
   %add.ptr1 = getelementptr inbounds i8, ptr addrspace(1) %main.addr.011, i32 %add.ptr.sum
   %2 = load i32, ptr addrspace(1) %add.ptr1, align 4
   %add.ptr2 = getelementptr inbounds i8, ptr addrspace(1) %main.addr.011, i32 %add.ptr1.sum
   %3 = load i32, ptr addrspace(1) %add.ptr2, align 4
   %add.ptr3 = getelementptr inbounds i8, ptr addrspace(1) %main.addr.011, i32 %add.ptr4.sum
   %4 = load i32, ptr addrspace(1) %add.ptr3, align 4
   %add.ptr6 = getelementptr inbounds i8, ptr addrspace(1) %main.addr.011, i32 undef
-  br i1 undef, label %for.end, label %for.body
+  br i1 %c1, label %for.end, label %for.body
 
 for.end:                                          ; preds = %for.body, %entry
   ret void
 }
 
 attributes #0 = { nounwind }

llvm/test/CodeGen/AMDGPU/uniform-cfg.ll

else:
   br label %done
 
 done:
   %value = phi i32 [0, %if], [1, %else]
   store i32 %value, ptr addrspace(1) %out
   ret void
 }
 
-define void @move_to_valu_vgpr_operand_phi(ptr addrspace(3) %out) {
+define void @move_to_valu_vgpr_operand_phi(ptr addrspace(3) %out, i1 %c0) {
 ; SI-LABEL: move_to_valu_vgpr_operand_phi:
 ; SI:       ; %bb.0: ; %bb0
 ; SI-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; SI-NEXT:    v_add_i32_e32 v0, vcc, 28, v0
 ; SI-NEXT:    v_mov_b32_e32 v1, 1
 ; SI-NEXT:    s_and_b64 vcc, exec, 0
 ; SI-NEXT:    s_mov_b32 m0, -1
 ; SI-NEXT:    s_branch .LBB20_2

 ; VI-NEXT:    s_setpc_b64 s[30:31]
 bb0:
   br label %bb1
 
 bb1:                                              ; preds = %bb3, %bb0
   %tmp0 = phi i32 [ 8, %bb0 ], [ %tmp4, %bb3 ]
   %tmp1 = add nsw i32 %tmp0, -1
   %tmp2 = getelementptr inbounds i32, ptr addrspace(3) %out, i32 %tmp1
-  br i1 undef, label %bb2, label %bb3
+  br i1 %c0, label %bb2, label %bb3
 
 bb2:                                              ; preds = %bb1
   store volatile i32 1, ptr addrspace(3) %tmp2, align 4
   br label %bb3
 
 bb3:                                              ; preds = %bb2, %bb1
   %tmp4 = add nsw i32 %tmp0, 2
   br label %bb1
 }
 
 declare i32 @llvm.amdgcn.workitem.id.x() #0
 
 attributes #0 = { nounwind readnone }
 ;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
 ; GCN: {{.*}}

llvm/test/CodeGen/AMDGPU/uniform-crash.ll

ENDIF:                                            ; preds = %IF, %main_body
   ret void
 }
 
 ; GCN-LABEL: {{^}}fix_sgpr_live_ranges_crash:
 ; GCN: s_cbranch_scc1 [[BB0:.L[A-Z0-9_]+]]
 ; GCN: {{^}}[[LOOP:.L[A-Z0-9_]+]]:
 ; GCN: s_cbranch_scc1 [[LOOP]]
 ; GCN: {{^}}[[BB0]]:
-define amdgpu_kernel void @fix_sgpr_live_ranges_crash(i32 %arg, i32 %arg1)  {
+define amdgpu_kernel void @fix_sgpr_live_ranges_crash(i32 %arg, i32 %arg1, i1 %c0)  {
 bb:
   %cnd = trunc i32 %arg to i1
   br i1 %cnd, label %bb2, label %bb5
 
 bb2:                                              ; preds = %bb
   %tmp = mul i32 10, %arg1
   br label %bb3
 
 bb3:                                              ; preds = %bb3, %bb2
   %val = load volatile i32, ptr addrspace(4) undef
   %tmp4 = icmp eq i32 %val, %arg1
   br i1 %tmp4, label %bb5, label %bb3
 
 bb5:                                              ; preds = %bb3, %bb
   %tmp6 = tail call i32 @llvm.amdgcn.workitem.id.y() #1
   %tmp10 = icmp ult i32 %tmp6, %arg
   br i1 %tmp10, label %bb11, label %bb12
 
 bb11:                                             ; preds = %bb11, %bb5
-  br i1 undef, label %bb11, label %bb12
+  br i1 %c0, label %bb11, label %bb12
 
 bb12:                                             ; preds = %bb11, %bb5
   ret void
 }
 
 ; Function Attrs: nounwind readnone
 declare i32 @llvm.amdgcn.workitem.id.y() #1
 
 attributes #1 = { nounwind readnone }

llvm/test/CodeGen/AMDGPU/v1024.ll

 ; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
 
 ; Check that we do not use AGPRs for v32i32 type
 
 ; GCN-LABEL: {{^}}test_v1024:
 ; GCN-NOT: v_accvgpr
 ; GCN-COUNT-32: v_mov_b32_e32
 ; GCN-NOT: v_accvgpr
-define amdgpu_kernel void @test_v1024() {
+define amdgpu_kernel void @test_v1024(i1 %c0) {
 entry:
   %alloca = alloca <32 x i32>, align 16, addrspace(5)
   call void @llvm.memset.p5.i32(ptr addrspace(5) %alloca, i8 0, i32 128, i1 false)
-  br i1 undef, label %if.then.i.i, label %if.else.i
+  br i1 %c0, label %if.then.i.i, label %if.else.i
 
 if.then.i.i:                                      ; preds = %entry
   call void @llvm.memcpy.p5.p5.i64(ptr addrspace(5) align 16 %alloca, ptr addrspace(5) align 4 undef, i64 128, i1 false)
   br label %if.then.i62.i
 
 if.else.i:                                        ; preds = %entry
   br label %if.then.i62.i
 

llvm/test/CodeGen/AMDGPU/vgpr-liverange-ir.ll

else:
   br label %end
 
 end:
   %r = phi float [ %v.if, %if ], [ %v.else, %else ]
   %r2 = fadd float %r, %v
   ret float %r2
 }
 
-define amdgpu_kernel void @livevariables_update_missed_block(ptr addrspace(1) %src1) {
+define amdgpu_kernel void @livevariables_update_missed_block(ptr addrspace(1) %src1, i1 %c0) {
   ; SI-LABEL: name: livevariables_update_missed_block
   ; SI: bb.0.entry:
   ; SI-NEXT:   successors: %bb.2(0x40000000), %bb.5(0x40000000)
   ; SI-NEXT:   liveins: $vgpr0, $sgpr0_sgpr1
   ; SI-NEXT: {{  $}}
   ; SI-NEXT:   [[COPY:%[0-9]+]]:sgpr_64(p4) = COPY killed $sgpr0_sgpr1
   ; SI-NEXT:   [[COPY1:%[0-9]+]]:vgpr_32(s32) = COPY killed $vgpr0
   ; SI-NEXT:   [[V_CMP_NE_U32_e64_:%[0-9]+]]:sreg_32 = V_CMP_NE_U32_e64 0, [[COPY1]](s32), implicit $exec
if.then:                                          ; preds = %entry
   %i13 = insertelement <3 x i8> %i12, i8 0, i64 1
   %i14 = insertelement <3 x i8> %i13, i8 0, i64 1
   %i15 = select <3 x i1> zeroinitializer, <3 x i8> zeroinitializer, <3 x i8> %i14
   %i16 = extractelement <3 x i8> %i15, i64 0
   store i8 %i16, ptr addrspace(1) null, align 1
   ret void
 
 if.then9:                                         ; preds = %entry
-  br i1 undef, label %sw.bb18, label %sw.bb
+  br i1 %c0, label %sw.bb18, label %sw.bb
 
 sw.bb:                                            ; preds = %if.then9
   %i17 = load i8, ptr addrspace(1) null, align 1
   %i18 = insertelement <4 x i8> zeroinitializer, i8 %i17, i64 0
   %a.sroa.0.0.vecblend = shufflevector <4 x i8> %i18, <4 x i8> zeroinitializer, <4 x i32> <i32 0, i32 0, i32 0, i32 undef>
   br label %sw.bb18
 
 sw.bb18:                                          ; preds = %sw.bb, %if.then9

llvm/test/CodeGen/AMDGPU/wave32.ll

define amdgpu_ps void @test_wqm_vote(float %a) {
   call void @llvm.amdgcn.kill(i1 %c2)
   call void @llvm.amdgcn.exp.f32(i32 0, i32 0, float 0.0, float 0.0, float 0.0, float 0.0, i1 false, i1 false)
   ret void
 }
 
 ; GCN-LABEL: {{^}}test_branch_true:
 ; GFX1032: s_mov_b32 vcc_lo, exec_lo
 ; GFX1064: s_mov_b64 vcc, exec
-define amdgpu_kernel void @test_branch_true() #2 {
+define amdgpu_kernel void @test_branch_true(i1 %c0) #2 {
 entry:
   br i1 true, label %for.end, label %for.body.lr.ph
 
 for.body.lr.ph:                                   ; preds = %entry
   br label %for.body
 
 for.body:                                         ; preds = %for.body, %for.body.lr.ph
-  br i1 undef, label %for.end, label %for.body
+  br i1 %c0, label %for.end, label %for.body
 
 for.end:                                          ; preds = %for.body, %entry
   ret void
 }
 
 ; GCN-LABEL: {{^}}test_ps_live:
 ; GFX1032: s_mov_b32 [[C:s[0-9]+]], exec_lo
 ; GFX1064: s_mov_b64 [[C:s\[[0-9:]+\]]], exec{{$}}