AMDGPU: Fix using deprecated buffer intrinsics in test

This commit is contained in:
Matt Arsenault 2022-01-18 18:41:44 -05:00
parent fa40c53960
commit e3dd47f987
1 changed files with 10 additions and 6 deletions

View File

@ -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 }