Skip to content

AMDGPU: Replace some test undef uses with poison #131103

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
arsenm merged 1 commit into from
Mar 14, 2025
Merged

AMDGPU: Replace some test undef uses with poison #131103

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions llvm/test/CodeGen/AMDGPU/GlobalISel/call-outgoing-stack-args.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -53,7 +53,7 @@ define amdgpu_kernel void @kernel_caller_stack() {
; FLATSCR-NEXT: scratch_store_dword off, v0, s2
; FLATSCR-NEXT: s_swappc_b64 s[30:31], s[0:1]
; FLATSCR-NEXT: s_endpgm
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>)
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>)
ret void
}

Expand Down Expand Up @@ -294,7 +294,7 @@ define void @func_caller_stack() {
; FLATSCR-NEXT: s_mov_b32 s33, s0
; FLATSCR-NEXT: s_waitcnt vmcnt(0)
; FLATSCR-NEXT: s_setpc_b64 s[30:31]
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>)
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>)
ret void
}

Expand Down
2 changes: 1 addition & 1 deletion llvm/test/CodeGen/AMDGPU/GlobalISel/divergent-control-flow.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -231,7 +231,7 @@ bb:
br label %bb1

bb1:
%lsr.iv = phi i32 [ undef, %bb ], [ %lsr.iv.next, %bb4 ]
%lsr.iv = phi i32 [ poison, %bb ], [ %lsr.iv.next, %bb4 ]
%lsr.iv.next = add i32 %lsr.iv, 1
%cmp0 = icmp slt i32 %lsr.iv.next, 0
br i1 %cmp0, label %bb4, label %bb9
Expand Down
2 changes: 1 addition & 1 deletion llvm/test/CodeGen/AMDGPU/diverge-interp-mov-lower.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,7 +21,7 @@ define dllexport amdgpu_ps void @_amdgpu_ps_main(i32 inreg %arg) local_unnamed_a
%tmp6 = load <4 x float>, ptr addrspace(4) %tmp5, align 16
%tmp7 = extractelement <4 x float> %tmp6, i32 3
%tmp8 = call <2 x half> @llvm.amdgcn.cvt.pkrtz(float poison, float %tmp7) #1
call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 15, <2 x half> undef, <2 x half> %tmp8, i1 true, i1 true) #2
call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 15, <2 x half> poison, <2 x half> %tmp8, i1 true, i1 true) #2
ret void
}

Expand Down
2 changes: 1 addition & 1 deletion llvm/test/CodeGen/AMDGPU/llvm.dbg.value.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,7 +29,7 @@ entry:
; OPT: s_endpgm
define amdgpu_kernel void @only_undef_dbg_value() #1 {
bb:
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
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
ret void, !dbg !14
}

