diff --git a/llvm/lib/Target/AMDGPU/VOP3PInstructions.td b/llvm/lib/Target/AMDGPU/VOP3PInstructions.td --- a/llvm/lib/Target/AMDGPU/VOP3PInstructions.td +++ b/llvm/lib/Target/AMDGPU/VOP3PInstructions.td @@ -427,6 +427,17 @@ } multiclass VOP3P_Real_MAI op> { + def _vi : VOP3P_Real(NAME), SIEncodingFamily.VI>, + VOP3Pe_MAI (NAME).Pfl> { + let AssemblerPredicate = HasMAIInsts; + let DecoderNamespace = "GFX8"; + let Inst{14} = 1; // op_sel_hi(2) default value + let Inst{59} = 1; // op_sel_hi(0) default value + let Inst{60} = 1; // op_sel_hi(1) default value + } +} + +multiclass VOP3P_Real_MFMA op> { def _vi : VOP3P_Real(NAME), SIEncodingFamily.VI>, VOP3Pe_MAI (NAME).Pfl> { let AssemblerPredicate = HasMAIInsts; @@ -495,26 +506,26 @@ defm V_ACCVGPR_READ_B32 : VOP3P_Real_MAI <0x58>; defm V_ACCVGPR_WRITE_B32 : VOP3P_Real_MAI <0x59>; -defm V_MFMA_F32_32X32X1F32 : VOP3P_Real_MAI <0x40>; -defm V_MFMA_F32_16X16X1F32 : VOP3P_Real_MAI <0x41>; -defm V_MFMA_F32_4X4X1F32 : VOP3P_Real_MAI <0x42>; -defm V_MFMA_F32_32X32X2F32 : VOP3P_Real_MAI <0x44>; -defm V_MFMA_F32_16X16X4F32 : VOP3P_Real_MAI <0x45>; -defm V_MFMA_F32_32X32X4F16 : VOP3P_Real_MAI <0x48>; -defm V_MFMA_F32_16X16X4F16 : VOP3P_Real_MAI <0x49>; -defm V_MFMA_F32_4X4X4F16 : VOP3P_Real_MAI <0x4a>; -defm V_MFMA_F32_32X32X8F16 : VOP3P_Real_MAI <0x4c>; -defm V_MFMA_F32_16X16X16F16 : VOP3P_Real_MAI <0x4d>; -defm V_MFMA_I32_32X32X4I8 : VOP3P_Real_MAI <0x50>; -defm V_MFMA_I32_16X16X4I8 : VOP3P_Real_MAI <0x51>; -defm V_MFMA_I32_4X4X4I8 : VOP3P_Real_MAI <0x52>; -defm V_MFMA_I32_32X32X8I8 : VOP3P_Real_MAI <0x54>; -defm V_MFMA_I32_16X16X16I8 : VOP3P_Real_MAI <0x55>; -defm V_MFMA_F32_32X32X2BF16 : VOP3P_Real_MAI <0x68>; -defm V_MFMA_F32_16X16X2BF16 : VOP3P_Real_MAI <0x69>; -defm V_MFMA_F32_4X4X2BF16 : VOP3P_Real_MAI <0x6b>; -defm V_MFMA_F32_32X32X4BF16 : VOP3P_Real_MAI <0x6c>; -defm V_MFMA_F32_16X16X8BF16 : VOP3P_Real_MAI <0x6d>; +defm V_MFMA_F32_32X32X1F32 : VOP3P_Real_MFMA <0x40>; +defm V_MFMA_F32_16X16X1F32 : VOP3P_Real_MFMA <0x41>; +defm V_MFMA_F32_4X4X1F32 : VOP3P_Real_MFMA <0x42>; +defm V_MFMA_F32_32X32X2F32 : VOP3P_Real_MFMA <0x44>; +defm V_MFMA_F32_16X16X4F32 : VOP3P_Real_MFMA <0x45>; +defm V_MFMA_F32_32X32X4F16 : VOP3P_Real_MFMA <0x48>; +defm V_MFMA_F32_16X16X4F16 : VOP3P_Real_MFMA <0x49>; +defm V_MFMA_F32_4X4X4F16 : VOP3P_Real_MFMA <0x4a>; +defm V_MFMA_F32_32X32X8F16 : VOP3P_Real_MFMA <0x4c>; +defm V_MFMA_F32_16X16X16F16 : VOP3P_Real_MFMA <0x4d>; +defm V_MFMA_I32_32X32X4I8 : VOP3P_Real_MFMA <0x50>; +defm V_MFMA_I32_16X16X4I8 : VOP3P_Real_MFMA <0x51>; +defm V_MFMA_I32_4X4X4I8 : VOP3P_Real_MFMA <0x52>; +defm V_MFMA_I32_16X16X16I8 : VOP3P_Real_MFMA <0x55>; +defm V_MFMA_I32_32X32X8I8 : VOP3P_Real_MFMA <0x54>; +defm V_MFMA_F32_32X32X2BF16 : VOP3P_Real_MFMA <0x68>; +defm V_MFMA_F32_16X16X2BF16 : VOP3P_Real_MFMA <0x69>; +defm V_MFMA_F32_4X4X2BF16 : VOP3P_Real_MFMA <0x6b>; +defm V_MFMA_F32_32X32X4BF16 : VOP3P_Real_MFMA <0x6c>; +defm V_MFMA_F32_16X16X8BF16 : VOP3P_Real_MFMA <0x6d>; } // End SubtargetPredicate = HasMAIInsts diff --git a/llvm/test/MC/AMDGPU/accvgpr-altnames.s b/llvm/test/MC/AMDGPU/accvgpr-altnames.s --- a/llvm/test/MC/AMDGPU/accvgpr-altnames.s +++ b/llvm/test/MC/AMDGPU/accvgpr-altnames.s @@ -1,10 +1,10 @@ // RUN: llvm-mc -arch=amdgcn -mcpu=gfx908 -show-encoding %s | FileCheck -check-prefix=GFX908 %s v_accvgpr_read_b32 v2, acc0 -// GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x00,0xd8,0xd3,0x00,0x01,0x00,0x08] +// GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x40,0xd8,0xd3,0x00,0x01,0x00,0x18] v_accvgpr_write_b32 acc2, -2.0 -// GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x00,0xd9,0xd3,0xf5,0x00,0x00,0x00] +// GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x40,0xd9,0xd3,0xf5,0x00,0x00,0x18] v_mfma_f32_32x32x1f32 acc[0:31], acc0, acc1, acc[1:32] // GFX908: v_mfma_f32_32x32x1f32 a[0:31], a0, a1, a[1:32] ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x06,0x1c] diff --git a/llvm/test/MC/AMDGPU/mai.s b/llvm/test/MC/AMDGPU/mai.s --- a/llvm/test/MC/AMDGPU/mai.s +++ b/llvm/test/MC/AMDGPU/mai.s @@ -2,28 +2,28 @@ // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx908 %s 2>&1 | FileCheck -check-prefix=NOGFX908 --implicit-check-not=error: %s v_accvgpr_read_b32 v2, a0 -// GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x00,0xd8,0xd3,0x00,0x01,0x00,0x08] +// GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x40,0xd8,0xd3,0x00,0x01,0x00,0x18] v_accvgpr_read_b32 v2, a1 -// GFX908: v_accvgpr_read_b32 v2, a1 ; encoding: [0x02,0x00,0xd8,0xd3,0x01,0x01,0x00,0x08] +// GFX908: v_accvgpr_read_b32 v2, a1 ; encoding: [0x02,0x40,0xd8,0xd3,0x01,0x01,0x00,0x18] v_accvgpr_read_b32 v2, a255 -// GFX908: v_accvgpr_read_b32 v2, a255 ; encoding: [0x02,0x00,0xd8,0xd3,0xff,0x01,0x00,0x08] +// GFX908: v_accvgpr_read_b32 v2, a255 ; encoding: [0x02,0x40,0xd8,0xd3,0xff,0x01,0x00,0x18] v_accvgpr_read v2, a10 -// GFX908: v_accvgpr_read_b32 v2, a10 ; encoding: [0x02,0x00,0xd8,0xd3,0x0a,0x01,0x00,0x08] +// GFX908: v_accvgpr_read_b32 v2, a10 ; encoding: [0x02,0x40,0xd8,0xd3,0x0a,0x01,0x00,0x18] v_accvgpr_write_b32 a2, -2.0 -// GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x00,0xd9,0xd3,0xf5,0x00,0x00,0x00] +// GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x40,0xd9,0xd3,0xf5,0x00,0x00,0x18] v_accvgpr_write_b32 a2, -2 -// GFX908: v_accvgpr_write_b32 a2, -2 ; encoding: [0x02,0x00,0xd9,0xd3,0xc2,0x00,0x00,0x00] +// GFX908: v_accvgpr_write_b32 a2, -2 ; encoding: [0x02,0x40,0xd9,0xd3,0xc2,0x00,0x00,0x18] v_accvgpr_write_b32 a2, v1 -// GFX908: v_accvgpr_write_b32 a2, v1 ; encoding: [0x02,0x00,0xd9,0xd3,0x01,0x01,0x00,0x00] +// GFX908: v_accvgpr_write_b32 a2, v1 ; encoding: [0x02,0x40,0xd9,0xd3,0x01,0x01,0x00,0x18] v_accvgpr_write a2, v255 -// GFX908: v_accvgpr_write_b32 a2, v255 ; encoding: [0x02,0x00,0xd9,0xd3,0xff,0x01,0x00,0x00] +// GFX908: v_accvgpr_write_b32 a2, v255 ; encoding: [0x02,0x40,0xd9,0xd3,0xff,0x01,0x00,0x18] v_accvgpr_write a2, 100 // NOGFX908: error: invalid operand for instruction diff --git a/llvm/test/MC/Disassembler/AMDGPU/mai.txt b/llvm/test/MC/Disassembler/AMDGPU/mai.txt --- a/llvm/test/MC/Disassembler/AMDGPU/mai.txt +++ b/llvm/test/MC/Disassembler/AMDGPU/mai.txt @@ -1,22 +1,22 @@ # RUN: llvm-mc -arch=amdgcn -mcpu=gfx908 -show-encoding -disassemble %s | FileCheck -check-prefix=GFX908 %s -# GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x00,0xd8,0xd3,0x00,0x01,0x00,0x08] -0x02,0x00,0xd8,0xd3,0x00,0x01,0x00,0x08 +# GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x40,0xd8,0xd3,0x00,0x01,0x00,0x18] +0x02,0x40,0xd8,0xd3,0x00,0x01,0x00,0x18 -# GFX908: v_accvgpr_read_b32 v2, a1 ; encoding: [0x02,0x00,0xd8,0xd3,0x01,0x01,0x00,0x08] -0x02,0x00,0xd8,0xd3,0x01,0x01,0x00,0x08 +# GFX908: v_accvgpr_read_b32 v2, a1 ; encoding: [0x02,0x40,0xd8,0xd3,0x01,0x01,0x00,0x18] +0x02,0x40,0xd8,0xd3,0x01,0x01,0x00,0x18 -# GFX908: v_accvgpr_read_b32 v2, a255 ; encoding: [0x02,0x00,0xd8,0xd3,0xff,0x01,0x00,0x08] -0x02,0x00,0xd8,0xd3,0xff,0x01,0x00,0x08 +# GFX908: v_accvgpr_read_b32 v2, a255 ; encoding: [0x02,0x40,0xd8,0xd3,0xff,0x01,0x00,0x18] +0x02,0x40,0xd8,0xd3,0xff,0x01,0x00,0x18 -# GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x00,0xd9,0xd3,0xf5,0x00,0x00,0x00] -0x02,0x00,0xd9,0xd3,0xf5,0x00,0x00,0x00 +# GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x40,0xd9,0xd3,0xf5,0x00,0x00,0x18] +0x02,0x40,0xd9,0xd3,0xf5,0x00,0x00,0x18 -# GFX908: v_accvgpr_write_b32 a2, -2 ; encoding: [0x02,0x00,0xd9,0xd3,0xc2,0x00,0x00,0x00] -0x02,0x00,0xd9,0xd3,0xc2,0x00,0x00,0x00 +# GFX908: v_accvgpr_write_b32 a2, -2 ; encoding: [0x02,0x40,0xd9,0xd3,0xc2,0x00,0x00,0x18] +0x02,0x40,0xd9,0xd3,0xc2,0x00,0x00,0x18 -# GFX908: v_accvgpr_write_b32 a2, v1 ; encoding: [0x02,0x00,0xd9,0xd3,0x01,0x01,0x00,0x00] -0x02,0x00,0xd9,0xd3,0x01,0x01,0x00,0x00 +# GFX908: v_accvgpr_write_b32 a2, v1 ; encoding: [0x02,0x40,0xd9,0xd3,0x01,0x01,0x00,0x18] +0x02,0x40,0xd9,0xd3,0x01,0x01,0x00,0x18 # GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[1:32] ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x06,0x04] 0x00,0x00,0xc0,0xd3,0x00,0x03,0x06,0x04