Skip to content

Commit fd4a1d8

Browse files
arsenmarsenm
committed
AMDGPU: Replace some test undef uses with poison
1 parent 7251260 commit fd4a1d8

13 files changed

+30
-30
lines changed

‎llvm/test/CodeGen/AMDGPU/GlobalISel/call-outgoing-stack-args.ll‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -53,7 +53,7 @@ define amdgpu_kernel void @kernel_caller_stack() {
5353
; FLATSCR-NEXT: scratch_store_dword off, v0, s2
5454
; FLATSCR-NEXT: s_swappc_b64 s[30:31], s[0:1]
5555
; FLATSCR-NEXT: s_endpgm
56-
call void @external_void_func_v16i32_v16i32_v4i32(<16 x i32> undef, <16 x i32> undef, <4 x i32> <i32 9, i32 10, i32 11, i32 12>)
56+
call void @external_void_func_v16i32_v16i32_v4i32(<16 x i32> poison, <16 x i32> poison, <4 x i32> <i32 9, i32 10, i32 11, i32 12>)
5757
ret void
5858
}
5959

@@ -294,7 +294,7 @@ define void @func_caller_stack() {
294294
; FLATSCR-NEXT: s_mov_b32 s33, s0
295295
; FLATSCR-NEXT: s_waitcnt vmcnt(0)
296296
; FLATSCR-NEXT: s_setpc_b64 s[30:31]
297-
call void @external_void_func_v16i32_v16i32_v4i32(<16 x i32> undef, <16 x i32> undef, <4 x i32> <i32 9, i32 10, i32 11, i32 12>)
297+
call void @external_void_func_v16i32_v16i32_v4i32(<16 x i32> poison, <16 x i32> poison, <4 x i32> <i32 9, i32 10, i32 11, i32 12>)
298298
ret void
299299
}
300300

‎llvm/test/CodeGen/AMDGPU/GlobalISel/divergent-control-flow.ll‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -231,7 +231,7 @@ bb:
231231
br label %bb1
232232

233233
bb1:
234-
%lsr.iv = phi i32 [ undef, %bb ], [ %lsr.iv.next, %bb4 ]
234+
%lsr.iv = phi i32 [ poison, %bb ], [ %lsr.iv.next, %bb4 ]
235235
%lsr.iv.next = add i32 %lsr.iv, 1
236236
%cmp0 = icmp slt i32 %lsr.iv.next, 0
237237
br i1 %cmp0, label %bb4, label %bb9

‎llvm/test/CodeGen/AMDGPU/diverge-interp-mov-lower.ll‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -21,7 +21,7 @@ define dllexport amdgpu_ps void @_amdgpu_ps_main(i32 inreg %arg) local_unnamed_a
2121
%tmp6 = load <4 x float>, ptr addrspace(4) %tmp5, align 16
2222
%tmp7 = extractelement <4 x float> %tmp6, i32 3
2323
%tmp8 = call <2 x half> @llvm.amdgcn.cvt.pkrtz(float poison, float %tmp7) #1
24-
call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 15, <2 x half> undef, <2 x half> %tmp8, i1 true, i1 true) #2
24+
call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 15, <2 x half> poison, <2 x half> %tmp8, i1 true, i1 true) #2
2525
ret void
2626
}
2727

‎llvm/test/CodeGen/AMDGPU/llvm.dbg.value.ll‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -29,7 +29,7 @@ entry:
2929
; OPT: s_endpgm
3030
define amdgpu_kernel void @only_undef_dbg_value() #1 {
3131
bb:
32-
call void @llvm.dbg.value(metadata <4 x float> undef, metadata !10, metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)) #2, !dbg !14
32+
call void @llvm.dbg.value(metadata <4 x float> poison, metadata !10, metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)) #2, !dbg !14
3333
ret void, !dbg !14
3434
}
3535

