[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