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 @@ -573,16 +573,56 @@ VOP3Pe_MAI (NAME # "_vgprcd" # "_e64").Pfl, 0>; } // End AssemblerPredicate = isGFX90AOnly, DecoderNamespace = "GFX90A" } +} + +multiclass VOP3P_Real_MFMA_gfx940_aliases(Op # "_e64"), + VOP3_Pseudo PS_VCD = !cast(Op # "_vgprcd" # "_e64"), + VOPProfile Pfl_ACD = PS_ACD.Pfl, + VOPProfile Pfl_VCD = PS_VCD.Pfl> { + let Predicates = [isGFX940Plus] in { + foreach _ = BoolToList.ret in { + def : InstAlias (Op # "_gfx940_acd") Pfl_ACD.DstRC:$vdst, + Pfl_ACD.Src0RC64:$src0, Pfl_ACD.Src1RC64:$src1, Pfl_ACD.Src2RC64:$src2, + cbsz:$cbsz, abid:$abid, blgp:$blgp)>, PredicateControl; + def : InstAlias (Op # "_gfx940_vcd") Pfl_VCD.DstRC:$vdst, + Pfl_VCD.Src0RC64:$src0, Pfl_VCD.Src1RC64:$src1, Pfl_VCD.Src2RC64:$src2, + cbsz:$cbsz, abid:$abid, blgp:$blgp)>, PredicateControl; + } + } // End Predicates = [isGFX940Plus] +} + +multiclass VOP3P_Real_MFMA_gfx940 op, string Name = !cast(NAME#"_e64").Mnemonic, + VOP3_Pseudo PS_ACD = !cast(NAME # "_e64"), + VOP3_Pseudo PS_VCD = !cast(NAME # "_vgprcd" # "_e64")> { + let SubtargetPredicate = isGFX940Plus, + AssemblerPredicate = isGFX940Plus, DecoderNamespace = "GFX9", + AsmString = Name # PS_ACD.AsmOperands, Constraints = "" in { + def _gfx940_acd : VOP3P_Real, + VOP3Pe_MAI ; + + def _gfx940_vcd : VOP3P_Real, + VOP3Pe_MAI ; + } // End AssemblerPredicate = isGFX940Plus, DecoderNamespace = "GFX9" + + defm : VOP3P_Real_MFMA_gfx940_aliases; -multiclass VOP3P_Real_MFMA op> : - VOP3P_Real_MFMA_gfx90a { + foreach _ = BoolToList.ret in + defm : VOP3P_Real_MFMA_gfx940_aliases; +} + +multiclass VOP3P_Real_MFMA op, string GFX940Name = !cast(NAME#"_e64").Mnemonic> : + VOP3P_Real_MFMA_gfx90a , + VOP3P_Real_MFMA_gfx940 { def _vi : VOP3P_Real(NAME#"_e64"), SIEncodingFamily.VI>, VOP3Pe_MAI (NAME#"_e64").Pfl, ?> { let AssemblerPredicate = HasMAIInsts; let DecoderNamespace = "GFX8"; + let Constraints = ""; } } -} defm V_PK_MAD_I16 : VOP3P_Real_vi <0x00>; defm V_PK_MUL_LO_U16 : VOP3P_Real_vi <0x01>; @@ -650,19 +690,20 @@ 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_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_F32_32X32X1F32 : VOP3P_Real_MFMA <0x40, "v_mfma_f32_32x32x1_2b_f32">; +defm V_MFMA_F32_16X16X1F32 : VOP3P_Real_MFMA <0x41, "v_mfma_f32_16x16x1_4b_f32">; +defm V_MFMA_F32_4X4X1F32 : VOP3P_Real_MFMA <0x42, "v_mfma_f32_4x4x1_16b_f32">; +defm V_MFMA_F32_32X32X2F32 : VOP3P_Real_MFMA <0x44, "v_mfma_f32_32x32x2_f32">; +defm V_MFMA_F32_16X16X4F32 : VOP3P_Real_MFMA <0x45, "v_mfma_f32_16x16x4_f32">; +defm V_MFMA_F32_32X32X4F16 : VOP3P_Real_MFMA <0x48, "v_mfma_f32_32x32x4_2b_f16">; +defm V_MFMA_F32_16X16X4F16 : VOP3P_Real_MFMA <0x49, "v_mfma_f32_16x16x4_4b_f16">; +defm V_MFMA_F32_4X4X4F16 : VOP3P_Real_MFMA <0x4a, "v_mfma_f32_4x4x4_16b_f16">; +defm V_MFMA_F32_32X32X8F16 : VOP3P_Real_MFMA <0x4c, "v_mfma_f32_32x32x8_f16">; +defm V_MFMA_F32_16X16X16F16 : VOP3P_Real_MFMA <0x4d, "v_mfma_f32_16x16x16_f16">; +defm V_MFMA_I32_32X32X4I8 : VOP3P_Real_MFMA <0x50, "v_mfma_i32_32x32x4_2b_i8">; +defm V_MFMA_I32_16X16X4I8 : VOP3P_Real_MFMA <0x51, "v_mfma_i32_16x16x4_4b_i8">; +defm V_MFMA_I32_4X4X4I8 : VOP3P_Real_MFMA <0x52, "v_mfma_i32_4x4x4_16b_i8">; + 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>; @@ -681,6 +722,15 @@ defm V_MFMA_F64_16X16X4F64 : VOP3P_Real_MFMA_gfx90a <0x6e>; defm V_MFMA_F64_4X4X4F64 : VOP3P_Real_MFMA_gfx90a <0x6f>; +defm V_MFMA_F32_32X32X4BF16_1K : VOP3P_Real_MFMA_gfx940 <0x5d, "v_mfma_f32_32x32x4_2b_bf16">; +defm V_MFMA_F32_16X16X4BF16_1K : VOP3P_Real_MFMA_gfx940 <0x5e, "v_mfma_f32_16x16x4_4b_bf16">; +defm V_MFMA_F32_4X4X4BF16_1K : VOP3P_Real_MFMA_gfx940 <0x5f, "v_mfma_f32_4x4x4_16b_bf16">; +defm V_MFMA_F32_32X32X8BF16_1K : VOP3P_Real_MFMA_gfx940 <0x60, "v_mfma_f32_32x32x8_bf16">; +defm V_MFMA_F32_16X16X16BF16_1K : VOP3P_Real_MFMA_gfx940 <0x61, "v_mfma_f32_16x16x16_bf16">; + +defm V_MFMA_F64_16X16X4F64 : VOP3P_Real_MFMA_gfx940 <0x6e, "v_mfma_f64_16x16x4_f64">; +defm V_MFMA_F64_4X4X4F64 : VOP3P_Real_MFMA_gfx940 <0x6f, "v_mfma_f64_4x4x4_4b_f64">; + let SubtargetPredicate = HasPackedFP32Ops in { defm V_PK_FMA_F32 : VOP3P_Real_vi <0x30>; defm V_PK_MUL_F32 : VOP3P_Real_vi <0x31>; diff --git a/llvm/test/MC/AMDGPU/mai-gfx940.s b/llvm/test/MC/AMDGPU/mai-gfx940.s --- a/llvm/test/MC/AMDGPU/mai-gfx940.s +++ b/llvm/test/MC/AMDGPU/mai-gfx940.s @@ -4,3 +4,339 @@ v_accvgpr_write_b32 a10, s20 // GFX940: v_accvgpr_write_b32 a10, s20 ; encoding: [0x0a,0x40,0xd9,0xd3,0x14,0x00,0x00,0x18] // GFX90A: error: source operand must be either a VGPR or an inline constant + +v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] +// GFX940: v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] ; encoding: [0x00,0x80,0xef,0xd3,0x00,0x05,0x0a,0x14] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], a[2:3], v[2:3] +// GFX940: v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], a[2:3], v[2:3] ; encoding: [0x00,0x00,0xef,0xd3,0x00,0x05,0x0a,0x14] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f64_4x4x4f64 a[0:1], v[0:1], a[2:3], a[2:3] +// GFX940: v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] ; encoding: [0x00,0x80,0xef,0xd3,0x00,0x05,0x0a,0x14] + +v_mfma_f64_4x4x4f64 v[0:1], v[0:1], a[2:3], v[2:3] +// GFX940: v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], a[2:3], v[2:3] ; encoding: [0x00,0x00,0xef,0xd3,0x00,0x05,0x0a,0x14] + +v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] +// GFX940: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] ; encoding: [0x00,0x80,0xee,0xd3,0x00,0x05,0x02,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7] +// GFX940: v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7] ; encoding: [0x00,0x00,0xee,0xd3,0x00,0x05,0x02,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7] +// GFX940: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] ; encoding: [0x00,0x80,0xee,0xd3,0x00,0x05,0x02,0x04] + +v_mfma_f64_16x16x4f64 v[0:7], v[0:1], v[2:3], v[0:7] +// GFX940: v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7] ; encoding: [0x00,0x00,0xee,0xd3,0x00,0x05,0x02,0x04] + +v_mfma_f32_16x16x1_4b_f32 a[0:15], v0, v1, a[18:33] +// GFX940: v_mfma_f32_16x16x1_4b_f32 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xc1,0xd3,0x00,0x03,0x4a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_16x16x1_4b_f32 v[0:15], v0, v1, v[18:33] +// GFX940: v_mfma_f32_16x16x1_4b_f32 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xc1,0xd3,0x00,0x03,0x4a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_16x16x1f32 a[0:15], v0, v1, a[18:33] +// GFX940: v_mfma_f32_16x16x1_4b_f32 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xc1,0xd3,0x00,0x03,0x4a,0x04] + +v_mfma_f32_16x16x1f32 v[0:15], v0, v1, v[18:33] +// GFX940: v_mfma_f32_16x16x1_4b_f32 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xc1,0xd3,0x00,0x03,0x4a,0x04] + +v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[2:5] +// GFX940: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xc2,0xd3,0x00,0x03,0x0a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_4x4x1_16b_f32 v[0:3], v0, v1, v[2:5] +// GFX940: v_mfma_f32_4x4x1_16b_f32 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xc2,0xd3,0x00,0x03,0x0a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_4x4x1f32 a[0:3], v0, v1, a[2:5] +// GFX940: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xc2,0xd3,0x00,0x03,0x0a,0x04] + +v_mfma_f32_4x4x1f32 v[0:3], v0, v1, v[2:5] +// GFX940: v_mfma_f32_4x4x1_16b_f32 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xc2,0xd3,0x00,0x03,0x0a,0x04] + +v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33] +// GFX940: v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xc4,0xd3,0x00,0x03,0x4a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33] +// GFX940: v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xc4,0xd3,0x00,0x03,0x4a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_32x32x2f32 a[0:15], v0, v1, a[18:33] +// GFX940: v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xc4,0xd3,0x00,0x03,0x4a,0x04] + +v_mfma_f32_32x32x2f32 v[0:15], v0, v1, v[18:33] +// GFX940: v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xc4,0xd3,0x00,0x03,0x4a,0x04] + +v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5] +// GFX940: v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xc5,0xd3,0x00,0x03,0x0a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5] +// GFX940: v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xc5,0xd3,0x00,0x03,0x0a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_16x16x4f32 a[0:3], v0, v1, a[2:5] +// GFX940: v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xc5,0xd3,0x00,0x03,0x0a,0x04] + +v_mfma_f32_16x16x4f32 v[0:3], v0, v1, v[2:5] +// GFX940: v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xc5,0xd3,0x00,0x03,0x0a,0x04] + +v_mfma_f32_32x32x4_2b_f16 a[0:31], v[0:1], v[2:3], a[34:65] +// GFX940: v_mfma_f32_32x32x4_2b_f16 a[0:31], v[0:1], v[2:3], a[34:65] ; encoding: [0x00,0x80,0xc8,0xd3,0x00,0x05,0x8a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_32x32x4_2b_f16 v[0:31], v[0:1], v[2:3], v[34:65] +// GFX940: v_mfma_f32_32x32x4_2b_f16 v[0:31], v[0:1], v[2:3], v[34:65] ; encoding: [0x00,0x00,0xc8,0xd3,0x00,0x05,0x8a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_32x32x4f16 a[0:31], v[0:1], v[2:3], a[34:65] +// GFX940: v_mfma_f32_32x32x4_2b_f16 a[0:31], v[0:1], v[2:3], a[34:65] ; encoding: [0x00,0x80,0xc8,0xd3,0x00,0x05,0x8a,0x04] + +v_mfma_f32_32x32x4f16 v[0:31], v[0:1], v[2:3], v[34:65] +// GFX940: v_mfma_f32_32x32x4_2b_f16 v[0:31], v[0:1], v[2:3], v[34:65] ; encoding: [0x00,0x00,0xc8,0xd3,0x00,0x05,0x8a,0x04] + +v_mfma_f32_16x16x4_4b_f16 a[0:15], v[0:1], v[2:3], a[18:33] +// GFX940: v_mfma_f32_16x16x4_4b_f16 a[0:15], v[0:1], v[2:3], a[18:33] ; encoding: [0x00,0x80,0xc9,0xd3,0x00,0x05,0x4a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_16x16x4_4b_f16 v[0:15], v[0:1], v[2:3], v[18:33] +// GFX940: v_mfma_f32_16x16x4_4b_f16 v[0:15], v[0:1], v[2:3], v[18:33] ; encoding: [0x00,0x00,0xc9,0xd3,0x00,0x05,0x4a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_16x16x4f16 a[0:15], v[0:1], v[2:3], a[18:33] +// GFX940: v_mfma_f32_16x16x4_4b_f16 a[0:15], v[0:1], v[2:3], a[18:33] ; encoding: [0x00,0x80,0xc9,0xd3,0x00,0x05,0x4a,0x04] + +v_mfma_f32_16x16x4f16 v[0:15], v[0:1], v[2:3], v[18:33] +// GFX940: v_mfma_f32_16x16x4_4b_f16 v[0:15], v[0:1], v[2:3], v[18:33] ; encoding: [0x00,0x00,0xc9,0xd3,0x00,0x05,0x4a,0x04] + +v_mfma_f32_4x4x4_16b_f16 a[0:3], v[0:1], v[2:3], a[2:5] +// GFX940: v_mfma_f32_4x4x4_16b_f16 a[0:3], v[0:1], v[2:3], a[2:5] ; encoding: [0x00,0x80,0xca,0xd3,0x00,0x05,0x0a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_4x4x4_16b_f16 v[0:3], v[0:1], v[2:3], v[2:5] +// GFX940: v_mfma_f32_4x4x4_16b_f16 v[0:3], v[0:1], v[2:3], v[2:5] ; encoding: [0x00,0x00,0xca,0xd3,0x00,0x05,0x0a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_4x4x4f16 a[0:3], v[0:1], v[2:3], a[2:5] +// GFX940: v_mfma_f32_4x4x4_16b_f16 a[0:3], v[0:1], v[2:3], a[2:5] ; encoding: [0x00,0x80,0xca,0xd3,0x00,0x05,0x0a,0x04] + +v_mfma_f32_4x4x4f16 v[0:3], v[0:1], v[2:3], v[2:5] +// GFX940: v_mfma_f32_4x4x4_16b_f16 v[0:3], v[0:1], v[2:3], v[2:5] ; encoding: [0x00,0x00,0xca,0xd3,0x00,0x05,0x0a,0x04] + +v_mfma_f32_32x32x8_f16 a[0:15], v[0:1], v[2:3], a[18:33] +// GFX940: v_mfma_f32_32x32x8_f16 a[0:15], v[0:1], v[2:3], a[18:33] ; encoding: [0x00,0x80,0xcc,0xd3,0x00,0x05,0x4a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_32x32x8_f16 v[0:15], v[0:1], v[2:3], v[18:33] +// GFX940: v_mfma_f32_32x32x8_f16 v[0:15], v[0:1], v[2:3], v[18:33] ; encoding: [0x00,0x00,0xcc,0xd3,0x00,0x05,0x4a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_32x32x8f16 a[0:15], v[0:1], v[2:3], a[18:33] +// GFX940: v_mfma_f32_32x32x8_f16 a[0:15], v[0:1], v[2:3], a[18:33] ; encoding: [0x00,0x80,0xcc,0xd3,0x00,0x05,0x4a,0x04] + +v_mfma_f32_32x32x8f16 v[0:15], v[0:1], v[2:3], v[18:33] +// GFX940: v_mfma_f32_32x32x8_f16 v[0:15], v[0:1], v[2:3], v[18:33] ; encoding: [0x00,0x00,0xcc,0xd3,0x00,0x05,0x4a,0x04] + +v_mfma_f32_16x16x16_f16 a[0:3], v[0:1], v[2:3], a[2:5] +// GFX940: v_mfma_f32_16x16x16_f16 a[0:3], v[0:1], v[2:3], a[2:5] ; encoding: [0x00,0x80,0xcd,0xd3,0x00,0x05,0x0a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_16x16x16_f16 v[0:3], v[0:1], v[2:3], v[2:5] +// GFX940: v_mfma_f32_16x16x16_f16 v[0:3], v[0:1], v[2:3], v[2:5] ; encoding: [0x00,0x00,0xcd,0xd3,0x00,0x05,0x0a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_16x16x16f16 a[0:3], v[0:1], v[2:3], a[2:5] +// GFX940: v_mfma_f32_16x16x16_f16 a[0:3], v[0:1], v[2:3], a[2:5] ; encoding: [0x00,0x80,0xcd,0xd3,0x00,0x05,0x0a,0x04] + +v_mfma_f32_16x16x16f16 v[0:3], v[0:1], v[2:3], v[2:5] +// GFX940: v_mfma_f32_16x16x16_f16 v[0:3], v[0:1], v[2:3], v[2:5] ; encoding: [0x00,0x00,0xcd,0xd3,0x00,0x05,0x0a,0x04] + +v_mfma_i32_32x32x4_2b_i8 a[0:31], v0, v1, a[34:65] +// GFX940: v_mfma_i32_32x32x4_2b_i8 a[0:31], v0, v1, a[34:65] ; encoding: [0x00,0x80,0xd0,0xd3,0x00,0x03,0x8a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_i32_32x32x4_2b_i8 v[0:31], v0, a1, v[34:65] +// GFX940: v_mfma_i32_32x32x4_2b_i8 v[0:31], v0, a1, v[34:65] ; encoding: [0x00,0x00,0xd0,0xd3,0x00,0x03,0x8a,0x14] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_i32_32x32x4i8 a[0:31], v0, v1, a[34:65] +// GFX940: v_mfma_i32_32x32x4_2b_i8 a[0:31], v0, v1, a[34:65] ; encoding: [0x00,0x80,0xd0,0xd3,0x00,0x03,0x8a,0x04] + +v_mfma_i32_32x32x4i8 v[0:31], v0, a1, v[34:65] +// GFX940: v_mfma_i32_32x32x4_2b_i8 v[0:31], v0, a1, v[34:65] ; encoding: [0x00,0x00,0xd0,0xd3,0x00,0x03,0x8a,0x14] + +v_mfma_i32_16x16x4_4b_i8 a[0:15], v0, v1, a[18:33] +// GFX940: v_mfma_i32_16x16x4_4b_i8 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xd1,0xd3,0x00,0x03,0x4a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_i32_16x16x4_4b_i8 v[0:15], v0, v1, v[18:33] +// GFX940: v_mfma_i32_16x16x4_4b_i8 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xd1,0xd3,0x00,0x03,0x4a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_i32_16x16x4i8 a[0:15], v0, v1, a[18:33] +// GFX940: v_mfma_i32_16x16x4_4b_i8 a[0:15], v0, v1, a[18:33] ; encoding: [0x00,0x80,0xd1,0xd3,0x00,0x03,0x4a,0x04] + +v_mfma_i32_16x16x4i8 v[0:15], v0, v1, v[18:33] +// GFX940: v_mfma_i32_16x16x4_4b_i8 v[0:15], v0, v1, v[18:33] ; encoding: [0x00,0x00,0xd1,0xd3,0x00,0x03,0x4a,0x04] + +v_mfma_i32_4x4x4_16b_i8 a[0:3], v0, v1, a[2:5] +// GFX940: v_mfma_i32_4x4x4_16b_i8 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xd2,0xd3,0x00,0x03,0x0a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_i32_4x4x4_16b_i8 v[0:3], v0, v1, v[2:5] +// GFX940: v_mfma_i32_4x4x4_16b_i8 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xd2,0xd3,0x00,0x03,0x0a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_i32_4x4x4i8 a[0:3], v0, v1, a[2:5] +// GFX940: v_mfma_i32_4x4x4_16b_i8 a[0:3], v0, v1, a[2:5] ; encoding: [0x00,0x80,0xd2,0xd3,0x00,0x03,0x0a,0x04] + +v_mfma_i32_4x4x4i8 v[0:3], v0, v1, v[2:5] +// GFX940: v_mfma_i32_4x4x4_16b_i8 v[0:3], v0, v1, v[2:5] ; encoding: [0x00,0x00,0xd2,0xd3,0x00,0x03,0x0a,0x04] + +v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7 +// GFX940: v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7 ; encoding: [0x00,0x80,0xc0,0xd3,0x00,0x03,0x8a,0xe4] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7 +// GFX940: v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7 ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x8a,0xe4] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7 +// GFX940: v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7 ; encoding: [0x00,0x80,0xc0,0xd3,0x00,0x03,0x8a,0xe4] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7 +// GFX940: v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7 ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x8a,0xe4] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[34:65] blgp:7 +// GFX940: v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7 ; encoding: [0x00,0x80,0xc0,0xd3,0x00,0x03,0x8a,0xe4] + +v_mfma_f32_32x32x1f32 v[0:31], v0, v1, v[34:65] blgp:7 +// GFX940: v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7 ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x8a,0xe4] + +v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65] +// GFX940: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65] ; encoding: [0x00,0x00,0xdd,0xd3,0x02,0x09,0x8a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65] +// GFX940: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65] ; encoding: [0x00,0x80,0xdd,0xd3,0x02,0x09,0x8a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65] +// GFX940: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65] ; encoding: [0x00,0x00,0xdd,0xd3,0x02,0x09,0x8a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65] +// GFX940: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65] ; encoding: [0x00,0x80,0xdd,0xd3,0x02,0x09,0x8a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_32x32x4bf16 v[0:31], v[2:3], v[4:5], v[34:65] blgp:5 +// GFX940: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65] blgp:5 ; encoding: [0x00,0x00,0xdd,0xd3,0x02,0x09,0x8a,0xa4] +// GFX90A: error: operands are not valid for this GPU or mode + +v_mfma_f32_32x32x4bf16 a[0:31], v[2:3], v[4:5], a[34:65] blgp:5 +// GFX940: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65] blgp:5 ; encoding: [0x00,0x80,0xdd,0xd3,0x02,0x09,0x8a,0xa4] +// GFX90A: error: operands are not valid for this GPU or mode + +v_mfma_f32_32x32x4bf16_1k v[0:31], v[2:3], v[4:5], v[34:65] blgp:5 +// GFX940: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[34:65] blgp:5 ; encoding: [0x00,0x00,0xdd,0xd3,0x02,0x09,0x8a,0xa4] + +v_mfma_f32_32x32x4bf16_1k a[0:31], v[2:3], v[4:5], a[34:65] blgp:5 +// GFX940: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[34:65] blgp:5 ; encoding: [0x00,0x80,0xdd,0xd3,0x02,0x09,0x8a,0xa4] + +v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33] +// GFX940: v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xde,0xd3,0x02,0x09,0x4a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33] +// GFX940: v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xde,0xd3,0x02,0x09,0x4a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_16x16x4bf16 v[0:15], v[2:3], v[4:5], v[18:33] blgp:5 +// GFX940: v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33] blgp:5 ; encoding: [0x00,0x00,0xde,0xd3,0x02,0x09,0x4a,0xa4] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_16x16x4bf16 a[0:15], v[2:3], v[4:5], a[18:33] blgp:5 +// GFX940: v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33] blgp:5 ; encoding: [0x00,0x80,0xde,0xd3,0x02,0x09,0x4a,0xa4] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_16x16x4bf16_1k v[0:15], v[2:3], v[4:5], v[18:33] blgp:5 +// GFX940: v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33] blgp:5 ; encoding: [0x00,0x00,0xde,0xd3,0x02,0x09,0x4a,0xa4] + +v_mfma_f32_16x16x4bf16_1k a[0:15], v[2:3], v[4:5], a[18:33] blgp:5 +// GFX940: v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33] blgp:5 ; encoding: [0x00,0x80,0xde,0xd3,0x02,0x09,0x4a,0xa4] + +v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[2:3], v[4:5], v[2:5] +// GFX940: v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xdf,0xd3,0x02,0x09,0x0a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[2:3], v[4:5], a[2:5] +// GFX940: v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xdf,0xd3,0x02,0x09,0x0a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_4x4x4bf16 v[0:3], v[2:3], v[4:5], v[2:5] +// GFX940: v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xdf,0xd3,0x02,0x09,0x0a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_4x4x4bf16 a[0:3], v[2:3], v[4:5], a[2:5] +// GFX940: v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xdf,0xd3,0x02,0x09,0x0a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_4x4x4bf16_1k v[0:3], v[2:3], v[4:5], v[2:5] +// GFX940: v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xdf,0xd3,0x02,0x09,0x0a,0x04] + +v_mfma_f32_4x4x4bf16_1k a[0:3], v[2:3], v[4:5], a[2:5] +// GFX940: v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xdf,0xd3,0x02,0x09,0x0a,0x04] + +v_mfma_f32_32x32x8_bf16 v[0:15], v[2:3], v[4:5], v[18:33] +// GFX940: v_mfma_f32_32x32x8_bf16 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xe0,0xd3,0x02,0x09,0x4a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[4:5], a[18:33] +// GFX940: v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xe0,0xd3,0x02,0x09,0x4a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_32x32x8bf16 v[0:15], v[2:3], v[4:5], v[18:33] +// GFX940: v_mfma_f32_32x32x8_bf16 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xe0,0xd3,0x02,0x09,0x4a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_32x32x8bf16 a[0:15], v[2:3], v[4:5], a[18:33] +// GFX940: v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xe0,0xd3,0x02,0x09,0x4a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_32x32x8bf16_1k v[0:15], v[2:3], v[4:5], v[18:33] +// GFX940: v_mfma_f32_32x32x8_bf16 v[0:15], v[2:3], v[4:5], v[18:33] ; encoding: [0x00,0x00,0xe0,0xd3,0x02,0x09,0x4a,0x04] + +v_mfma_f32_32x32x8bf16_1k a[0:15], v[2:3], v[4:5], a[18:33] +// GFX940: v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[4:5], a[18:33] ; encoding: [0x00,0x80,0xe0,0xd3,0x02,0x09,0x4a,0x04] + +v_mfma_f32_16x16x16_bf16 v[0:3], v[2:3], v[4:5], v[2:5] +// GFX940: v_mfma_f32_16x16x16_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xe1,0xd3,0x02,0x09,0x0a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_16x16x16_bf16 a[0:3], v[2:3], v[4:5], a[2:5] +// GFX940: v_mfma_f32_16x16x16_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xe1,0xd3,0x02,0x09,0x0a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_16x16x16bf16 v[0:3], v[2:3], v[4:5], v[2:5] +// GFX940: v_mfma_f32_16x16x16_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xe1,0xd3,0x02,0x09,0x0a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_16x16x16bf16 a[0:3], v[2:3], v[4:5], a[2:5] +// GFX940: v_mfma_f32_16x16x16_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xe1,0xd3,0x02,0x09,0x0a,0x04] +// GFX90A: error: instruction not supported on this GPU + +v_mfma_f32_16x16x16bf16_1k v[0:3], v[2:3], v[4:5], v[2:5] +// GFX940: v_mfma_f32_16x16x16_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xe1,0xd3,0x02,0x09,0x0a,0x04] + +v_mfma_f32_16x16x16bf16_1k a[0:3], v[2:3], v[4:5], a[2:5] +// GFX940: v_mfma_f32_16x16x16_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xe1,0xd3,0x02,0x09,0x0a,0x04] diff --git a/llvm/test/MC/Disassembler/AMDGPU/mai-gfx940.txt b/llvm/test/MC/Disassembler/AMDGPU/mai-gfx940.txt --- a/llvm/test/MC/Disassembler/AMDGPU/mai-gfx940.txt +++ b/llvm/test/MC/Disassembler/AMDGPU/mai-gfx940.txt @@ -2,3 +2,33 @@ # GFX940: v_accvgpr_write_b32 a10, s20 ; encoding: [0x0a,0x40,0xd9,0xd3,0x14,0x00,0x00,0x18] 0x0a,0x40,0xd9,0xd3,0x14,0x00,0x00,0x18 + +# GFX940: v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[2:3], v[4:5], v[2:33] ; encoding: [0x00,0x00,0xdd,0xd3,0x02,0x09,0x0a,0x04] +0x00,0x00,0xdd,0xd3,0x02,0x09,0x0a,0x04 + +# GFX940: v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[2:3], v[4:5], a[2:33] ; encoding: [0x00,0x80,0xdd,0xd3,0x02,0x09,0x0a,0x04] +0x00,0x80,0xdd,0xd3,0x02,0x09,0x0a,0x04 + +# GFX940: v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[2:17] ; encoding: [0x00,0x00,0xde,0xd3,0x02,0x09,0x0a,0x04] +0x00,0x00,0xde,0xd3,0x02,0x09,0x0a,0x04 + +# GFX940: v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[2:17] ; encoding: [0x00,0x80,0xde,0xd3,0x02,0x09,0x0a,0x04] +0x00,0x80,0xde,0xd3,0x02,0x09,0x0a,0x04 + +# GFX940: v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xdf,0xd3,0x02,0x09,0x0a,0x04] +0x00,0x00,0xdf,0xd3,0x02,0x09,0x0a,0x04 + +# GFX940: v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xdf,0xd3,0x02,0x09,0x0a,0x04] +0x00,0x80,0xdf,0xd3,0x02,0x09,0x0a,0x04 + +# GFX940: v_mfma_f32_32x32x8_bf16 v[0:15], v[2:3], v[4:5], v[2:17] ; encoding: [0x00,0x00,0xe0,0xd3,0x02,0x09,0x0a,0x04] +0x00,0x00,0xe0,0xd3,0x02,0x09,0x0a,0x04 + +# GFX940: v_mfma_f32_32x32x8_bf16 a[0:15], v[2:3], v[4:5], a[2:17] ; encoding: [0x00,0x80,0xe0,0xd3,0x02,0x09,0x0a,0x04] +0x00,0x80,0xe0,0xd3,0x02,0x09,0x0a,0x04 + +# GFX940: v_mfma_f32_16x16x16_bf16 v[0:3], v[2:3], v[4:5], v[2:5] ; encoding: [0x00,0x00,0xe1,0xd3,0x02,0x09,0x0a,0x04] +0x00,0x00,0xe1,0xd3,0x02,0x09,0x0a,0x04 + +# GFX940: v_mfma_f32_16x16x16_bf16 a[0:3], v[2:3], v[4:5], a[2:5] ; encoding: [0x00,0x80,0xe1,0xd3,0x02,0x09,0x0a,0x04] +0x00,0x80,0xe1,0xd3,0x02,0x09,0x0a,0x04