# RUN: llvm-mc -triple=amdgcn -mcpu=gfx950 -show-encoding -disassemble %s | FileCheck -check-prefix=GFX950 %s

# Check behavior of truncated instruction just to the ld_scale

# GFX950: v_mfma_ld_scale_b32 v20, v21 op_sel_hi:[0,0] ; encoding: [0x00,0x40,0xac,0xd3,0x14,0x2b,0x02,0x00]
0x00,0x00,0xac,0xd3,0x14,0x2b,0x02,0x00

# GFX950: v_mfma_ld_scale_b32 v20, v21 op_sel_hi:[0,0] ; encoding: [0x00,0x40,0xac,0xd3,0x14,0x2b,0x02,0x00]
0x00,0x00,0xac,0xd3,0x14,0x2b,0x02,0x00

# GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3] ; encoding: [0x00,0x80,0xad,0xd3,0x00,0x11,0x02,0x04]
0x00,0x00,0xac,0xd3,0x14,0x2b,0x02,0x00,0x00,0x80,0xad,0xd3,0x00,0x11,0x02,0x04

# GFX950: v_mfma_ld_scale_b32 v20, v21 op_sel_hi:[0,0] ; encoding: [0x00,0x40,0xac,0xd3,0x14,0x2b,0x02,0x00]
0x00,0x00,0xac,0xd3,0x14,0x2b,0x02,0x00

# GFX950: s_nop 0 ; encoding: [0x00,0x00,0x80,0xbf]
0x00,0x00,0x80,0xbf

# GFX950: v_mfma_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3] ; encoding: [0x00,0x80,0xad,0xd3,0x00,0x11,0x02,0x04]
0x00,0x80,0xad,0xd3,0x00,0x11,0x02,0x04


# GFX950: v_mfma_ld_scale_b32 v20, v21 op_sel_hi:[0,0] ; encoding: [0x00,0x40,0xac,0xd3,0x14,0x2b,0x02,0x00]
0x00,0x00,0xac,0xd3,0x14,0x2b,0x02,0x00

# GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], v[0:7], v[8:15], a[0:3], v20, v21 op_sel_hi:[0,0,0] ; encoding: [0x00,0x00,0xac,0xd3,0x14,0x2b,0x02,0x00,0x00,0x88,0xad,0xd3,0x00,0x11,0x02,0x04]
0x00,0x00,0xac,0xd3,0x14,0x2b,0x02,0x00,0x00,0x88,0xad,0xd3,0x00,0x11,0x02,0x04