Expand Down
18 changes: 9 additions & 9 deletions llvm/test/CodeGen/AMDGPU/mfma-bf16-vgpr-cd-select.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,7 +33,7 @@ bb:
define amdgpu_kernel void @test_mfma_f32_16x16x2bf16(ptr addrspace(1) %arg) #0 {
bb:
%in.1 = load <16 x float>, ptr addrspace(1) %arg
%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)
%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)
store <16 x float> %mai.1, ptr addrspace(1) %arg
ret void
}
Expand All @@ -43,7 +43,7 @@ bb:
define amdgpu_kernel void @test_mfma_f32_4x4x2bf16(ptr addrspace(1) %arg) #0 {
bb:
%in.1 = load <4 x float>, ptr addrspace(1) %arg
%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)
%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)
store <4 x float> %mai.1, ptr addrspace(1) %arg
ret void
}
Expand All @@ -53,7 +53,7 @@ bb:
define amdgpu_kernel void @test_mfma_f32_32x32x4bf16(ptr addrspace(1) %arg) #0 {
bb:
%in.1 = load <16 x float>, ptr addrspace(1) %arg
%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)
%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)
store <16 x float> %mai.1, ptr addrspace(1) %arg
ret void
}
Expand All @@ -63,7 +63,7 @@ bb:
define amdgpu_kernel void @test_mfma_f32_16x16x8bf16(ptr addrspace(1) %arg) #0 {
bb:
%in.1 = load <4 x float>, ptr addrspace(1) %arg
%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)
%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)
store <4 x float> %mai.1, ptr addrspace(1) %arg
ret void
}
Expand All @@ -73,7 +73,7 @@ bb:
define amdgpu_kernel void @test_mfma_f32_32x32x4bf16_1k(ptr addrspace(1) %arg) #0 {
bb:
%in.1 = load <32 x float>, ptr addrspace(1) %arg
%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)
%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)
store <32 x float> %mai.1, ptr addrspace(1) %arg
ret void
}
Expand All @@ -83,7 +83,7 @@ bb:
define amdgpu_kernel void @test_mfma_f32_16x16x4bf16_1k(ptr addrspace(1) %arg) #0 {
bb:
%in.1 = load <16 x float>, ptr addrspace(1) %arg
%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)
%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)
store <16 x float> %mai.1, ptr addrspace(1) %arg
ret void
}
Expand All @@ -93,7 +93,7 @@ bb:
define amdgpu_kernel void @test_mfma_f32_4x4x4bf16_1k(ptr addrspace(1) %arg) #0 {
bb:
%in.1 = load <4 x float>, ptr addrspace(1) %arg
%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)
%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)
store <4 x float> %mai.1, ptr addrspace(1) %arg
ret void
}
Expand All @@ -103,7 +103,7 @@ bb:
define amdgpu_kernel void @test_mfma_f32_32x32x8bf16_1k(ptr addrspace(1) %arg) #0 {
bb:
%in.1 = load <16 x float>, ptr addrspace(1) %arg
%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)
%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)
store <16 x float> %mai.1, ptr addrspace(1) %arg
ret void
}
Expand All @@ -113,7 +113,7 @@ bb:
define amdgpu_kernel void @test_mfma_f32_16x16x16bf16_1k(ptr addrspace(1) %arg) #0 {
bb:
%in.1 = load <4 x float>, ptr addrspace(1) %arg
%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)
%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)
store <4 x float> %mai.1, ptr addrspace(1) %arg
ret void
}
Expand Down
2 changes: 1 addition & 1 deletion llvm/test/CodeGen/AMDGPU/mfma-cd-select.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,7 +42,7 @@ 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_inline_asm_phys_agpr(ptr addrspace(1) %arg) {
bb:
call void asm sideeffect "; use $0", "{a[100:131]}"(<32 x float> undef)
call void asm sideeffect "; use $0", "{a[100:131]}"(<32 x float> poison)
%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 0, i32 0, i32 0)
store <32 x float> %mai.1, ptr addrspace(1) %arg
Expand Down
10 changes: 5 additions & 5 deletions llvm/test/CodeGen/AMDGPU/mfma-vgpr-cd-select.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -72,7 +72,7 @@ bb:
define amdgpu_kernel void @test_mfma_f32_32x32x4f16(ptr addrspace(1) %arg) #0 {
bb:
%in.1 = load <32 x float>, ptr addrspace(1) %arg
%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)
%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)
store <32 x float> %mai.1, ptr addrspace(1) %arg
ret void
}
Expand All @@ -82,7 +82,7 @@ bb:
define amdgpu_kernel void @test_mfma_f32_16x16x4f16(ptr addrspace(1) %arg) #0 {
bb:
%in.1 = load <16 x float>, ptr addrspace(1) %arg
%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)
%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)
store <16 x float> %mai.1, ptr addrspace(1) %arg
ret void
}
Expand All @@ -92,7 +92,7 @@ bb:
define amdgpu_kernel void @test_mfma_f32_4x4x4f16(ptr addrspace(1) %arg) #0 {
bb:
%in.1 = load <4 x float>, ptr addrspace(1) %arg
%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)
%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)
store <4 x float> %mai.1, ptr addrspace(1) %arg
ret void
}
Expand All @@ -102,7 +102,7 @@ bb:
define amdgpu_kernel void @test_mfma_f32_32x32x8f16(ptr addrspace(1) %arg) #0 {
bb:
%in.1 = load <16 x float>, ptr addrspace(1) %arg
%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)
%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)
store <16 x float> %mai.1, ptr addrspace(1) %arg
ret void
}
Expand All @@ -112,7 +112,7 @@ bb:
define amdgpu_kernel void @test_mfma_f32_16x16x16f16(ptr addrspace(1) %arg) #0 {
bb:
%in.1 = load <4 x float>, ptr addrspace(1) %arg
%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)
%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)
store <4 x float> %mai.1, ptr addrspace(1) %arg
ret void
}
Expand Down
4 changes: 2 additions & 2 deletions llvm/test/CodeGen/AMDGPU/ret_jump.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,7 +50,7 @@ unreachable.bb: ; preds = %else
unreachable

ret.bb: ; preds = %else, %main_body
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
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
}

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

ret.bb: ; preds = %else, %main_body
store volatile i32 11, ptr addrspace(1) poison
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
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
}

; Function Attrs: nounwind readnone
Expand Down
4 changes: 2 additions & 2 deletions llvm/test/CodeGen/AMDGPU/sgpr-copy.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -371,7 +371,7 @@ bb:
%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)
%tmp10 = extractelement <4 x float> %tmp, i32 0
%tmp12 = call <2 x half> @llvm.amdgcn.cvt.pkrtz(float poison, float %tmp10)
call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 15, <2 x half> %tmp12, <2 x half> undef, i1 true, i1 true) #0
call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 15, <2 x half> %tmp12, <2 x half> poison, i1 true, i1 true) #0
ret void
}

