Skip to content

Commit 1073e90

Browse files
arsenmarsenm
authored
AMDGPU: Clean up more real instruction predicate overrides (#116868)
In general real instructions should not have manually specified predicates.
1 parent e3ff649 commit 1073e90

File tree

1 file changed

+16
-26
lines changed

1 file changed

+16
-26
lines changed

‎llvm/lib/Target/AMDGPU/VOP3PInstructions.td

Lines changed: 16 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -669,8 +669,6 @@ class VgprMAIFrag<SDPatternOperator Op> : MAIFrag<Op, MayNotNeedAGPRs> {
669669
let GISelPredicateCode = MayNotNeedAGPRs_gisel;
670670
}
671671

672-
let SubtargetPredicate = HasMAIInsts in {
673-
674672
let isAsCheapAsAMove = 1, isReMaterializable = 1 in {
675673
defm V_ACCVGPR_READ_B32 : VOP3Inst<"v_accvgpr_read_b32", VOPProfileAccRead>;
676674
let isMoveImm = 1 in {
@@ -680,6 +678,7 @@ let isAsCheapAsAMove = 1, isReMaterializable = 1 in {
680678

681679
class MAIInst<string OpName, VOPProfile P, SDPatternOperator node>
682680
: VOP3InstBase<OpName, P, node> {
681+
let SubtargetPredicate = HasMAIInsts;
683682
Instruction Opcode = !cast<Instruction>(NAME);
684683
bit is_dgemm = 0;
685684
bit is_gfx940_xdl = 0;
@@ -695,7 +694,7 @@ multiclass MAIInst<string OpName, string P, SDPatternOperator node> {
695694
!if(!or(NoDstOverlap, !eq(node, null_frag)), null_frag, AgprMAIFrag<node>)>,
696695
MFMATable<0, NAME # "_e64">;
697696

698-
let SubtargetPredicate = isGFX90APlus, Mnemonic = OpName in
697+
let OtherPredicates = [isGFX90APlus], Mnemonic = OpName in
699698
def _vgprcd_e64 : MAIInst<OpName # "_vgprcd", !cast<VOPProfileMAI>("VOPProfileMAI_" # P # "_VCD"),
700699
!if(!or(NoDstOverlap, !eq(node, null_frag)), null_frag, VgprMAIFrag<node>)>,
701700
MFMATable<0, NAME # "_vgprcd_e64">;
@@ -709,7 +708,7 @@ multiclass MAIInst<string OpName, string P, SDPatternOperator node> {
709708
!if(!eq(node, null_frag), null_frag, AgprMAIFrag<node>)>,
710709
MFMATable<1, NAME # "_e64">;
711710

712-
let SubtargetPredicate = isGFX90APlus in
711+
let OtherPredicates = [isGFX90APlus] in
713712
def _mac_vgprcd_e64 : MAIInst<OpName # "_mac_vgprcd", !cast<VOPProfileMAI>("VOPProfileMAI_" # P # "_VCD"),
714713
!if(!eq(node, null_frag), null_frag, VgprMAIFrag<node>)>,
715714
MFMATable<1, NAME # "_vgprcd_e64">;
@@ -735,7 +734,7 @@ defm V_MFMA_F32_32X32X8F16 : MAIInst<"v_mfma_f32_32x32x8f16", "F32_V4F16_X16",
735734
defm V_MFMA_I32_32X32X4I8 : MAIInst<"v_mfma_i32_32x32x4i8", "I32_I32_X32", int_amdgcn_mfma_i32_32x32x4i8>;
736735
}
737736

738-
let Predicates = [isGFX908orGFX90A] in {
737+
let SubtargetPredicate = isGFX908orGFX90A in {
739738
defm V_MFMA_I32_16X16X16I8 : MAIInst<"v_mfma_i32_16x16x16i8", "I32_I32_X4", int_amdgcn_mfma_i32_16x16x16i8>;
740739
defm V_MFMA_I32_32X32X8I8 : MAIInst<"v_mfma_i32_32x32x8i8", "I32_I32_X16", int_amdgcn_mfma_i32_32x32x8i8>;
741740
defm V_MFMA_F32_4X4X2BF16 : MAIInst<"v_mfma_f32_4x4x2bf16", "F32_V2I16_X4", int_amdgcn_mfma_f32_4x4x2bf16>;
@@ -745,15 +744,13 @@ defm V_MFMA_F32_32X32X2BF16 : MAIInst<"v_mfma_f32_32x32x2bf16", "F32_V2I16_X32",
745744
defm V_MFMA_F32_32X32X4BF16 : MAIInst<"v_mfma_f32_32x32x4bf16", "F32_V2I16_X16", int_amdgcn_mfma_f32_32x32x4bf16>;
746745
}
747746

748-
} // End SubtargetPredicate = HasMAIInsts
749-
750747
let SubtargetPredicate = HasGFX950Insts, is_gfx940_xdl = 1 in {
751748
defm V_MFMA_F32_16X16X32_F16 : MAIInst<"v_mfma_f32_16x16x32f16", "F32_V8F16_X32", int_amdgcn_mfma_f32_16x16x32_f16>;
752749
defm V_MFMA_F32_32X32X16_F16 : MAIInst<"v_mfma_f32_32x32x16f16", "F32_V8F16_X16", int_amdgcn_mfma_f32_32x32x16_f16>;
753750
defm V_MFMA_F32_32X32X16_BF16 : MAIInst<"v_mfma_f32_32x32x16bf16", "F32_V8BF16_X16", int_amdgcn_mfma_f32_32x32x16_bf16>;
754751
}
755752

756-
let Predicates = [isGFX90APlus] in {
753+
let SubtargetPredicate = isGFX90APlus in {
757754
let is_gfx940_xdl = 1 in {
758755
defm V_MFMA_F32_32X32X4BF16_1K : MAIInst<"v_mfma_f32_32x32x4bf16_1k", "F32_V4I16_X32", int_amdgcn_mfma_f32_32x32x4bf16_1k>;
759756
defm V_MFMA_F32_16X16X4BF16_1K : MAIInst<"v_mfma_f32_16x16x4bf16_1k", "F32_V4I16_X16", int_amdgcn_mfma_f32_16x16x4bf16_1k>;
@@ -766,7 +763,7 @@ let Predicates = [isGFX90APlus] in {
766763
defm V_MFMA_F64_16X16X4F64 : MAIInst<"v_mfma_f64_16x16x4f64", "F64_16X16X4F64", int_amdgcn_mfma_f64_16x16x4f64>;
767764
defm V_MFMA_F64_4X4X4F64 : MAIInst<"v_mfma_f64_4x4x4f64", "F64_4X4X4F64", int_amdgcn_mfma_f64_4x4x4f64>;
768765
}
769-
} // End Predicates = [isGFX90APlus]
766+
} // End SubtargetPredicate = isGFX90APlus
770767

771768
let SubtargetPredicate = isGFX940Plus, is_gfx940_xdl = 1 in {
772769
defm V_MFMA_I32_32X32X16I8 : MAIInst<"v_mfma_i32_32x32x16i8", "I32_I64_X32", int_amdgcn_mfma_i32_32x32x16_i8>;
@@ -1632,14 +1629,17 @@ multiclass VOP3P_Real_MFMA_gfx940_aliases<string NameFrom, string NameTo, string
16321629
VOPProfile Pfl_ACD = PS_ACD.Pfl,
16331630
VOPProfile Pfl_VCD = PS_VCD.Pfl> {
16341631
if !ne(NameFrom, NameTo) then {
1635-
def : InstAlias <NameTo # " " # PS_ACD.AsmOperands,
1632+
let SubtargetPredicate = PS_ACD.SubtargetPredicate,
1633+
OtherPredicates = PS_ACD.OtherPredicates in {
1634+
def : InstAlias <NameTo # " " # PS_ACD.AsmOperands,
16361635
(!cast<VOP3P_Real>(Op # "_gfx940_acd") Pfl_ACD.DstRC:$vdst,
16371636
Pfl_ACD.Src0RC64:$src0, Pfl_ACD.Src1RC64:$src1, Pfl_ACD.Src2RC64:$src2,
16381637
CBSZ:$cbsz, ABID:$abid, blgp:$blgp)>, PredicateControl;
1639-
def : InstAlias <NameTo # " " # PS_VCD.AsmOperands,
1638+
def : InstAlias <NameTo # " " # PS_VCD.AsmOperands,
16401639
(!cast<VOP3P_Real>(Op # "_gfx940_vcd") Pfl_VCD.DstRC:$vdst,
16411640
Pfl_VCD.Src0RC64:$src0, Pfl_VCD.Src1RC64:$src1, Pfl_VCD.Src2RC64:$src2,
16421641
CBSZ:$cbsz, ABID:$abid, blgp:$blgp)>, PredicateControl;
1642+
}
16431643
}
16441644
}
16451645

@@ -1656,7 +1656,10 @@ multiclass VOP3P_Real_MFMA_gfx940<bits<7> op, string Name = !cast<VOP3_Pseudo>(N
16561656
VOP3Pe_MAI <op, PS_VCD.Pfl, 0>;
16571657
} // End AssemblerPredicate = isGFX940Plus, DecoderNamespace = "GFX940"
16581658

1659-
let SubtargetPredicate = isGFX940Plus in {
1659+
let SubtargetPredicate = PS_ACD.SubtargetPredicate,
1660+
OtherPredicates = PS_ACD.OtherPredicates,
1661+
AssemblerPredicate = isGFX940Plus
1662+
in {
16601663
defm : VOP3P_Real_MFMA_gfx940_aliases<Name, PS_ACD.Mnemonic, NAME>;
16611664

16621665
if !ne(!subst("_1k", "", PS_ACD.Mnemonic), PS_ACD.Mnemonic) then
@@ -1703,7 +1706,6 @@ multiclass VOP3P_Real_SMFMAC<bits<7> op, string alias> {
17031706
}
17041707
}
17051708

1706-
let SubtargetPredicate = isGFX8GFX9 in {
17071709
defm V_PK_MAD_I16 : VOP3P_Real_vi <0x00>;
17081710
defm V_PK_MUL_LO_U16 : VOP3P_Real_vi <0x01>;
17091711
defm V_PK_ADD_I16 : VOP3P_Real_vi <0x02>;
@@ -1725,11 +1727,9 @@ defm V_PK_MUL_F16 : VOP3P_Real_vi <0x10>;
17251727
defm V_PK_MIN_F16 : VOP3P_Real_vi <0x11>;
17261728
defm V_PK_MAX_F16 : VOP3P_Real_vi <0x12>;
17271729

1728-
let OtherPredicates = [HasMadMixInsts] in {
17291730
defm V_MAD_MIX_F32 : VOP3P_Real_vi <0x20>;
17301731
defm V_MAD_MIXLO_F16 : VOP3P_Real_vi <0x21>;
17311732
defm V_MAD_MIXHI_F16 : VOP3P_Real_vi <0x22>;
1732-
}
17331733

17341734
let OtherPredicates = [HasFmaMixInsts],
17351735
DecoderNamespace = "GFX9_DL" in {
@@ -1750,9 +1750,6 @@ defm V_DOT8_U32_U4 : VOP3P_Real_vi <0x2b>;
17501750

17511751
defm V_DOT4_I32_I8 : VOP3P_Real_vi <0x28>;
17521752
defm V_DOT8_I32_I4 : VOP3P_Real_vi <0x2a>;
1753-
} // End SubtargetPredicate = isGFX8GFX9
1754-
1755-
let OtherPredicates = [HasMAIInsts] in {
17561753

17571754
defm V_ACCVGPR_READ_B32 : VOP3P_Real_MAI <0x58>;
17581755
defm V_ACCVGPR_WRITE_B32 : VOP3P_Real_MAI <0x59>;
@@ -1778,8 +1775,6 @@ defm V_MFMA_F32_4X4X2BF16 : VOP3P_Real_MFMA_vi_gfx90a <0x6b>;
17781775
defm V_MFMA_F32_32X32X4BF16 : VOP3P_Real_MFMA_vi_gfx90a <0x6c>;
17791776
defm V_MFMA_F32_16X16X8BF16 : VOP3P_Real_MFMA_vi_gfx90a <0x6d>;
17801777

1781-
} // End OtherPredicates = [HasMAIInsts]
1782-
17831778
defm V_MFMA_F32_32X32X4BF16_1K : VOP3P_Real_MFMA_gfx90a <0x63>;
17841779
defm V_MFMA_F32_16X16X4BF16_1K : VOP3P_Real_MFMA_gfx90a <0x64>;
17851780
defm V_MFMA_F32_4X4X4BF16_1K : VOP3P_Real_MFMA_gfx90a <0x65>;
@@ -1794,11 +1789,9 @@ defm V_MFMA_F32_32X32X16_BF16 : VOP3P_Real_MFMA_gfx950 <0x37, "v_mfma_f32_32x
17941789

17951790
defm V_MFMA_I32_32X32X16I8 : VOP3P_Real_MFMA_gfx940 <0x56, "v_mfma_i32_32x32x16_i8">;
17961791
defm V_MFMA_I32_16X16X32I8 : VOP3P_Real_MFMA_gfx940 <0x57, "v_mfma_i32_16x16x32_i8">;
1797-
let SubtargetPredicate = HasXF32Insts in {
17981792
defm V_MFMA_F32_16X16X8XF32 : VOP3P_Real_MFMA_gfx940 <0x3e, "v_mfma_f32_16x16x8_xf32">;
17991793
defm V_MFMA_F32_32X32X4XF32 : VOP3P_Real_MFMA_gfx940 <0x3f, "v_mfma_f32_32x32x4_xf32">;
1800-
} // End SubtargetPredicate = HasXF32Insts
1801-
let SubtargetPredicate = HasFP8Insts in {
1794+
18021795
defm V_MFMA_F32_16X16X32_BF8_BF8 : VOP3P_Real_MFMA_gfx940 <0x70>;
18031796
defm V_MFMA_F32_16X16X32_BF8_FP8 : VOP3P_Real_MFMA_gfx940 <0x71>;
18041797
defm V_MFMA_F32_16X16X32_FP8_BF8 : VOP3P_Real_MFMA_gfx940 <0x72>;
@@ -1807,7 +1800,6 @@ defm V_MFMA_F32_32X32X16_BF8_BF8 : VOP3P_Real_MFMA_gfx940 <0x74>;
18071800
defm V_MFMA_F32_32X32X16_BF8_FP8 : VOP3P_Real_MFMA_gfx940 <0x75>;
18081801
defm V_MFMA_F32_32X32X16_FP8_BF8 : VOP3P_Real_MFMA_gfx940 <0x76>;
18091802
defm V_MFMA_F32_32X32X16_FP8_FP8 : VOP3P_Real_MFMA_gfx940 <0x77>;
1810-
} // End SubtargetPredicate = HasFP8Insts
18111803

18121804
defm V_MFMA_F32_32X32X4BF16_1K : VOP3P_Real_MFMA_gfx940 <0x5d, "v_mfma_f32_32x32x4_2b_bf16">;
18131805
defm V_MFMA_F32_16X16X4BF16_1K : VOP3P_Real_MFMA_gfx940 <0x5e, "v_mfma_f32_16x16x4_4b_bf16">;
@@ -1824,7 +1816,6 @@ defm V_SMFMAC_F32_16X16X32_BF16 : VOP3P_Real_SMFMAC <0x66, "v_smfmac_f32_16x1
18241816
defm V_SMFMAC_F32_32X32X16_BF16 : VOP3P_Real_SMFMAC <0x68, "v_smfmac_f32_32x32x16bf16">;
18251817
defm V_SMFMAC_I32_16X16X64_I8 : VOP3P_Real_SMFMAC <0x6a, "v_smfmac_i32_16x16x64i8">;
18261818
defm V_SMFMAC_I32_32X32X32_I8 : VOP3P_Real_SMFMAC <0x6c, "v_smfmac_i32_32x32x32i8">;
1827-
let SubtargetPredicate = HasFP8Insts in {
18281819
defm V_SMFMAC_F32_16X16X64_BF8_BF8 : VOP3P_Real_SMFMAC <0x78, "v_smfmac_f32_16x16x64bf8bf8">;
18291820
defm V_SMFMAC_F32_16X16X64_BF8_FP8 : VOP3P_Real_SMFMAC <0x79, "v_smfmac_f32_16x16x64bf8fp8">;
18301821
defm V_SMFMAC_F32_16X16X64_FP8_BF8 : VOP3P_Real_SMFMAC <0x7a, "v_smfmac_f32_16x16x64fp8bf8">;
@@ -1833,7 +1824,6 @@ defm V_SMFMAC_F32_32X32X32_BF8_BF8 : VOP3P_Real_SMFMAC <0x7c, "v_smfmac_f32_32x3
18331824
defm V_SMFMAC_F32_32X32X32_BF8_FP8 : VOP3P_Real_SMFMAC <0x7d, "v_smfmac_f32_32x32x32bf8fp8">;
18341825
defm V_SMFMAC_F32_32X32X32_FP8_BF8 : VOP3P_Real_SMFMAC <0x7e, "v_smfmac_f32_32x32x32fp8bf8">;
18351826
defm V_SMFMAC_F32_32X32X32_FP8_FP8 : VOP3P_Real_SMFMAC <0x7f, "v_smfmac_f32_32x32x32fp8fp8">;
1836-
} // End SubtargetPredicate = HasFP8Insts
18371827

18381828
defm V_PK_FMA_F32 : VOP3P_Real_vi <0x30>;
18391829
defm V_PK_MUL_F32 : VOP3P_Real_vi <0x31>;

0 commit comments

Comments
 (0)