From 047efda60db6a3a5b614dab900bd7e432c32db1a Mon Sep 17 00:00:00 2001 From: Nuno Lopes Date: Tue, 4 Apr 2023 13:45:56 +0100 Subject: [PATCH] Revert "[test] Remove occurences of br undef in CodeGen/AMDGPU tests" This reverts commit 18c594c176036b7ffcd8439ed9c4b08d2085a244. Build bots broke --- llvm/test/CodeGen/AMDGPU/coalescer_distribute.ll | 12 +++++----- llvm/test/CodeGen/AMDGPU/extract-subvector.ll | 28 +++++++++++----------- llvm/test/CodeGen/AMDGPU/i1-copy-phi.ll | 4 ++-- llvm/test/CodeGen/AMDGPU/implicit-def-muse.ll | 4 ++-- llvm/test/CodeGen/AMDGPU/inline-asm.ll | 4 ++-- .../CodeGen/AMDGPU/large-constant-initializer.ll | 4 ++-- llvm/test/CodeGen/AMDGPU/mdt-preserving-crash.ll | 4 ++-- llvm/test/CodeGen/AMDGPU/mfma-cd-select.ll | 4 ++-- .../CodeGen/AMDGPU/promote-alloca-to-lds-phi.ll | 16 ++++++------- .../CodeGen/AMDGPU/promote-alloca-to-lds-select.ll | 4 ++-- .../CodeGen/AMDGPU/reg-coalescer-sched-crash.ll | 4 ++-- .../test/CodeGen/AMDGPU/rename-disconnected-bug.ll | 8 +++---- .../CodeGen/AMDGPU/si-annotate-cf-unreachable.ll | 4 ++-- .../AMDGPU/simplifydemandedbits-recursion.ll | 14 +++++------ .../AMDGPU/switch-default-block-unreachable.ll | 4 ++-- llvm/test/CodeGen/AMDGPU/uniform-crash.ll | 4 ++-- llvm/test/CodeGen/AMDGPU/v1024.ll | 4 ++-- llvm/test/CodeGen/AMDGPU/wave32.ll | 4 ++-- 18 files changed, 65 insertions(+), 65 deletions(-) diff --git a/llvm/test/CodeGen/AMDGPU/coalescer_distribute.ll b/llvm/test/CodeGen/AMDGPU/coalescer_distribute.ll index 2d8ef33..c622349 100644 --- a/llvm/test/CodeGen/AMDGPU/coalescer_distribute.ll +++ b/llvm/test/CodeGen/AMDGPU/coalescer_distribute.ll @@ -2,13 +2,13 @@ ; This testcase produces a situation with unused value numbers in subregister ; liveranges that get distributed by ConnectedVNInfoEqClasses. -define amdgpu_kernel void @hoge(i1 %c0, i1 %c1, i1 %c2, i1 %c3, i1 %c4) { +define amdgpu_kernel void @hoge() { bb: %tmp = tail call i32 @llvm.amdgcn.workitem.id.x() - br i1 %c0, label %bb2, label %bb23 + br i1 undef, label %bb2, label %bb23 bb2: - br i1 %c1, label %bb6, label %bb8 + br i1 undef, label %bb6, label %bb8 bb6: %tmp7 = or i64 undef, undef @@ -20,7 +20,7 @@ bb8: br i1 %tmp10, label %bb11, label %bb23 bb11: - br i1 %c2, label %bb20, label %bb17 + br i1 undef, label %bb20, label %bb17 bb17: br label %bb20 @@ -36,10 +36,10 @@ bb23: bb25: %tmp26 = phi i32 [ %tmp24, %bb23 ], [ undef, %bb25 ] - br i1 %c3, label %bb25, label %bb30 + br i1 undef, label %bb25, label %bb30 bb30: - br i1 %c4, label %bb32, label %bb34 + br i1 undef, label %bb32, label %bb34 bb32: %tmp33 = zext i32 %tmp26 to i64 diff --git a/llvm/test/CodeGen/AMDGPU/extract-subvector.ll b/llvm/test/CodeGen/AMDGPU/extract-subvector.ll index 3a1c9f1..7d1a37a 100644 --- a/llvm/test/CodeGen/AMDGPU/extract-subvector.ll +++ b/llvm/test/CodeGen/AMDGPU/extract-subvector.ll @@ -20,8 +20,8 @@ ; GCN: v_bfe_i32 ; GCN: v_bfe_i32 -define <2 x i16> @extract_2xi16(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) { - br i1 %c0, label %T, label %F +define <2 x i16> @extract_2xi16(ptr addrspace(1) %p0, ptr addrspace(1) %p1) { + br i1 undef, label %T, label %F T: %t = load volatile <8 x i16>, ptr addrspace(1) %p0 @@ -41,8 +41,8 @@ exit: ; GCN-LABEL: extract_2xi64 ; GCN-COUNT-2: v_cndmask_b32 -define <2 x i64> @extract_2xi64(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) { - br i1 %c0, label %T, label %F +define <2 x i64> @extract_2xi64(ptr addrspace(1) %p0, ptr addrspace(1) %p1) { + br i1 undef, label %T, label %F T: %t = load volatile <8 x i64>, ptr addrspace(1) %p0 @@ -62,8 +62,8 @@ exit: ; GCN-LABEL: extract_4xi64 ; GCN-COUNT-4: v_cndmask_b32 -define <4 x i64> @extract_4xi64(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) { - br i1 %c0, label %T, label %F +define <4 x i64> @extract_4xi64(ptr addrspace(1) %p0, ptr addrspace(1) %p1) { + br i1 undef, label %T, label %F T: %t = load volatile <8 x i64>, ptr addrspace(1) %p0 @@ -83,8 +83,8 @@ exit: ; GCN-LABEL: extract_8xi64 ; GCN-COUNT-8: v_cndmask_b32 -define <8 x i64> @extract_8xi64(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) { - br i1 %c0, label %T, label %F +define <8 x i64> @extract_8xi64(ptr addrspace(1) %p0, ptr addrspace(1) %p1) { + br i1 undef, label %T, label %F T: %t = load volatile <16 x i64>, ptr addrspace(1) %p0 @@ -104,8 +104,8 @@ exit: ; GCN-LABEL: extract_2xf64 ; GCN-COUNT-2: v_cndmask_b32 -define <2 x double> @extract_2xf64(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) { - br i1 %c0, label %T, label %F +define <2 x double> @extract_2xf64(ptr addrspace(1) %p0, ptr addrspace(1) %p1) { + br i1 undef, label %T, label %F T: %t = load volatile <8 x double>, ptr addrspace(1) %p0 @@ -125,8 +125,8 @@ exit: ; GCN-LABEL: extract_4xf64 ; GCN-COUNT-4: v_cndmask_b32 -define <4 x double> @extract_4xf64(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) { - br i1 %c0, label %T, label %F +define <4 x double> @extract_4xf64(ptr addrspace(1) %p0, ptr addrspace(1) %p1) { + br i1 undef, label %T, label %F T: %t = load volatile <8 x double>, ptr addrspace(1) %p0 @@ -146,8 +146,8 @@ exit: ; GCN-LABEL: extract_8xf64 ; GCN-COUNT-8: v_cndmask_b32 -define <8 x double> @extract_8xf64(ptr addrspace(1) %p0, ptr addrspace(1) %p1, i1 %c0) { - br i1 %c0, label %T, label %F +define <8 x double> @extract_8xf64(ptr addrspace(1) %p0, ptr addrspace(1) %p1) { + br i1 undef, label %T, label %F T: %t = load volatile <16 x double>, ptr addrspace(1) %p0 diff --git a/llvm/test/CodeGen/AMDGPU/i1-copy-phi.ll b/llvm/test/CodeGen/AMDGPU/i1-copy-phi.ll index 9fc7530..82887c7 100644 --- a/llvm/test/CodeGen/AMDGPU/i1-copy-phi.ll +++ b/llvm/test/CodeGen/AMDGPU/i1-copy-phi.ll @@ -42,9 +42,9 @@ attributes #0 = { nounwind readnone } ; 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, i1 %c0) { +define <2 x float> @vcopy_i1_undef(ptr addrspace(1) %p) { entry: - br i1 %c0, label %exit, label %false + br i1 undef, label %exit, label %false false: %x = load <2 x float>, ptr addrspace(1) %p diff --git a/llvm/test/CodeGen/AMDGPU/implicit-def-muse.ll b/llvm/test/CodeGen/AMDGPU/implicit-def-muse.ll index be6f3f7..5e13fa2 100644 --- a/llvm/test/CodeGen/AMDGPU/implicit-def-muse.ll +++ b/llvm/test/CodeGen/AMDGPU/implicit-def-muse.ll @@ -7,9 +7,9 @@ ; CHECK-NOT: COPY [[IMPDEF0]] ; CHECK-NOT: COPY [[IMPDEF1]] ; CHECK: .false: -define <2 x float> @vcopy_i1_undef(ptr addrspace(1) %p, i1 %c0) { +define <2 x float> @vcopy_i1_undef(ptr addrspace(1) %p) { entry: - br i1 %c0, label %exit, label %false + br i1 undef, label %exit, label %false false: %x = load <2 x float>, ptr addrspace(1) %p diff --git a/llvm/test/CodeGen/AMDGPU/inline-asm.ll b/llvm/test/CodeGen/AMDGPU/inline-asm.ll index 173e4f6..b9637ab 100644 --- a/llvm/test/CodeGen/AMDGPU/inline-asm.ll +++ b/llvm/test/CodeGen/AMDGPU/inline-asm.ll @@ -299,9 +299,9 @@ entry: ; Check aggregate types are handled properly. ; CHECK-LABEL: mad_u64 ; CHECK: v_mad_u64_u32 -define void @mad_u64(i32 %x, i1 %c0) { +define void @mad_u64(i32 %x) { entry: - br i1 %c0, label %exit, label %false + br i1 undef, 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) diff --git a/llvm/test/CodeGen/AMDGPU/large-constant-initializer.ll b/llvm/test/CodeGen/AMDGPU/large-constant-initializer.ll index 65bf419..ff91860 100644 --- a/llvm/test/CodeGen/AMDGPU/large-constant-initializer.ll +++ b/llvm/test/CodeGen/AMDGPU/large-constant-initializer.ll @@ -4,10 +4,10 @@ @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, i1 %c0) nounwind { +define amdgpu_kernel void @opencv_cvtfloat_crash(ptr addrspace(1) %out, i32 %x) 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 %c0, label %exit, label %bb + br i1 undef, label %exit, label %bb bb: %cmp = icmp slt i32 %x, 0 diff --git a/llvm/test/CodeGen/AMDGPU/mdt-preserving-crash.ll b/llvm/test/CodeGen/AMDGPU/mdt-preserving-crash.ll index fc80ff2..2e015a9 100644 --- a/llvm/test/CodeGen/AMDGPU/mdt-preserving-crash.ll +++ b/llvm/test/CodeGen/AMDGPU/mdt-preserving-crash.ll @@ -4,7 +4,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__________________________________(i1 %c0) local_unnamed_addr #0 { +define protected amdgpu_kernel void @_RSENC_PRInit__________________________________() local_unnamed_addr #0 { entry: %runtimeVersionCopy = alloca [128 x i8], align 16, addrspace(5) %licenseVersionCopy = alloca [128 x i8], align 16, addrspace(5) @@ -18,7 +18,7 @@ if.end: ; preds = %entry br i1 %cmp13, label %cleanup.cont, label %if.end15 if.end15: ; preds = %if.end - br i1 %c0, label %cleanup.cont, label %lor.lhs.false17 + br i1 undef, label %cleanup.cont, label %lor.lhs.false17 lor.lhs.false17: ; preds = %if.end15 br label %while.cond.i diff --git a/llvm/test/CodeGen/AMDGPU/mfma-cd-select.ll b/llvm/test/CodeGen/AMDGPU/mfma-cd-select.ll index 7c1784f..96030a4 100644 --- a/llvm/test/CodeGen/AMDGPU/mfma-cd-select.ll +++ b/llvm/test/CodeGen/AMDGPU/mfma-cd-select.ll @@ -78,12 +78,12 @@ bb: ; 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, i1 %c0) #0 { +define amdgpu_kernel void @test_mfma_f32_32x32x1f32_call_multi_bb(ptr addrspace(1) %arg) #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 %c0, label %bb2, label %bb3 + br i1 undef, label %bb2, label %bb3 br label %bb2 bb2: diff --git a/llvm/test/CodeGen/AMDGPU/promote-alloca-to-lds-phi.ll b/llvm/test/CodeGen/AMDGPU/promote-alloca-to-lds-phi.ll index 8f34612..a982bc8 100644 --- a/llvm/test/CodeGen/AMDGPU/promote-alloca-to-lds-phi.ll +++ b/llvm/test/CodeGen/AMDGPU/promote-alloca-to-lds-phi.ll @@ -13,10 +13,10 @@ ; 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, i1 %c0) #0 { +define amdgpu_kernel void @branch_ptr_var_same_alloca(i32 %a, i32 %b) #0 { entry: %alloca = alloca [64 x i32], align 4, addrspace(5) - br i1 %c0, label %if, label %else + br i1 undef, label %if, label %else if: %arrayidx0 = getelementptr inbounds [64 x i32], ptr addrspace(5) %alloca, i32 0, i32 %a @@ -34,10 +34,10 @@ endif: ; 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, i1 %c0) #0 { +define amdgpu_kernel void @branch_ptr_phi_alloca_null_0(i32 %a, i32 %b) #0 { entry: %alloca = alloca [64 x i32], align 4, addrspace(5) - br i1 %c0, label %if, label %endif + br i1 undef, label %if, label %endif if: %arrayidx0 = getelementptr inbounds [64 x i32], ptr addrspace(5) %alloca, i32 0, i32 %a @@ -51,10 +51,10 @@ endif: ; 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, i1 %c0) #0 { +define amdgpu_kernel void @branch_ptr_phi_alloca_null_1(i32 %a, i32 %b) #0 { entry: %alloca = alloca [64 x i32], align 4, addrspace(5) - br i1 %c0, label %if, label %endif + br i1 undef, label %if, label %endif if: %arrayidx0 = getelementptr inbounds [64 x i32], ptr addrspace(5) %alloca, i32 0, i32 %a @@ -97,10 +97,10 @@ exit: ; 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, i1 %c0) #0 { +define amdgpu_kernel void @branch_ptr_alloca_unknown_obj(i32 %a, i32 %b) #0 { entry: %alloca = alloca [64 x i32], align 4, addrspace(5) - br i1 %c0, label %if, label %else + br i1 undef, label %if, label %else if: %arrayidx0 = getelementptr inbounds [64 x i32], ptr addrspace(5) %alloca, i32 0, i32 %a diff --git a/llvm/test/CodeGen/AMDGPU/promote-alloca-to-lds-select.ll b/llvm/test/CodeGen/AMDGPU/promote-alloca-to-lds-select.ll index 8fcbccc..b32faec 100644 --- a/llvm/test/CodeGen/AMDGPU/promote-alloca-to-lds-select.ll +++ b/llvm/test/CodeGen/AMDGPU/promote-alloca-to-lds-select.ll @@ -77,13 +77,13 @@ define amdgpu_kernel void @lds_promoted_alloca_select_input_select(i32 %a, i32 % ret void } -define amdgpu_kernel void @lds_promoted_alloca_select_input_phi(i32 %a, i32 %b, i32 %c, i1 %c0) #0 { +define amdgpu_kernel void @lds_promoted_alloca_select_input_phi(i32 %a, i32 %b, i32 %c) #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 %c0, label %bb1, label %bb2 + br i1 undef, label %bb1, label %bb2 bb1: %ptr2 = getelementptr inbounds [16 x i32], ptr addrspace(5) %alloca, i32 0, i32 %c diff --git a/llvm/test/CodeGen/AMDGPU/reg-coalescer-sched-crash.ll b/llvm/test/CodeGen/AMDGPU/reg-coalescer-sched-crash.ll index fb39653..ceec966 100644 --- a/llvm/test/CodeGen/AMDGPU/reg-coalescer-sched-crash.ll +++ b/llvm/test/CodeGen/AMDGPU/reg-coalescer-sched-crash.ll @@ -6,7 +6,7 @@ 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, i1 %c0) #1 { +define amdgpu_kernel void @reg_coalescer_breaks_dead(ptr addrspace(1) nocapture readonly %arg, i32 %arg1, i32 %arg2, i32 %arg3) #1 { bb: %id.x = call i32 @llvm.amdgcn.workitem.id.x() %cmp0 = icmp eq i32 %id.x, 0 @@ -18,7 +18,7 @@ bb3: ; preds = %bb bb4: ; preds = %bb6, %bb %tmp5 = phi <2 x i32> [ zeroinitializer, %bb ], [ %tmp13, %bb6 ] - br i1 %c0, label %bb15, label %bb16 + br i1 undef, label %bb15, label %bb16 bb6: ; preds = %bb6, %bb3 %tmp7 = phi <2 x i32> [ zeroinitializer, %bb3 ], [ %tmp13, %bb6 ] diff --git a/llvm/test/CodeGen/AMDGPU/rename-disconnected-bug.ll b/llvm/test/CodeGen/AMDGPU/rename-disconnected-bug.ll index a6867d4..9839b71 100644 --- a/llvm/test/CodeGen/AMDGPU/rename-disconnected-bug.ll +++ b/llvm/test/CodeGen/AMDGPU/rename-disconnected-bug.ll @@ -3,23 +3,23 @@ ; definition on every path (there should at least be IMPLICIT_DEF instructions). target triple = "amdgcn--" -define amdgpu_kernel void @func(i1 %c0, i1 %c1, i1 %c2) { +define amdgpu_kernel void @func() { B0: - br i1 %c0, label %B1, label %B2 + br i1 undef, label %B1, label %B2 B1: br label %B2 B2: %v0 = phi <4 x float> [ zeroinitializer, %B1 ], [ , %B0 ] - br i1 %c1, label %B20.1, label %B20.2 + br i1 undef, 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 %c2, label %B30.1, label %B30.2 + br i1 undef, label %B30.1, label %B30.2 B30.1: %sub = fsub <4 x float> %v2, undef diff --git a/llvm/test/CodeGen/AMDGPU/si-annotate-cf-unreachable.ll b/llvm/test/CodeGen/AMDGPU/si-annotate-cf-unreachable.ll index db8b09d..f8fde0a 100644 --- a/llvm/test/CodeGen/AMDGPU/si-annotate-cf-unreachable.ll +++ b/llvm/test/CodeGen/AMDGPU/si-annotate-cf-unreachable.ll @@ -11,7 +11,7 @@ ; 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, i1 %c0) #0 { +define amdgpu_kernel void @annotate_unreachable(ptr addrspace(1) noalias nocapture readonly %arg) #0 { bb: %tmp = tail call i32 @llvm.amdgcn.workitem.id.x() br label %bb1 @@ -20,7 +20,7 @@ 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 %c0, label %bb3, label %bb5 ; label order reversed + br i1 undef, label %bb3, label %bb5 ; label order reversed bb3: ; preds = %bb1 %tmp6 = extractelement <4 x float> %tmp4, i32 2 diff --git a/llvm/test/CodeGen/AMDGPU/simplifydemandedbits-recursion.ll b/llvm/test/CodeGen/AMDGPU/simplifydemandedbits-recursion.ll index 0b22625..ed07486 100644 --- a/llvm/test/CodeGen/AMDGPU/simplifydemandedbits-recursion.ll +++ b/llvm/test/CodeGen/AMDGPU/simplifydemandedbits-recursion.ll @@ -19,7 +19,7 @@ declare i32 @llvm.amdgcn.workitem.id.x() #0 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, i1 %c0, i1 %c1, i1 %c2, i1 %c3, i1 %c4, i1 %c5) 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) 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() @@ -31,7 +31,7 @@ bb: br label %bb12 bb11: ; preds = %bb30 - br i1 %c0, label %bb37, label %bb38 + br i1 undef, label %bb37, label %bb38 bb12: ; preds = %bb30, %bb br i1 false, label %.preheader, label %.loopexit145 @@ -43,7 +43,7 @@ bb13: ; preds = %.loopexit, %.loopex %tmp14 = phi i32 [ %tmp5, %.loopexit145 ], [ %tmp20, %.loopexit ] %tmp15 = add nsw i32 %tmp14, -3 %tmp16 = mul i32 %tmp14, 21 - br i1 %c1, label %bb17, label %.loopexit + br i1 undef, label %bb17, label %.loopexit bb17: ; preds = %bb13 %tmp18 = mul i32 %tmp15, 224 @@ -52,7 +52,7 @@ bb17: ; preds = %bb13 .loopexit: ; preds = %bb21, %bb13 %tmp20 = add nuw nsw i32 %tmp14, 16 - br i1 %c2, label %bb13, label %bb26 + br i1 undef, label %bb13, label %bb26 bb21: ; preds = %bb21, %bb17 %tmp22 = phi i32 [ %tmp4, %bb17 ], [ %tmp25, %bb21 ] @@ -60,7 +60,7 @@ bb21: ; preds = %bb21, %bb17 %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 %c3, label %bb21, label %.loopexit + br i1 undef, label %bb21, label %.loopexit bb26: ; preds = %.loopexit br label %bb31 @@ -72,7 +72,7 @@ bb26: ; preds = %.loopexit br i1 %tmp29, label %.preheader, label %.loopexit145 bb30: ; preds = %bb31 - br i1 %c4, label %bb11, label %bb12 + br i1 undef, label %bb11, label %bb12 bb31: ; preds = %bb31, %bb26 %tmp32 = phi i32 [ %tmp9, %bb26 ], [ undef, %bb31 ] @@ -80,7 +80,7 @@ bb31: ; preds = %bb31, %bb26 %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 %c5, label %bb30, label %bb31 + br i1 undef, label %bb30, label %bb31 bb37: ; preds = %bb11 br label %bb38 diff --git a/llvm/test/CodeGen/AMDGPU/switch-default-block-unreachable.ll b/llvm/test/CodeGen/AMDGPU/switch-default-block-unreachable.ll index a5e1423..ad22c1b 100644 --- a/llvm/test/CodeGen/AMDGPU/switch-default-block-unreachable.ll +++ b/llvm/test/CodeGen/AMDGPU/switch-default-block-unreachable.ll @@ -1,5 +1,5 @@ ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -verify-machineinstrs -stop-after=amdgpu-isel -o - %s | FileCheck -check-prefix=GCN %s -define void @test(i1 %c0) #1 { +define void @test() #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 @@ -22,7 +22,7 @@ define void @test(i1 %c0) #1 { ; GCN: bb.{{[0-9]+}}.UnifiedReturnBlock: entry: %idx = tail call i32 @llvm.amdgcn.workitem.id.x() #0 - br i1 %c0, label %entry.true.blk, label %entry.false.blk + br i1 undef, label %entry.true.blk, label %entry.false.blk entry.true.blk: ; preds = %entry %exit.cmp = icmp ult i32 %idx, 3 diff --git a/llvm/test/CodeGen/AMDGPU/uniform-crash.ll b/llvm/test/CodeGen/AMDGPU/uniform-crash.ll index c57f346..c00ef9a 100644 --- a/llvm/test/CodeGen/AMDGPU/uniform-crash.ll +++ b/llvm/test/CodeGen/AMDGPU/uniform-crash.ll @@ -25,7 +25,7 @@ ENDIF: ; preds = %IF, %main_body ; 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, i1 %c0) { +define amdgpu_kernel void @fix_sgpr_live_ranges_crash(i32 %arg, i32 %arg1) { bb: %cnd = trunc i32 %arg to i1 br i1 %cnd, label %bb2, label %bb5 @@ -45,7 +45,7 @@ bb5: ; preds = %bb3, %bb br i1 %tmp10, label %bb11, label %bb12 bb11: ; preds = %bb11, %bb5 - br i1 %c0, label %bb11, label %bb12 + br i1 undef, label %bb11, label %bb12 bb12: ; preds = %bb11, %bb5 ret void diff --git a/llvm/test/CodeGen/AMDGPU/v1024.ll b/llvm/test/CodeGen/AMDGPU/v1024.ll index 5f3d926..03a3685 100644 --- a/llvm/test/CodeGen/AMDGPU/v1024.ll +++ b/llvm/test/CodeGen/AMDGPU/v1024.ll @@ -6,11 +6,11 @@ ; GCN-NOT: v_accvgpr ; GCN-COUNT-8: global_store_dwordx4 ; GCN-NOT: v_accvgpr -define amdgpu_kernel void @test_v1024(i1 %c0) { +define amdgpu_kernel void @test_v1024() { 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 %c0, label %if.then.i.i, label %if.else.i + br i1 undef, 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) diff --git a/llvm/test/CodeGen/AMDGPU/wave32.ll b/llvm/test/CodeGen/AMDGPU/wave32.ll index 8607aa6..0d4144a 100644 --- a/llvm/test/CodeGen/AMDGPU/wave32.ll +++ b/llvm/test/CodeGen/AMDGPU/wave32.ll @@ -2290,7 +2290,7 @@ define amdgpu_ps void @test_wqm_vote(float %a) { ret void } -define amdgpu_kernel void @test_branch_true(i1 %c0) #2 { +define amdgpu_kernel void @test_branch_true() #2 { ; GFX1032-LABEL: test_branch_true: ; GFX1032: ; %bb.0: ; %entry ; GFX1032-NEXT: s_mov_b32 vcc_lo, exec_lo @@ -2327,7 +2327,7 @@ for.body.lr.ph: ; preds = %entry br label %for.body for.body: ; preds = %for.body, %for.body.lr.ph - br i1 %c0, label %for.end, label %for.body + br i1 undef, label %for.end, label %for.body for.end: ; preds = %for.body, %entry ret void -- 2.7.4