Expand All @@ -386,7 +386,7 @@ bb:
%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)
%tmp10 = extractelement <4 x float> %tmp, i32 0
%tmp12 = call <2 x half> @llvm.amdgcn.cvt.pkrtz(float %tmp10, float poison)
call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 15, <2 x half> %tmp12, <2 x half> undef, i1 true, i1 true) #0
call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 15, <2 x half> %tmp12, <2 x half> poison, i1 true, i1 true) #0
ret void
}

Expand Down
2 changes: 1 addition & 1 deletion llvm/test/CodeGen/AMDGPU/split-smrd.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,7 @@ bb3: ; preds = %bb
%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)
%tmp10 = extractelement <4 x float> %tmp9, i32 0
%tmp12 = call <2 x half> @llvm.amdgcn.cvt.pkrtz(float %tmp10, float poison)
call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 15, <2 x half> %tmp12, <2 x half> undef, i1 true, i1 true) #0
call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 15, <2 x half> %tmp12, <2 x half> poison, i1 true, i1 true) #0
ret void
}

Expand Down
4 changes: 2 additions & 2 deletions llvm/test/CodeGen/AMDGPU/v1024.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,14 +13,14 @@ entry:
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)
call void @llvm.memcpy.p5.p5.i64(ptr addrspace(5) align 16 %alloca, ptr addrspace(5) align 4 poison, i64 128, i1 false)
br label %if.then.i62.i

if.else.i: ; preds = %entry
br label %if.then.i62.i

if.then.i62.i: ; preds = %if.else.i, %if.then.i.i
call void @llvm.memcpy.p1.p5.i64(ptr addrspace(1) align 4 undef, ptr addrspace(5) align 16 %alloca, i64 128, i1 false)
call void @llvm.memcpy.p1.p5.i64(ptr addrspace(1) align 4 poison, ptr addrspace(5) align 16 %alloca, i64 128, i1 false)
ret void
}

Expand Down
2 changes: 1 addition & 1 deletion llvm/test/CodeGen/AMDGPU/wmma_modifiers.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@ define amdgpu_cs void @xyz () {
loop:
%ld = load <8 x float>, ptr addrspace(5) null, align 32
%in_shuffle = shufflevector <8 x float> %ld, <8 x float> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 3>
%wmma = call <4 x float> @llvm.amdgcn.wmma.f32.16x16x16.f16.v4f32.v16f16(<16 x half> undef, <16 x half> undef, <4 x float> %in_shuffle)
%wmma = call <4 x float> @llvm.amdgcn.wmma.f32.16x16x16.f16.v4f32.v16f16(<16 x half> poison, <16 x half> poison, <4 x float> %in_shuffle)
%out_shuffle = shufflevector <4 x float> %wmma, <4 x float> poison, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 poison, i32 poison, i32 poison, i32 poison>
store <8 x float> %out_shuffle, ptr addrspace(5) null, align 32
br i1 false, label %.exit, label %loop
Expand Down
4 changes: 2 additions & 2 deletions llvm/test/CodeGen/AMDGPU/wqm.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -209,7 +209,7 @@ define amdgpu_ps <4 x float> @test4(<8 x i32> inreg %rsrc, <4 x i32> inreg %samp
main_body:
%c.1 = mul i32 %c, %d

call void @llvm.amdgcn.struct.buffer.store.v4f32(<4 x float> undef, <4 x i32> poison, i32 %c.1, i32 0, i32 0, i32 0)
call void @llvm.amdgcn.struct.buffer.store.v4f32(<4 x float> poison, <4 x i32> poison, i32 %c.1, i32 0, i32 0, i32 0)
%c.1.bc = bitcast i32 %c.1 to float
%tex = call <4 x float> @llvm.amdgcn.image.sample.1d.v4f32.f32(i32 15, float %c.1.bc, <8 x i32> %rsrc, <4 x i32> %sampler, i1 false, i32 0, i32 0) #0
%tex0 = extractelement <4 x float> %tex, i32 0
Expand Down Expand Up @@ -247,7 +247,7 @@ define amdgpu_ps <4 x float> @test4_ptr_buf(<8 x i32> inreg %rsrc, <4 x i32> inr
main_body:
%c.1 = mul i32 %c, %d

call void @llvm.amdgcn.struct.ptr.buffer.store.v4f32(<4 x float> undef, ptr addrspace(8) poison, i32 %c.1, i32 0, i32 0, i32 0)
call void @llvm.amdgcn.struct.ptr.buffer.store.v4f32(<4 x float> poison, ptr addrspace(8) poison, i32 %c.1, i32 0, i32 0, i32 0)
%c.1.bc = bitcast i32 %c.1 to float
%tex = call <4 x float> @llvm.amdgcn.image.sample.1d.v4f32.f32(i32 15, float %c.1.bc, <8 x i32> %rsrc, <4 x i32> %sampler, i1 false, i32 0, i32 0) #0
%tex0 = extractelement <4 x float> %tex, i32 0
Expand Down
Loading