‎llvm/test/CodeGen/AMDGPU/mfma-bf16-vgpr-cd-select.ll‎

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -33,7 +33,7 @@ bb:
3333
define amdgpu_kernel void @test_mfma_f32_16x16x2bf16(ptr addrspace(1) %arg) #0 {
3434
bb:
3535
%in.1 = load <16 x float>, ptr addrspace(1) %arg
36-
%mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x2bf16(<2 x i16> undef, <2 x i16> undef, <16 x float> %in.1, i32 0, i32 0, i32 0)
36+
%mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x2bf16(<2 x i16> poison, <2 x i16> poison, <16 x float> %in.1, i32 0, i32 0, i32 0)
3737
store <16 x float> %mai.1, ptr addrspace(1) %arg
3838
ret void
3939
}
@@ -43,7 +43,7 @@ bb:
4343
define amdgpu_kernel void @test_mfma_f32_4x4x2bf16(ptr addrspace(1) %arg) #0 {
4444
bb:
4545
%in.1 = load <4 x float>, ptr addrspace(1) %arg
46-
%mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> undef, <2 x i16> undef, <4 x float> %in.1, i32 0, i32 0, i32 0)
46+
%mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> poison, <2 x i16> poison, <4 x float> %in.1, i32 0, i32 0, i32 0)
4747
store <4 x float> %mai.1, ptr addrspace(1) %arg
4848
ret void
4949
}
@@ -53,7 +53,7 @@ bb:
5353
define amdgpu_kernel void @test_mfma_f32_32x32x4bf16(ptr addrspace(1) %arg) #0 {
5454
bb:
5555
%in.1 = load <16 x float>, ptr addrspace(1) %arg
56-
%mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x4bf16(<2 x i16> undef, <2 x i16> undef, <16 x float> %in.1, i32 0, i32 0, i32 0)
56+
%mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x4bf16(<2 x i16> poison, <2 x i16> poison, <16 x float> %in.1, i32 0, i32 0, i32 0)
5757
store <16 x float> %mai.1, ptr addrspace(1) %arg
5858
ret void
5959
}
@@ -63,7 +63,7 @@ bb:
6363
define amdgpu_kernel void @test_mfma_f32_16x16x8bf16(ptr addrspace(1) %arg) #0 {
6464
bb:
6565
%in.1 = load <4 x float>, ptr addrspace(1) %arg
66-
%mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x8bf16(<2 x i16> undef, <2 x i16> undef, <4 x float> %in.1, i32 0, i32 0, i32 0)
66+
%mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x8bf16(<2 x i16> poison, <2 x i16> poison, <4 x float> %in.1, i32 0, i32 0, i32 0)
6767
store <4 x float> %mai.1, ptr addrspace(1) %arg
6868
ret void
6969
}
@@ -73,7 +73,7 @@ bb:
7373
define amdgpu_kernel void @test_mfma_f32_32x32x4bf16_1k(ptr addrspace(1) %arg) #0 {
7474
bb:
7575
%in.1 = load <32 x float>, ptr addrspace(1) %arg
76-
%mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x4bf16.1k(<4 x i16> undef, <4 x i16> undef, <32 x float> %in.1, i32 0, i32 0, i32 0)
76+
%mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x4bf16.1k(<4 x i16> poison, <4 x i16> poison, <32 x float> %in.1, i32 0, i32 0, i32 0)
7777
store <32 x float> %mai.1, ptr addrspace(1) %arg
7878
ret void
7979
}
@@ -83,7 +83,7 @@ bb:
8383
define amdgpu_kernel void @test_mfma_f32_16x16x4bf16_1k(ptr addrspace(1) %arg) #0 {
8484
bb:
8585
%in.1 = load <16 x float>, ptr addrspace(1) %arg
86-
%mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x4bf16.1k(<4 x i16> undef, <4 x i16> undef, <16 x float> %in.1, i32 0, i32 0, i32 0)
86+
%mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x4bf16.1k(<4 x i16> poison, <4 x i16> poison, <16 x float> %in.1, i32 0, i32 0, i32 0)
8787
store <16 x float> %mai.1, ptr addrspace(1) %arg
8888
ret void
8989
}
@@ -93,7 +93,7 @@ bb:
9393
define amdgpu_kernel void @test_mfma_f32_4x4x4bf16_1k(ptr addrspace(1) %arg) #0 {
9494
bb:
9595
%in.1 = load <4 x float>, ptr addrspace(1) %arg
96-
%mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x4bf16.1k(<4 x i16> undef, <4 x i16> undef, <4 x float> %in.1, i32 0, i32 0, i32 0)
96+
%mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x4bf16.1k(<4 x i16> poison, <4 x i16> poison, <4 x float> %in.1, i32 0, i32 0, i32 0)
9797
store <4 x float> %mai.1, ptr addrspace(1) %arg
9898
ret void
9999
}
@@ -103,7 +103,7 @@ bb:
103103
define amdgpu_kernel void @test_mfma_f32_32x32x8bf16_1k(ptr addrspace(1) %arg) #0 {
104104
bb:
105105
%in.1 = load <16 x float>, ptr addrspace(1) %arg
106-
%mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8bf16.1k(<4 x i16> undef, <4 x i16> undef, <16 x float> %in.1, i32 0, i32 0, i32 0)
106+
%mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8bf16.1k(<4 x i16> poison, <4 x i16> poison, <16 x float> %in.1, i32 0, i32 0, i32 0)
107107
store <16 x float> %mai.1, ptr addrspace(1) %arg
108108
ret void
109109
}
@@ -113,7 +113,7 @@ bb:
113113
define amdgpu_kernel void @test_mfma_f32_16x16x16bf16_1k(ptr addrspace(1) %arg) #0 {
114114
bb:
115115
%in.1 = load <4 x float>, ptr addrspace(1) %arg
116-
%mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> undef, <4 x i16> undef, <4 x float> %in.1, i32 0, i32 0, i32 0)
116+
%mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16bf16.1k(<4 x i16> poison, <4 x i16> poison, <4 x float> %in.1, i32 0, i32 0, i32 0)
117117
store <4 x float> %mai.1, ptr addrspace(1) %arg
118118
ret void
119119
}

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -42,7 +42,7 @@ bb:
4242
; GCN: v_mfma_f32_32x32x1{{.*}} a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9:]+}}, a[{{[0-9:]+}}]
4343
define amdgpu_kernel void @test_mfma_f32_32x32x1f32_inline_asm_phys_agpr(ptr addrspace(1) %arg) {
4444
bb:
45-
call void asm sideeffect "; use $0", "{a[100:131]}"(<32 x float> undef)
45+
call void asm sideeffect "; use $0", "{a[100:131]}"(<32 x float> poison)
4646
%in.1 = load <32 x float>, ptr addrspace(1) %arg
4747
%mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %in.1, i32 0, i32 0, i32 0)
4848
store <32 x float> %mai.1, ptr addrspace(1) %arg

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

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -72,7 +72,7 @@ bb:
7272
define amdgpu_kernel void @test_mfma_f32_32x32x4f16(ptr addrspace(1) %arg) #0 {
7373
bb:
7474
%in.1 = load <32 x float>, ptr addrspace(1) %arg
75-
%mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x4f16(<4 x half> undef, <4 x half> undef, <32 x float> %in.1, i32 0, i32 0, i32 0)
75+
%mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x4f16(<4 x half> poison, <4 x half> poison, <32 x float> %in.1, i32 0, i32 0, i32 0)
7676
store <32 x float> %mai.1, ptr addrspace(1) %arg
7777
ret void
7878
}
@@ -82,7 +82,7 @@ bb:
8282
define amdgpu_kernel void @test_mfma_f32_16x16x4f16(ptr addrspace(1) %arg) #0 {
8383
bb:
8484
%in.1 = load <16 x float>, ptr addrspace(1) %arg
85-
%mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x4f16(<4 x half> undef, <4 x half> undef, <16 x float> %in.1, i32 0, i32 0, i32 0)
85+
%mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x4f16(<4 x half> poison, <4 x half> poison, <16 x float> %in.1, i32 0, i32 0, i32 0)
8686
store <16 x float> %mai.1, ptr addrspace(1) %arg
8787
ret void
8888
}
@@ -92,7 +92,7 @@ bb:
9292
define amdgpu_kernel void @test_mfma_f32_4x4x4f16(ptr addrspace(1) %arg) #0 {
9393
bb:
9494
%in.1 = load <4 x float>, ptr addrspace(1) %arg
95-
%mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x4f16(<4 x half> undef, <4 x half> undef, <4 x float> %in.1, i32 0, i32 0, i32 0)
95+
%mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x4f16(<4 x half> poison, <4 x half> poison, <4 x float> %in.1, i32 0, i32 0, i32 0)
9696
store <4 x float> %mai.1, ptr addrspace(1) %arg
9797
ret void
9898
}
@@ -102,7 +102,7 @@ bb:
102102
define amdgpu_kernel void @test_mfma_f32_32x32x8f16(ptr addrspace(1) %arg) #0 {
103103
bb:
104104
%in.1 = load <16 x float>, ptr addrspace(1) %arg
105-
%mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> undef, <4 x half> undef, <16 x float> %in.1, i32 0, i32 0, i32 0)
105+
%mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> poison, <4 x half> poison, <16 x float> %in.1, i32 0, i32 0, i32 0)
106106
store <16 x float> %mai.1, ptr addrspace(1) %arg
107107
ret void
108108
}
@@ -112,7 +112,7 @@ bb:
112112
define amdgpu_kernel void @test_mfma_f32_16x16x16f16(ptr addrspace(1) %arg) #0 {
113113
bb:
114114
%in.1 = load <4 x float>, ptr addrspace(1) %arg
115-
%mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> undef, <4 x half> undef, <4 x float> %in.1, i32 0, i32 0, i32 0)
115+
%mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> poison, <4 x half> poison, <4 x float> %in.1, i32 0, i32 0, i32 0)
116116
store <4 x float> %mai.1, ptr addrspace(1) %arg
117117
ret void
118118
}

