[llvm] e3dd47f - AMDGPU: Fix using deprecated buffer intrinsics in test
Matt Arsenault via llvm-commits
llvm-commits at lists.llvm.org
Tue Jan 18 16:02:53 PST 2022
Author: Matt Arsenault
Date: 2022-01-18T19:02:47-05:00
New Revision: e3dd47f987e71568e3dab0a779c0520388eb6cd9
URL: https://github.com/llvm/llvm-project/commit/e3dd47f987e71568e3dab0a779c0520388eb6cd9
DIFF: https://github.com/llvm/llvm-project/commit/e3dd47f987e71568e3dab0a779c0520388eb6cd9.diff
LOG: AMDGPU: Fix using deprecated buffer intrinsics in test
Added:
Modified:
llvm/test/CodeGen/AMDGPU/gfx90a-enc.ll
Removed:
################################################################################
diff --git a/llvm/test/CodeGen/AMDGPU/gfx90a-enc.ll b/llvm/test/CodeGen/AMDGPU/gfx90a-enc.ll
index 8a3119ada24a9..f8f805fb9e26d 100644
--- a/llvm/test/CodeGen/AMDGPU/gfx90a-enc.ll
+++ b/llvm/test/CodeGen/AMDGPU/gfx90a-enc.ll
@@ -7,16 +7,20 @@
; GFX90A-DAG: v_mfma_i32_4x4x4i8 a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9]+}}, a[{{[0-9:]+}}] ; encoding: [{{0x..,0x8.,}}
define amdgpu_kernel void @test(<4 x i32> %x) {
%id = tail call i32 @llvm.amdgcn.workitem.id.x()
- %r1 = tail call <4 x float> @llvm.amdgcn.buffer.load.format.v4f32(<4 x i32> %x, i32 %id, i32 0, i1 zeroext false, i1 zeroext false)
+ %r1 = tail call <4 x float> @llvm.amdgcn.struct.buffer.load.format.v4f32(<4 x i32> %x, i32 %id, i32 0, i32 0, i32 0)
store volatile <4 x float> %r1, <4 x float>* undef
- %r2 = tail call <4 x half> @llvm.amdgcn.buffer.load.format.v4f16(<4 x i32> %x, i32 %id, i32 0, i1 zeroext false, i1 zeroext false)
+ %r2 = tail call <4 x half> @llvm.amdgcn.struct.buffer.load.format.v4f16(<4 x i32> %x, i32 %id, i32 0, i32 0, i32 0)
store volatile <4 x half> %r2, <4 x half>* undef
%r3 = tail call <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32 1, i32 2, <4 x i32> %x, i32 0, i32 0, i32 0)
store <4 x i32> %r3, <4 x i32>* undef
ret void
}
-declare i32 @llvm.amdgcn.workitem.id.x()
-declare <4 x float> @llvm.amdgcn.buffer.load.format.v4f32(<4 x i32>, i32, i32, i1 immarg, i1 immarg)
-declare <4 x half> @llvm.amdgcn.buffer.load.format.v4f16(<4 x i32>, i32, i32, i1 immarg, i1 immarg)
-declare <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32, i32, <4 x i32>, i32, i32, i32)
+declare i32 @llvm.amdgcn.workitem.id.x() #0
+declare <4 x float> @llvm.amdgcn.struct.buffer.load.format.v4f32(<4 x i32>, i32, i32, i32, i32 immarg) #1
+declare <4 x half> @llvm.amdgcn.struct.buffer.load.format.v4f16(<4 x i32>, i32, i32, i32, i32 immarg) #1
+declare <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32, i32, <4 x i32>, i32 immarg, i32 immarg, i32 immarg) #2
+
+attributes #0 = { nounwind readnone speculatable willreturn }
+attributes #1 = { nounwind readonly willreturn }
+attributes #2 = { convergent nounwind readnone willreturn }
More information about the llvm-commits
mailing list