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

[kritgpt]

Summary

This patch removes trivial occurences of 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.

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/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-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/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/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/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-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/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/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/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-8: global_store_dwordx4
 ; 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