‎llvm/test/CodeGen/AMDGPU/ret_jump.ll‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -50,7 +50,7 @@ unreachable.bb: ; preds = %else
5050
unreachable
5151

5252
ret.bb: ; preds = %else, %main_body
53-
ret <{ i32, i32, i32, i32, i32, i32, i32, i32, i32, float, float, float, float, float, float, float, float, float, float, float, float, float, float }> undef
53+
ret <{ i32, i32, i32, i32, i32, i32, i32, i32, i32, float, float, float, float, float, float, float, float, float, float, float, float, float, float }> poison
5454
}
5555

5656
; GCN-LABEL: {{^}}uniform_br_nontrivial_ret_divergent_br_nontrivial_unreachable:
@@ -103,7 +103,7 @@ unreachable.bb: ; preds = %else
103103

104104
ret.bb: ; preds = %else, %main_body
105105
store volatile i32 11, ptr addrspace(1) poison
106-
ret <{ i32, i32, i32, i32, i32, i32, i32, i32, i32, float, float, float, float, float, float, float, float, float, float, float, float, float, float }> undef
106+
ret <{ i32, i32, i32, i32, i32, i32, i32, i32, i32, float, float, float, float, float, float, float, float, float, float, float, float, float, float }> poison
107107
}
108108

109109
; Function Attrs: nounwind readnone

