1; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs -show-mc-encoding < %s | FileCheck -check-prefixes=GFX9,GFX908 %s
2; RUN: llc -march=amdgcn -mcpu=gfx90a -verify-machineinstrs -show-mc-encoding < %s | FileCheck -check-prefixes=GFX9,GFX90A %s
3
4; GFX9-DAG:   buffer_load_format_xyzw v[{{[0-9:]+}}], v{{[0-9]+}}, s[{{[0-9:]+}}], 0 idxen ; encoding:
5; GFX9-DAG:   buffer_load_format_d16_xyzw v[{{[0-9:]+}}], v{{[0-9]+}}, s[{{[0-9:]+}}], 0 idxen ; encoding:
6; GFX908-DAG: v_mfma_i32_4x4x4i8 a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9]+}}, a[{{[0-9:]+}}] ; encoding: [{{0x..,0x0.,}}
7; GFX90A-DAG: v_mfma_i32_4x4x4i8 a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9]+}}, a[{{[0-9:]+}}] ; encoding: [{{0x..,0x8.,}}
8define amdgpu_kernel void @test(<4 x i32> %x)  {
9  %id = tail call i32 @llvm.amdgcn.workitem.id.x()
10  %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)
11  store volatile <4 x float> %r1, <4 x float>* undef
12  %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)
13  store volatile <4 x half> %r2, <4 x half>* undef
14  %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)
15  store <4 x i32> %r3, <4 x i32>* undef
16  ret void
17}
18
19declare i32 @llvm.amdgcn.workitem.id.x()
20declare <4 x float> @llvm.amdgcn.buffer.load.format.v4f32(<4 x i32>, i32, i32, i1 immarg, i1 immarg)
21declare <4 x half> @llvm.amdgcn.buffer.load.format.v4f16(<4 x i32>, i32, i32, i1 immarg, i1 immarg)
22declare <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32, i32, <4 x i32>, i32, i32, i32)
23