‎llvm/test/CodeGen/AMDGPU/sgpr-copy.ll‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -371,7 +371,7 @@ bb:
371371
%tmp = call <4 x float> @llvm.amdgcn.image.sample.2d.v4f32.f32(i32 15, float 7.500000e-01, float 2.500000e-01, <8 x i32> %tmp8, <4 x i32> poison, i1 0, i32 0, i32 0)
372372
%tmp10 = extractelement <4 x float> %tmp, i32 0
373373
%tmp12 = call <2 x half> @llvm.amdgcn.cvt.pkrtz(float poison, float %tmp10)
374-
call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 15, <2 x half> %tmp12, <2 x half> undef, i1 true, i1 true) #0
374+
call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 15, <2 x half> %tmp12, <2 x half> poison, i1 true, i1 true) #0
375375
ret void
376376
}
377377

@@ -386,7 +386,7 @@ bb:
386386
%tmp = call <4 x float> @llvm.amdgcn.image.sample.2d.v4f32.f32(i32 15, float 7.500000e-01, float 2.500000e-01, <8 x i32> poison, <4 x i32> %tmp8, i1 0, i32 0, i32 0)
387387
%tmp10 = extractelement <4 x float> %tmp, i32 0
388388
%tmp12 = call <2 x half> @llvm.amdgcn.cvt.pkrtz(float %tmp10, float poison)
389-
call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 15, <2 x half> %tmp12, <2 x half> undef, i1 true, i1 true) #0
389+
call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 15, <2 x half> %tmp12, <2 x half> poison, i1 true, i1 true) #0
390390
ret void
391391
}
392392

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -24,7 +24,7 @@ bb3: ; preds = %bb
2424
%tmp9 = call <4 x float> @llvm.amdgcn.image.sample.2d.v4f32.f32(i32 15, float bitcast (i32 1061158912 to float), float bitcast (i32 1048576000 to float), <8 x i32> %tmp8, <4 x i32> poison, i1 0, i32 0, i32 0)
2525
%tmp10 = extractelement <4 x float> %tmp9, i32 0
2626
%tmp12 = call <2 x half> @llvm.amdgcn.cvt.pkrtz(float %tmp10, float poison)
27-
call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 15, <2 x half> %tmp12, <2 x half> undef, i1 true, i1 true) #0
27+
call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 15, <2 x half> %tmp12, <2 x half> poison, i1 true, i1 true) #0
2828
ret void
2929
}
3030

0 commit comments

Comments
 (0)