2914 lines · plain
1//===-- VOP3PInstructions.td - Vector Instruction Definitions -------------===//2//3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.4// See https://llvm.org/LICENSE.txt for license information.5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception6//7//===----------------------------------------------------------------------===//8 9//===----------------------------------------------------------------------===//10// VOP3P Classes11//===----------------------------------------------------------------------===//12 13class VOP3P_Profile<VOPProfile P, VOP3Features Features = VOP3_REGULAR,14 bit HasDPP = 0> : VOP3_Profile<P, Features> {15 let IsVOP3P = 1;16 let HasExtVOP3DPP = HasDPP;17 // We do not want to print src modifiers for vop3p because the bits are18 // overloaded in meaning and the logic in printOperandAndFPInputMods is19 // wrong for vop3p20 let AsmVOP3Base = AsmVOP3P;21 bit IsSMFMAC = false;22}23 24def VOP_MFMA_LD_SCALE : VOP3P_Profile<VOPProfile<[untyped, i32, i32, untyped]>, VOP3P_LD_SCALE> {25 let HasModifiers = 1;26 let HasNeg = 0;27}28 29// Used for FMA_MIX* and MAD_MIX* insts30// Their operands are only sort of f16 operands. Depending on31// op_sel_hi, these may be interpreted as f32. The inline immediate32// values are really f16 converted to f32, so we treat these as f1633// operands.34class VOP3P_Mix_Profile<VOPProfile P, VOP3Features Features = VOP3_REGULAR,35 bit useTiedOutput = 0> : VOP3P_Profile<P, Features, 1> {36 bit UseTiedOutput = useTiedOutput;37 38 defvar Src0RC = getVCSrcForVT<P.Src0VT>.ret;39 defvar Src1RC = getVCSrcForVT<P.Src1VT>.ret;40 defvar Src2RC = getVCSrcForVT<P.Src2VT>.ret;41 42 dag srcs =43 (ins FP16InputMods:$src0_modifiers, Src0RC:$src0,44 FP16InputMods:$src1_modifiers, Src1RC:$src1,45 FP16InputMods:$src2_modifiers, Src2RC:$src2);46 dag dpp_srcs =47 (ins FPVRegInputMods:$src0_modifiers, VGPROp_32:$src0,48 FPVRegInputMods:$src1_modifiers, VRegSrc_32:$src1,49 FP16InputMods:$src2_modifiers, Src2RC:$src2);50 51 // FIXME: Clamp0 misbehaves with the non-default vdst_in52 // following it. For now workaround this by requiring clamp53 // in tied patterns. This should use undef_tied_input, but it54 // seems underdeveloped and doesn't apply the right register55 // class constraints.56 dag mods = !con(!if(UseTiedOutput, (ins Clamp:$clamp, VGPR_32:$vdst_in),57 (ins Clamp0:$clamp)),58 (ins op_sel0:$op_sel, op_sel_hi0:$op_sel_hi));59 // We use Ins64 because that is the one which populates InOperandList60 // due to the logic in class VOP3_Pseudo61 let Ins64 = !con(srcs, mods);62 let InsVOP3Base = !con(dpp_srcs, mods);63 let AsmVOP3Base =64 "$vdst, $src0_modifiers, $src1_modifiers, $src2_modifiers$op_sel$op_sel_hi$clamp";65}66 67class VOP3P_Mix_Profile_t16<VOPProfile P, VOP3Features Features = VOP3_REGULAR>68 : VOP3P_Mix_Profile<P, Features, 0> {69 let IsTrue16 = 1;70 let IsRealTrue16 = 1;71 let DstRC64 = getVALUDstForVT<P.DstVT, 1 /*IsTrue16*/, 1 /*IsVOP3Encoding*/>.ret;72}73 74multiclass VOP3PInst<string OpName, VOPProfile P,75 SDPatternOperator node = null_frag, bit IsDOT = 0> {76 def NAME : VOP3P_Pseudo<OpName, P,77 !if (P.HasModifiers,78 getVOP3PModPat<P, node, !or(P.EnableClamp, IsDOT), IsDOT>.ret,79 getVOP3Pat<P, node>.ret)>;80 let SubtargetPredicate = isGFX11Plus in {81 if P.HasExtVOP3DPP then82 def _dpp : VOP3_DPP_Pseudo<OpName, P> {83 let VOP3P = 1;84 let PseudoInstr = OpName #"_dpp";85 }86 } // end SubtargetPredicate = isGFX11Plus87}88 89// Non-packed instructions that use the VOP3P encoding.90// VOP3 neg/abs and VOP3P opsel/opsel_hi modifiers are allowed.91multiclass VOP3_VOP3PInst<string OpName, VOP3P_Mix_Profile P> {92 def NAME : VOP3P_Pseudo<OpName, P> {93 let Constraints = !if(P.UseTiedOutput, "$vdst = $vdst_in", "");94 }95 let SubtargetPredicate = isGFX11Plus in {96 if P.HasExtVOP3DPP then97 def _dpp : VOP3_DPP_Pseudo<OpName, P> {98 let VOP3P = 1;99 let PseudoInstr = OpName#"_dpp";100 let Constraints = !if(P.UseTiedOutput, "$vdst = $vdst_in", "");101 }102 } // end SubtargetPredicate = isGFX11Plus103}104 105multiclass VOP3_VOP3PInst_t16<string OpName, VOP3P_Mix_Profile P> {106 def NAME : VOP3P_Pseudo<OpName, P>;107 108 if P.HasExtVOP3DPP then109 def _dpp : VOP3_DPP_Pseudo<OpName, P> {110 let VOP3P = 1;111 let PseudoInstr = OpName#"_dpp";112 }113}114 115let isReMaterializable = 1 in {116let isCommutable = 1 in {117defm V_PK_MAD_I16 : VOP3PInst<"v_pk_mad_i16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16_V2I16>>;118defm V_PK_MAD_U16 : VOP3PInst<"v_pk_mad_u16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16_V2I16>, imad>;119 120let FPDPRounding = 1 in {121defm V_PK_FMA_F16 : VOP3PInst<"v_pk_fma_f16", VOP3P_Profile<VOP_V2F16_V2F16_V2F16_V2F16>, any_fma>;122defm V_PK_ADD_F16 : VOP3PInst<"v_pk_add_f16", VOP3P_Profile<VOP_V2F16_V2F16_V2F16>, any_fadd>;123defm V_PK_MUL_F16 : VOP3PInst<"v_pk_mul_f16", VOP3P_Profile<VOP_V2F16_V2F16_V2F16>, any_fmul>;124} // End FPDPRounding = 1125defm V_PK_MAX_F16 : VOP3PInst<"v_pk_max_f16", VOP3P_Profile<VOP_V2F16_V2F16_V2F16>, fmaxnum_like>;126defm V_PK_MIN_F16 : VOP3PInst<"v_pk_min_f16", VOP3P_Profile<VOP_V2F16_V2F16_V2F16>, fminnum_like>;127 128defm V_PK_ADD_U16 : VOP3PInst<"v_pk_add_u16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, add>;129defm V_PK_ADD_I16 : VOP3PInst<"v_pk_add_i16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>>;130defm V_PK_MUL_LO_U16 : VOP3PInst<"v_pk_mul_lo_u16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, mul>;131 132defm V_PK_MIN_I16 : VOP3PInst<"v_pk_min_i16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, smin>;133defm V_PK_MIN_U16 : VOP3PInst<"v_pk_min_u16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, umin>;134defm V_PK_MAX_I16 : VOP3PInst<"v_pk_max_i16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, smax>;135defm V_PK_MAX_U16 : VOP3PInst<"v_pk_max_u16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, umax>;136 137let SubtargetPredicate = HasIEEEMinimumMaximumInsts, ReadsModeReg = 0 in {138defm V_PK_MAXIMUM_F16 : VOP3PInst<"v_pk_maximum_f16", VOP3P_Profile<VOP_V2F16_V2F16_V2F16, VOP3_PACKED>, fmaximum>;139defm V_PK_MINIMUM_F16 : VOP3PInst<"v_pk_minimum_f16", VOP3P_Profile<VOP_V2F16_V2F16_V2F16, VOP3_PACKED>, fminimum>;140} // End SubtargetPredicate = HasIEEEMinimumMaximumInsts, ReadsModeReg = 0141}142 143defm V_PK_SUB_U16 : VOP3PInst<"v_pk_sub_u16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>>;144defm V_PK_SUB_I16 : VOP3PInst<"v_pk_sub_i16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, sub>;145 146defm V_PK_LSHLREV_B16 : VOP3PInst<"v_pk_lshlrev_b16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, clshl_rev_16>;147defm V_PK_ASHRREV_I16 : VOP3PInst<"v_pk_ashrrev_i16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, cashr_rev_16>;148defm V_PK_LSHRREV_B16 : VOP3PInst<"v_pk_lshrrev_b16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, clshr_rev_16>;149} // End isReMaterializable = 1150 151let SubtargetPredicate = HasVOP3PInsts in {152 153// Integer operations with clamp bit set.154class VOP3PSatPat<SDPatternOperator pat, Instruction inst> : GCNPat<155 (pat (v2i16 (VOP3PMods v2i16:$src0, i32:$src0_modifiers)),156 (v2i16 (VOP3PMods v2i16:$src1, i32:$src1_modifiers))),157 (inst $src0_modifiers, $src0, $src1_modifiers, $src1, DSTCLAMP.ENABLE)158>;159 160def : VOP3PSatPat<uaddsat, V_PK_ADD_U16>;161def : VOP3PSatPat<saddsat, V_PK_ADD_I16>;162def : VOP3PSatPat<usubsat, V_PK_SUB_U16>;163def : VOP3PSatPat<ssubsat, V_PK_SUB_I16>;164} // End SubtargetPredicate = HasVOP3PInsts165 166let isCommutable = 1, FPDPRounding = 1 in {167let SubtargetPredicate = HasMin3Max3PKF16 in {168defm V_PK_MIN3_NUM_F16 : VOP3PInst<"v_pk_min3_num_f16", VOP3P_Profile<VOP_V2F16_V2F16_V2F16_V2F16>, AMDGPUfmin3>;169defm V_PK_MAX3_NUM_F16 : VOP3PInst<"v_pk_max3_num_f16", VOP3P_Profile<VOP_V2F16_V2F16_V2F16_V2F16>, AMDGPUfmax3>;170}171 172let SubtargetPredicate = HasMinimum3Maximum3PKF16 in {173defm V_PK_MINIMUM3_F16 : VOP3PInst<"v_pk_minimum3_f16", VOP3P_Profile<VOP_V2F16_V2F16_V2F16_V2F16>, AMDGPUfminimum3>;174defm V_PK_MAXIMUM3_F16 : VOP3PInst<"v_pk_maximum3_f16", VOP3P_Profile<VOP_V2F16_V2F16_V2F16_V2F16>, AMDGPUfmaximum3>;175}176} // End isCommutable = 1, FPDPRounding = 1177 178// TODO: Make sure we're doing the right thing with denormals. Note179// that FMA and MAD will differ.180multiclass MadFmaMixFP32Pats<SDPatternOperator fma_like,181 Instruction mix_inst,182 ValueType VT = f16> {183 defvar VOP3PMadMixModsPat = !if (!eq(VT, bf16), VOP3PMadMixBF16Mods, VOP3PMadMixMods);184 defvar VOP3PMadMixModsExtPat = !if (!eq(VT, bf16), VOP3PMadMixBF16ModsExt, VOP3PMadMixModsExt);185 // At least one of the operands needs to be an fpextend of an f16186 // for this to be worthwhile, so we need three patterns here.187 // TODO: Could we use a predicate to inspect src1/2/3 instead?188 def : GCNPat <189 (f32 (fma_like (f32 (VOP3PMadMixModsExtPat VT:$src0, i32:$src0_mods)),190 (f32 (VOP3PMadMixModsPat VT:$src1, i32:$src1_mods)),191 (f32 (VOP3PMadMixModsPat VT:$src2, i32:$src2_mods)))),192 (mix_inst $src0_mods, $src0, $src1_mods, $src1, $src2_mods, $src2,193 DSTCLAMP.NONE)>;194 def : GCNPat <195 (f32 (fma_like (f32 (VOP3PMadMixModsPat VT:$src0, i32:$src0_mods)),196 (f32 (VOP3PMadMixModsExtPat VT:$src1, i32:$src1_mods)),197 (f32 (VOP3PMadMixModsPat f32:$src2, i32:$src2_mods)))),198 (mix_inst $src0_mods, $src0, $src1_mods, $src1, $src2_mods, $src2,199 DSTCLAMP.NONE)>;200 def : GCNPat <201 (f32 (fma_like (f32 (VOP3PMadMixModsPat VT:$src0, i32:$src0_mods)),202 (f32 (VOP3PMadMixModsPat f32:$src1, i32:$src1_mods)),203 (f32 (VOP3PMadMixModsExtPat VT:$src2, i32:$src2_mods)))),204 (mix_inst $src0_mods, $src0, $src1_mods, $src1, $src2_mods, $src2,205 DSTCLAMP.NONE)>;206}207 208multiclass MadFmaMixFP16Pats<SDPatternOperator fma_like,209 Instruction mixlo_inst,210 Instruction mixhi_inst,211 ValueType VT = f16,212 ValueType vecVT = v2f16> {213 defvar VOP3PMadMixModsPat = !if (!eq(VT, bf16), VOP3PMadMixBF16Mods, VOP3PMadMixMods);214 def : GCNPat <215 (AMDGPUclamp (build_vector216 (VT (fpround (fma_like (f32 (VOP3PMadMixModsPat VT:$lo_src0, i32:$lo_src0_modifiers)),217 (f32 (VOP3PMadMixModsPat VT:$lo_src1, i32:$lo_src1_modifiers)),218 (f32 (VOP3PMadMixModsPat VT:$lo_src2, i32:$lo_src2_modifiers))))),219 (VT (fpround (fma_like (f32 (VOP3PMadMixModsPat VT:$hi_src0, i32:$hi_src0_modifiers)),220 (f32 (VOP3PMadMixModsPat VT:$hi_src1, i32:$hi_src1_modifiers)),221 (f32 (VOP3PMadMixModsPat VT:$hi_src2, i32:$hi_src2_modifiers))))))),222 (vecVT (mixhi_inst $hi_src0_modifiers, $hi_src0,223 $hi_src1_modifiers, $hi_src1,224 $hi_src2_modifiers, $hi_src2,225 DSTCLAMP.ENABLE,226 (mixlo_inst $lo_src0_modifiers, $lo_src0,227 $lo_src1_modifiers, $lo_src1,228 $lo_src2_modifiers, $lo_src2,229 DSTCLAMP.ENABLE,230 (i32 (IMPLICIT_DEF)))))231 >;232 233 def : GCNPat <234 (VT (fpround (fmul (f32 (VOP3PMadMixModsPat f32:$src0, i32:$src0_modifiers)),235 (f32 (VOP3PMadMixModsPat f32:$src1, i32:$src1_modifiers))))),236 (mixlo_inst $src0_modifiers, $src0,237 $src1_modifiers, $src1,238 (i32 0), (i32 0),239 DSTCLAMP.NONE,240 (i32 (IMPLICIT_DEF)))241 >;242 243 def : GCNPat <244 (build_vector VT:$elt0, (VT (fpround (fmul (f32 (VOP3PMadMixModsPat f32:$src0, i32:$src0_modifiers)),245 (f32 (VOP3PMadMixModsPat f32:$src1, i32:$src1_modifiers)))))),246 (vecVT (mixhi_inst $src0_modifiers, $src0,247 $src1_modifiers, $src1,248 (i32 0), (i32 0),249 DSTCLAMP.NONE,250 VGPR_32:$elt0))251 >;252 253 def : GCNPat <254 (VT (fpround (fma_like (f32 (VOP3PMadMixModsPat VT:$src0, i32:$src0_modifiers)),255 (f32 (VOP3PMadMixModsPat VT:$src1, i32:$src1_modifiers)),256 (f32 (VOP3PMadMixModsPat VT:$src2, i32:$src2_modifiers))))),257 (mixlo_inst $src0_modifiers, $src0,258 $src1_modifiers, $src1,259 $src2_modifiers, $src2,260 DSTCLAMP.NONE,261 (i32 (IMPLICIT_DEF)))262 >;263 264 // FIXME: Special case handling for maxhi (especially for clamp)265 // because dealing with the write to high half of the register is266 // difficult.267 def : GCNPat <268 (build_vector VT:$elt0, (VT (fpround (fma_like (f32 (VOP3PMadMixModsPat VT:$src0, i32:$src0_modifiers)),269 (f32 (VOP3PMadMixModsPat VT:$src1, i32:$src1_modifiers)),270 (f32 (VOP3PMadMixModsPat VT:$src2, i32:$src2_modifiers)))))),271 (vecVT (mixhi_inst $src0_modifiers, $src0,272 $src1_modifiers, $src1,273 $src2_modifiers, $src2,274 DSTCLAMP.NONE,275 VGPR_32:$elt0))276 >;277 278 def : GCNPat <279 (build_vector280 VT:$elt0,281 (AMDGPUclamp (VT (fpround (fma_like (f32 (VOP3PMadMixModsPat VT:$src0, i32:$src0_modifiers)),282 (f32 (VOP3PMadMixModsPat VT:$src1, i32:$src1_modifiers)),283 (f32 (VOP3PMadMixModsPat VT:$src2, i32:$src2_modifiers))))))),284 (vecVT (mixhi_inst $src0_modifiers, $src0,285 $src1_modifiers, $src1,286 $src2_modifiers, $src2,287 DSTCLAMP.ENABLE,288 VGPR_32:$elt0))289 >;290}291 292multiclass MadFmaMixFP16Pats_t16<SDPatternOperator fma_like,293 Instruction mix_inst_16,294 ValueType VT = f16,295 ValueType vecVT = v2f16> {296 defvar VOP3PMadMixModsPat = !if (!eq(VT, bf16), VOP3PMadMixBF16Mods, VOP3PMadMixMods);297 def : GCNPat <298 (VT (fpround (fmul (f32 (VOP3PMadMixModsPat f32:$src0, i32:$src0_modifiers)),299 (f32 (VOP3PMadMixModsPat f32:$src1, i32:$src1_modifiers))))),300 (mix_inst_16 $src0_modifiers, $src0,301 $src1_modifiers, $src1,302 (i32 0), (i32 0),303 DSTCLAMP.NONE)304 >;305 306 def : GCNPat <307 (VT (fpround (fma_like (f32 (VOP3PMadMixModsPat VT:$src0, i32:$src0_modifiers)),308 (f32 (VOP3PMadMixModsPat VT:$src1, i32:$src1_modifiers)),309 (f32 (VOP3PMadMixModsPat VT:$src2, i32:$src2_modifiers))))),310 (mix_inst_16 $src0_modifiers, $src0,311 $src1_modifiers, $src1,312 $src2_modifiers, $src2,313 DSTCLAMP.NONE)314 >;315 316 317 def : GCNPat <318 (AMDGPUclamp (VT (fpround (fma_like (f32 (VOP3PMadMixModsPat VT:$src0, i32:$src0_modifiers)),319 (f32 (VOP3PMadMixModsPat VT:$src1, i32:$src1_modifiers)),320 (f32 (VOP3PMadMixModsPat VT:$src2, i32:$src2_modifiers)))))),321 (mix_inst_16 $src0_modifiers, $src0,322 $src1_modifiers, $src1,323 $src2_modifiers, $src2,324 DSTCLAMP.ENABLE)325 >;326 327 def : GCNPat <328 (AMDGPUclamp (build_vector329 (VT (fpround (fma_like (f32 (VOP3PMadMixModsPat VT:$lo_src0, i32:$lo_src0_modifiers)),330 (f32 (VOP3PMadMixModsPat VT:$lo_src1, i32:$lo_src1_modifiers)),331 (f32 (VOP3PMadMixModsPat VT:$lo_src2, i32:$lo_src2_modifiers))))),332 (VT (fpround (fma_like (f32 (VOP3PMadMixModsPat VT:$hi_src0, i32:$hi_src0_modifiers)),333 (f32 (VOP3PMadMixModsPat VT:$hi_src1, i32:$hi_src1_modifiers)),334 (f32 (VOP3PMadMixModsPat VT:$hi_src2, i32:$hi_src2_modifiers))))))),335 (vecVT (REG_SEQUENCE VGPR_32, (mix_inst_16 $lo_src0_modifiers, $lo_src0,336 $lo_src1_modifiers, $lo_src1,337 $lo_src2_modifiers, $lo_src2,338 DSTCLAMP.ENABLE), lo16,339 (mix_inst_16 $hi_src0_modifiers, $hi_src0,340 $hi_src1_modifiers, $hi_src1,341 $hi_src2_modifiers, $hi_src2,342 DSTCLAMP.ENABLE), hi16))343 >;344}345 346class MinimumMaximumByMinimum3Maximum3VOP3P<SDPatternOperator node,347 Instruction inst> : GCNPat<348 (v2f16 (node (VOP3PMods v2f16:$src0, i32:$src0_mods), (VOP3PMods v2f16:$src1, i32:$src1_mods))),349 (inst $src0_mods, $src0, $src1_mods, $src1, $src1_mods, $src1)350>;351 352let SubtargetPredicate = HasMinimum3Maximum3PKF16 in {353def : MinimumMaximumByMinimum3Maximum3VOP3P<fminimum, V_PK_MINIMUM3_F16>;354def : MinimumMaximumByMinimum3Maximum3VOP3P<fmaximum, V_PK_MAXIMUM3_F16>;355}356 357let SubtargetPredicate = HasMadMixInsts in {358let OtherPredicates = [NoFP32Denormals] in {359 360// These are VOP3a-like opcodes which accept no omod.361// Size of src arguments (16/32) is controlled by op_sel.362// For 16-bit src arguments their location (hi/lo) are controlled by op_sel_hi.363let isCommutable = 1, mayRaiseFPException = 0 in {364let isReMaterializable = 1 in365defm V_MAD_MIX_F32 : VOP3_VOP3PInst<"v_mad_mix_f32", VOP3P_Mix_Profile<VOP_F32_F16_F16_F16, VOP3_OPSEL>>;366 367let FPDPRounding = 1 in {368// Clamp modifier is applied after conversion to f16.369defm V_MAD_MIXLO_F16 : VOP3_VOP3PInst<"v_mad_mixlo_f16", VOP3P_Mix_Profile<VOP_F16_F16_F16_F16, VOP3_OPSEL, 1>>;370 371let ClampLo = 0, ClampHi = 1 in {372defm V_MAD_MIXHI_F16 : VOP3_VOP3PInst<"v_mad_mixhi_f16", VOP3P_Mix_Profile<VOP_F16_F16_F16_F16, VOP3_OPSEL, 1>>;373}374} // End FPDPRounding = 1375}376 377defm : MadFmaMixFP32Pats<fmad, V_MAD_MIX_F32>;378defm : MadFmaMixFP16Pats<fmad, V_MAD_MIXLO_F16, V_MAD_MIXHI_F16>;379} // OtherPredicates = [NoFP32Denormals]380} // End SubtargetPredicate = HasMadMixInsts381 382 383// Essentially the same as the mad_mix versions384let SubtargetPredicate = HasFmaMixInsts in {385let isCommutable = 1 in {386 387let isReMaterializable = 1 in388defm V_FMA_MIX_F32 : VOP3_VOP3PInst<"v_fma_mix_f32", VOP3P_Mix_Profile<VOP_F32_F16_F16_F16, VOP3_OPSEL>>;389 390let FPDPRounding = 1 in {391// Clamp modifier is applied after conversion to f16.392defm V_FMA_MIXLO_F16 : VOP3_VOP3PInst<"v_fma_mixlo_f16", VOP3P_Mix_Profile<VOP_F16_F16_F16_F16, VOP3_OPSEL, 1>>;393 394let ClampLo = 0, ClampHi = 1 in {395defm V_FMA_MIXHI_F16 : VOP3_VOP3PInst<"v_fma_mixhi_f16", VOP3P_Mix_Profile<VOP_F16_F16_F16_F16, VOP3_OPSEL, 1>>;396}397 398// Pseudo true16 inst for v_fma_mixlo/hi_f16399defm V_FMA_MIX_F16_t16 : VOP3_VOP3PInst_t16<"v_fma_mix_f16_t16", VOP3P_Mix_Profile_t16<VOP_F16_F16_F16_F16, VOP3_OPSEL>>;400} // End FPDPRounding = 1401}402 403defm : MadFmaMixFP32Pats<fma, V_FMA_MIX_F32>;404 405let True16Predicate = NotUseRealTrue16Insts in406defm : MadFmaMixFP16Pats<fma, V_FMA_MIXLO_F16, V_FMA_MIXHI_F16>;407let True16Predicate = UseRealTrue16Insts in408defm : MadFmaMixFP16Pats_t16<fma, V_FMA_MIX_F16_t16>;409}410 411let SubtargetPredicate = HasFmaMixBF16Insts in {412let isCommutable = 1 in {413 414let isReMaterializable = 1 in415defm V_FMA_MIX_F32_BF16 : VOP3_VOP3PInst<"v_fma_mix_f32_bf16", VOP3P_Mix_Profile<VOP_F32_BF16_BF16_BF16, VOP3_OPSEL>>;416 417let FPDPRounding = 1 in {418defm V_FMA_MIXLO_BF16 : VOP3_VOP3PInst<"v_fma_mixlo_bf16", VOP3P_Mix_Profile<VOP_BF16_BF16_BF16_BF16, VOP3_OPSEL, 1>>;419 420let ClampLo = 0, ClampHi = 1 in {421defm V_FMA_MIXHI_BF16 : VOP3_VOP3PInst<"v_fma_mixhi_bf16", VOP3P_Mix_Profile<VOP_BF16_BF16_BF16_BF16, VOP3_OPSEL, 1>>;422}423 424// Pseudo true16 inst for v_fma_mixlo/hi_bf16425defm V_FMA_MIX_BF16_t16 : VOP3_VOP3PInst_t16<"v_fma_mix_bf16_t16", VOP3P_Mix_Profile_t16<VOP_F16_F16_F16_F16, VOP3_OPSEL>>;426} // End FPDPRounding = 1427} // End isCommutable = 1428 429defm : MadFmaMixFP32Pats<fma, V_FMA_MIX_F32_BF16, bf16>;430let True16Predicate = NotUseRealTrue16Insts in431defm : MadFmaMixFP16Pats<fma, V_FMA_MIXLO_BF16, V_FMA_MIXHI_BF16, bf16, v2bf16>;432let True16Predicate = UseRealTrue16Insts in433defm : MadFmaMixFP16Pats_t16<fma, V_FMA_MIX_BF16_t16>;434} // End SubtargetPredicate = HasFmaMixBF16Insts435 436def PK_ADD_MINMAX_Profile : VOP3P_Profile<VOP_V2I16_V2I16_V2I16_V2I16, VOP3_PACKED> {437 let HasNeg = 0;438 let EnableClamp = 1;439}440 441let isCommutable = 1, isReMaterializable = 1 in {442let SubtargetPredicate = HasPkAddMinMaxInsts in {443defm V_PK_ADD_MAX_I16 : VOP3PInst<"v_pk_add_max_i16", PK_ADD_MINMAX_Profile, int_amdgcn_pk_add_max_i16>;444defm V_PK_ADD_MAX_U16 : VOP3PInst<"v_pk_add_max_u16", PK_ADD_MINMAX_Profile, int_amdgcn_pk_add_max_u16>;445defm V_PK_ADD_MIN_I16 : VOP3PInst<"v_pk_add_min_i16", PK_ADD_MINMAX_Profile, int_amdgcn_pk_add_min_i16>;446defm V_PK_ADD_MIN_U16 : VOP3PInst<"v_pk_add_min_u16", PK_ADD_MINMAX_Profile, int_amdgcn_pk_add_min_u16>;447}448let SubtargetPredicate = HasPkMinMax3Insts in {449defm V_PK_MAX3_I16 : VOP3PInst<"v_pk_max3_i16", PK_ADD_MINMAX_Profile>;450defm V_PK_MAX3_U16 : VOP3PInst<"v_pk_max3_u16", PK_ADD_MINMAX_Profile>;451defm V_PK_MIN3_I16 : VOP3PInst<"v_pk_min3_i16", PK_ADD_MINMAX_Profile>;452defm V_PK_MIN3_U16 : VOP3PInst<"v_pk_min3_u16", PK_ADD_MINMAX_Profile>;453}454} // End isCommutable = 1, isReMaterializable = 1455 456// TODO: Extend pattern to select op_sel and op_sel_hi.457class ThreeOp_OpSelClampPats <SDPatternOperator op1, SDPatternOperator op2,458 VOP3P_Pseudo inst,459 ValueType vt = inst.Pfl.Src0VT,460 RegisterOperand RC = getVCSrcForVT<vt>.ret> : GCNPat <461 (ThreeOpFrag<op1, op2> vt:$src0, vt:$src1, vt:$src2),462 (inst SRCMODS.OP_SEL_1, RC:$src0, SRCMODS.OP_SEL_1, RC:$src1,463 SRCMODS.OP_SEL_1, RC:$src2, DSTCLAMP.NONE, 0)464>;465 466let SubtargetPredicate = HasPkAddMinMaxInsts in {467def : ThreeOp_OpSelClampPats<saddsat, smax, V_PK_ADD_MAX_I16>;468def : ThreeOp_OpSelClampPats<uaddsat, umax, V_PK_ADD_MAX_U16>;469def : ThreeOp_OpSelClampPats<saddsat, smin, V_PK_ADD_MIN_I16>;470def : ThreeOp_OpSelClampPats<uaddsat, umin, V_PK_ADD_MIN_U16>;471}472 473let SubtargetPredicate = HasPkMinMax3Insts in {474def : ThreeOp_OpSelClampPats<smax, smax, V_PK_MAX3_I16>;475def : ThreeOp_OpSelClampPats<umax, umax, V_PK_MAX3_U16>;476def : ThreeOp_OpSelClampPats<smin, smin, V_PK_MIN3_I16>;477def : ThreeOp_OpSelClampPats<umin, umin, V_PK_MIN3_U16>;478}479 480// Defines patterns that extract signed 4bit from each Idx[0].481foreach Idx = [[0,28],[4,24],[8,20],[12,16],[16,12],[20,8],[24,4]] in482 def ExtractSigned4bit_#Idx[0] : PatFrag<(ops node:$src),483 (sra (shl node:$src, (i32 Idx[1])), (i32 28))>;484 485// Defines code pattern that extracts U(unsigned/signed) 4/8bit from FromBitIndex.486class Extract<int FromBitIndex, int BitMask, bit U>: PatFrag<487 (ops node:$src),488 !if (!or (!and (!eq (BitMask, 255), !eq (FromBitIndex, 24)), !eq (FromBitIndex, 28)), // last element489 !if (U, (srl node:$src, (i32 FromBitIndex)), (sra node:$src, (i32 FromBitIndex))),490 !if (!eq (FromBitIndex, 0), // first element491 !if (U, (and node:$src, (i32 BitMask)),492 !if (!eq (BitMask, 15), (!cast<PatFrag>("ExtractSigned4bit_"#FromBitIndex) node:$src),493 (sext_inreg node:$src, i8))),494 !if (U, (and (srl node:$src, (i32 FromBitIndex)), (i32 BitMask)),495 !if (!eq (BitMask, 15), (!cast<PatFrag>("ExtractSigned4bit_"#FromBitIndex) node:$src),496 (sext_inreg (srl node:$src, (i32 FromBitIndex)), i8)))))>;497 498 499foreach Type = ["I", "U"] in500 foreach Index = 0-3 in {501 // Defines patterns that extract each Index'ed 8bit from an unsigned502 // 32bit scalar value;503 def Type#Index#"_8bit" : Extract<!shl(Index, 3), 255, !eq (Type, "U")>;504 505 // Defines multiplication patterns where the multiplication is happening on each506 // Index'ed 8bit of a 32bit scalar value.507 508 def Mul#Type#_Elt#Index : PatFrag<509 (ops node:$src0, node:$src1),510 (!cast<HasOneUseBinOp>(!if (!eq (Type, "I"), AMDGPUmul_i24_oneuse, AMDGPUmul_u24_oneuse))511 (!cast<Extract>(Type#Index#"_8bit") node:$src0),512 (!cast<Extract>(Type#Index#"_8bit") node:$src1))>;513 }514 515// Different variants of dot8 patterns cause a huge increase in the compile time.516// Define non-associative/commutative add/mul to prevent permutation in the dot8517// pattern.518def NonACAdd : SDNode<"ISD::ADD" , SDTIntBinOp>;519def NonACAdd_oneuse : HasOneUseBinOp<NonACAdd>;520 521def NonACAMDGPUmul_u24 : SDNode<"AMDGPUISD::MUL_U24" , SDTIntBinOp>;522def NonACAMDGPUmul_u24_oneuse : HasOneUseBinOp<NonACAMDGPUmul_u24>;523 524def NonACAMDGPUmul_i24 : SDNode<"AMDGPUISD::MUL_I24" , SDTIntBinOp>;525def NonACAMDGPUmul_i24_oneuse : HasOneUseBinOp<NonACAMDGPUmul_i24>;526 527foreach Type = ["I", "U"] in528 foreach Index = 0-7 in {529 // Defines patterns that extract each Index'ed 4bit from an unsigned530 // 32bit scalar value;531 def Type#Index#"_4bit" : Extract<!shl(Index, 2), 15, !eq (Type, "U")>;532 533 // Defines multiplication patterns where the multiplication is happening on each534 // Index'ed 8bit of a 32bit scalar value.535 def Mul#Type#Index#"_4bit" : PatFrag<536 (ops node:$src0, node:$src1),537 (!cast<HasOneUseBinOp>(!if (!eq (Type, "I"), NonACAMDGPUmul_i24_oneuse, NonACAMDGPUmul_u24_oneuse))538 (!cast<Extract>(Type#Index#"_4bit") node:$src0),539 (!cast<Extract>(Type#Index#"_4bit") node:$src1))>;540 }541 542class UDot2Pat<VOP_Pseudo Inst> : GCNPat <543 (add (add_oneuse (AMDGPUmul_u24_oneuse (srl i32:$src0, (i32 16)),544 (srl i32:$src1, (i32 16))), i32:$src2),545 (AMDGPUmul_u24_oneuse (and i32:$src0, (i32 65535)),546 (and i32:$src1, (i32 65535)))547 ),548 (Inst (i32 8), $src0, (i32 8), $src1, (i32 8), $src2, (i1 0))> {549 let Predicates = Inst.Predicates;550}551 552class SDot2Pat<VOP_Pseudo Inst> : GCNPat <553 (add (add_oneuse (AMDGPUmul_i24_oneuse (sra i32:$src0, (i32 16)),554 (sra i32:$src1, (i32 16))), i32:$src2),555 (AMDGPUmul_i24_oneuse (sext_inreg i32:$src0, i16),556 (sext_inreg i32:$src1, i16))),557 (Inst (i32 8), $src0, (i32 8), $src1, (i32 8), $src2, (i1 0))> {558 let Predicates = Inst.Predicates;559}560 561let IsDOT = 1 in {562let OtherPredicates = [HasDot2Insts] in {563defm V_DOT2_I32_I16 : VOP3PInst<"v_dot2_i32_i16",564 VOP3P_Profile<VOP_I32_V2I16_V2I16_I32>, int_amdgcn_sdot2, 1>;565defm V_DOT2_U32_U16 : VOP3PInst<"v_dot2_u32_u16",566 VOP3P_Profile<VOP_I32_V2I16_V2I16_I32>, int_amdgcn_udot2, 1>;567} // End OtherPredicates = [HasDot2Insts]568 569let OtherPredicates = [HasDot10Insts] in570defm V_DOT2_F32_F16 : VOP3PInst<"v_dot2_f32_f16",571 VOP3P_Profile<VOP_F32_V2F16_V2F16_F32, VOP3_REGULAR, /*HasDPP*/ 1>,572 AMDGPUfdot2, 1/*ExplicitClamp*/>;573 574let OtherPredicates = [HasDot7Insts] in {575defm V_DOT4_U32_U8 : VOP3PInst<"v_dot4_u32_u8",576 VOP3P_Profile<VOP_I32_I32_I32_I32, VOP3_PACKED_NO_OPSEL>, int_amdgcn_udot4, 1>;577defm V_DOT8_U32_U4 : VOP3PInst<"v_dot8_u32_u4",578 VOP3P_Profile<VOP_I32_I32_I32_I32, VOP3_PACKED_NO_OPSEL>, int_amdgcn_udot8, 1>;579} // End OtherPredicates = [HasDot7Insts]580 581let OtherPredicates = [HasDot1Insts] in {582defm V_DOT4_I32_I8 : VOP3PInst<"v_dot4_i32_i8",583 VOP3P_Profile<VOP_I32_I32_I32_I32, VOP3_PACKED_NO_OPSEL>, int_amdgcn_sdot4, 1>;584defm V_DOT8_I32_I4 : VOP3PInst<"v_dot8_i32_i4",585 VOP3P_Profile<VOP_I32_I32_I32_I32, VOP3_PACKED_NO_OPSEL>, int_amdgcn_sdot8, 1>;586} // End OtherPredicates = [HasDot1Insts]587 588def DOT2_BF16_Profile589 : VOP3P_Profile<VOP_F32_V2BF16_V2BF16_F32, VOP3_REGULAR, /*HasDPP*/ 1> {590 let HasSrc1Mods = 1;591}592 593let SubtargetPredicate = HasDot12Insts in {594 595defm V_DOT2_F32_BF16 : VOP3PInst<"v_dot2_f32_bf16", DOT2_BF16_Profile,596 int_amdgcn_fdot2_f32_bf16, 1>;597 598} // End SubtargetPredicate = HasDot12Insts599 600} // End let IsDOT = 1601 602multiclass VOP3PDOTIUInst <string OpName, SDPatternOperator intrinsic_node> {603 let IsDOT = 1 in604 defm NAME : VOP3PInst<OpName, VOP3P_Profile<VOP_I32_I32_I32_I32, VOP3_PACKED_NO_OPSEL>,605 null_frag, 1>;606 // Dot-iu instructions consider input as signed if imod neg bits are set. Thus607 // Dot-iu Intrinsics have extra operands and require separate codegen pattern.608 def : GCNPat < (intrinsic_node timm:$src0_mods, i32:$src0,609 timm:$src1_mods, i32:$src1,610 i32:$src2, (i1 timm:$clamp)),611 (!cast<Instruction>(NAME) (VOP3PModsNeg $src0_mods), i32:$src0,612 (VOP3PModsNeg $src1_mods), i32:$src1,613 (i32 8), i32:$src2, i1:$clamp)614 >;615}616 617let SubtargetPredicate = HasDot8Insts in {618defm V_DOT4_I32_IU8 : VOP3PDOTIUInst<"v_dot4_i32_iu8", int_amdgcn_sudot4>;619defm V_DOT8_I32_IU4 : VOP3PDOTIUInst<"v_dot8_i32_iu4", int_amdgcn_sudot8>;620 621def : GCNPat < (int_amdgcn_sdot8 i32:$src0,622 i32:$src1,623 i32:$src2, (i1 timm:$clamp)),624 (V_DOT8_I32_IU4 (i32 9), i32:$src0,625 (i32 9), i32:$src1, (i32 8), i32:$src2, i1:$clamp)626>;627 628def : GCNPat < (int_amdgcn_sdot4 i32:$src0,629 i32:$src1,630 i32:$src2, (i1 timm:$clamp)),631 (V_DOT4_I32_IU8 (i32 9), i32:$src0,632 (i32 9), i32:$src1, (i32 8), i32:$src2, i1:$clamp)633>;634} // End SubtargetPredicate = HasDot8Insts635 636// Does not use opsel, no src_modifiers on src0 and src1.637// src_modifiers on src2(f32) are f32 fneg(neg_lo[2]) and f32 fabs(neg_hi[2]).638def VOP3P_DOTF8_Profile : VOP3P_Profile<VOPProfile <[f32, i32, i32, f32]>,639 VOP3_PACKED, 1> {640 let HasClamp = 0;641 let HasOpSel = 0;642 let HasOMod = 0;643 let IsDOT = 1;644 let HasSrc0Mods = 0;645 let HasSrc1Mods = 0;646 let HasSrc2Mods = 1;647 648 let InsVOP3P = (ins VSrc_b32:$src0, VSrc_b32:$src1,649 PackedF16InputMods:$src2_modifiers, VSrc_f32:$src2,650 neg_lo0:$neg_lo, neg_hi0:$neg_hi);651 652 let InsVOP3DPP8 = (ins DstRC:$old, VGPR_32:$src0, VRegSrc_32:$src1,653 PackedF16InputMods:$src2_modifiers, VRegSrc_32:$src2,654 neg_lo0:$neg_lo, neg_hi0:$neg_hi, dpp8:$dpp8, Dpp8FI:$fi);655 656 let InsVOP3DPP16 = (ins DstRC:$old, VGPR_32:$src0, VRegSrc_32:$src1,657 PackedF16InputMods:$src2_modifiers, VRegSrc_32:$src2,658 neg_lo0:$neg_lo, neg_hi0:$neg_hi, dpp_ctrl:$dpp_ctrl,659 DppRowMask:$row_mask, DppBankMask:$bank_mask,660 DppBoundCtrl:$bound_ctrl, Dpp16FI:$fi);661}662 663multiclass VOP3PDOTF8Inst <string OpName, SDPatternOperator intrinsic_node> {664 defm NAME : VOP3PInst<OpName, VOP3P_DOTF8_Profile, null_frag, 1>;665 666 let SubtargetPredicate = isGFX12Plus in667 def : GCNPat <(intrinsic_node i32:$src0, i32:$src1,668 (VOP3Mods f32:$src2, i32:$src2_modifiers)),669 (!cast<Instruction>(NAME) i32:$src0, i32:$src1,670 i32:$src2_modifiers, f32:$src2)>;671}672 673let OtherPredicates = [HasDot11Insts] in {674defm V_DOT4_F32_FP8_BF8 : VOP3PDOTF8Inst<"v_dot4_f32_fp8_bf8", int_amdgcn_dot4_f32_fp8_bf8>;675defm V_DOT4_F32_BF8_FP8 : VOP3PDOTF8Inst<"v_dot4_f32_bf8_fp8", int_amdgcn_dot4_f32_bf8_fp8>;676defm V_DOT4_F32_FP8_FP8 : VOP3PDOTF8Inst<"v_dot4_f32_fp8_fp8", int_amdgcn_dot4_f32_fp8_fp8>;677defm V_DOT4_F32_BF8_BF8 : VOP3PDOTF8Inst<"v_dot4_f32_bf8_bf8", int_amdgcn_dot4_f32_bf8_bf8>;678}679 680def : UDot2Pat<V_DOT2_U32_U16>;681def : SDot2Pat<V_DOT2_I32_I16>;682 683foreach Type = ["U", "I"] in684 let Predicates = !cast<VOP_Pseudo>("V_DOT4_"#Type#"32_"#Type#8).Predicates in685 def : GCNPat <686 !cast<dag>(!foldl((i32 i32:$src2), [0, 1, 2, 3], lhs, y,687 (add_oneuse lhs, (!cast<PatFrag>("Mul"#Type#"_Elt"#y) i32:$src0, i32:$src1)))),688 (!cast<VOP3P_Pseudo>("V_DOT4_"#Type#"32_"#Type#8) (i32 8), $src0, (i32 8), $src1, (i32 8), $src2, (i1 0))>;689 690foreach Type = ["U", "I"] in691 let Predicates = !cast<VOP_Pseudo>("V_DOT8_"#Type#"32_"#Type#4).Predicates in692 def : GCNPat <693 !cast<dag>(!foldl((add_oneuse i32:$src2, (!cast<PatFrag>("Mul"#Type#"0_4bit") i32:$src0, i32:$src1)),694 [1, 2, 3, 4, 5, 6, 7], lhs, y,695 (NonACAdd_oneuse lhs, (!cast<PatFrag>("Mul"#Type#y#"_4bit") i32:$src0, i32:$src1)))),696 (!cast<VOP3P_Pseudo>("V_DOT8_"#Type#"32_"#Type#4) (i32 8), $src0, (i32 8), $src1, (i32 8), $src2, (i1 0))>;697 698// Different variants of dot8 code-gen dag patterns are not generated through table-gen due to a huge increase699// in the compile time. Directly handle the pattern generated by the FE here.700foreach Type = ["U", "I"] in701 let Predicates = !cast<VOP_Pseudo>("V_DOT8_"#Type#"32_"#Type#4).Predicates in702 def : GCNPat <703 !cast<dag>(!foldl((add_oneuse i32:$src2, (!cast<PatFrag>("Mul"#Type#"0_4bit") i32:$src0, i32:$src1)),704 [7, 1, 2, 3, 4, 5, 6], lhs, y,705 (NonACAdd_oneuse lhs, (!cast<PatFrag>("Mul"#Type#y#"_4bit") i32:$src0, i32:$src1)))),706 (!cast<VOP3P_Pseudo>("V_DOT8_"#Type#"32_"#Type#4) (i32 8), $src0, (i32 8), $src1, (i32 8), $src2, (i1 0))>;707 708def ADst_32 : VOPDstOperand<AGPR_32>;709def ADst_64 : VOPDstOperand<AReg_64_AlignTarget>;710def ADst_128 : VOPDstOperand<AReg_128_AlignTarget>;711def ADst_256 : VOPDstOperand<AReg_256_AlignTarget>;712def ADst_512 : VOPDstOperand<AReg_512_AlignTarget>;713def ADst_1024 : VOPDstOperand<AReg_1024_AlignTarget>;714def VDst_64 : VOPDstOperand<VReg_64_AlignTarget>;715def VDst_128 : VOPDstOperand<VReg_128_AlignTarget>;716def VDst_256 : VOPDstOperand<VReg_256_AlignTarget>;717def VDst_512 : VOPDstOperand<VReg_512_AlignTarget>;718def VDst_1024 : VOPDstOperand<VReg_1024_AlignTarget>;719 720def VOPProfileAccRead : VOP3P_Profile<VOP_I32_I32, VOP3_MAI> {721 let Src0RC64 = ARegSrc_32;722}723 724def VOPProfileAccWrite : VOP3P_Profile<VOP_I32_I32, VOP3_MAI> {725 let DstRC = ADst_32;726 let Src0RC64 = VCSrc_b32;727}728 729class VOPProfileMAI<VOPProfile P, RegisterOperand _SrcRC, RegisterOperand _DstRC,730 RegisterOperand SrcARC = AVSrc_32, RegisterOperand SrcBRC = SrcARC>731 : VOP3P_Profile<P, VOP3_MAI> {732 bit HasAbid = true;733 let DstRC = _DstRC;734 let Src0RC64 = SrcARC;735 let Src1RC64 = SrcBRC;736 let Src2RC64 = _SrcRC;737 let HasOpSel = 0;738 let HasClamp = 0;739 let HasIntClamp = 0;740 let HasOMod = 0;741 let HasModifiers = 0;742 let AsmVOP3Base = "$vdst, $src0, $src1, $src2$cbsz"#!if(HasAbid,"$abid","")#"$blgp";743 let Ins64 = !con(744 (ins Src0RC64:$src0, Src1RC64:$src1, Src2RC64:$src2, CBSZ:$cbsz),745 !if(HasAbid, (ins ABID:$abid), (ins)),746 (ins blgp:$blgp));747 let InsVOP3Base = Ins64;748 // Dst and SrcC cannot partially overlap if SrcC/Dst is bigger than 4 VGPRs.749 // We then create two versions of the instruction: with tied dst and src2750 // and with the earlyclobber flag on the dst. This is stricter than the751 // actual HW restriction. In particular earlyclobber also affects src0 and752 // src1 allocation which is not required.753 bit NoDstOverlap = !gt(DstVT.Size, 128);754}755 756class VOPProfileSMFMAC<VOPProfile P, RegisterOperand _DstRC,757 RegisterOperand _SrcARC, RegisterOperand _SrcBRC>758 : VOPProfileMAI<P, _DstRC, _DstRC, _SrcARC> {759 let Src1RC64 = _SrcBRC;760 let Src2VT = DstVT;761 let Asm64 = " $vdst, $src0, $src1, $idx$cbsz$abid";762 let Outs64 = (outs DstRC:$vdst);763 let Ins64 = (ins Src0RC64:$src0, Src1RC64:$src1, VRegSrc_32:$idx, CBSZ:$cbsz, ABID:$abid, Src2RC64:$src2);764 let IsSMFMAC = true;765}766 767def VOPProfileMAI_F32_F32_X4 : VOPProfileMAI<VOP_V4F32_F32_F32_V4F32, AISrc_128_f32, ADst_128>;768def VOPProfileMAI_F32_F32_X16 : VOPProfileMAI<VOP_V16F32_F32_F32_V16F32, AISrc_512_f32, ADst_512>;769def VOPProfileMAI_F32_F32_X32 : VOPProfileMAI<VOP_V32F32_F32_F32_V32F32, AISrc_1024_f32, ADst_1024>;770def VOPProfileMAI_I32_I32_X4 : VOPProfileMAI<VOP_V4I32_I32_I32_V4I32, AISrc_128_b32, ADst_128>;771def VOPProfileMAI_I32_I32_X16 : VOPProfileMAI<VOP_V16I32_I32_I32_V16I32, AISrc_512_b32, ADst_512>;772def VOPProfileMAI_I32_I32_X32 : VOPProfileMAI<VOP_V32I32_I32_I32_V32I32, AISrc_1024_b32, ADst_1024>;773def VOPProfileMAI_F32_V2I16_X4 : VOPProfileMAI<VOP_V4F32_V2I16_V2I16_V4F32, AISrc_128_b32, ADst_128>;774def VOPProfileMAI_F32_V2I16_X16 : VOPProfileMAI<VOP_V16F32_V2I16_V2I16_V16F32, AISrc_512_b32, ADst_512>;775def VOPProfileMAI_F32_V2I16_X32 : VOPProfileMAI<VOP_V32F32_V2I16_V2I16_V32F32, AISrc_1024_b32, ADst_1024>;776def VOPProfileMAI_F32_V4F16_X4 : VOPProfileMAI<VOP_V4F32_V4F16_V4F16_V4F32, AISrc_128_b32, ADst_128, AVSrc_64>;777def VOPProfileMAI_F32_V4F16_X16 : VOPProfileMAI<VOP_V16F32_V4F16_V4F16_V16F32, AISrc_512_b32, ADst_512, AVSrc_64>;778def VOPProfileMAI_F32_V4F16_X32 : VOPProfileMAI<VOP_V32F32_V4F16_V4F16_V32F32, AISrc_1024_b32, ADst_1024, AVSrc_64>;779def VOPProfileMAI_F32_V4I16_X4 : VOPProfileMAI<VOP_V4F32_V4I16_V4I16_V4F32, AISrc_128_b32, ADst_128, AVSrc_64>;780def VOPProfileMAI_F32_V4I16_X16 : VOPProfileMAI<VOP_V16F32_V4I16_V4I16_V16F32, AISrc_512_b32, ADst_512, AVSrc_64>;781def VOPProfileMAI_F32_V4I16_X32 : VOPProfileMAI<VOP_V32F32_V4I16_V4I16_V32F32, AISrc_1024_b32, ADst_1024, AVSrc_64>;782def VOPProfileMAI_F64_16X16X4F64 : VOPProfileMAI<VOP_V4F64_F64_F64_V4F64, AISrc_256_f64, ADst_256, AVSrc_64>;783def VOPProfileMAI_F64_4X4X4F64 : VOPProfileMAI<VOP_F64_F64_F64_F64, AISrc_64_f64, ADst_64, AVSrc_64>;784def VOPProfileMAI_I32_I64_X16 : VOPProfileMAI<VOP_V4I32_I64_I64_V4I32, AISrc_128_b32, ADst_128, AVSrc_64>;785def VOPProfileMAI_I32_I64_X32 : VOPProfileMAI<VOP_V16I32_I64_I64_V16I32, AISrc_512_b32, ADst_512, AVSrc_64>;786def VOPProfileMAI_F32_V2F32_X16 : VOPProfileMAI<VOP_V4F32_V2F32_V2F32_V4F32, AISrc_128_b32, ADst_128, AVSrc_64>;787def VOPProfileMAI_F32_V2F32_X32 : VOPProfileMAI<VOP_V16F32_V2F32_V2F32_V16F32, AISrc_512_b32, ADst_512, AVSrc_64>;788def VOPProfileMAI_F32_I64_X32 : VOPProfileMAI<VOP_V4F32_I64_I64_V4F32, AISrc_128_b32, ADst_128, AVSrc_64>;789def VOPProfileMAI_F32_I64_X16 : VOPProfileMAI<VOP_V16F32_I64_I64_V16F32, AISrc_512_b32, ADst_512, AVSrc_64>;790 791def VOPProfileMAI_F32_F32_X4_VCD : VOPProfileMAI<VOP_V4F32_F32_F32_V4F32, VISrc_128_f32, VDst_128>;792def VOPProfileMAI_F32_F32_X16_VCD : VOPProfileMAI<VOP_V16F32_F32_F32_V16F32, VISrc_512_f32, VDst_512>;793def VOPProfileMAI_F32_F32_X32_VCD : VOPProfileMAI<VOP_V32F32_F32_F32_V32F32, VISrc_1024_f32, VDst_1024>;794def VOPProfileMAI_I32_I32_X4_VCD : VOPProfileMAI<VOP_V4I32_I32_I32_V4I32, VISrc_128_b32, VDst_128>;795def VOPProfileMAI_I32_I32_X16_VCD : VOPProfileMAI<VOP_V16I32_I32_I32_V16I32, VISrc_512_b32, VDst_512>;796def VOPProfileMAI_I32_I32_X32_VCD : VOPProfileMAI<VOP_V32I32_I32_I32_V32I32, VISrc_1024_b32, VDst_1024>;797def VOPProfileMAI_F32_V2I16_X4_VCD : VOPProfileMAI<VOP_V4F32_V2I16_V2I16_V4F32, VISrc_128_b32, VDst_128>;798def VOPProfileMAI_F32_V2I16_X16_VCD : VOPProfileMAI<VOP_V16F32_V2I16_V2I16_V16F32, VISrc_512_b32, VDst_512>;799def VOPProfileMAI_F32_V2I16_X32_VCD : VOPProfileMAI<VOP_V32F32_V2I16_V2I16_V32F32, VISrc_1024_b32, VDst_1024>;800def VOPProfileMAI_F32_V4F16_X4_VCD : VOPProfileMAI<VOP_V4F32_V4F16_V4F16_V4F32, VISrc_128_b32, VDst_128, AVSrc_64>;801def VOPProfileMAI_F32_V4F16_X16_VCD : VOPProfileMAI<VOP_V16F32_V4F16_V4F16_V16F32, VISrc_512_b32, VDst_512, AVSrc_64>;802def VOPProfileMAI_F32_V4F16_X32_VCD : VOPProfileMAI<VOP_V32F32_V4F16_V4F16_V32F32, VISrc_1024_b32, VDst_1024, AVSrc_64>;803def VOPProfileMAI_F32_V4I16_X4_VCD : VOPProfileMAI<VOP_V4F32_V4I16_V4I16_V4F32, VISrc_128_b32, VDst_128, AVSrc_64>;804def VOPProfileMAI_F32_V4I16_X16_VCD : VOPProfileMAI<VOP_V16F32_V4I16_V4I16_V16F32, VISrc_512_b32, VDst_512, AVSrc_64>;805def VOPProfileMAI_F32_V4I16_X32_VCD : VOPProfileMAI<VOP_V32F32_V4I16_V4I16_V32F32, VISrc_1024_b32, VDst_1024, AVSrc_64>;806def VOPProfileMAI_F64_16X16X4F64_VCD : VOPProfileMAI<VOP_V4F64_F64_F64_V4F64, VISrc_256_f64, VDst_256, AVSrc_64>;807def VOPProfileMAI_F64_4X4X4F64_VCD : VOPProfileMAI<VOP_F64_F64_F64_F64, VISrc_64_f64, VDst_64, AVSrc_64>;808def VOPProfileMAI_I32_I64_X16_VCD : VOPProfileMAI<VOP_V4I32_I64_I64_V4I32, VISrc_128_b32, VDst_128, AVSrc_64>;809def VOPProfileMAI_I32_I64_X32_VCD : VOPProfileMAI<VOP_V16I32_I64_I64_V16I32, VISrc_512_b32, VDst_512, AVSrc_64>;810def VOPProfileMAI_F32_V2F32_X16_VCD : VOPProfileMAI<VOP_V4F32_V2F32_V2F32_V4F32, VISrc_128_b32, VDst_128, AVSrc_64>;811def VOPProfileMAI_F32_V2F32_X32_VCD : VOPProfileMAI<VOP_V16F32_V2F32_V2F32_V16F32, VISrc_512_b32, VDst_512, AVSrc_64>;812def VOPProfileMAI_F32_I64_X32_VCD : VOPProfileMAI<VOP_V4F32_I64_I64_V4F32, VISrc_128_b32, VDst_128, AVSrc_64>;813def VOPProfileMAI_F32_I64_X16_VCD : VOPProfileMAI<VOP_V16F32_I64_I64_V16F32, VISrc_512_b32, VDst_512, AVSrc_64>;814 815def VOPProfileSMFMAC_F32_16X16X32_F16 : VOPProfileSMFMAC<VOP_V4F32_V4F16_V8F16_I32, AVDst_128_Align2, AVSrc_64_Align2, AVSrc_128_Align2>;816def VOPProfileSMFMAC_F32_16X16X64_F16 : VOPProfileSMFMAC<VOP_V4F32_V8F16_V16F16_I32, AVDst_128_Align2, AVSrc_128_Align2, AVSrc_256_Align2>;817def VOPProfileSMFMAC_F32_32X32X32_F16 : VOPProfileSMFMAC<VOP_V16F32_V8F16_V16F16_I32, AVDst_512_Align2, AVSrc_128_Align2, AVSrc_256_Align2>;818def VOPProfileSMFMAC_F32_16X16X64_BF16 : VOPProfileSMFMAC<VOP_V4F32_V8BF16_V16BF16_I32, AVDst_128_Align2, AVSrc_128_Align2, AVSrc_256_Align2>;819def VOPProfileSMFMAC_F32_32X32X32_BF16 : VOPProfileSMFMAC<VOP_V16F32_V8BF16_V16BF16_I32, AVDst_512_Align2, AVSrc_128_Align2, AVSrc_256_Align2>;820def VOPProfileSMFMAC_F32_32X32X16_F16 : VOPProfileSMFMAC<VOP_V16F32_V4F16_V8F16_I32, AVDst_512_Align2, AVSrc_64_Align2, AVSrc_128_Align2>;821def VOPProfileSMFMAC_F32_16X16X32_I16 : VOPProfileSMFMAC<VOP_V4F32_V4I16_V8I16_I32, AVDst_128_Align2, AVSrc_64_Align2, AVSrc_128_Align2>;822def VOPProfileSMFMAC_F32_32X32X16_I16 : VOPProfileSMFMAC<VOP_V16F32_V4I16_V8I16_I32, AVDst_512_Align2, AVSrc_64_Align2, AVSrc_128_Align2>;823def VOPProfileSMFMAC_I32_16X16X64_I8 : VOPProfileSMFMAC<VOP_V4I32_V2I32_V4I32_I32, AVDst_128_Align2, AVSrc_64_Align2, AVSrc_128_Align2>;824def VOPProfileSMFMAC_I32_32X32X32_I8 : VOPProfileSMFMAC<VOP_V16I32_V2I32_V4I32_I32, AVDst_512_Align2, AVSrc_64_Align2, AVSrc_128_Align2>;825def VOPProfileSMFMAC_F32_16X16X64_F8 : VOPProfileSMFMAC<VOP_V4F32_V2I32_V4I32_I32, AVDst_128_Align2, AVSrc_64_Align2, AVSrc_128_Align2>;826def VOPProfileSMFMAC_F32_32X32X32_F8 : VOPProfileSMFMAC<VOP_V16F32_V2I32_V4I32_I32, AVDst_512_Align2, AVSrc_64_Align2, AVSrc_128_Align2>;827def VOPProfileSMFMAC_I32_16X16X128_I8 : VOPProfileSMFMAC<VOP_V4I32_V4I32_V8I32_I32, AVDst_128_Align2, AVSrc_128_Align2, AVSrc_256_Align2>;828def VOPProfileSMFMAC_I32_32X32X64_I8 : VOPProfileSMFMAC<VOP_V16I32_V4I32_V8I32_I32, AVDst_512_Align2, AVSrc_128_Align2, AVSrc_256_Align2>;829 830def VOPProfileSMFMAC_F32_16X16X128_F8 : VOPProfileSMFMAC<VOP_V4F32_V4I32_V8I32_I32, AVDst_128_Align2, AVSrc_128_Align2, AVSrc_256_Align2>;831def VOPProfileSMFMAC_F32_32X32X64_F8 : VOPProfileSMFMAC<VOP_V16F32_V4I32_V8I32_I32, AVDst_512_Align2, AVSrc_128_Align2, AVSrc_256_Align2>;832 833def VOPProfileMAI_F32_V8F16_X32 : VOPProfileMAI<VOP_V4F32_V8F16_V8F16_V4F32, AISrc_128_f32, ADst_128, AVSrc_128>;834def VOPProfileMAI_F32_V8F16_X32_VCD : VOPProfileMAI<VOP_V4F32_V8F16_V8F16_V4F32, VISrc_128_f32, VDst_128, AVSrc_128>;835def VOPProfileMAI_F32_V8F16_X16 : VOPProfileMAI<VOP_V16F32_V8F16_V8F16_V16F32, AISrc_512_f32, ADst_512, AVSrc_128>;836def VOPProfileMAI_F32_V8F16_X16_VCD : VOPProfileMAI<VOP_V16F32_V8F16_V8F16_V16F32, VISrc_512_f32, VDst_512, AVSrc_128>;837 838def VOPProfileMAI_F32_V8BF16_X16 : VOPProfileMAI<VOP_V16F32_V8BF16_V8BF16_V16F32, AISrc_512_f32, ADst_512, AVSrc_128>;839def VOPProfileMAI_F32_V8BF16_X16_VCD : VOPProfileMAI<VOP_V16F32_V8BF16_V8BF16_V16F32, VISrc_512_f32, VDst_512, AVSrc_128>;840 841def VOPProfileMAI_F32_V8BF16_X4 : VOPProfileMAI<VOP_V4F32_V8BF16_V8BF16_V4F32, AISrc_128_f32, ADst_128, AVSrc_128>;842def VOPProfileMAI_F32_V8BF16_X4_VCD : VOPProfileMAI<VOP_V4F32_V8BF16_V8BF16_V4F32, VISrc_128_f32, VDst_128, AVSrc_128>;843 844 845let HasAbid = false in {846// For f32_16x16x128_f8f6f4 - f8 x f8 case847def VOPProfileMAI_F32_V8I32_V8I32_X128 : VOPProfileMAI<VOP_V4F32_V8I32_V8I32_V4F32, AISrc_128_f32, ADst_128, AVSrc_256>;848def VOPProfileMAI_F32_V8I32_V8I32_X128_VCD : VOPProfileMAI<VOP_V4F32_V8I32_V8I32_V4F32, VISrc_128_f32, VDst_128, AVSrc_256>;849 850// For f32_16x16x128_f8f6f4 - f8 x f6 case851def VOPProfileMAI_F32_V8I32_V6I32_X128 : VOPProfileMAI<VOP_V4F32_V8I32_V6I32_V4F32, AISrc_128_f32, ADst_128, AVSrc_256, AVSrc_192>;852def VOPProfileMAI_F32_V8I32_V6I32_X128_VCD : VOPProfileMAI<VOP_V4F32_V8I32_V6I32_V4F32, VISrc_128_f32, VDst_128, AVSrc_256, AVSrc_192>;853 854// For f32_16x16x128_f8f6f4 - f6 x f8 case855def VOPProfileMAI_F32_V6I32_V8I32_X128 : VOPProfileMAI<VOP_V4F32_V6I32_V8I32_V4F32, AISrc_128_f32, ADst_128, AVSrc_192, AVSrc_256>;856def VOPProfileMAI_F32_V6I32_V8I32_X128_VCD : VOPProfileMAI<VOP_V4F32_V6I32_V8I32_V4F32, VISrc_128_f32, VDst_128, AVSrc_192, AVSrc_256>;857 858// For f32_16x16x128_f8f6f4 - f6 x f6 case859def VOPProfileMAI_F32_V6I32_V6I32_X128 : VOPProfileMAI<VOP_V4F32_V6I32_V6I32_V4F32, AISrc_128_f32, ADst_128, AVSrc_192, AVSrc_192>;860def VOPProfileMAI_F32_V6I32_V6I32_X128_VCD : VOPProfileMAI<VOP_V4F32_V6I32_V6I32_V4F32, VISrc_128_f32, VDst_128, AVSrc_192, AVSrc_192>;861 862// For f32_16x16x128_f8f6f4 - f6 x f4 case863def VOPProfileMAI_F32_V6I32_V4I32_X128 : VOPProfileMAI<VOP_V4F32_V6I32_V4I32_V4F32, AISrc_128_f32, ADst_128, AVSrc_192, AVSrc_128>;864def VOPProfileMAI_F32_V6I32_V4I32_X128_VCD : VOPProfileMAI<VOP_V4F32_V6I32_V4I32_V4F32, VISrc_128_f32, VDst_128, AVSrc_192, AVSrc_128>;865 866// For f32_16x16x128_f8f6f4 - f4 x f6 case867def VOPProfileMAI_F32_V4I32_V6I32_X128 : VOPProfileMAI<VOP_V4F32_V4I32_V6I32_V4F32, AISrc_128_f32, ADst_128, AVSrc_128, AVSrc_192>;868def VOPProfileMAI_F32_V4I32_V6I32_X128_VCD : VOPProfileMAI<VOP_V4F32_V4I32_V6I32_V4F32, VISrc_128_f32, VDst_128, AVSrc_128, AVSrc_192>;869 870// For f32_16x16x128_f8f6f4 - f8 x f4 case871def VOPProfileMAI_F32_V8I32_V4I32_X128 : VOPProfileMAI<VOP_V4F32_V8I32_V4I32_V4F32, AISrc_128_f32, ADst_128, AVSrc_256, AVSrc_128>;872def VOPProfileMAI_F32_V8I32_V4I32_X128_VCD : VOPProfileMAI<VOP_V4F32_V8I32_V4I32_V4F32, VISrc_128_f32, VDst_128, AVSrc_256, AVSrc_128>;873 874// For f32_16x16x128_f8f6f4 - f4 x f8 case875def VOPProfileMAI_F32_V4I32_V8I32_X128 : VOPProfileMAI<VOP_V4F32_V4I32_V8I32_V4F32, AISrc_128_f32, ADst_128, AVSrc_128, AVSrc_256>;876def VOPProfileMAI_F32_V4I32_V8I32_X128_VCD : VOPProfileMAI<VOP_V4F32_V4I32_V8I32_V4F32, VISrc_128_f32, VDst_128, AVSrc_128, AVSrc_256>;877 878// For f32_16x16x128_f8f6f4 - f4 x f4 case879def VOPProfileMAI_F32_V4I32_V4I32_X128 : VOPProfileMAI<VOP_V4F32_V4I32_V4I32_V4F32, AISrc_128_f32, ADst_128, AVSrc_128, AVSrc_128>;880def VOPProfileMAI_F32_V4I32_V4I32_X128_VCD : VOPProfileMAI<VOP_V4F32_V4I32_V4I32_V4F32, VISrc_128_f32, VDst_128, AVSrc_128, AVSrc_128>;881 882// For f32_32x32x64_f8f6f4 - f8 x f8 case883def VOPProfileMAI_F32_V8I32_V8I32_X512 : VOPProfileMAI<VOP_V16F32_V8I32_V8I32_V16F32, AISrc_512_f32, ADst_512, AVSrc_256>;884def VOPProfileMAI_F32_V8I32_V8I32_X512_VCD : VOPProfileMAI<VOP_V16F32_V8I32_V8I32_V16F32, VISrc_512_f32, VDst_512, AVSrc_256>;885 886// For f32_32x32x64_f8f6f4 - f8 x f6 case887def VOPProfileMAI_F32_V8I32_V6I32_X512 : VOPProfileMAI<VOP_V16F32_V8I32_V6I32_V16F32, AISrc_512_f32, ADst_512, AVSrc_256, AVSrc_192>;888def VOPProfileMAI_F32_V8I32_V6I32_X512_VCD : VOPProfileMAI<VOP_V16F32_V8I32_V6I32_V16F32, VISrc_512_f32, VDst_512, AVSrc_256, AVSrc_192>;889 890// For f32_32x32x64_f8f6f4 - f8 x f4 case891def VOPProfileMAI_F32_V8I32_V4I32_X512 : VOPProfileMAI<VOP_V16F32_V8I32_V4I32_V16F32, AISrc_512_f32, ADst_512, AVSrc_256, AVSrc_128>;892def VOPProfileMAI_F32_V8I32_V4I32_X512_VCD : VOPProfileMAI<VOP_V16F32_V8I32_V4I32_V16F32, VISrc_512_f32, VDst_512, AVSrc_256, AVSrc_128>;893 894// For f32_32x32x64_f8f6f4 - f4 x f8 case895def VOPProfileMAI_F32_V4I32_V8I32_X512 : VOPProfileMAI<VOP_V16F32_V4I32_V8I32_V16F32, AISrc_512_f32, ADst_512, AVSrc_128, AVSrc_256>;896def VOPProfileMAI_F32_V4I32_V8I32_X512_VCD : VOPProfileMAI<VOP_V16F32_V4I32_V8I32_V16F32, VISrc_512_f32, VDst_512, AVSrc_128, AVSrc_256>;897 898// For f32_32x32x64_f8f6f4 - f6 x f8 case899def VOPProfileMAI_F32_V6I32_V8I32_X512 : VOPProfileMAI<VOP_V16F32_V6I32_V8I32_V16F32, AISrc_512_f32, ADst_512, AVSrc_192, AVSrc_256>;900def VOPProfileMAI_F32_V6I32_V8I32_X512_VCD : VOPProfileMAI<VOP_V16F32_V6I32_V8I32_V16F32, VISrc_512_f32, VDst_512, AVSrc_192, AVSrc_256>;901 902// For f32_32x32x64_f8f6f4 - f6 x f6 case903def VOPProfileMAI_F32_V6I32_V6I32_X512 : VOPProfileMAI<VOP_V16F32_V6I32_V6I32_V16F32, AISrc_512_f32, ADst_512, AVSrc_192, AVSrc_192>;904def VOPProfileMAI_F32_V6I32_V6I32_X512_VCD : VOPProfileMAI<VOP_V16F32_V6I32_V6I32_V16F32, VISrc_512_f32, VDst_512, AVSrc_192, AVSrc_192>;905 906// For f32_32x32x64_f8f6f4 - f6 x f4 case907def VOPProfileMAI_F32_V6I32_V4I32_X512 : VOPProfileMAI<VOP_V16F32_V6I32_V4I32_V16F32, AISrc_512_f32, ADst_512, AVSrc_192, AVSrc_128>;908def VOPProfileMAI_F32_V6I32_V4I32_X512_VCD : VOPProfileMAI<VOP_V16F32_V6I32_V4I32_V16F32, VISrc_512_f32, VDst_512, AVSrc_192, AVSrc_128>;909 910// For f32_32x32x64_f8f6f4 - f4 x f6 case911def VOPProfileMAI_F32_V4I32_V6I32_X512 : VOPProfileMAI<VOP_V16F32_V4I32_V6I32_V16F32, AISrc_512_f32, ADst_512, AVSrc_128, AVSrc_192>;912def VOPProfileMAI_F32_V4I32_V6I32_X512_VCD : VOPProfileMAI<VOP_V16F32_V4I32_V6I32_V16F32, VISrc_512_f32, VDst_512, AVSrc_128, AVSrc_192>;913 914// For f32_32x32x64_f8f6f4 - f4 x f4 case915def VOPProfileMAI_F32_V4I32_V4I32_X512 : VOPProfileMAI<VOP_V16F32_V4I32_V4I32_V16F32, AISrc_512_f32, ADst_512, AVSrc_128, AVSrc_128>;916def VOPProfileMAI_F32_V4I32_V4I32_X512_VCD : VOPProfileMAI<VOP_V16F32_V4I32_V4I32_V16F32, VISrc_512_f32, VDst_512, AVSrc_128, AVSrc_128>;917}918 919 920// For i32_16x16x64_i8921def VOPProfileMAI_I32_V4I32_X128 : VOPProfileMAI<VOP_V4I32_V4I32_V4I32_V4I32, AISrc_128_f32, ADst_128, AVSrc_128>;922def VOPProfileMAI_I32_V4I32_X128_VCD : VOPProfileMAI<VOP_V4I32_V4I32_V4I32_V4I32, VISrc_128_f32, VDst_128, AVSrc_128>;923 924// For i32_32x32x32_i8925def VOPProfileMAI_I32_V4I32_X16 : VOPProfileMAI<VOP_V16I32_V4I32_V4I32_V16I32, AISrc_512_b32, ADst_512, AVSrc_128>;926def VOPProfileMAI_I32_V4I32_X16_VCD : VOPProfileMAI<VOP_V16I32_V4I32_V4I32_V16I32, VISrc_512_b32, VDst_512, AVSrc_128>;927 928 929class MFMATable <bit is_mac, string Kind, string Name,930 string AGPROpName = NAME> {931 bit IsMac = is_mac;932 string FMAOp = Name;933 string AGPROp = AGPROpName;934 935 // Does this MFMA use "AGPR" or "VGPR" for srcC/vdst936 string MFMAKind = Kind;937}938 939class MFMA_F8F6F4_WithSizeTable<int A, int B, Instruction ThisVariant, Instruction F8F8Variant> {940 Instruction F8F8Opcode = F8F8Variant;941 Instruction Opcode = ThisVariant;942 bits<8> NumRegsSrcA = A;943 bits<8> NumRegsSrcB = B;944}945 946class MFMA_F8F6F4_WithSizeTable_Helper<VOP3_Pseudo ps, string F8F8Op> :947 MFMA_F8F6F4_WithSizeTable<!srl(ps.Pfl.Src0VT.Size, 5),948 !srl(ps.Pfl.Src1VT.Size, 5),949 !cast<Instruction>(NAME),950 !cast<Instruction>(F8F8Op)> {951}952 953// Currently assumes scaled instructions never have abid954class MAIFrag<SDPatternOperator Op, bit HasAbid = true, bit Scaled = false> : PatFrag <955 !if(Scaled, (ops node:$src0, node:$src1, node:$src2, node:$cbsz, node:$blgp,956 node:$src0_modifiers, node:$scale_src0,957 node:$src1_modifiers, node:$scale_src1),958 !con((ops node:$src0, node:$src1, node:$src2, node:$cbsz),959 !if(HasAbid, (ops node:$abid), (ops)),960 (ops node:$blgp))),961 !if(Scaled, (Op $src0, $src1, $src2, $cbsz, $blgp, $src0_modifiers, $scale_src0, $src1_modifiers, $scale_src1),962 !if(HasAbid, (Op $src0, $src1, $src2, $cbsz, $abid, $blgp),963 (Op $src0, $src1, $src2, $cbsz, $blgp)))>;964 965class CanUseAGPR_MAI<ValueType vt> {966 code PredicateCode = [{967 return !Subtarget->hasGFX90AInsts() ||968 MF->getInfo<SIMachineFunctionInfo>()->selectAGPRFormMFMA(969 }] # !srl(vt.Size, 5) # ");";970 971 code GISelPredicateCode = [{972 return !Subtarget->hasGFX90AInsts() ||973 MF.getInfo<SIMachineFunctionInfo>()->selectAGPRFormMFMA(974 }] # !srl(vt.Size, 5) # ");";975}976 977class AgprMAIFrag<SDPatternOperator Op, ValueType vt, bit HasAbid = true,978 bit Scaled = false> :979 MAIFrag<Op, HasAbid, Scaled>,980 CanUseAGPR_MAI<vt>;981 982class VgprMAIFrag<SDPatternOperator Op, bit HasAbid = true,983 bit Scaled = false> :984 MAIFrag<Op, HasAbid, Scaled>;985 986let isAsCheapAsAMove = 1, isReMaterializable = 1 in {987 defm V_ACCVGPR_READ_B32 : VOP3Inst<"v_accvgpr_read_b32", VOPProfileAccRead>;988 let isMoveImm = 1 in {989 defm V_ACCVGPR_WRITE_B32 : VOP3Inst<"v_accvgpr_write_b32", VOPProfileAccWrite>;990 } // End isMoveImm = 1991} // End isAsCheapAsAMove = 1, isReMaterializable = 1992 993class MAIInst<string OpName, VOPProfile P, SDPatternOperator node, bit Scaled = false>994 : VOP3InstBase<OpName, P, node, /*IsVOP2=*/0, Scaled> {995 let SubtargetPredicate = HasMAIInsts;996 Instruction Opcode = !cast<Instruction>(NAME);997 bit is_dgemm = 0;998 bit is_gfx940_xdl = 0;999 let PseudoInstr = NAME; // FIXME: Why is this not the default1000}1001 1002// FIXME: Intrinsic should probably not have op_sel operands, we can1003// pattern match byte select patterns into op_sel.1004// FIXME: Missing neg and clamp modifiers1005//1006// FIXME: Usual syntax for op_sel is quite hostile here.1007class ScaledMAIInst<string OpName, MAIInst BaseInst, SDPatternOperator node> :1008 MAIInst<OpName, BaseInst.Pfl, node, /*Scaled=*/true> {1009 // Append operands from V_MFMA_LD_SCALE_B32, but we need to rename them.1010 // Restrict to VGPR only (VRegSrc_32) for the scale operands to workaround a1011 // hardware design defect: For all Inline/SGPR constants, SP HW use bits1012 // [30:23] as the scale.1013 // TODO: We may still be able to allow Inline Constants/SGPR, with a proper1014 // shift, to obtain a potentially better performance.1015 let InOperandList = !con(BaseInst.InOperandList,1016 (ins VRegSrc_32:$scale_src0,1017 VRegSrc_32:$scale_src1,1018 op_sel0:$src0_modifiers,1019 op_sel_hi0:$src1_modifiers));1020 let AsmOperands =1021 "$vdst, $src0, $src1, $src2, $scale_src0, $scale_src1"1022 "$src0_modifiers$src1_modifiers$cbsz$blgp";1023 let AsmMatchConverter = "cvtScaledMFMA";1024 let FixedSize = 1;1025 let Size = 16;1026}1027 1028multiclass MAIInst<string OpName, string P, SDPatternOperator node = null_frag,1029 bit HasAbid = true,1030 bit Scaled = false> {1031 defvar NoDstOverlap = !cast<VOPProfileMAI>("VOPProfileMAI_" # P).NoDstOverlap;1032 defvar ProfileAGPR = !cast<VOPProfileMAI>("VOPProfileMAI_" # P);1033 defvar ProfileVGPR = !cast<VOPProfileMAI>("VOPProfileMAI_" # P # "_VCD");1034 1035 1036 let isConvergent = 1, mayRaiseFPException = 0, ReadsModeReg = 1 in {1037 // FP32 denorm mode is respected, rounding mode is not. Exceptions are not supported.1038 let Constraints = !if(NoDstOverlap, "@earlyclobber $vdst", "") in {1039 def _e64 : MAIInst<OpName, ProfileAGPR,1040 !if(!or(NoDstOverlap, !eq(node, null_frag)), null_frag, AgprMAIFrag<node, ProfileAGPR.DstVT, HasAbid, Scaled>), Scaled>,1041 MFMATable<0, "AGPR", NAME # "_e64">;1042 1043 let OtherPredicates = [isGFX90APlus], Mnemonic = OpName in1044 def _vgprcd_e64 : MAIInst<OpName # "_vgprcd", ProfileVGPR,1045 !if(!or(NoDstOverlap, !eq(node, null_frag)), null_frag, VgprMAIFrag<node, HasAbid, Scaled>), Scaled>,1046 MFMATable<0, "VGPR", NAME # "_vgprcd_e64", NAME # "_e64">;1047 }1048 1049 if NoDstOverlap then {1050 let Constraints = !if(NoDstOverlap, "$vdst = $src2", ""),1051 isConvertibleToThreeAddress = NoDstOverlap,1052 Mnemonic = OpName in {1053 def "_mac_e64" : MAIInst<OpName # "_mac", ProfileAGPR,1054 !if(!eq(node, null_frag), null_frag, AgprMAIFrag<node, ProfileAGPR.DstVT, HasAbid, Scaled>), Scaled>,1055 MFMATable<1, "AGPR", NAME # "_e64", NAME # "_mac_e64">;1056 1057 let OtherPredicates = [isGFX90APlus] in1058 def _mac_vgprcd_e64 : MAIInst<OpName # "_mac_vgprcd", ProfileVGPR,1059 !if(!eq(node, null_frag), null_frag, VgprMAIFrag<node, HasAbid, Scaled>), Scaled>,1060 MFMATable<1, "VGPR", NAME # "_vgprcd_e64", NAME # "_mac_e64">;1061 }1062 }1063 } // End isConvergent = 1, mayRaiseFPException = 0, ReadsModeReg = 11064}1065 1066// Provide a wrapper around MAIInst that provides the appended operands from V_MFMA_LD_SCALE_B321067multiclass ScaledMAIInst_mc<string OpName, string UnscaledOpName_, SDPatternOperator node> {1068 defvar VariantSuffix = !subst(!toupper(OpName), "", NAME); // Drop the main opcode name prefix to get the "_fN_fM" suffix.1069 defvar UnscaledOpName = UnscaledOpName_#VariantSuffix;1070 1071 defvar HasAbid = false;1072 defvar Profile = !cast<VOPProfileMAI>(!cast<MAIInst>(UnscaledOpName#"_e64").Pfl);1073 defvar NoDstOverlap = Profile.NoDstOverlap;1074 1075 def _e64 : ScaledMAIInst<OpName,1076 !cast<MAIInst>(UnscaledOpName#"_e64"), !if(NoDstOverlap, null_frag, AgprMAIFrag<node, Profile.DstVT, HasAbid, true>)>,1077 MFMATable<0, "AGPR", NAME # "_e64">;1078 1079 def _vgprcd_e64 : ScaledMAIInst<OpName # "_vgprcd",1080 !cast<MAIInst>(UnscaledOpName#"_vgprcd_e64"), !if(NoDstOverlap, null_frag, VgprMAIFrag<node, HasAbid, true>)>,1081 MFMATable<0, "VGPR", NAME # "_vgprcd_e64", NAME # "_e64">;1082 1083 if NoDstOverlap then {1084 let Constraints = !if(NoDstOverlap, "$vdst = $src2", ""),1085 isConvertibleToThreeAddress = NoDstOverlap,1086 Mnemonic = UnscaledOpName_ in {1087 def _mac_e64 : ScaledMAIInst<OpName # "_mac",1088 !cast<MAIInst>(UnscaledOpName # "_mac_e64"), AgprMAIFrag<node, Profile.DstVT, HasAbid, true>>,1089 MFMATable<1, "AGPR", NAME # "_e64">;1090 1091 def _mac_vgprcd_e64 : ScaledMAIInst<OpName # " _mac_vgprcd",1092 !cast<MAIInst>(UnscaledOpName # "_mac_vgprcd_e64"), VgprMAIFrag<node, HasAbid, true>>,1093 MFMATable<1, "VGPR", NAME # "_vgprcd_e64", NAME # "_mac_e64">;1094 }1095 }1096}1097 1098// Each of SrcA and SrcB can be encoded using 3 different sizes, so1099// define 9 permutations of register classes.1100multiclass MAIInst_SrcFormats_mc<string OpName, string ProfileSuffix, SDPatternOperator node> {1101 defvar HasAbid = false;1102 defm _f8_f8 : MAIInst<OpName, "F32_V8I32_V8I32"#ProfileSuffix, node, HasAbid>;1103 defm _f8_f6 : MAIInst<OpName, "F32_V8I32_V6I32"#ProfileSuffix, node, HasAbid>;1104 defm _f6_f8 : MAIInst<OpName, "F32_V6I32_V8I32"#ProfileSuffix, node, HasAbid>;1105 defm _f6_f6 : MAIInst<OpName, "F32_V6I32_V6I32"#ProfileSuffix, node, HasAbid>;1106 defm _f8_f4 : MAIInst<OpName, "F32_V8I32_V4I32"#ProfileSuffix, node, HasAbid>;1107 defm _f4_f8 : MAIInst<OpName, "F32_V4I32_V8I32"#ProfileSuffix, node, HasAbid>;1108 defm _f6_f4 : MAIInst<OpName, "F32_V6I32_V4I32"#ProfileSuffix, node, HasAbid>;1109 defm _f4_f6 : MAIInst<OpName, "F32_V4I32_V6I32"#ProfileSuffix, node, HasAbid>;1110 defm _f4_f4 : MAIInst<OpName, "F32_V4I32_V4I32"#ProfileSuffix, node, HasAbid>;1111}1112 1113multiclass MAIInst_SrcFormats_Scaled_mc<string OpName, string UnscaledOpName, SDPatternOperator node> {1114 defm _f8_f8 : ScaledMAIInst_mc<OpName, UnscaledOpName, node>;1115 defm _f8_f6 : ScaledMAIInst_mc<OpName, UnscaledOpName, node>;1116 defm _f6_f8 : ScaledMAIInst_mc<OpName, UnscaledOpName, node>;1117 defm _f6_f6 : ScaledMAIInst_mc<OpName, UnscaledOpName, node>;1118 defm _f8_f4 : ScaledMAIInst_mc<OpName, UnscaledOpName, node>;1119 defm _f4_f8 : ScaledMAIInst_mc<OpName, UnscaledOpName, node>;1120 defm _f6_f4 : ScaledMAIInst_mc<OpName, UnscaledOpName, node>;1121 defm _f4_f6 : ScaledMAIInst_mc<OpName, UnscaledOpName, node>;1122 defm _f4_f4 : ScaledMAIInst_mc<OpName, UnscaledOpName, node>;1123}1124 1125defm V_MFMA_F32_4X4X1F32 : MAIInst<"v_mfma_f32_4x4x1f32", "F32_F32_X4", int_amdgcn_mfma_f32_4x4x1f32>;1126defm V_MFMA_F32_16X16X1F32 : MAIInst<"v_mfma_f32_16x16x1f32", "F32_F32_X16", int_amdgcn_mfma_f32_16x16x1f32>;1127defm V_MFMA_F32_16X16X4F32 : MAIInst<"v_mfma_f32_16x16x4f32", "F32_F32_X4", int_amdgcn_mfma_f32_16x16x4f32>;1128defm V_MFMA_F32_32X32X1F32 : MAIInst<"v_mfma_f32_32x32x1f32", "F32_F32_X32", int_amdgcn_mfma_f32_32x32x1f32>;1129defm V_MFMA_F32_32X32X2F32 : MAIInst<"v_mfma_f32_32x32x2f32", "F32_F32_X16", int_amdgcn_mfma_f32_32x32x2f32>;1130 1131let is_gfx940_xdl = 1 in {1132defm V_MFMA_F32_4X4X4F16 : MAIInst<"v_mfma_f32_4x4x4f16", "F32_V4F16_X4", int_amdgcn_mfma_f32_4x4x4f16>;1133defm V_MFMA_I32_4X4X4I8 : MAIInst<"v_mfma_i32_4x4x4i8", "I32_I32_X4", int_amdgcn_mfma_i32_4x4x4i8>;1134defm V_MFMA_F32_16X16X4F16 : MAIInst<"v_mfma_f32_16x16x4f16", "F32_V4F16_X16", int_amdgcn_mfma_f32_16x16x4f16>;1135defm V_MFMA_F32_16X16X16F16 : MAIInst<"v_mfma_f32_16x16x16f16", "F32_V4F16_X4", int_amdgcn_mfma_f32_16x16x16f16>;1136defm V_MFMA_I32_16X16X4I8 : MAIInst<"v_mfma_i32_16x16x4i8", "I32_I32_X16", int_amdgcn_mfma_i32_16x16x4i8>;1137defm V_MFMA_F32_32X32X4F16 : MAIInst<"v_mfma_f32_32x32x4f16", "F32_V4F16_X32", int_amdgcn_mfma_f32_32x32x4f16>;1138defm V_MFMA_F32_32X32X8F16 : MAIInst<"v_mfma_f32_32x32x8f16", "F32_V4F16_X16", int_amdgcn_mfma_f32_32x32x8f16>;1139defm V_MFMA_I32_32X32X4I8 : MAIInst<"v_mfma_i32_32x32x4i8", "I32_I32_X32", int_amdgcn_mfma_i32_32x32x4i8>;1140}1141 1142let SubtargetPredicate = isGFX908orGFX90A in {1143defm V_MFMA_I32_16X16X16I8 : MAIInst<"v_mfma_i32_16x16x16i8", "I32_I32_X4", int_amdgcn_mfma_i32_16x16x16i8>;1144defm V_MFMA_I32_32X32X8I8 : MAIInst<"v_mfma_i32_32x32x8i8", "I32_I32_X16", int_amdgcn_mfma_i32_32x32x8i8>;1145defm V_MFMA_F32_4X4X2BF16 : MAIInst<"v_mfma_f32_4x4x2bf16", "F32_V2I16_X4", int_amdgcn_mfma_f32_4x4x2bf16>;1146defm V_MFMA_F32_16X16X2BF16 : MAIInst<"v_mfma_f32_16x16x2bf16", "F32_V2I16_X16", int_amdgcn_mfma_f32_16x16x2bf16>;1147defm V_MFMA_F32_16X16X8BF16 : MAIInst<"v_mfma_f32_16x16x8bf16", "F32_V2I16_X4", int_amdgcn_mfma_f32_16x16x8bf16>;1148defm V_MFMA_F32_32X32X2BF16 : MAIInst<"v_mfma_f32_32x32x2bf16", "F32_V2I16_X32", int_amdgcn_mfma_f32_32x32x2bf16>;1149defm V_MFMA_F32_32X32X4BF16 : MAIInst<"v_mfma_f32_32x32x4bf16", "F32_V2I16_X16", int_amdgcn_mfma_f32_32x32x4bf16>;1150}1151 1152let SubtargetPredicate = HasGFX950Insts, is_gfx940_xdl = 1 in {1153defm V_MFMA_F32_16X16X32_F16 : MAIInst<"v_mfma_f32_16x16x32f16", "F32_V8F16_X32", int_amdgcn_mfma_f32_16x16x32_f16>;1154defm V_MFMA_F32_32X32X16_F16 : MAIInst<"v_mfma_f32_32x32x16f16", "F32_V8F16_X16", int_amdgcn_mfma_f32_32x32x16_f16>;1155defm V_MFMA_F32_16X16X32_BF16 : MAIInst<"v_mfma_f32_16x16x32bf16", "F32_V8BF16_X4", int_amdgcn_mfma_f32_16x16x32_bf16>;1156defm V_MFMA_I32_16X16X64_I8 : MAIInst<"v_mfma_i32_16x16x64i8", "I32_V4I32_X128", int_amdgcn_mfma_i32_16x16x64_i8>;1157defm V_MFMA_F32_32X32X16_BF16 : MAIInst<"v_mfma_f32_32x32x16bf16", "F32_V8BF16_X16", int_amdgcn_mfma_f32_32x32x16_bf16>;1158defm V_MFMA_I32_32X32X32_I8 : MAIInst<"v_mfma_i32_32x32x32i8", "I32_V4I32_X16", int_amdgcn_mfma_i32_32x32x32_i8>;1159 1160defm V_MFMA_F32_16X16X128_F8F6F4 : MAIInst_SrcFormats_mc<"v_mfma_f32_16x16x128f8f6f4",1161 "_X128", mfma_f32_16x16x128_f8f6f4>;1162defm V_MFMA_F32_32X32X64_F8F6F4 : MAIInst_SrcFormats_mc<"v_mfma_f32_32x32x64f8f6f4",1163 "_X512", mfma_f32_32x32x64_f8f6f4>;1164 1165defm V_MFMA_SCALE_F32_16X16X128_F8F6F4 : MAIInst_SrcFormats_Scaled_mc<1166 "v_mfma_scale_f32_16x16x128_f8f6f4", "V_MFMA_F32_16X16X128_F8F6F4",1167 int_amdgcn_mfma_scale_f32_16x16x128_f8f6f4>;1168 1169defm V_MFMA_SCALE_F32_32X32X64_F8F6F4 : MAIInst_SrcFormats_Scaled_mc<1170 "v_mfma_scale_f32_32x32x64_f8f6f4",1171 "V_MFMA_F32_32X32X64_F8F6F4",1172 int_amdgcn_mfma_scale_f32_32x32x64_f8f6f4>;1173}1174 1175let SubtargetPredicate = HasGFX950Insts in {1176defm V_MFMA_LD_SCALE_B32 : VOP3PInst<"v_mfma_ld_scale_b32", VOP_MFMA_LD_SCALE>;1177}1178 1179let SubtargetPredicate = isGFX90APlus in {1180 let is_gfx940_xdl = 1 in {1181 defm V_MFMA_F32_32X32X4BF16_1K : MAIInst<"v_mfma_f32_32x32x4bf16_1k", "F32_V4I16_X32", int_amdgcn_mfma_f32_32x32x4bf16_1k>;1182 defm V_MFMA_F32_16X16X4BF16_1K : MAIInst<"v_mfma_f32_16x16x4bf16_1k", "F32_V4I16_X16", int_amdgcn_mfma_f32_16x16x4bf16_1k>;1183 defm V_MFMA_F32_4X4X4BF16_1K : MAIInst<"v_mfma_f32_4x4x4bf16_1k", "F32_V4I16_X4", int_amdgcn_mfma_f32_4x4x4bf16_1k>;1184 defm V_MFMA_F32_32X32X8BF16_1K : MAIInst<"v_mfma_f32_32x32x8bf16_1k", "F32_V4I16_X16", int_amdgcn_mfma_f32_32x32x8bf16_1k>;1185 defm V_MFMA_F32_16X16X16BF16_1K : MAIInst<"v_mfma_f32_16x16x16bf16_1k", "F32_V4I16_X4", int_amdgcn_mfma_f32_16x16x16bf16_1k>;1186 }1187 1188 let is_dgemm = 1 in {1189 defm V_MFMA_F64_16X16X4F64 : MAIInst<"v_mfma_f64_16x16x4f64", "F64_16X16X4F64", int_amdgcn_mfma_f64_16x16x4f64>;1190 defm V_MFMA_F64_4X4X4F64 : MAIInst<"v_mfma_f64_4x4x4f64", "F64_4X4X4F64", int_amdgcn_mfma_f64_4x4x4f64>;1191 }1192} // End SubtargetPredicate = isGFX90APlus1193 1194let SubtargetPredicate = isGFX940Plus, is_gfx940_xdl = 1 in {1195 defm V_MFMA_I32_32X32X16I8 : MAIInst<"v_mfma_i32_32x32x16i8", "I32_I64_X32", int_amdgcn_mfma_i32_32x32x16_i8>;1196 defm V_MFMA_I32_16X16X32I8 : MAIInst<"v_mfma_i32_16x16x32i8", "I32_I64_X16", int_amdgcn_mfma_i32_16x16x32_i8>;1197} // End SubtargetPredicate = isGFX940Plus, is_gfx940_xdl = 11198 1199let SubtargetPredicate = HasXF32Insts, is_gfx940_xdl = 1 in {1200 defm V_MFMA_F32_16X16X8XF32 : MAIInst<"v_mfma_f32_16x16x8xf32", "F32_V2F32_X16", int_amdgcn_mfma_f32_16x16x8_xf32>;1201 defm V_MFMA_F32_32X32X4XF32 : MAIInst<"v_mfma_f32_32x32x4xf32", "F32_V2F32_X32", int_amdgcn_mfma_f32_32x32x4_xf32>;1202} // End SubtargetPredicate = HasXF32Insts, is_gfx940_xdl = 11203 1204let SubtargetPredicate = HasFP8Insts, is_gfx940_xdl = 1 in {1205 defm V_MFMA_F32_16X16X32_BF8_BF8 : MAIInst<"v_mfma_f32_16x16x32_bf8_bf8", "F32_I64_X32", int_amdgcn_mfma_f32_16x16x32_bf8_bf8>;1206 defm V_MFMA_F32_16X16X32_BF8_FP8 : MAIInst<"v_mfma_f32_16x16x32_bf8_fp8", "F32_I64_X32", int_amdgcn_mfma_f32_16x16x32_bf8_fp8>;1207 defm V_MFMA_F32_16X16X32_FP8_BF8 : MAIInst<"v_mfma_f32_16x16x32_fp8_bf8", "F32_I64_X32", int_amdgcn_mfma_f32_16x16x32_fp8_bf8>;1208 defm V_MFMA_F32_16X16X32_FP8_FP8 : MAIInst<"v_mfma_f32_16x16x32_fp8_fp8", "F32_I64_X32", int_amdgcn_mfma_f32_16x16x32_fp8_fp8>;1209 defm V_MFMA_F32_32X32X16_BF8_BF8 : MAIInst<"v_mfma_f32_32x32x16_bf8_bf8", "F32_I64_X16", int_amdgcn_mfma_f32_32x32x16_bf8_bf8>;1210 defm V_MFMA_F32_32X32X16_BF8_FP8 : MAIInst<"v_mfma_f32_32x32x16_bf8_fp8", "F32_I64_X16", int_amdgcn_mfma_f32_32x32x16_bf8_fp8>;1211 defm V_MFMA_F32_32X32X16_FP8_BF8 : MAIInst<"v_mfma_f32_32x32x16_fp8_bf8", "F32_I64_X16", int_amdgcn_mfma_f32_32x32x16_fp8_bf8>;1212 defm V_MFMA_F32_32X32X16_FP8_FP8 : MAIInst<"v_mfma_f32_32x32x16_fp8_fp8", "F32_I64_X16", int_amdgcn_mfma_f32_32x32x16_fp8_fp8>;1213} // End SubtargetPredicate = HasFP8Insts, is_gfx940_xdl = 11214 1215multiclass SMFMACInst<string OpName, string P, SDPatternOperator node> {1216 let Constraints = "$vdst = $src2",1217 isConvergent = 1, mayRaiseFPException = 0, ReadsModeReg = 1, is_gfx940_xdl = 1 in {1218 def _e64 : MAIInst<OpName, !cast<VOPProfileSMFMAC>("VOPProfileSMFMAC_" # P), node>;1219 }1220}1221 1222let SubtargetPredicate = isGFX940Plus in {1223defm V_SMFMAC_F32_16X16X32_F16 : SMFMACInst<"v_smfmac_f32_16x16x32_f16", "F32_16X16X32_F16", int_amdgcn_smfmac_f32_16x16x32_f16>;1224defm V_SMFMAC_F32_32X32X16_F16 : SMFMACInst<"v_smfmac_f32_32x32x16_f16", "F32_32X32X16_F16", int_amdgcn_smfmac_f32_32x32x16_f16>;1225defm V_SMFMAC_F32_16X16X32_BF16 : SMFMACInst<"v_smfmac_f32_16x16x32_bf16", "F32_16X16X32_I16", int_amdgcn_smfmac_f32_16x16x32_bf16>;1226defm V_SMFMAC_F32_32X32X16_BF16 : SMFMACInst<"v_smfmac_f32_32x32x16_bf16", "F32_32X32X16_I16", int_amdgcn_smfmac_f32_32x32x16_bf16>;1227defm V_SMFMAC_I32_16X16X64_I8 : SMFMACInst<"v_smfmac_i32_16x16x64_i8", "I32_16X16X64_I8", int_amdgcn_smfmac_i32_16x16x64_i8>;1228defm V_SMFMAC_I32_32X32X32_I8 : SMFMACInst<"v_smfmac_i32_32x32x32_i8", "I32_32X32X32_I8", int_amdgcn_smfmac_i32_32x32x32_i8>;1229}1230 1231let SubtargetPredicate = HasFP8Insts, is_gfx940_xdl = 1 in {1232defm V_SMFMAC_F32_16X16X64_BF8_BF8 : SMFMACInst<"v_smfmac_f32_16x16x64_bf8_bf8", "F32_16X16X64_F8", int_amdgcn_smfmac_f32_16x16x64_bf8_bf8>;1233defm V_SMFMAC_F32_16X16X64_BF8_FP8 : SMFMACInst<"v_smfmac_f32_16x16x64_bf8_fp8", "F32_16X16X64_F8", int_amdgcn_smfmac_f32_16x16x64_bf8_fp8>;1234defm V_SMFMAC_F32_16X16X64_FP8_BF8 : SMFMACInst<"v_smfmac_f32_16x16x64_fp8_bf8", "F32_16X16X64_F8", int_amdgcn_smfmac_f32_16x16x64_fp8_bf8>;1235defm V_SMFMAC_F32_16X16X64_FP8_FP8 : SMFMACInst<"v_smfmac_f32_16x16x64_fp8_fp8", "F32_16X16X64_F8", int_amdgcn_smfmac_f32_16x16x64_fp8_fp8>;1236defm V_SMFMAC_F32_32X32X32_BF8_BF8 : SMFMACInst<"v_smfmac_f32_32x32x32_bf8_bf8", "F32_32X32X32_F8", int_amdgcn_smfmac_f32_32x32x32_bf8_bf8>;1237defm V_SMFMAC_F32_32X32X32_BF8_FP8 : SMFMACInst<"v_smfmac_f32_32x32x32_bf8_fp8", "F32_32X32X32_F8", int_amdgcn_smfmac_f32_32x32x32_bf8_fp8>;1238defm V_SMFMAC_F32_32X32X32_FP8_BF8 : SMFMACInst<"v_smfmac_f32_32x32x32_fp8_bf8", "F32_32X32X32_F8", int_amdgcn_smfmac_f32_32x32x32_fp8_bf8>;1239defm V_SMFMAC_F32_32X32X32_FP8_FP8 : SMFMACInst<"v_smfmac_f32_32x32x32_fp8_fp8", "F32_32X32X32_F8", int_amdgcn_smfmac_f32_32x32x32_fp8_fp8>;1240} // End SubtargetPredicate = HasFP8Insts, is_gfx940_xdl = 11241 1242let SubtargetPredicate = HasGFX950Insts in {1243defm V_SMFMAC_F32_16X16X64_F16 : SMFMACInst<"v_smfmac_f32_16x16x64_f16", "F32_16X16X64_F16", int_amdgcn_smfmac_f32_16x16x64_f16>;1244defm V_SMFMAC_F32_32X32X32_F16 : SMFMACInst<"v_smfmac_f32_32x32x32_f16", "F32_32X32X32_F16", int_amdgcn_smfmac_f32_32x32x32_f16>;1245defm V_SMFMAC_F32_16X16X64_BF16 : SMFMACInst<"v_smfmac_f32_16x16x64_bf16", "F32_16X16X64_BF16", int_amdgcn_smfmac_f32_16x16x64_bf16>;1246defm V_SMFMAC_F32_32X32X32_BF16 : SMFMACInst<"v_smfmac_f32_32x32x32_bf16", "F32_32X32X32_BF16", int_amdgcn_smfmac_f32_32x32x32_bf16>;1247defm V_SMFMAC_I32_16X16X128_I8 : SMFMACInst<"v_smfmac_i32_16x16x128_i8", "I32_16X16X128_I8", int_amdgcn_smfmac_i32_16x16x128_i8>;1248defm V_SMFMAC_I32_32X32X64_I8 : SMFMACInst<"v_smfmac_i32_32x32x64_i8", "I32_32X32X64_I8", int_amdgcn_smfmac_i32_32x32x64_i8>;1249defm V_SMFMAC_F32_16X16X128_BF8_BF8 : SMFMACInst<"v_smfmac_f32_16x16x128_bf8_bf8", "F32_16X16X128_F8", int_amdgcn_smfmac_f32_16x16x128_bf8_bf8>;1250defm V_SMFMAC_F32_16X16X128_BF8_FP8 : SMFMACInst<"v_smfmac_f32_16x16x128_bf8_fp8", "F32_16X16X128_F8", int_amdgcn_smfmac_f32_16x16x128_bf8_fp8>;1251defm V_SMFMAC_F32_16X16X128_FP8_BF8 : SMFMACInst<"v_smfmac_f32_16x16x128_fp8_bf8", "F32_16X16X128_F8", int_amdgcn_smfmac_f32_16x16x128_fp8_bf8>;1252defm V_SMFMAC_F32_16X16X128_FP8_FP8 : SMFMACInst<"v_smfmac_f32_16x16x128_fp8_fp8", "F32_16X16X128_F8", int_amdgcn_smfmac_f32_16x16x128_fp8_fp8>;1253defm V_SMFMAC_F32_32X32X64_BF8_BF8 : SMFMACInst<"v_smfmac_f32_32x32x64_bf8_bf8", "F32_32X32X64_F8", int_amdgcn_smfmac_f32_32x32x64_bf8_bf8>;1254defm V_SMFMAC_F32_32X32X64_BF8_FP8 : SMFMACInst<"v_smfmac_f32_32x32x64_bf8_fp8", "F32_32X32X64_F8", int_amdgcn_smfmac_f32_32x32x64_bf8_fp8>;1255defm V_SMFMAC_F32_32X32X64_FP8_BF8 : SMFMACInst<"v_smfmac_f32_32x32x64_fp8_bf8", "F32_32X32X64_F8", int_amdgcn_smfmac_f32_32x32x64_fp8_bf8>;1256defm V_SMFMAC_F32_32X32X64_FP8_FP8 : SMFMACInst<"v_smfmac_f32_32x32x64_fp8_fp8", "F32_32X32X64_F8", int_amdgcn_smfmac_f32_32x32x64_fp8_fp8>;1257}1258 1259def MAIInstInfoTable : GenericTable {1260 let FilterClass = "MAIInst";1261 let CppTypeName = "MAIInstInfo";1262 let Fields = [1263 "Opcode", "is_dgemm", "is_gfx940_xdl"1264 ];1265 1266 let PrimaryKey = ["Opcode"];1267 let PrimaryKeyName = "getMAIInstInfoHelper";1268}1269 1270let isCommutable = 1, isReMaterializable = 1 in {1271 let SubtargetPredicate = HasPackedFP32Ops in {1272 defm V_PK_FMA_F32 : VOP3PInst<"v_pk_fma_f32", VOP3P_Profile<VOP_V2F32_V2F32_V2F32_V2F32, VOP3_PACKED>, any_fma>;1273 defm V_PK_MUL_F32 : VOP3PInst<"v_pk_mul_f32", VOP3P_Profile<VOP_V2F32_V2F32_V2F32, VOP3_PACKED>, any_fmul>;1274 defm V_PK_ADD_F32 : VOP3PInst<"v_pk_add_f32", VOP3P_Profile<VOP_V2F32_V2F32_V2F32, VOP3_PACKED>, any_fadd>;1275 } // End SubtargetPredicate = HasPackedFP32Ops1276 1277 let SubtargetPredicate = HasPkMovB32, isAsCheapAsAMove = 1 in1278 defm V_PK_MOV_B32 : VOP3PInst<"v_pk_mov_b32", VOP3P_Profile<VOP_V2I32_V2I32_V2I32, VOP3_PACKED>>;1279 1280 let SubtargetPredicate = HasBF16PackedInsts in {1281 defm V_PK_ADD_BF16 : VOP3PInst<"v_pk_add_bf16", VOP3P_Profile<VOP_V2BF16_V2BF16_V2BF16, VOP3_PACKED>, any_fadd>;1282 defm V_PK_MUL_BF16 : VOP3PInst<"v_pk_mul_bf16", VOP3P_Profile<VOP_V2BF16_V2BF16_V2BF16, VOP3_PACKED>, any_fmul>;1283 defm V_PK_MIN_NUM_BF16 : VOP3PInst<"v_pk_min_num_bf16", VOP3P_Profile<VOP_V2BF16_V2BF16_V2BF16, VOP3_PACKED>, fminnum_like>;1284 defm V_PK_MAX_NUM_BF16 : VOP3PInst<"v_pk_max_num_bf16", VOP3P_Profile<VOP_V2BF16_V2BF16_V2BF16, VOP3_PACKED>, fmaxnum_like>;1285 defm V_PK_FMA_BF16 : VOP3PInst<"v_pk_fma_bf16", VOP3P_Profile<VOP_V2BF16_V2BF16_V2BF16_V2BF16, VOP3_PACKED>, any_fma>;1286 1287 // Scalar pseudo used to emulate AMDGPUClamp.1288 // Expanded to V_PK_MAX_NUM_BF16 with unused high half.1289 // FIXME-TRUE16: Pseudo expansion of this won't work with True16.1290 let True16Predicate = UseFakeTrue16Insts in1291 defm V_MAX_BF16_PSEUDO : VOP3Inst <"v_max_bf16", VOP_BF16_BF16_BF16>;1292 }1293} // End isCommutable = 1, isReMaterializable = 11294 1295def : AMDGPUMnemonicAlias<"v_accvgpr_read", "v_accvgpr_read_b32">;1296def : AMDGPUMnemonicAlias<"v_accvgpr_write", "v_accvgpr_write_b32">;1297 1298class VOPProfileWMMA<VOPProfile P, string Suffix, RegisterOperand _Src01RC64, bit _HasClamp, bit _HasOpSel> : VOP3P_Profile<P> {1299 let DstRC = !if(!eq(Suffix, "_w32"), VDst_256, VDst_128);1300 let Src0RC64 = _Src01RC64;1301 let Src1RC64 = _Src01RC64;1302 let Src2RC64 = !if(!eq(Suffix, "_w32"), VISrc_256_f64, VISrc_128_f32);1303 let HasClamp = _HasClamp;1304 let HasOpSel = _HasOpSel;1305 let IsPacked = 1;1306 let IsWMMA = 1;1307}1308 1309def VOP_V8F32_V16F16_V16F16_V8F32 : VOPProfile <[v8f32, v16f16, v16f16, v8f32]>;1310def VOP_V8F32_V16I16_V16I16_V8F32 : VOPProfile <[v8f32, v16i16, v16i16, v8f32]>;1311def VOP_V16F16_V16F16_V16F16_V16F16 : VOPProfile <[v16f16, v16f16, v16f16, v16f16]>;1312def VOP_V16I16_V16I16_V16I16_V16I16 : VOPProfile <[v16i16, v16i16, v16i16, v16i16]>;1313def VOP_V8I32_V4I32_V4I32_V8I32 : VOPProfile <[v8i32, v4i32, v4i32, v8i32]>;1314def VOP_V8I32_V2I32_V2I32_V8I32 : VOPProfile <[v8i32, v2i32, v2i32, v8i32]>;1315 1316def VOP_V4F32_V16F16_V16F16_V4F32 : VOPProfile <[v4f32, v16f16, v16f16, v4f32]>;1317def VOP_V4F32_V16I16_V16I16_V4F32 : VOPProfile <[v4f32, v16i16, v16i16, v4f32]>;1318def VOP_V8F16_V16F16_V16F16_V8F16 : VOPProfile <[v8f16, v16f16, v16f16, v8f16]>;1319def VOP_V8I16_V16I16_V16I16_V8I16 : VOPProfile <[v8i16, v16i16, v16i16, v8i16]>;1320def VOP_V4I32_V2I32_V2I32_V4I32 : VOPProfile <[v4i32, v2i32, v2i32, v4i32]>;1321 1322 1323class WMMAType <bits<2> val> {1324 bit hasClamp = val{0};1325 bit hasOpsel = val{1};1326}1327 1328def WMMARegular : WMMAType<0b00>;1329def WMMAUIClamp : WMMAType<0b01>;1330def WMMAOpSel : WMMAType<0b10>;1331 1332class WMMARegularPat<Instruction Inst, SDPatternOperator node, VOPProfile P> :1333 GCNPat < (P.DstVT (node1334 (P.Src0VT (VOP3PMods P.Src0VT:$src0, i32:$src0_modifiers)),1335 (P.Src1VT (VOP3PMods P.Src1VT:$src1, i32:$src1_modifiers)),1336 (P.Src2VT (VOP3PMods P.Src2VT:$src2, i32:$src2_modifiers))1337 )),1338 (P.DstVT (Inst i32:$src0_modifiers, P.Src0VT:$src0, i32:$src1_modifiers, P.Src1VT:$src1, $src2_modifiers, P.Src2VT:$src2))1339>;1340 1341class WMMAOpSelPat<Instruction Inst, SDPatternOperator node, VOPProfile P> :1342 GCNPat < (P.DstVT (node1343 (P.Src0VT P.Src0VT:$src0),1344 (P.Src1VT P.Src1VT:$src1),1345 (P.Src2VT P.Src2VT:$src2), (WMMAOpSelVOP3PMods i32:$src2_modifiers)1346 )),1347 (P.DstVT (Inst (i32 8), P.Src0VT:$src0, (i32 8), P.Src1VT:$src1, i32:$src2_modifiers, P.Src2VT:$src2))1348>;1349 1350class WMMAUIClampPat<Instruction Inst, SDPatternOperator node, VOPProfile P> :1351 GCNPat < (P.DstVT (node1352 timm:$src0_modifiers, (P.Src0VT P.Src0VT:$src0),1353 timm:$src1_modifiers, (P.Src1VT P.Src1VT:$src1),1354 (P.Src2VT P.Src2VT:$src2), (i1 timm:$clamp)1355 )),1356 (P.DstVT (Inst (VOP3PModsNeg $src0_modifiers), P.Src0VT:$src0, (VOP3PModsNeg $src1_modifiers), P.Src1VT:$src1, (i32 8), P.Src2VT:$src2, i1:$clamp))1357>;1358 1359class WMMAOpcodeMapping<Instruction TwoAddr, Instruction ThreeAddr> {1360 Instruction Opcode2Addr = TwoAddr;1361 Instruction Opcode3Addr = ThreeAddr;1362 Predicate WaveSizePredicate;1363 Predicate SubtargetPredicate;1364 field bit is_wmma_xdl;1365}1366 1367def WMMAOpcode : GenericEnum {1368 let FilterClass = "VOP3P_Pseudo";1369}1370 1371class WMMAMappingTable : GenericTable {1372 let FilterClass = "WMMAOpcodeMapping";1373 let CppTypeName = "WMMAOpcodeMappingInfo";1374 let Fields = ["Opcode2Addr", "Opcode3Addr"];1375 string TypeOf_Opcode2Addr = "WMMAOpcode";1376 string TypeOf_Opcode3Addr = "WMMAOpcode";1377}1378 1379def WMMAOpcode2AddrMappingTable : WMMAMappingTable {1380 let PrimaryKey = ["Opcode2Addr"];1381 let PrimaryKeyName = "getWMMAMappingInfoFrom2AddrOpcode";1382}1383 1384def WMMAOpcode3AddrMappingTable : WMMAMappingTable {1385 let PrimaryKey = ["Opcode3Addr"];1386 let PrimaryKeyName = "getWMMAMappingInfoFrom3AddrOpcode";1387}1388 1389// The WMMA instruction has extra constraints:1390// Matrices A and B cannot overlap with D. C cannot partially overlap with D,1391// but it is OK for them to be the same (which is a typical case).1392//1393// We implement it as follows:1394// 1) Map the intrinsic to the pseudo where D is tied to C ($vdst = $src2).1395// 2) The pass twoaddressinstruction checks if src2 is live and if that is the case1396// it converts the default pseudo to the pseudo where src2 is not the same as vdst.1397// 3) @earlyclobber on the destination satisfies the constraint during RA.1398 1399multiclass WMMAInst<string Suffix, string Instr, VOPProfile P, SDPatternOperator node = null_frag, RegisterOperand _Src01RC64 = VRegSrc_256, WMMAType Type, bit convertibleTo3Addr> {1400 1401 defvar WMMAConstraints2Addr = "@earlyclobber $vdst,$vdst = $src2";1402 defvar WMMAConstraints3Addr = "@earlyclobber $vdst";1403 1404 defvar WMMAProfile = VOPProfileWMMA<P, Suffix, _Src01RC64, Type.hasClamp, Type.hasOpsel>;1405 let Mnemonic = Instr, mayRaiseFPException = 0, ReadsModeReg = 0 in {1406 let Constraints = WMMAConstraints2Addr, isConvertibleToThreeAddress = convertibleTo3Addr in {1407 def _twoaddr # Suffix : VOP3P_Pseudo<Instr # Suffix, WMMAProfile>;1408 }1409 }1410 if convertibleTo3Addr then {1411 let Mnemonic = Instr, mayRaiseFPException = 0, ReadsModeReg = 0 in {1412 let Constraints = WMMAConstraints3Addr, SchedRW = [Write32Bit, Write32Bit] in {1413 def _threeaddr # Suffix : VOP3P_Pseudo<Instr # Suffix, WMMAProfile>;1414 }1415 }1416 def : WMMAOpcodeMapping<!cast<Instruction>(NAME # _twoaddr # Suffix),1417 !cast<Instruction>(NAME # _threeaddr # Suffix)>;1418 }1419 1420 let SubtargetPredicate = isGFX11Only in {1421 if !eq(Type, WMMAOpSel) then {1422 def : WMMAOpSelPat<!cast<Instruction>(NAME # _twoaddr # Suffix), node, P>;1423 } else if !eq(Type, WMMAUIClamp) then {1424 def : WMMAUIClampPat<!cast<Instruction>(NAME # _twoaddr # Suffix), node, P>;1425 } else {1426 def : WMMARegularPat<!cast<Instruction>(NAME # _twoaddr # Suffix), node, P>;1427 }1428 }1429}1430 1431 1432 1433let WaveSizePredicate = isWave32 in {1434 defm V_WMMA_F32_16X16X16_F16 : WMMAInst<"_w32", "v_wmma_f32_16x16x16_f16", VOP_V8F32_V16F16_V16F16_V8F32, int_amdgcn_wmma_f32_16x16x16_f16, VRegSrc_256, WMMARegular, 1>;1435 defm V_WMMA_F32_16X16X16_BF16 : WMMAInst<"_w32", "v_wmma_f32_16x16x16_bf16", VOP_V8F32_V16I16_V16I16_V8F32, int_amdgcn_wmma_f32_16x16x16_bf16, VRegSrc_256, WMMARegular, 1>;1436 defm V_WMMA_F16_16X16X16_F16 : WMMAInst<"_w32", "v_wmma_f16_16x16x16_f16", VOP_V16F16_V16F16_V16F16_V16F16, int_amdgcn_wmma_f16_16x16x16_f16, VRegSrc_256, WMMAOpSel, 1>;1437 defm V_WMMA_BF16_16X16X16_BF16 : WMMAInst<"_w32", "v_wmma_bf16_16x16x16_bf16", VOP_V16I16_V16I16_V16I16_V16I16, int_amdgcn_wmma_bf16_16x16x16_bf16, VRegSrc_256, WMMAOpSel, 1>;1438 defm V_WMMA_F16_16X16X16_F16_TIED : WMMAInst<"_w32", "v_wmma_f16_16x16x16_f16", VOP_V16F16_V16F16_V16F16_V16F16, int_amdgcn_wmma_f16_16x16x16_f16_tied, VRegSrc_256, WMMAOpSel, 0>;1439 defm V_WMMA_BF16_16X16X16_BF16_TIED : WMMAInst<"_w32", "v_wmma_bf16_16x16x16_bf16", VOP_V16I16_V16I16_V16I16_V16I16, int_amdgcn_wmma_bf16_16x16x16_bf16_tied, VRegSrc_256, WMMAOpSel, 0>;1440 defm V_WMMA_I32_16X16X16_IU8 : WMMAInst<"_w32", "v_wmma_i32_16x16x16_iu8", VOP_V8I32_V4I32_V4I32_V8I32, int_amdgcn_wmma_i32_16x16x16_iu8, VRegSrc_128, WMMAUIClamp, 1>;1441 defm V_WMMA_I32_16X16X16_IU4 : WMMAInst<"_w32", "v_wmma_i32_16x16x16_iu4", VOP_V8I32_V2I32_V2I32_V8I32, int_amdgcn_wmma_i32_16x16x16_iu4, VRegSrc_64, WMMAUIClamp, 1>;1442}1443 1444let WaveSizePredicate = isWave64 in {1445 defm V_WMMA_F32_16X16X16_F16 : WMMAInst<"_w64", "v_wmma_f32_16x16x16_f16", VOP_V4F32_V16F16_V16F16_V4F32, int_amdgcn_wmma_f32_16x16x16_f16, VRegSrc_256, WMMARegular, 1>;1446 defm V_WMMA_F32_16X16X16_BF16 : WMMAInst<"_w64", "v_wmma_f32_16x16x16_bf16", VOP_V4F32_V16I16_V16I16_V4F32, int_amdgcn_wmma_f32_16x16x16_bf16, VRegSrc_256, WMMARegular, 1>;1447 defm V_WMMA_F16_16X16X16_F16 : WMMAInst<"_w64", "v_wmma_f16_16x16x16_f16", VOP_V8F16_V16F16_V16F16_V8F16, int_amdgcn_wmma_f16_16x16x16_f16, VRegSrc_256, WMMAOpSel, 1>;1448 defm V_WMMA_BF16_16X16X16_BF16 : WMMAInst<"_w64", "v_wmma_bf16_16x16x16_bf16", VOP_V8I16_V16I16_V16I16_V8I16, int_amdgcn_wmma_bf16_16x16x16_bf16, VRegSrc_256, WMMAOpSel, 1>;1449 defm V_WMMA_F16_16X16X16_F16_TIED : WMMAInst<"_w64", "v_wmma_f16_16x16x16_f16", VOP_V8F16_V16F16_V16F16_V8F16, int_amdgcn_wmma_f16_16x16x16_f16_tied, VRegSrc_256, WMMAOpSel, 0>;1450 defm V_WMMA_BF16_16X16X16_BF16_TIED : WMMAInst<"_w64", "v_wmma_bf16_16x16x16_bf16", VOP_V8I16_V16I16_V16I16_V8I16, int_amdgcn_wmma_bf16_16x16x16_bf16_tied, VRegSrc_256, WMMAOpSel, 0>;1451 defm V_WMMA_I32_16X16X16_IU8 : WMMAInst<"_w64", "v_wmma_i32_16x16x16_iu8", VOP_V4I32_V4I32_V4I32_V4I32, int_amdgcn_wmma_i32_16x16x16_iu8, VRegSrc_128, WMMAUIClamp, 1>;1452 defm V_WMMA_I32_16X16X16_IU4 : WMMAInst<"_w64", "v_wmma_i32_16x16x16_iu4", VOP_V4I32_V2I32_V2I32_V4I32, int_amdgcn_wmma_i32_16x16x16_iu4, VRegSrc_64, WMMAUIClamp, 1>;1453 1454}1455 1456class VOP3PWMMA_Profile<list<ValueType> ArgTy, bit _IsSWMMAC, int _IndexType,1457 bit _IsIU, bit _IsFP8BF8XF32, bit _Has_ImodOp = 0,1458 bit _HasMatrixFMT = 0, bit _HasMatrixScale = 0,1459 bit _Scale16 = 0, bit _HasMatrixReuse = 0, bit _IsF4 = 0>1460 : VOP3P_Profile<VOPProfile<ArgTy>> {1461 bit IsIU = _IsIU;1462 bit NoABMods = !or(_IsFP8BF8XF32, _IsF4); // No IMOD support for A and B1463 bit IsXF32 = !and(_IsFP8BF8XF32, !eq(ArgTy[1], v8f32));1464 1465 int IndexType = _IndexType;1466 let HasMatrixFMT = _HasMatrixFMT;1467 let HasMatrixScale = _HasMatrixScale;1468 bit Scale16 = _Scale16;1469 let HasMatrixReuse = _HasMatrixReuse;1470 1471 bit HasIModOp = _Has_ImodOp;1472 let HasClamp = !and(IsIU, !not(HasIModOp));1473 let IsPacked = 1;1474 let IsWMMA = !not(_IsSWMMAC);1475 let IsSWMMAC = _IsSWMMAC;1476 1477 bit IsAB_F64 = !or(!eq(ArgTy[1], v2f64), !eq(ArgTy[1], v4f64));1478 bit IsAB_F32 = !eq(ArgTy[1], v2f32);1479 bit IsAB_F16 = !or(!eq(ArgTy[1], v16f16), !eq(ArgTy[1], v8f16), !eq(ArgTy[1], v4f16));1480 bit IsAB_BF16 = !or(!eq(ArgTy[1], v16i16), !eq(ArgTy[1], v8i16), !eq(ArgTy[1], v4i16),1481 !eq(ArgTy[1], v16bf16), !eq(ArgTy[1], v8bf16), !eq(ArgTy[1], v4bf16));1482 bit IsF16BF16 = !or(IsAB_F16, IsAB_BF16);1483 1484 bit IsC_F64 = !eq(ArgTy[3], v8f64);1485 bit IsC_F32 = !or(!eq(ArgTy[3], v8f32), !eq(ArgTy[3], v4f32));1486 bit IsC_BF16 = !or(!eq(ArgTy[3], v8i16), !eq(ArgTy[3], v4i16),1487 !eq(ArgTy[3], v8bf16), !eq(ArgTy[3], v4bf16));1488 bit IsC_F16 = !or(!eq(ArgTy[3], v8f16), !eq(ArgTy[3], v4f16));1489 1490 bit NegLo01 = !not(NoABMods);1491 bit NegLo2 = !and(!not(IsIU), !not(IsXF32), IsWMMA);1492 bit NegHi01 = IsF16BF16; // Only F16BF16 can have neg_hi[0:1]1493 bit NegHi2 = !and(!not(IsIU), !not(IsXF32), IsWMMA);1494 bit NegLoAny = !or(NegLo01, NegLo2);1495 bit NegHiAny = !or(NegHi01, NegHi2);1496 1497 let DstRC = !cast<RegisterOperand>("VDst_"#ArgTy[0].Size);1498 let Src0RC64 = !cast<RegisterOperand>("VRegSrc_"#ArgTy[1].Size);1499 let Src1RC64 = !cast<RegisterOperand>("VRegSrc_"#ArgTy[2].Size);1500 let Src2RC64 = !if(IsSWMMAC, DstRC,1501 !cast<RegisterOperand>("VISrc_"#ArgTy[3].Size#1502 !cond(IsC_F64: "_f64",1503 IsC_F32: "_f32",1504 IsC_F16: "_f16",1505 IsC_BF16: "_bf16",1506 1: "_b32")));1507 ValueType ScaleTy = !if(Scale16, i64, i32);1508 1509 // For f16 and bf16 matrices A and B, each element can be modified by1510 // fneg(neg_lo,neg_hi = 1). For f32 and f64, neg_lo[0:1] is allowed, but1511 // neg_hi[0:1] is ignored. For iu4 and iu8 matrices A and B neg_lo is1512 // overloaded to mean unsigned/signed: neg_lo = 0 (u4 and u8) unsigned(zext)1513 // neg_lo = 1 (i4 and i8) signed(sext). For f16, bf16, f32 and f64 matrix C1514 // each element can be modified by fneg(neg_lo = 1) or fabs(neg_hi = 1).1515 1516 // Opcode | src0/src1 - matrix A/B | src2 - matrix C or Index1517 // ---------------------------------------------------------------------------1518 // wmma f64_f64 | neg_lo for neg A/B | neg_lo = 1 neg C(f64)1519 // | neg_hi ignored | neg_hi = 1 abs C(f64)1520 // ---------------------------------------------------------------------------1521 // wmma f32_f32 | neg_lo for neg A/B | neg_lo = 1 neg C(f32)1522 // | neg_hi ignored | neg_hi = 1 abs C(f32)1523 // ---------------------------------------------------------------------------1524 // wmma f32_xf32 | not allowed for xf32 | not allowed1525 // ---------------------------------------------------------------------------1526 // wmma f32_f16 | both neg_lo,neg_hi = 1 | neg_lo = 1 neg C(f32)1527 // wmma f32_bf16 | neg A/B (f16 or bf16) | neg_hi = 1 abs C(f32)1528 // ---------------------------------------------------------------------------1529 // wmma f16_f16 | both neg_lo,neg_hi = 1 | neg_lo = 1 neg C(f16 or bf16)1530 // wmma bf16_bf16 | neg A/B (f16 or bf16) | neg_hi = 1 abs C(f16 or bf16)1531 // ---------------------------------------------------------------------------1532 // wmma i32_iu8/iu4 | neg_lo = 0 u4/u8(zext) | not allowed for1533 // | neg_lo = 1 i4/i8(sext) | i32 matrices1534 // ---------------------------------------------------------------------------1535 // wmma f32_fp8/bf8 | not allowed for | neg_lo = 1 neg C(f32)1536 // | fp8 and bf8 matrices | neg_hi = 1 abs C(f32)1537 // ---------------------------------------------------------------------------1538 // wmma f16_fp8/bf8 | not allowed for | neg_lo = 1 neg C(f16)1539 // | fp8 and bf8 matrices | neg_hi = 1 abs C(f16)1540 // ---------------------------------------------------------------------------1541 // swmmac f32_f16 | both neg_lo,neg_hi = 1 | not allowed for sparse matrix1542 // swmmac f32_bf16 | neg A/B (f16 or bf16) | A Index - matrix C is in dst1543 // ---------------------------------------------------------------------------1544 // swmmac f16_f16 | both neg_lo,neg_hi = 1 | not allowed for sparse matrix1545 // swmmac bf16_bf16 | neg A/B (f16 or bf16) | A Index - matrix C is in dst1546 // ---------------------------------------------------------------------------1547 // swmmac i32_iu8/iu4 | neg_lo = 0 u4/u8(zext) | not allowed for sparse matrix1548 // | neg_lo = 1 i4/i8(sext) | A Index - matrix C is in dst1549 // ---------------------------------------------------------------------------1550 // swmmac f32_fp8/bf8 | not allowed for | not allowed for sparse matrix1551 // swmmac f16_fp8/bf8 | f8 and bf8 matrices | A Index - matrix C is in dst1552 // ---------------------------------------------------------------------------1553 1554 // pseudo1555 1556 // fp8bf8 and xf32 wmmas don't use src (0 and 1) modifiers, iu use neg_lo, f16 and bf161557 // use neg_lo and neg_hi. iu wmmas (C is i32) don't use src 2 modifiers,1558 // remaining wmmas(f16, bf16 and f8bf8) use neg_lo and neg_hi for C (C is f321559 // f16 or bf16). swmmac use index_key and don't use src 2 modifiers.1560 dag Src0Mods = !if(NoABMods, (ins), (ins PackedF16InputMods:$src0_modifiers));1561 dag Src1Mods = !if(NoABMods, (ins), (ins PackedF16InputMods:$src1_modifiers));1562 dag Src2Mods = !if(!or(IsIU, IsXF32, IsSWMMAC), (ins), (ins PackedF16InputMods:$src2_modifiers));1563 dag IndexKey = !cond(!eq(IndexType, 0) : (ins),1564 !eq(IndexType, 8) : (ins IndexKey8bit:$index_key_8bit),1565 !eq(IndexType, 16): (ins IndexKey16bit:$index_key_16bit),1566 !eq(IndexType, 32): (ins IndexKey32bit:$index_key_32bit));1567 dag MatrixFMT = !if(HasMatrixFMT, (ins MatrixAFMT:$matrix_a_fmt, MatrixBFMT:$matrix_b_fmt),1568 (ins));1569 dag MatrixScaleSrc = !if(HasMatrixScale,1570 !if(Scale16, (ins VCSrc_b64_Lo256:$scale_src0, VCSrc_b64_Lo256:$scale_src1),1571 (ins VCSrc_b32_Lo256:$scale_src0, VCSrc_b32_Lo256:$scale_src1)),1572 (ins));1573 dag MatrixScale = !if(HasMatrixScale, (ins MatrixAScale:$matrix_a_scale, MatrixBScale:$matrix_b_scale,1574 MatrixAScaleFmt:$matrix_a_scale_fmt, MatrixBScaleFmt:$matrix_b_scale_fmt),1575 (ins));1576 dag MatrixReuse = !if(HasMatrixReuse, (ins MatrixAReuse:$matrix_a_reuse, MatrixBReuse:$matrix_b_reuse), (ins));1577 dag Clamp = !if(HasClamp, (ins Clamp0:$clamp), (ins));1578 dag Neg = !cond(!and(NegLoAny, NegHiAny) : (ins neg_lo0:$neg_lo, neg_hi0:$neg_hi),1579 !and(NegLoAny, !not(NegHiAny)) : (ins neg_lo0:$neg_lo),1580 !and(!not(NegLoAny), !not(NegHiAny)) : (ins));1581 1582 let InsVOP3P = !con(Src0Mods, (ins Src0RC64:$src0), Src1Mods, (ins Src1RC64:$src1),1583 !cond(IsWMMA : !con(Src2Mods, (ins Src2RC64:$src2)),1584 IsSWMMAC : !con((ins DstRC:$srcTiedDef),1585 !if(!eq(IndexType, 32),1586 (ins VRegSrc_64:$src2),1587 (ins VRegSrc_32:$src2)),1588 IndexKey)),1589 MatrixScaleSrc, MatrixFMT, MatrixScale, MatrixReuse, Clamp, Neg);1590 1591 // asm1592 1593 string IndexKeyAsm = !cond(!eq(IndexType, 0) : "",1594 !eq(IndexType, 8) : "$index_key_8bit",1595 !eq(IndexType, 16) : "$index_key_16bit",1596 !eq(IndexType, 32) : "$index_key_32bit");1597 string MatrxFMTAsm = !if(HasMatrixFMT, "$matrix_a_fmt$matrix_b_fmt", "");1598 string MatrixScaleSrcAsm = !if(HasMatrixScale, ", $scale_src0, $scale_src1", "");1599 string MatrixScaleAsm = !if(HasMatrixScale, "$matrix_a_scale$matrix_b_scale$matrix_a_scale_fmt$matrix_b_scale_fmt", "");1600 string MatrixReuseAsm = !if(HasMatrixReuse, "$matrix_a_reuse$matrix_b_reuse", "");1601 string ClampAsm = !if(HasClamp, "$clamp", "");1602 string NegAsm = !cond(!and(NegLoAny, NegHiAny) : "$neg_lo$neg_hi",1603 !and(NegLoAny, !not(NegHiAny)) : "$neg_lo",1604 !and(!not(NegLoAny), !not(NegHiAny)) : "");1605 1606 let AsmVOP3P = "$vdst, $src0, $src1, $src2"#IndexKeyAsm#MatrixScaleSrcAsm#MatrxFMTAsm#MatrixScaleAsm#MatrixReuseAsm#NegAsm#ClampAsm;1607 1608 // isel patterns1609 bit IsAB_BF16_IMod0 = !and(IsAB_BF16, !not(HasIModOp));1610 bit IsAB_F16_IMod0 = !and(IsAB_F16, !not(HasIModOp));1611 bit IsAB_F32F64_IMod1 = !and(!or(IsAB_F64, IsAB_F32), HasIModOp);1612 bit IsAB_F16BF16_IMod1 = !and(!or(IsAB_F16, IsAB_BF16), HasIModOp);1613 dag Src0InPat = !cond(IsAB_F32F64_IMod1 : (ins timm:$src0_modifiers, Src0VT:$src0),1614 IsAB_F16BF16_IMod1 : (ins timm:$src0_modifiers, Src0VT:$src0),1615 IsAB_F16_IMod0 : (ins (Src0VT (WMMAModsF16Neg Src0VT:$src0, i32:$src0_modifiers))),1616 IsAB_BF16_IMod0 : (ins Src0VT:$src0),1617 IsIU : (ins timm:$src0_modifiers, Src0VT:$src0),1618 HasMatrixFMT : (ins timm:$matrix_a_fmt, Src0VT:$src0),1619 NoABMods : (ins Src0VT:$src0));1620 dag Src0OutPat = !cond(IsAB_F32F64_IMod1 : (ins (VOP3PModsNeg $src0_modifiers), Src0VT:$src0),1621 IsAB_F16BF16_IMod1 : (ins (VOP3PModsNegs $src0_modifiers), Src0VT:$src0),1622 IsAB_F16_IMod0 : (ins i32:$src0_modifiers, Src0VT:$src0),1623 IsAB_BF16_IMod0 : (ins (i32 8), Src0VT:$src0),1624 IsIU : (ins (VOP3PModsNeg $src0_modifiers), Src0VT:$src0),1625 NoABMods : (ins Src0VT:$src0));1626 dag Src1InPat = !cond(IsAB_F32F64_IMod1 : (ins timm:$src1_modifiers, Src1VT:$src1),1627 IsAB_F16BF16_IMod1 : (ins timm:$src1_modifiers, Src1VT:$src1),1628 IsAB_F16_IMod0 : (ins (Src1VT (WMMAModsF16Neg Src1VT:$src1, i32:$src1_modifiers))),1629 IsAB_BF16_IMod0 : (ins Src1VT:$src1),1630 IsIU : (ins timm:$src1_modifiers, Src1VT:$src1),1631 HasMatrixFMT : (ins timm:$matrix_b_fmt, Src1VT:$src1),1632 NoABMods : (ins Src1VT:$src1));1633 dag Src1OutPat = !cond(IsAB_F32F64_IMod1 : (ins (VOP3PModsNeg $src1_modifiers), Src1VT:$src1),1634 IsAB_F16BF16_IMod1 : (ins (VOP3PModsNegs $src1_modifiers), Src1VT:$src1),1635 IsAB_F16_IMod0 : (ins i32:$src1_modifiers, Src1VT:$src1),1636 IsAB_BF16_IMod0 : (ins (i32 8), Src1VT:$src1),1637 IsIU : (ins (VOP3PModsNeg $src1_modifiers), Src1VT:$src1),1638 NoABMods : (ins Src1VT:$src1));1639 bit IsC_IMod1 = !and(HasIModOp, IsWMMA, !not(IsIU), !not(IsXF32));1640 bit IsC_F32_IMod0 = !and(IsC_F32, !not(HasIModOp));1641 bit IsC_F16_IMod0 = !and(IsC_F16, !not(HasIModOp));1642 bit IsC_BF16_IMod0 = !and(IsC_BF16, !not(HasIModOp));1643 bit IsIUXF32 = !or(IsIU, IsXF32);1644 dag Src2InPatWmma = !cond(IsC_IMod1 : (ins timm:$src2_modifiers, Src2VT:$src2),1645 IsC_F32_IMod0 : (ins (Src2VT (WMMAModsF32NegAbs Src2VT:$src2, i32:$src2_modifiers))),1646 IsC_F16_IMod0 : (ins (Src2VT (WMMAModsF16NegAbs Src2VT:$src2, i32:$src2_modifiers))),1647 IsC_BF16_IMod0 : (ins Src2VT:$src2),1648 IsIUXF32 : (ins Src2VT:$src2),1649 IsSWMMAC : (ins));1650 dag Src2OutPatWmma = !cond(IsC_IMod1 : (ins (VOP3PModsNegAbs $src2_modifiers), Src2VT:$src2),1651 IsC_F32_IMod0 : (ins i32:$src2_modifiers, Src2VT:$src2),1652 IsC_F16_IMod0 : (ins i32:$src2_modifiers, Src2VT:$src2),1653 IsC_BF16_IMod0 : (ins (i32 8), Src2VT:$src2),1654 IsIUXF32 : (ins Src2VT:$src2),1655 IsSWMMAC : (ins));1656 dag ClampPat = !if(HasClamp, (ins i1:$clamp), (ins));1657 dag IndexInPat = !cond(!eq(IndexType, 0) : (ins i32:$src2),1658 !eq(IndexType, 8) : (ins (i32 (SWMMACIndex8 i32:$src2, i32:$index_key_8bit))),1659 !eq(IndexType, 16): (ins (i32 (SWMMACIndex16 i32:$src2, i32:$index_key_16bit))),1660 !eq(IndexType, 32): (ins (i64 (SWMMACIndex32 i64:$src2, i32:$index_key_32bit))));1661 dag IndexOutPat = !cond(!eq(IndexType, 0) : (ins i32:$src2),1662 !eq(IndexType, 8) : (ins i32:$src2, i32:$index_key_8bit),1663 !eq(IndexType, 16): (ins i32:$src2, i32:$index_key_16bit),1664 !eq(IndexType, 32): (ins i64:$src2, i32:$index_key_32bit));1665 dag MatrixFMTOutPat = !if(HasMatrixFMT, (ins i32:$matrix_a_fmt, i32:$matrix_b_fmt), (ins));1666 dag Src2InlineInPat = !con(!if(IsC_IMod1, (ins timm:$src2_modifiers), (ins)), (ins (Src2VT (WMMAVISrc Src2VT:$src2))));1667 dag Src2InlineOutPat = !con(!if(IsIUXF32, (ins), !if(IsC_IMod1, (ins (VOP3PModsNegAbs $src2_modifiers)), (ins (i32 8)))), (ins Src2VT:$src2));1668 dag MatrixScaleInPat = !if(HasMatrixScale, (ins timm:$matrix_a_scale, timm:$matrix_a_scale_fmt, ScaleTy:$scale_src0,1669 timm:$matrix_b_scale, timm:$matrix_b_scale_fmt, ScaleTy:$scale_src1),1670 (ins));1671 1672 dag MatrixReuseInPat = !if(HasMatrixReuse, (ins timm:$matrix_a_reuse, timm:$matrix_b_reuse), (ins));1673 dag MatrixScaleOutSrcPat = !if(HasMatrixScale, (ins ScaleTy:$scale_src0, ScaleTy:$scale_src1), (ins));1674 dag MatrixScaleOutModPat = !if(HasMatrixScale, (ins i32:$matrix_a_scale, i32:$matrix_b_scale, i32:$matrix_a_scale_fmt, i32:$matrix_b_scale_fmt), (ins));1675 dag MatrixReuseOutModPat = !if(HasMatrixReuse, (ins i1:$matrix_a_reuse, i1:$matrix_b_reuse), (ins));1676 1677 dag WmmaInPat = !con(Src0InPat, Src1InPat, Src2InPatWmma, MatrixScaleInPat, MatrixReuseInPat, ClampPat);1678 dag WmmaOutPat = !con(Src0OutPat, Src1OutPat, Src2OutPatWmma, MatrixScaleOutSrcPat, MatrixFMTOutPat,1679 MatrixScaleOutModPat, MatrixReuseOutModPat, ClampPat);1680 1681 dag SwmmacInPat = !con(Src0InPat, Src1InPat, (ins Src2VT:$srcTiedDef), IndexInPat, MatrixReuseInPat, ClampPat);1682 dag SwmmacOutPat = !con(Src0OutPat, Src1OutPat, (ins Src2VT:$srcTiedDef), IndexOutPat, MatrixReuseOutModPat, ClampPat);1683 1684 // wmma pattern where src2 is inline imm uses _threeaddr pseudo,1685 // can't use _twoaddr since it would violate src2 tied to vdst constraint.1686 dag WmmaInlineInPat = !con(Src0InPat, Src1InPat, Src2InlineInPat, MatrixScaleInPat, MatrixReuseInPat, ClampPat);1687 dag WmmaInlineOutPat = !con(Src0OutPat, Src1OutPat, Src2InlineOutPat, MatrixScaleOutSrcPat,1688 MatrixFMTOutPat, MatrixScaleOutModPat, MatrixReuseOutModPat, ClampPat);1689}1690 1691def WMMAInstInfoTable : GenericTable {1692 let FilterClass = "WMMAInstInfo";1693 let CppTypeName = "WMMAInstInfo";1694 let Fields = ["Opcode", "is_wmma_xdl"];1695 1696 let PrimaryKey = ["Opcode"];1697 let PrimaryKeyName = "getWMMAInstInfoHelper";1698}1699 1700class WMMAInstInfo {1701 Instruction Opcode = !cast<Instruction>(NAME);1702 bit is_wmma_xdl = 0;1703}1704 1705multiclass WMMAInstGFX12<string Instr, VOP3PWMMA_Profile WMMAProfile, string PseudoInstrSuffix, bit DiffVdstSrc2 = 0> {1706 1707 defvar WMMAConstraints2Addr = !if(DiffVdstSrc2, "@earlyclobber $vdst", "@earlyclobber $vdst,$vdst = $src2");1708 defvar WMMAConstraints3Addr = "@earlyclobber $vdst";1709 1710 let Mnemonic = Instr, mayRaiseFPException = 0, ReadsModeReg = 0, isConvergent = 1 in {1711 let Constraints = WMMAConstraints2Addr, isConvertibleToThreeAddress = 1 in1712 def _twoaddr : VOP3P_Pseudo<Instr, WMMAProfile>, WMMAInstInfo {1713 let PseudoInstr = Instr#PseudoInstrSuffix;1714 let FixedSize = WMMAProfile.HasMatrixScale;1715 let Size = !if(WMMAProfile.HasMatrixScale, 16, 8);1716 }1717 1718 let Constraints = WMMAConstraints3Addr, SchedRW = [Write32Bit, Write32Bit] in1719 def _threeaddr : VOP3P_Pseudo<Instr, WMMAProfile>, WMMAInstInfo {1720 let PseudoInstr = Instr#PseudoInstrSuffix;1721 let FixedSize = WMMAProfile.HasMatrixScale;1722 let Size = !if(WMMAProfile.HasMatrixScale, 16, 8);1723 }1724 1725 }1726 def : WMMAOpcodeMapping<!cast<Instruction>(NAME # _twoaddr),1727 !cast<Instruction>(NAME # _threeaddr)>;1728}1729 1730multiclass SWMMACInstGFX12<string Instr, VOP3PWMMA_Profile WMMAProfile, string PseudoInstrSuffix> {1731 def _twoaddr : VOP3P_Pseudo<Instr, WMMAProfile>, WMMAInstInfo {1732 let Mnemonic = Instr;1733 let PseudoInstr = Instr#PseudoInstrSuffix;1734 let mayRaiseFPException = 0;1735 let ReadsModeReg = 0;1736 let AsmMatchConverter = "cvtSWMMAC";1737 let isConvergent = 1;1738 let Constraints = "@earlyclobber $vdst,$vdst = $srcTiedDef";1739 }1740}1741 1742// First argument in Profile is types for matrices D, A, B and C (D = A * B + C)1743// as used by llvm ir, types are vectors(with matrix elements)1744// wave32:1745// For 16x16 matrices, lanes 0 to 31 will have 8 matrix elts,1746// for 16 x 32 16 elts and for 16 x 64 lanes have 32 elts.1747// wave64:1748// lanes will have half the size of elements in lanes compared to wave32 with1749// exception of 16x16_iu4: lanes0-31 will have 8xi4, remaining lanes are ignored1750 1751// general idea on element distribution differences:1752// wave32: lane n has 8 matrix elements1753// wave64: lane n has first 4, lane n+32 has other 4 elements1754 1755// index size, for each 2 elements in lane you need 4bits in index1756 1757// Non-standard types (iu8, iu4, fp8, bf8) will be packed in vectors of i32s.1758// Original type for them is in comment on the right and refers to A and B.1759 1760def F32_F16_WMMA_w32 : VOP3PWMMA_Profile<[v8f32, v8f16, v8f16, v8f32], 0, 0, 0, 0>;1761def F32_BF16_WMMA_w32 : VOP3PWMMA_Profile<[v8f32, v8i16, v8i16, v8f32], 0, 0, 0, 0>;1762def F16_F16_WMMA_w32 : VOP3PWMMA_Profile<[v8f16, v8f16, v8f16, v8f16], 0, 0, 0, 0>;1763def BF16_BF16_WMMA_w32 : VOP3PWMMA_Profile<[v8i16, v8i16, v8i16, v8i16], 0, 0, 0, 0>;1764def I32_IU8_WMMA_w32 : VOP3PWMMA_Profile<[v8i32, v2i32, v2i32, v8i32], 0, 0, 1, 0>; // 8xi81765def I32_IU4X16_WMMA_w32 : VOP3PWMMA_Profile<[v8i32, i32, i32, v8i32], 0, 0, 1, 0>; // 8xi41766def F32_FP8BF8_WMMA_w32 : VOP3PWMMA_Profile<[v8f32, v2i32, v2i32, v8f32], 0, 0, 0, 1>; // 8xf81767def I32_IU4X32_WMMA_w32 : VOP3PWMMA_Profile<[v8i32, v2i32, v2i32, v8i32], 0, 0, 1, 0>; // 16xi41768 1769def F32_F16_WMMA_w64 : VOP3PWMMA_Profile<[v4f32, v4f16, v4f16, v4f32], 0, 0, 0, 0>;1770def F32_BF16_WMMA_w64 : VOP3PWMMA_Profile<[v4f32, v4i16, v4i16, v4f32], 0, 0, 0, 0>;1771def F16_F16_WMMA_w64 : VOP3PWMMA_Profile<[v4f16, v4f16, v4f16, v4f16], 0, 0, 0, 0>;1772def BF16_BF16_WMMA_w64 : VOP3PWMMA_Profile<[v4i16, v4i16, v4i16, v4i16], 0, 0, 0, 0>;1773def I32_IU8_WMMA_w64 : VOP3PWMMA_Profile<[v4i32, i32, i32, v4i32], 0, 0, 1, 0>; // 4xi81774def I32_IU4X16_WMMA_w64 : VOP3PWMMA_Profile<[v4i32, i32, i32, v4i32], 0, 0, 1, 0>; // 8xi4 *1775def F32_FP8BF8_WMMA_w64 : VOP3PWMMA_Profile<[v4f32, i32, i32, v4f32], 0, 0, 0, 1>; // 4xf81776def I32_IU4X32_WMMA_w64 : VOP3PWMMA_Profile<[v4i32, i32, i32, v4i32], 0, 0, 1, 0>; // 8xi41777 1778def F32_F16_SWMMAC_w32 : VOP3PWMMA_Profile<[v8f32, v8f16, v16f16, v8f32], 1, 16, 0, 0>;1779def F32_BF16_SWMMAC_w32 : VOP3PWMMA_Profile<[v8f32, v8i16, v16i16, v8f32], 1, 16, 0, 0>;1780def F16_F16_SWMMAC_w32 : VOP3PWMMA_Profile<[v8f16, v8f16, v16f16, v8f16], 1, 16, 0, 0>;1781def BF16_BF16_SWMMAC_w32 : VOP3PWMMA_Profile<[v8i16, v8i16, v16i16, v8i16], 1, 16, 0, 0>;1782def I32_IU8_SWMMAC_w32 : VOP3PWMMA_Profile<[v8i32, v2i32, v4i32, v8i32], 1, 16, 1, 0>; // 8xi8, 16xi81783def I32_IU4X32_SWMMAC_w32 : VOP3PWMMA_Profile<[v8i32, i32, v2i32, v8i32], 1, 16, 1, 0>; // 8xi4, 16xi41784def I32_IU4X64_SWMMAC_w32 : VOP3PWMMA_Profile<[v8i32, v2i32, v4i32, v8i32], 1, 0, 1, 0>; // 16xi4, 32xi4 **1785def F32_FP8BF8_SWMMAC_w32 : VOP3PWMMA_Profile<[v8f32, v2i32, v4i32, v8f32], 1, 16, 0, 1>; // 8xf8, 16xf81786 1787def F32_F16_SWMMAC_w64 : VOP3PWMMA_Profile<[v4f32, v4f16, v8f16, v4f32], 1, 8, 0, 0>;1788def F32_BF16_SWMMAC_w64 : VOP3PWMMA_Profile<[v4f32, v4i16, v8i16, v4f32], 1, 8, 0, 0>;1789def F16_F16_SWMMAC_w64 : VOP3PWMMA_Profile<[v4f16, v4f16, v8f16, v4f16], 1, 8, 0, 0>;1790def BF16_BF16_SWMMAC_w64 : VOP3PWMMA_Profile<[v4i16, v4i16, v8i16, v4i16], 1, 8, 0, 0>;1791def I32_IU8_SWMMAC_w64 : VOP3PWMMA_Profile<[v4i32, i32, v2i32, v4i32], 1, 8, 1, 0>; // 4xi8, 8xi81792def I32_IU4X32_SWMMAC_w64 : VOP3PWMMA_Profile<[v4i32, i32, i32, v4i32], 1, 16, 1, 0>; // 8xi4, 8xi4 ***1793def I32_IU4X64_SWMMAC_w64 : VOP3PWMMA_Profile<[v4i32, i32, v2i32, v4i32], 1, 16, 1, 0>; // 8xi4, 16xi41794def F32_FP8BF8_SWMMAC_w64 : VOP3PWMMA_Profile<[v4f32, i32, v2i32, v4f32], 1, 8, 0, 1>; // 4xf8, 8xf81795 1796// * IU4X16_WMMA_w64 lanes 0-31 will have 8xi4, remaining lanes are ignored1797// ** IU4X64_SWMMAC_w32 index is i32, index_key is not used1798// *** IU4X32_SWMMAC_w64 lanes 0-31 will have 8xi4 remaining lanes are ignored1799// for matrix A, index is i16; Matrix B uses all lanes1800 1801def F32_F32_WMMA_w32 : VOP3PWMMA_Profile<[v8f32, v2f32, v2f32, v8f32], 0, 0, 0, 0, 1, 0, 0, 0, 1>;1802def F32_BF16X32_WMMA_w32 : VOP3PWMMA_Profile<[v8f32, v16bf16, v16bf16, v8f32], 0, 0, 0, 0, 1, 0, 0, 0, 1>;1803def F32_F16X32_WMMA_w32 : VOP3PWMMA_Profile<[v8f32, v16f16, v16f16, v8f32], 0, 0, 0, 0, 1, 0, 0, 0, 1>;1804def F16_F16X32_WMMA_w32 : VOP3PWMMA_Profile<[v8f16, v16f16, v16f16, v8f16], 0, 0, 0, 0, 1, 0, 0, 0, 1>;1805def BF16_BF16X32_WMMA_w32 : VOP3PWMMA_Profile<[v8bf16, v16bf16, v16bf16, v8bf16], 0, 0, 0, 0, 1, 0, 0, 0, 1>;1806def BF16F32_BF16_WMMA_w32 : VOP3PWMMA_Profile<[v8bf16, v16bf16, v16bf16, v8f32], 0, 0, 0, 0, 1, 0, 0, 0, 1>;1807def F32_FP8BF8X64_WMMA_w32 : VOP3PWMMA_Profile<[v8f32, v8i32, v8i32, v8f32], 0, 0, 0, 1, 1, 0, 0, 0, 1>;1808def F32_FP8BF8X128_WMMA_w32 : VOP3PWMMA_Profile<[v8f32, v16i32, v16i32, v8f32], 0, 0, 0, 1, 1, 0, 0, 0, 1>;1809def F16_FP8BF8X64_WMMA_w32 : VOP3PWMMA_Profile<[v8f16, v8i32, v8i32, v8f16], 0, 0, 0, 1, 1, 0, 0, 0, 1>;1810def F16_FP8BF8X128_WMMA_w32 : VOP3PWMMA_Profile<[v8f16, v16i32, v16i32, v8f16], 0, 0, 0, 1, 1, 0, 0, 0, 1>;1811def F32_32X16X128_F4_WMMA_w32 : VOP3PWMMA_Profile<[v16f32, v16i32, v8i32, v16f32], 0, 0, 0, 0, 1, 0, 0, 0, 0, 1>;1812def I32_IU8X64_WMMA_w32 : VOP3PWMMA_Profile<[v8i32, v8i32, v8i32, v8i32], 0, 0, 1, 0, 1, 0, 0, 0, 1>;1813def F32_32X16X128_F4_SCALE_w32 : VOP3PWMMA_Profile<[v16f32, v16i32, v8i32, v16f32], 0, 0, 0, 1, 1, 0, 1, 0, 1>;1814def F32_32X16X128_F4_SCALE16_w32 : VOP3PWMMA_Profile<[v16f32, v16i32, v8i32, v16f32], 0, 0, 0, 1, 1, 0, 1, 1, 1>;1815def F32_F16X64_SWMMAC_w32 : VOP3PWMMA_Profile<[v8f32, v16f16, v32f16, v8f32], 1, 16, 0, 0, 1, 0, 0, 0, 1>;1816def F32_BF16X64_SWMMAC_w32 : VOP3PWMMA_Profile<[v8f32, v16bf16, v32bf16, v8f32], 1, 16, 0, 0, 1, 0, 0, 0, 1>;1817def F16_F16X64_SWMMAC_w32 : VOP3PWMMA_Profile<[v8f16, v16f16, v32f16, v8f16], 1, 16, 0, 0, 1, 0, 0, 0, 1>;1818def BF16_BF16X64_SWMMAC_w32 : VOP3PWMMA_Profile<[v8bf16, v16bf16, v32bf16, v8bf16], 1, 16, 0, 0, 1, 0, 0, 0, 1>;1819def F32_FP8BF8X128_SWMMAC_w32 : VOP3PWMMA_Profile<[v8f32, v8i32, v16i32, v8f32], 1, 32, 0, 1, 1, 0, 0, 0, 1>;1820def F16_FP8BF8X128_SWMMAC_w32 : VOP3PWMMA_Profile<[v8f16, v8i32, v16i32, v8f16], 1, 32, 0, 1, 1, 0, 0, 0, 1>;1821def I32_IU8X128_SWMMAC_w32 : VOP3PWMMA_Profile<[v8i32, v8i32, v16i32, v8i32], 1, 32, 1, 0, 1, 0, 0, 0, 1>;1822 1823multiclass WMMA_F8F6F4_Profiles<bit HasMatrixScale, bit Scale16, bit HasMatrixReuse> {1824 def _f8_f8_w32 : VOP3PWMMA_Profile<[v8f32, v16i32, v16i32, v8f32], 0, 0, 0, 1, 1, 1, HasMatrixScale, Scale16, HasMatrixReuse>;1825 def _f8_f6_w32 : VOP3PWMMA_Profile<[v8f32, v16i32, v12i32, v8f32], 0, 0, 0, 1, 1, 1, HasMatrixScale, Scale16, HasMatrixReuse>;1826 def _f8_f4_w32 : VOP3PWMMA_Profile<[v8f32, v16i32, v8i32, v8f32], 0, 0, 0, 1, 1, 1, HasMatrixScale, Scale16, HasMatrixReuse>;1827 def _f6_f8_w32 : VOP3PWMMA_Profile<[v8f32, v12i32, v16i32, v8f32], 0, 0, 0, 1, 1, 1, HasMatrixScale, Scale16, HasMatrixReuse>;1828 def _f6_f6_w32 : VOP3PWMMA_Profile<[v8f32, v12i32, v12i32, v8f32], 0, 0, 0, 1, 1, 1, HasMatrixScale, Scale16, HasMatrixReuse>;1829 def _f6_f4_w32 : VOP3PWMMA_Profile<[v8f32, v12i32, v8i32, v8f32], 0, 0, 0, 1, 1, 1, HasMatrixScale, Scale16, HasMatrixReuse>;1830 def _f4_f8_w32 : VOP3PWMMA_Profile<[v8f32, v8i32, v16i32, v8f32], 0, 0, 0, 1, 1, 1, HasMatrixScale, Scale16, HasMatrixReuse>;1831 def _f4_f6_w32 : VOP3PWMMA_Profile<[v8f32, v8i32, v12i32, v8f32], 0, 0, 0, 1, 1, 1, HasMatrixScale, Scale16, HasMatrixReuse>;1832 def _f4_f4_w32 : VOP3PWMMA_Profile<[v8f32, v8i32, v8i32, v8f32], 0, 0, 0, 1, 1, 1, HasMatrixScale, Scale16, HasMatrixReuse>;1833}1834 1835defm F32_16X16X128_F8F6F4 : WMMA_F8F6F4_Profiles<0, 0, 0>;1836defm F32_16X16X128_F8F6F4_SCALE : WMMA_F8F6F4_Profiles<1, 0, 1>;1837defm F32_16X16X128_F8F6F4_SCALE16 : WMMA_F8F6F4_Profiles<1, 1, 1>;1838 1839class VOP_WMMA_LD_SCALE<ValueType vt, RegisterOperand RC> : VOP3P_Profile<VOPProfile<[untyped, vt, vt, untyped]>> {1840 let HasMatrixScale = 1;1841 let HasMatrixReuse = 1;1842 let HasNeg = 0;1843 let Src0RC64 = RC;1844 let Src1RC64 = RC;1845 let Ins64 = (ins Src0RC64:$src0, Src1RC64:$src1, MatrixAScale:$matrix_a_scale, MatrixBScale:$matrix_b_scale,1846 MatrixAScaleFmt:$matrix_a_scale_fmt, MatrixBScaleFmt:$matrix_b_scale_fmt,1847 MatrixAReuse:$matrix_a_reuse, MatrixBReuse:$matrix_b_reuse);1848 let AsmVOP3P = " $src0, $src1$matrix_a_scale$matrix_b_scale$matrix_a_scale_fmt$matrix_b_scale_fmt$matrix_a_reuse$matrix_b_reuse";1849}1850 1851multiclass WMMAInst_SrcFormats_mc<string OpName, string Profile> {1852 foreach I = ["f8_f8", "f8_f6", "f8_f4", "f6_f8", "f6_f6", "f6_f4", "f4_f8", "f4_f6", "f4_f4"] in {1853 defm _#I#_w32 : WMMAInstGFX12<OpName # "_" # I # "_w32", !cast<VOP3PWMMA_Profile>(Profile # "_" # I # "_w32"), "_w32">;1854 }1855}1856 1857let WaveSizePredicate = isWave32 in {1858let SubtargetPredicate = isGFX125xOnly in {1859defm V_WMMA_F32_16X16X4_F32_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x4_f32", F32_F32_WMMA_w32, "_w32">;1860 1861let is_wmma_xdl = 1 in {1862defm V_WMMA_F32_16X16X32_BF16_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x32_bf16", F32_BF16X32_WMMA_w32, "_w32">;1863defm V_WMMA_BF16_16X16X32_BF16_w32 : WMMAInstGFX12<"v_wmma_bf16_16x16x32_bf16", BF16_BF16X32_WMMA_w32, "_w32">;1864defm V_WMMA_BF16F32_16X16X32_BF16_w32 : WMMAInstGFX12<"v_wmma_bf16f32_16x16x32_bf16", BF16F32_BF16_WMMA_w32, "_w32", 1>;1865defm V_WMMA_F32_16X16X64_FP8_FP8_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x64_fp8_fp8", F32_FP8BF8X64_WMMA_w32, "_w32">;1866defm V_WMMA_F32_16X16X64_FP8_BF8_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x64_fp8_bf8", F32_FP8BF8X64_WMMA_w32, "_w32">;1867defm V_WMMA_F32_16X16X64_BF8_FP8_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x64_bf8_fp8", F32_FP8BF8X64_WMMA_w32, "_w32">;1868defm V_WMMA_F32_16X16X64_BF8_BF8_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x64_bf8_bf8", F32_FP8BF8X64_WMMA_w32, "_w32">;1869defm V_WMMA_F16_16X16X64_FP8_FP8_w32 : WMMAInstGFX12<"v_wmma_f16_16x16x64_fp8_fp8", F16_FP8BF8X64_WMMA_w32, "_w32">;1870defm V_WMMA_F16_16X16X64_FP8_BF8_w32 : WMMAInstGFX12<"v_wmma_f16_16x16x64_fp8_bf8", F16_FP8BF8X64_WMMA_w32, "_w32">;1871defm V_WMMA_F16_16X16X64_BF8_FP8_w32 : WMMAInstGFX12<"v_wmma_f16_16x16x64_bf8_fp8", F16_FP8BF8X64_WMMA_w32, "_w32">;1872defm V_WMMA_F16_16X16X64_BF8_BF8_w32 : WMMAInstGFX12<"v_wmma_f16_16x16x64_bf8_bf8", F16_FP8BF8X64_WMMA_w32, "_w32">;1873defm V_WMMA_I32_16X16X64_IU8_w32 : WMMAInstGFX12<"v_wmma_i32_16x16x64_iu8", I32_IU8X64_WMMA_w32, "_w32">;1874defm V_WMMA_F32_16X16X32_F16_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x32_f16", F32_F16X32_WMMA_w32, "_w32">;1875defm V_WMMA_F16_16X16X32_F16_w32 : WMMAInstGFX12<"v_wmma_f16_16x16x32_f16", F16_F16X32_WMMA_w32, "_w32">;1876defm V_WMMA_F16_16X16X128_FP8_FP8_w32 : WMMAInstGFX12<"v_wmma_f16_16x16x128_fp8_fp8", F16_FP8BF8X128_WMMA_w32, "_w32">;1877defm V_WMMA_F16_16X16X128_FP8_BF8_w32 : WMMAInstGFX12<"v_wmma_f16_16x16x128_fp8_bf8", F16_FP8BF8X128_WMMA_w32, "_w32">;1878defm V_WMMA_F16_16X16X128_BF8_FP8_w32 : WMMAInstGFX12<"v_wmma_f16_16x16x128_bf8_fp8", F16_FP8BF8X128_WMMA_w32, "_w32">;1879defm V_WMMA_F16_16X16X128_BF8_BF8_w32 : WMMAInstGFX12<"v_wmma_f16_16x16x128_bf8_bf8", F16_FP8BF8X128_WMMA_w32, "_w32">;1880defm V_WMMA_F32_16X16X128_FP8_FP8_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x128_fp8_fp8", F32_FP8BF8X128_WMMA_w32, "_w32">;1881defm V_WMMA_F32_16X16X128_FP8_BF8_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x128_fp8_bf8", F32_FP8BF8X128_WMMA_w32, "_w32">;1882defm V_WMMA_F32_16X16X128_BF8_FP8_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x128_bf8_fp8", F32_FP8BF8X128_WMMA_w32, "_w32">;1883defm V_WMMA_F32_16X16X128_BF8_BF8_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x128_bf8_bf8", F32_FP8BF8X128_WMMA_w32, "_w32">;1884defm V_WMMA_F32_32X16X128_F4_w32 : WMMAInstGFX12<"v_wmma_f32_32x16x128_f4", F32_32X16X128_F4_WMMA_w32, "_w32">;1885 1886defm V_SWMMAC_F32_16X16X64_BF16_w32 : SWMMACInstGFX12<"v_swmmac_f32_16x16x64_bf16", F32_BF16X64_SWMMAC_w32, "_w32">;1887defm V_SWMMAC_BF16_16X16X64_BF16_w32 : SWMMACInstGFX12<"v_swmmac_bf16_16x16x64_bf16", BF16_BF16X64_SWMMAC_w32, "_w32">;1888defm V_SWMMAC_BF16F32_16X16X64_BF16_w32 : SWMMACInstGFX12<"v_swmmac_bf16f32_16x16x64_bf16", F32_BF16X64_SWMMAC_w32, "_w32">;1889defm V_SWMMAC_F32_16X16X128_FP8_FP8_w32 : SWMMACInstGFX12<"v_swmmac_f32_16x16x128_fp8_fp8", F32_FP8BF8X128_SWMMAC_w32, "_w32">;1890defm V_SWMMAC_F32_16X16X128_FP8_BF8_w32 : SWMMACInstGFX12<"v_swmmac_f32_16x16x128_fp8_bf8", F32_FP8BF8X128_SWMMAC_w32, "_w32">;1891defm V_SWMMAC_F32_16X16X128_BF8_FP8_w32 : SWMMACInstGFX12<"v_swmmac_f32_16x16x128_bf8_fp8", F32_FP8BF8X128_SWMMAC_w32, "_w32">;1892defm V_SWMMAC_F32_16X16X128_BF8_BF8_w32 : SWMMACInstGFX12<"v_swmmac_f32_16x16x128_bf8_bf8", F32_FP8BF8X128_SWMMAC_w32, "_w32">;1893defm V_SWMMAC_F16_16X16X128_FP8_FP8_w32 : SWMMACInstGFX12<"v_swmmac_f16_16x16x128_fp8_fp8", F16_FP8BF8X128_SWMMAC_w32, "_w32">;1894defm V_SWMMAC_F16_16X16X128_FP8_BF8_w32 : SWMMACInstGFX12<"v_swmmac_f16_16x16x128_fp8_bf8", F16_FP8BF8X128_SWMMAC_w32, "_w32">;1895defm V_SWMMAC_F16_16X16X128_BF8_FP8_w32 : SWMMACInstGFX12<"v_swmmac_f16_16x16x128_bf8_fp8", F16_FP8BF8X128_SWMMAC_w32, "_w32">;1896defm V_SWMMAC_F16_16X16X128_BF8_BF8_w32 : SWMMACInstGFX12<"v_swmmac_f16_16x16x128_bf8_bf8", F16_FP8BF8X128_SWMMAC_w32, "_w32">;1897defm V_SWMMAC_I32_16X16X128_IU8_w32 : SWMMACInstGFX12<"v_swmmac_i32_16x16x128_iu8", I32_IU8X128_SWMMAC_w32, "_w32">;1898defm V_SWMMAC_F32_16X16X64_F16_w32 : SWMMACInstGFX12<"v_swmmac_f32_16x16x64_f16", F32_F16X64_SWMMAC_w32, "_w32">;1899defm V_SWMMAC_F16_16X16X64_F16_w32 : SWMMACInstGFX12<"v_swmmac_f16_16x16x64_f16", F16_F16X64_SWMMAC_w32, "_w32">;1900 1901defm V_WMMA_F32_16X16X128_F8F6F4 : WMMAInst_SrcFormats_mc<"v_wmma_f32_16x16x128_f8f6f4", "F32_16X16X128_F8F6F4">;1902defm V_WMMA_SCALE_F32_16X16X128_F8F6F4 : WMMAInst_SrcFormats_mc<"v_wmma_scale_f32_16x16x128_f8f6f4", "F32_16X16X128_F8F6F4_SCALE">;1903defm V_WMMA_SCALE16_F32_16X16X128_F8F6F4 : WMMAInst_SrcFormats_mc<"v_wmma_scale16_f32_16x16x128_f8f6f4", "F32_16X16X128_F8F6F4_SCALE16">;1904 1905defm V_WMMA_SCALE_F32_32X16X128_F4_w32 : WMMAInstGFX12<"v_wmma_scale_f32_32x16x128_f4", F32_32X16X128_F4_SCALE_w32, "_w32">;1906defm V_WMMA_SCALE16_F32_32X16X128_F4_w32 : WMMAInstGFX12<"v_wmma_scale16_f32_32x16x128_f4", F32_32X16X128_F4_SCALE16_w32, "_w32">;1907} // End is_wmma_xdl = 1.1908 1909let isConvergent = 1 in {1910 defm V_WMMA_LD_SCALE_PAIRED_B32 : VOP3PInst<"v_wmma_ld_scale_paired_b32", VOP_WMMA_LD_SCALE<i32, VCSrc_b32_Lo256>>;1911 defm V_WMMA_LD_SCALE16_PAIRED_B64 : VOP3PInst<"v_wmma_ld_scale16_paired_b64", VOP_WMMA_LD_SCALE<i64, VCSrc_b64_Lo256>>;1912}1913} // End SubtargetPredicate = isGFX125xOnly1914} // End WaveSizePredicate = isWave321915 1916let WaveSizePredicate = isWave32 in {1917defm V_WMMA_F32_16X16X16_F16_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x16_f16", F32_F16_WMMA_w32, "_w32">;1918defm V_WMMA_F32_16X16X16_BF16_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x16_bf16", F32_BF16_WMMA_w32, "_w32">;1919defm V_WMMA_F16_16X16X16_F16_w32 : WMMAInstGFX12<"v_wmma_f16_16x16x16_f16", F16_F16_WMMA_w32, "_w32">;1920defm V_WMMA_BF16_16X16X16_BF16_w32 : WMMAInstGFX12<"v_wmma_bf16_16x16x16_bf16", BF16_BF16_WMMA_w32, "_w32">;1921defm V_WMMA_I32_16X16X16_IU8_w32 : WMMAInstGFX12<"v_wmma_i32_16x16x16_iu8", I32_IU8_WMMA_w32, "_w32">;1922defm V_WMMA_I32_16X16X16_IU4_w32 : WMMAInstGFX12<"v_wmma_i32_16x16x16_iu4", I32_IU4X16_WMMA_w32, "_w32">;1923defm V_WMMA_F32_16X16X16_FP8_FP8_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x16_fp8_fp8", F32_FP8BF8_WMMA_w32, "_w32">;1924defm V_WMMA_F32_16X16X16_FP8_BF8_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x16_fp8_bf8", F32_FP8BF8_WMMA_w32, "_w32">;1925defm V_WMMA_F32_16X16X16_BF8_FP8_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x16_bf8_fp8", F32_FP8BF8_WMMA_w32, "_w32">;1926defm V_WMMA_F32_16X16X16_BF8_BF8_w32 : WMMAInstGFX12<"v_wmma_f32_16x16x16_bf8_bf8", F32_FP8BF8_WMMA_w32, "_w32">;1927defm V_WMMA_I32_16X16X32_IU4_w32 : WMMAInstGFX12<"v_wmma_i32_16x16x32_iu4", I32_IU4X32_WMMA_w32, "_w32">;1928 1929defm V_SWMMAC_F32_16X16X32_F16_w32 : SWMMACInstGFX12<"v_swmmac_f32_16x16x32_f16", F32_F16_SWMMAC_w32, "_w32">;1930defm V_SWMMAC_F32_16X16X32_BF16_w32 : SWMMACInstGFX12<"v_swmmac_f32_16x16x32_bf16", F32_BF16_SWMMAC_w32, "_w32">;1931defm V_SWMMAC_F16_16X16X32_F16_w32 : SWMMACInstGFX12<"v_swmmac_f16_16x16x32_f16", F16_F16_SWMMAC_w32, "_w32">;1932defm V_SWMMAC_BF16_16X16X32_BF16_w32 : SWMMACInstGFX12<"v_swmmac_bf16_16x16x32_bf16", BF16_BF16_SWMMAC_w32, "_w32">;1933defm V_SWMMAC_I32_16X16X32_IU8_w32 : SWMMACInstGFX12<"v_swmmac_i32_16x16x32_iu8", I32_IU8_SWMMAC_w32, "_w32">;1934defm V_SWMMAC_I32_16X16X32_IU4_w32 : SWMMACInstGFX12<"v_swmmac_i32_16x16x32_iu4", I32_IU4X32_SWMMAC_w32, "_w32">;1935defm V_SWMMAC_I32_16X16X64_IU4_w32 : SWMMACInstGFX12<"v_swmmac_i32_16x16x64_iu4", I32_IU4X64_SWMMAC_w32, "_w32">;1936defm V_SWMMAC_F32_16X16X32_FP8_FP8_w32 : SWMMACInstGFX12<"v_swmmac_f32_16x16x32_fp8_fp8", F32_FP8BF8_SWMMAC_w32, "_w32">;1937defm V_SWMMAC_F32_16X16X32_FP8_BF8_w32 : SWMMACInstGFX12<"v_swmmac_f32_16x16x32_fp8_bf8", F32_FP8BF8_SWMMAC_w32, "_w32">;1938defm V_SWMMAC_F32_16X16X32_BF8_FP8_w32 : SWMMACInstGFX12<"v_swmmac_f32_16x16x32_bf8_fp8", F32_FP8BF8_SWMMAC_w32, "_w32">;1939defm V_SWMMAC_F32_16X16X32_BF8_BF8_w32 : SWMMACInstGFX12<"v_swmmac_f32_16x16x32_bf8_bf8", F32_FP8BF8_SWMMAC_w32, "_w32">;1940}1941 1942let WaveSizePredicate = isWave64 in {1943defm V_WMMA_F32_16X16X16_F16_w64 : WMMAInstGFX12<"v_wmma_f32_16x16x16_f16", F32_F16_WMMA_w64, "_w64">;1944defm V_WMMA_F32_16X16X16_BF16_w64 : WMMAInstGFX12<"v_wmma_f32_16x16x16_bf16", F32_BF16_WMMA_w64, "_w64">;1945defm V_WMMA_F16_16X16X16_F16_w64 : WMMAInstGFX12<"v_wmma_f16_16x16x16_f16", F16_F16_WMMA_w64, "_w64">;1946defm V_WMMA_BF16_16X16X16_BF16_w64 : WMMAInstGFX12<"v_wmma_bf16_16x16x16_bf16", BF16_BF16_WMMA_w64, "_w64">;1947defm V_WMMA_I32_16X16X16_IU8_w64 : WMMAInstGFX12<"v_wmma_i32_16x16x16_iu8", I32_IU8_WMMA_w64, "_w64">;1948defm V_WMMA_I32_16X16X16_IU4_w64 : WMMAInstGFX12<"v_wmma_i32_16x16x16_iu4", I32_IU4X16_WMMA_w64, "_w64">;1949defm V_WMMA_F32_16X16X16_FP8_FP8_w64 : WMMAInstGFX12<"v_wmma_f32_16x16x16_fp8_fp8", F32_FP8BF8_WMMA_w64, "_w64">;1950defm V_WMMA_F32_16X16X16_FP8_BF8_w64 : WMMAInstGFX12<"v_wmma_f32_16x16x16_fp8_bf8", F32_FP8BF8_WMMA_w64, "_w64">;1951defm V_WMMA_F32_16X16X16_BF8_FP8_w64 : WMMAInstGFX12<"v_wmma_f32_16x16x16_bf8_fp8", F32_FP8BF8_WMMA_w64, "_w64">;1952defm V_WMMA_F32_16X16X16_BF8_BF8_w64 : WMMAInstGFX12<"v_wmma_f32_16x16x16_bf8_bf8", F32_FP8BF8_WMMA_w64, "_w64">;1953defm V_WMMA_I32_16X16X32_IU4_w64 : WMMAInstGFX12<"v_wmma_i32_16x16x32_iu4", I32_IU4X32_WMMA_w64, "_w64">;1954 1955defm V_SWMMAC_F32_16X16X32_F16_w64 : SWMMACInstGFX12<"v_swmmac_f32_16x16x32_f16", F32_F16_SWMMAC_w64, "_w64">;1956defm V_SWMMAC_F32_16X16X32_BF16_w64 : SWMMACInstGFX12<"v_swmmac_f32_16x16x32_bf16", F32_BF16_SWMMAC_w64, "_w64">;1957defm V_SWMMAC_F16_16X16X32_F16_w64 : SWMMACInstGFX12<"v_swmmac_f16_16x16x32_f16", F16_F16_SWMMAC_w64, "_w64">;1958defm V_SWMMAC_BF16_16X16X32_BF16_w64 : SWMMACInstGFX12<"v_swmmac_bf16_16x16x32_bf16", BF16_BF16_SWMMAC_w64, "_w64">;1959defm V_SWMMAC_I32_16X16X32_IU8_w64 : SWMMACInstGFX12<"v_swmmac_i32_16x16x32_iu8", I32_IU8_SWMMAC_w64, "_w64">;1960defm V_SWMMAC_I32_16X16X32_IU4_w64 : SWMMACInstGFX12<"v_swmmac_i32_16x16x32_iu4", I32_IU4X32_SWMMAC_w64, "_w64">;1961defm V_SWMMAC_I32_16X16X64_IU4_w64 : SWMMACInstGFX12<"v_swmmac_i32_16x16x64_iu4", I32_IU4X64_SWMMAC_w64, "_w64">;1962defm V_SWMMAC_F32_16X16X32_FP8_FP8_w64 : SWMMACInstGFX12<"v_swmmac_f32_16x16x32_fp8_fp8", F32_FP8BF8_SWMMAC_w64, "_w64">;1963defm V_SWMMAC_F32_16X16X32_FP8_BF8_w64 : SWMMACInstGFX12<"v_swmmac_f32_16x16x32_fp8_bf8", F32_FP8BF8_SWMMAC_w64, "_w64">;1964defm V_SWMMAC_F32_16X16X32_BF8_FP8_w64 : SWMMACInstGFX12<"v_swmmac_f32_16x16x32_bf8_fp8", F32_FP8BF8_SWMMAC_w64, "_w64">;1965defm V_SWMMAC_F32_16X16X32_BF8_BF8_w64 : SWMMACInstGFX12<"v_swmmac_f32_16x16x32_bf8_bf8", F32_FP8BF8_SWMMAC_w64, "_w64">;1966}1967 1968// IsGFX11OpselIntrinsic: f16_f16 and bf16_bf16 Intrinsics have imm operand that1969// controls opsel. Used by gfx11, removed in gfx12 (operand must be 0).1970multiclass WMMAPat<string Inst, SDPatternOperator node, VOP3PWMMA_Profile P, bit IsGFX11OpselIntrinsic = 0> {1971 def : GCNPat <(P.DstVT !setdagop(!con(P.WmmaInPat, !if(IsGFX11OpselIntrinsic, (ins 0), (ins))), node)),1972 (P.DstVT !setdagop(P.WmmaOutPat, !cast<Instruction>(Inst#"_twoaddr")))>;1973 let AddedComplexity = 4 in1974 def : GCNPat <(P.DstVT !setdagop(!con(P.WmmaInlineInPat, !if(IsGFX11OpselIntrinsic, (ins 0), (ins))), node)),1975 (P.DstVT !setdagop(P.WmmaInlineOutPat, !cast<Instruction>(Inst#"_threeaddr")))>;1976}1977 1978class SWMMACPat<Instruction Inst, SDPatternOperator node, VOP3PWMMA_Profile P> :1979 GCNPat <(P.DstVT !setdagop(P.SwmmacInPat, node)),1980 (P.DstVT !setdagop(P.SwmmacOutPat, Inst))>;1981 1982class SWMMACPat_w64<Instruction Inst, SDPatternOperator node, VOP3PWMMA_Profile P> :1983 GCNPat <(P.DstVT !setdagop(P.SwmmacInPat, node)),1984 (P.DstVT !setdagop(P.SwmmacOutPat, Inst))>{1985 let WaveSizePredicate = isWave64;1986 }1987 1988let WaveSizePredicate = isWave32, SubtargetPredicate = isGFX12PlusNot12_50 in {1989 defm : WMMAPat<"V_WMMA_F32_16X16X16_F16_w32", int_amdgcn_wmma_f32_16x16x16_f16, F32_F16_WMMA_w32>;1990 defm : WMMAPat<"V_WMMA_F32_16X16X16_BF16_w32", int_amdgcn_wmma_f32_16x16x16_bf16, F32_BF16_WMMA_w32>;1991 defm : WMMAPat<"V_WMMA_F16_16X16X16_F16_w32", int_amdgcn_wmma_f16_16x16x16_f16, F16_F16_WMMA_w32,1>;1992 defm : WMMAPat<"V_WMMA_BF16_16X16X16_BF16_w32", int_amdgcn_wmma_bf16_16x16x16_bf16, BF16_BF16_WMMA_w32,1>;1993 defm : WMMAPat<"V_WMMA_I32_16X16X16_IU8_w32", int_amdgcn_wmma_i32_16x16x16_iu8, I32_IU8_WMMA_w32>;1994 defm : WMMAPat<"V_WMMA_I32_16X16X16_IU4_w32", int_amdgcn_wmma_i32_16x16x16_iu4, I32_IU4X16_WMMA_w32>;1995 defm : WMMAPat<"V_WMMA_F32_16X16X16_FP8_FP8_w32", int_amdgcn_wmma_f32_16x16x16_fp8_fp8, F32_FP8BF8_WMMA_w32>;1996 defm : WMMAPat<"V_WMMA_F32_16X16X16_FP8_BF8_w32", int_amdgcn_wmma_f32_16x16x16_fp8_bf8, F32_FP8BF8_WMMA_w32>;1997 defm : WMMAPat<"V_WMMA_F32_16X16X16_BF8_FP8_w32", int_amdgcn_wmma_f32_16x16x16_bf8_fp8, F32_FP8BF8_WMMA_w32>;1998 defm : WMMAPat<"V_WMMA_F32_16X16X16_BF8_BF8_w32", int_amdgcn_wmma_f32_16x16x16_bf8_bf8, F32_FP8BF8_WMMA_w32>;1999 defm : WMMAPat<"V_WMMA_I32_16X16X32_IU4_w32", int_amdgcn_wmma_i32_16x16x32_iu4, I32_IU4X32_WMMA_w32>;2000 2001 def : SWMMACPat<V_SWMMAC_F32_16X16X32_F16_w32_twoaddr, int_amdgcn_swmmac_f32_16x16x32_f16, F32_F16_SWMMAC_w32>;2002 def : SWMMACPat<V_SWMMAC_F32_16X16X32_BF16_w32_twoaddr, int_amdgcn_swmmac_f32_16x16x32_bf16, F32_BF16_SWMMAC_w32>;2003 def : SWMMACPat<V_SWMMAC_F16_16X16X32_F16_w32_twoaddr, int_amdgcn_swmmac_f16_16x16x32_f16, F16_F16_SWMMAC_w32>;2004 def : SWMMACPat<V_SWMMAC_BF16_16X16X32_BF16_w32_twoaddr, int_amdgcn_swmmac_bf16_16x16x32_bf16, BF16_BF16_SWMMAC_w32>;2005 def : SWMMACPat<V_SWMMAC_I32_16X16X32_IU8_w32_twoaddr, int_amdgcn_swmmac_i32_16x16x32_iu8, I32_IU8_SWMMAC_w32>;2006 def : SWMMACPat<V_SWMMAC_I32_16X16X32_IU4_w32_twoaddr, int_amdgcn_swmmac_i32_16x16x32_iu4, I32_IU4X32_SWMMAC_w32>;2007 def : GCNPat <(I32_IU4X64_SWMMAC_w32.DstVT !setdagop(I32_IU4X64_SWMMAC_w32.SwmmacInPat, int_amdgcn_swmmac_i32_16x16x64_iu4)),2008 (I32_IU4X64_SWMMAC_w32.DstVT !setdagop(I32_IU4X64_SWMMAC_w32.SwmmacOutPat, V_SWMMAC_I32_16X16X64_IU4_w32_twoaddr))>;2009 def : SWMMACPat<V_SWMMAC_F32_16X16X32_FP8_FP8_w32_twoaddr, int_amdgcn_swmmac_f32_16x16x32_fp8_fp8, F32_FP8BF8_SWMMAC_w32>;2010 def : SWMMACPat<V_SWMMAC_F32_16X16X32_FP8_BF8_w32_twoaddr, int_amdgcn_swmmac_f32_16x16x32_fp8_bf8, F32_FP8BF8_SWMMAC_w32>;2011 def : SWMMACPat<V_SWMMAC_F32_16X16X32_BF8_FP8_w32_twoaddr, int_amdgcn_swmmac_f32_16x16x32_bf8_fp8, F32_FP8BF8_SWMMAC_w32>;2012 def : SWMMACPat<V_SWMMAC_F32_16X16X32_BF8_BF8_w32_twoaddr, int_amdgcn_swmmac_f32_16x16x32_bf8_bf8, F32_FP8BF8_SWMMAC_w32>;2013}2014 2015let WaveSizePredicate = isWave64, SubtargetPredicate = isGFX12PlusNot12_50 in {2016 defm : WMMAPat<"V_WMMA_F32_16X16X16_F16_w64", int_amdgcn_wmma_f32_16x16x16_f16, F32_F16_WMMA_w64>;2017 defm : WMMAPat<"V_WMMA_F32_16X16X16_BF16_w64", int_amdgcn_wmma_f32_16x16x16_bf16, F32_BF16_WMMA_w64>;2018 defm : WMMAPat<"V_WMMA_F16_16X16X16_F16_w64", int_amdgcn_wmma_f16_16x16x16_f16, F16_F16_WMMA_w64,1>;2019 defm : WMMAPat<"V_WMMA_BF16_16X16X16_BF16_w64", int_amdgcn_wmma_bf16_16x16x16_bf16, BF16_BF16_WMMA_w64,1>;2020 defm : WMMAPat<"V_WMMA_I32_16X16X16_IU8_w64", int_amdgcn_wmma_i32_16x16x16_iu8, I32_IU8_WMMA_w64>;2021 defm : WMMAPat<"V_WMMA_I32_16X16X16_IU4_w64", int_amdgcn_wmma_i32_16x16x16_iu4, I32_IU4X16_WMMA_w64>;2022 defm : WMMAPat<"V_WMMA_F32_16X16X16_FP8_FP8_w64", int_amdgcn_wmma_f32_16x16x16_fp8_fp8, F32_FP8BF8_WMMA_w64>;2023 defm : WMMAPat<"V_WMMA_F32_16X16X16_FP8_BF8_w64", int_amdgcn_wmma_f32_16x16x16_fp8_bf8, F32_FP8BF8_WMMA_w64>;2024 defm : WMMAPat<"V_WMMA_F32_16X16X16_BF8_FP8_w64", int_amdgcn_wmma_f32_16x16x16_bf8_fp8, F32_FP8BF8_WMMA_w64>;2025 defm : WMMAPat<"V_WMMA_F32_16X16X16_BF8_BF8_w64", int_amdgcn_wmma_f32_16x16x16_bf8_bf8, F32_FP8BF8_WMMA_w64>;2026 defm : WMMAPat<"V_WMMA_I32_16X16X32_IU4_w64", int_amdgcn_wmma_i32_16x16x32_iu4, I32_IU4X32_WMMA_w64>;2027 2028 def : SWMMACPat<V_SWMMAC_F32_16X16X32_F16_w64_twoaddr, int_amdgcn_swmmac_f32_16x16x32_f16, F32_F16_SWMMAC_w64>;2029 def : SWMMACPat<V_SWMMAC_F32_16X16X32_BF16_w64_twoaddr, int_amdgcn_swmmac_f32_16x16x32_bf16, F32_BF16_SWMMAC_w64>;2030 def : SWMMACPat<V_SWMMAC_F16_16X16X32_F16_w64_twoaddr, int_amdgcn_swmmac_f16_16x16x32_f16, F16_F16_SWMMAC_w64>;2031 def : SWMMACPat<V_SWMMAC_BF16_16X16X32_BF16_w64_twoaddr, int_amdgcn_swmmac_bf16_16x16x32_bf16, BF16_BF16_SWMMAC_w64>;2032 def : SWMMACPat<V_SWMMAC_I32_16X16X32_IU8_w64_twoaddr, int_amdgcn_swmmac_i32_16x16x32_iu8, I32_IU8_SWMMAC_w64>;2033 def : SWMMACPat<V_SWMMAC_I32_16X16X32_IU4_w64_twoaddr, int_amdgcn_swmmac_i32_16x16x32_iu4, I32_IU4X32_SWMMAC_w64>;2034 def : SWMMACPat<V_SWMMAC_I32_16X16X64_IU4_w64_twoaddr, int_amdgcn_swmmac_i32_16x16x64_iu4, I32_IU4X64_SWMMAC_w64>;2035 def : SWMMACPat<V_SWMMAC_F32_16X16X32_FP8_FP8_w64_twoaddr, int_amdgcn_swmmac_f32_16x16x32_fp8_fp8, F32_FP8BF8_SWMMAC_w64>;2036 def : SWMMACPat<V_SWMMAC_F32_16X16X32_FP8_BF8_w64_twoaddr, int_amdgcn_swmmac_f32_16x16x32_fp8_bf8, F32_FP8BF8_SWMMAC_w64>;2037 def : SWMMACPat<V_SWMMAC_F32_16X16X32_BF8_FP8_w64_twoaddr, int_amdgcn_swmmac_f32_16x16x32_bf8_fp8, F32_FP8BF8_SWMMAC_w64>;2038 def : SWMMACPat<V_SWMMAC_F32_16X16X32_BF8_BF8_w64_twoaddr, int_amdgcn_swmmac_f32_16x16x32_bf8_bf8, F32_FP8BF8_SWMMAC_w64>;2039}2040 2041let WaveSizePredicate = isWave32 in {2042let SubtargetPredicate = isGFX125xOnly in {2043 defm : WMMAPat<"V_WMMA_F32_16X16X4_F32_w32", int_amdgcn_wmma_f32_16x16x4_f32, F32_F32_WMMA_w32>;2044 defm : WMMAPat<"V_WMMA_F32_16X16X32_BF16_w32", int_amdgcn_wmma_f32_16x16x32_bf16, F32_BF16X32_WMMA_w32>;2045 defm : WMMAPat<"V_WMMA_BF16_16X16X32_BF16_w32", int_amdgcn_wmma_bf16_16x16x32_bf16, BF16_BF16X32_WMMA_w32>;2046 defm : WMMAPat<"V_WMMA_BF16F32_16X16X32_BF16_w32", int_amdgcn_wmma_bf16f32_16x16x32_bf16, BF16F32_BF16_WMMA_w32>;2047 defm : WMMAPat<"V_WMMA_F32_16X16X64_FP8_FP8_w32", int_amdgcn_wmma_f32_16x16x64_fp8_fp8, F32_FP8BF8X64_WMMA_w32>;2048 defm : WMMAPat<"V_WMMA_F32_16X16X64_FP8_BF8_w32", int_amdgcn_wmma_f32_16x16x64_fp8_bf8, F32_FP8BF8X64_WMMA_w32>;2049 defm : WMMAPat<"V_WMMA_F32_16X16X64_BF8_FP8_w32", int_amdgcn_wmma_f32_16x16x64_bf8_fp8, F32_FP8BF8X64_WMMA_w32>;2050 defm : WMMAPat<"V_WMMA_F32_16X16X64_BF8_BF8_w32", int_amdgcn_wmma_f32_16x16x64_bf8_bf8, F32_FP8BF8X64_WMMA_w32>;2051 defm : WMMAPat<"V_WMMA_F16_16X16X64_FP8_FP8_w32", int_amdgcn_wmma_f16_16x16x64_fp8_fp8, F16_FP8BF8X64_WMMA_w32>;2052 defm : WMMAPat<"V_WMMA_F16_16X16X64_FP8_BF8_w32", int_amdgcn_wmma_f16_16x16x64_fp8_bf8, F16_FP8BF8X64_WMMA_w32>;2053 defm : WMMAPat<"V_WMMA_F16_16X16X64_BF8_FP8_w32", int_amdgcn_wmma_f16_16x16x64_bf8_fp8, F16_FP8BF8X64_WMMA_w32>;2054 defm : WMMAPat<"V_WMMA_F16_16X16X64_BF8_BF8_w32", int_amdgcn_wmma_f16_16x16x64_bf8_bf8, F16_FP8BF8X64_WMMA_w32>;2055 defm : WMMAPat<"V_WMMA_I32_16X16X64_IU8_w32", int_amdgcn_wmma_i32_16x16x64_iu8, I32_IU8X64_WMMA_w32>;2056 defm : WMMAPat<"V_WMMA_F32_16X16X32_F16_w32", int_amdgcn_wmma_f32_16x16x32_f16, F32_F16X32_WMMA_w32>;2057 defm : WMMAPat<"V_WMMA_F16_16X16X32_F16_w32", int_amdgcn_wmma_f16_16x16x32_f16, F16_F16X32_WMMA_w32>;2058 defm : WMMAPat<"V_WMMA_F16_16X16X128_FP8_FP8_w32", int_amdgcn_wmma_f16_16x16x128_fp8_fp8, F16_FP8BF8X128_WMMA_w32>;2059 defm : WMMAPat<"V_WMMA_F16_16X16X128_FP8_BF8_w32", int_amdgcn_wmma_f16_16x16x128_fp8_bf8, F16_FP8BF8X128_WMMA_w32>;2060 defm : WMMAPat<"V_WMMA_F16_16X16X128_BF8_FP8_w32", int_amdgcn_wmma_f16_16x16x128_bf8_fp8, F16_FP8BF8X128_WMMA_w32>;2061 defm : WMMAPat<"V_WMMA_F16_16X16X128_BF8_BF8_w32", int_amdgcn_wmma_f16_16x16x128_bf8_bf8, F16_FP8BF8X128_WMMA_w32>;2062 defm : WMMAPat<"V_WMMA_F32_16X16X128_FP8_FP8_w32", int_amdgcn_wmma_f32_16x16x128_fp8_fp8, F32_FP8BF8X128_WMMA_w32>;2063 defm : WMMAPat<"V_WMMA_F32_16X16X128_FP8_BF8_w32", int_amdgcn_wmma_f32_16x16x128_fp8_bf8, F32_FP8BF8X128_WMMA_w32>;2064 defm : WMMAPat<"V_WMMA_F32_16X16X128_BF8_FP8_w32", int_amdgcn_wmma_f32_16x16x128_bf8_fp8, F32_FP8BF8X128_WMMA_w32>;2065 defm : WMMAPat<"V_WMMA_F32_16X16X128_BF8_BF8_w32", int_amdgcn_wmma_f32_16x16x128_bf8_bf8, F32_FP8BF8X128_WMMA_w32>;2066 defm : WMMAPat<"V_WMMA_F32_32X16X128_F4_w32", int_amdgcn_wmma_f32_32x16x128_f4, F32_32X16X128_F4_WMMA_w32>;2067 defm : WMMAPat<"V_WMMA_SCALE_F32_32X16X128_F4_w32", int_amdgcn_wmma_scale_f32_32x16x128_f4, F32_32X16X128_F4_SCALE_w32>;2068 defm : WMMAPat<"V_WMMA_SCALE16_F32_32X16X128_F4_w32", int_amdgcn_wmma_scale16_f32_32x16x128_f4, F32_32X16X128_F4_SCALE16_w32>;2069 2070 foreach I = ["f8_f8", "f8_f6", "f8_f4", "f6_f8", "f6_f6", "f6_f4", "f4_f8", "f4_f6", "f4_f4"] in {2071 defm : WMMAPat<"V_WMMA_F32_16X16X128_F8F6F4_" # I # "_w32", int_amdgcn_wmma_f32_16x16x128_f8f6f4, !cast<VOP3PWMMA_Profile>("F32_16X16X128_F8F6F4_" # I # "_w32")>;2072 defm : WMMAPat<"V_WMMA_SCALE_F32_16X16X128_F8F6F4_" # I # "_w32", int_amdgcn_wmma_scale_f32_16x16x128_f8f6f4, !cast<VOP3PWMMA_Profile>("F32_16X16X128_F8F6F4_SCALE_" # I # "_w32")>;2073 defm : WMMAPat<"V_WMMA_SCALE16_F32_16X16X128_F8F6F4_" # I # "_w32", int_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4, !cast<VOP3PWMMA_Profile>("F32_16X16X128_F8F6F4_SCALE16_" # I # "_w32")>;2074 }2075 2076 def : SWMMACPat<V_SWMMAC_F32_16X16X64_BF16_w32_twoaddr, int_amdgcn_swmmac_f32_16x16x64_bf16, F32_BF16X64_SWMMAC_w32>;2077 def : SWMMACPat<V_SWMMAC_BF16_16X16X64_BF16_w32_twoaddr, int_amdgcn_swmmac_bf16_16x16x64_bf16, BF16_BF16X64_SWMMAC_w32>;2078 def : SWMMACPat<V_SWMMAC_BF16F32_16X16X64_BF16_w32_twoaddr, int_amdgcn_swmmac_bf16f32_16x16x64_bf16, F32_BF16X64_SWMMAC_w32>;2079 def : SWMMACPat<V_SWMMAC_F32_16X16X128_FP8_FP8_w32_twoaddr, int_amdgcn_swmmac_f32_16x16x128_fp8_fp8, F32_FP8BF8X128_SWMMAC_w32>;2080 def : SWMMACPat<V_SWMMAC_F32_16X16X128_FP8_BF8_w32_twoaddr, int_amdgcn_swmmac_f32_16x16x128_fp8_bf8, F32_FP8BF8X128_SWMMAC_w32>;2081 def : SWMMACPat<V_SWMMAC_F32_16X16X128_BF8_FP8_w32_twoaddr, int_amdgcn_swmmac_f32_16x16x128_bf8_fp8, F32_FP8BF8X128_SWMMAC_w32>;2082 def : SWMMACPat<V_SWMMAC_F32_16X16X128_BF8_BF8_w32_twoaddr, int_amdgcn_swmmac_f32_16x16x128_bf8_bf8, F32_FP8BF8X128_SWMMAC_w32>;2083 def : SWMMACPat<V_SWMMAC_F16_16X16X128_FP8_FP8_w32_twoaddr, int_amdgcn_swmmac_f16_16x16x128_fp8_fp8, F16_FP8BF8X128_SWMMAC_w32>;2084 def : SWMMACPat<V_SWMMAC_F16_16X16X128_FP8_BF8_w32_twoaddr, int_amdgcn_swmmac_f16_16x16x128_fp8_bf8, F16_FP8BF8X128_SWMMAC_w32>;2085 def : SWMMACPat<V_SWMMAC_F16_16X16X128_BF8_FP8_w32_twoaddr, int_amdgcn_swmmac_f16_16x16x128_bf8_fp8, F16_FP8BF8X128_SWMMAC_w32>;2086 def : SWMMACPat<V_SWMMAC_F16_16X16X128_BF8_BF8_w32_twoaddr, int_amdgcn_swmmac_f16_16x16x128_bf8_bf8, F16_FP8BF8X128_SWMMAC_w32>;2087 def : SWMMACPat<V_SWMMAC_I32_16X16X128_IU8_w32_twoaddr, int_amdgcn_swmmac_i32_16x16x128_iu8, I32_IU8X128_SWMMAC_w32>;2088 def : SWMMACPat<V_SWMMAC_F32_16X16X64_F16_w32_twoaddr, int_amdgcn_swmmac_f32_16x16x64_f16, F32_F16X64_SWMMAC_w32>;2089 def : SWMMACPat<V_SWMMAC_F16_16X16X64_F16_w32_twoaddr, int_amdgcn_swmmac_f16_16x16x64_f16, F16_F16X64_SWMMAC_w32>;2090} // End SubtargetPredicate = isGFX125xOnly2091} // End WaveSizePredicate = isWave322092 2093//===----------------------------------------------------------------------===//2094// Begin Real Encodings2095//===----------------------------------------------------------------------===//2096 2097class VOP3P_DPP16<bits<8> op, VOP_DPP_Pseudo ps, int subtarget,2098 string opName = ps.OpName>2099 : VOP3P_DPP<op, opName, ps.Pfl, 1>, SIMCInstr<ps.PseudoInstr, subtarget> {2100 let hasSideEffects = ps.hasSideEffects;2101 let Defs = ps.Defs;2102 let SchedRW = ps.SchedRW;2103 let Uses = ps.Uses;2104 let AssemblerPredicate = HasDPP16;2105 let SubtargetPredicate = ps.SubtargetPredicate;2106 let OtherPredicates = ps.OtherPredicates;2107 let IsPacked = ps.IsPacked;2108}2109 2110class VOP3P_DPP8_Base<bits<8> op, VOP_Pseudo ps, string opName = ps.OpName>2111 : VOP3P_DPP8<op, opName, ps.Pfl> {2112 let hasSideEffects = ps.hasSideEffects;2113 let Defs = ps.Defs;2114 let SchedRW = ps.SchedRW;2115 let Uses = ps.Uses;2116 let SubtargetPredicate = ps.SubtargetPredicate;2117 let OtherPredicates = ps.OtherPredicates;2118 let IsPacked = ps.IsPacked;2119}2120 2121//===----------------------------------------------------------------------===//2122// GFX11, GFX122123//===----------------------------------------------------------------------===//2124 2125multiclass VOP3P_Real_Base<GFXGen Gen, bits<8> op, string backing_ps_name = NAME,2126 string asmName = !cast<VOP3P_Pseudo>(NAME).Mnemonic> {2127 def Gen.Suffix :2128 VOP3P_Real_Gen<!cast<VOP3P_Pseudo>(backing_ps_name), Gen, asmName>,2129 VOP3Pe_gfx11_gfx12<op, !cast<VOP3P_Pseudo>(backing_ps_name).Pfl>;2130}2131 2132class VOP3PeWmma<bits<8> op, VOPProfile P, VOP3PWMMA_Profile WMMAP>2133 : VOP3Pe_gfx11_gfx12<op, P>{2134 2135 // opsel2136 let Inst{11} = !cond(WMMAP.HasMatrixFMT : matrix_a_fmt{0},2137 !eq(WMMAP.IndexType, 0) : 0,2138 !eq(WMMAP.IndexType, 8) : index_key_8bit{0},2139 !eq(WMMAP.IndexType, 16) : index_key_16bit{0},2140 !eq(WMMAP.IndexType, 32) : index_key_32bit{0});2141 let Inst{12} = !if(WMMAP.HasMatrixFMT, matrix_a_fmt{1},2142 !if(!eq(WMMAP.IndexType, 8), index_key_8bit{1}, 0));2143 let Inst{13} = !if (WMMAP.HasMatrixFMT, matrix_a_fmt{2},2144 !if(WMMAP.HasMatrixReuse, matrix_a_reuse, 0));2145 // opsel_hi2146 let Inst{59} = !if (WMMAP.HasMatrixFMT, matrix_b_fmt{0}, 1);2147 let Inst{60} = !if (WMMAP.HasMatrixFMT, matrix_b_fmt{1}, 1);2148 let Inst{14} = !if (WMMAP.HasMatrixFMT, matrix_b_fmt{2},2149 !if(WMMAP.HasMatrixReuse, matrix_b_reuse, 1));2150 // neg_lo2151 let Inst{61} = !if(WMMAP.NegLo01, src0_modifiers{0}, 0);2152 let Inst{62} = !if(WMMAP.NegLo01, src1_modifiers{0}, 0);2153 let Inst{63} = !if(WMMAP.NegLo2, src2_modifiers{0}, 0);2154 // neg_hi2155 let Inst{8} = !if(WMMAP.NegHi01, src0_modifiers{1}, 0);2156 let Inst{9} = !if(WMMAP.NegHi01, src1_modifiers{1}, 0);2157 let Inst{10} = !if(WMMAP.NegHi2, src2_modifiers{1}, 0);2158 // clamp2159 let Inst{15} = !if(WMMAP.HasClamp, clamp{0}, 0);2160}2161 2162multiclass VOP3P_WMMA_Real_Base<GFXGen Gen, bits<8> op, VOP3PWMMA_Profile WMMAP,2163 string backing_ps_name = NAME,2164 string asmName = !cast<VOP3P_Pseudo>(NAME).Mnemonic> {2165 def Gen.Suffix :2166 VOP3P_Real_Gen<!cast<VOP3P_Pseudo>(backing_ps_name), Gen, asmName>,2167 VOP3PeWmma<op, !cast<VOP3P_Pseudo>(backing_ps_name).Pfl, WMMAP>;2168}2169 2170multiclass VOP3P_Real_WMMA_gfx12 <bits<8> op, VOP3PWMMA_Profile WMMAP> {2171 let WaveSizePredicate = isWave32, DecoderNamespace = "GFX12" in {2172 defm _twoaddr : VOP3P_WMMA_Real_Base <GFX12Gen, op, WMMAP>;2173 }2174}2175 2176multiclass VOP3P_Real_WMMA_gfx12w64 <bits<8> op, VOP3PWMMA_Profile WMMAP> {2177 let WaveSizePredicate = isWave64, DecoderNamespace = "GFX12W64" in {2178 defm _twoaddr : VOP3P_WMMA_Real_Base <GFX12Gen, op, WMMAP>;2179 }2180}2181 2182multiclass VOP3P_Real_WMMA_gfx1250 <bits<8> op, VOP3PWMMA_Profile WMMAP> {2183 let WaveSizePredicate = isWave32, DecoderNamespace = "GFX12" in {2184 defm _twoaddr : VOP3P_WMMA_Real_Base <GFX1250Gen, op, WMMAP>;2185 }2186}2187 2188multiclass VOP3P_Real_WMMA_F8F6F4_gfx1250<bits<8> op, VOP3PWMMA_Profile WMMAP> {2189 defvar PS = !cast<VOP3P_Pseudo>(NAME # "_twoaddr");2190 defvar asmName = !substr(PS.Mnemonic, 0, !sub(!size(PS.Mnemonic), !size("_f8_f8_w32")));2191 defvar psName = !substr(NAME, 0, !sub(!size(PS.Mnemonic), !size("_f8_f8_w32")));2192 let AsmString = asmName # PS.AsmOperands in2193 defm NAME : VOP3P_Real_WMMA_gfx1250<op, WMMAP>,2194 MFMA_F8F6F4_WithSizeTable_Helper<PS, psName # "_f8_f8_w32_twoaddr_gfx1250">;2195}2196 2197multiclass VOP3P_Real_WMMA_gfx1250_SrcFormats<bits<8> op, string WMMAP> {2198 defm _f8_f8_w32 : VOP3P_Real_WMMA_F8F6F4_gfx1250<op, !cast<VOP3PWMMA_Profile>(WMMAP # "_f8_f8_w32")>;2199 foreach I = ["f8_f6", "f8_f4", "f6_f8", "f6_f6", "f6_f4", "f4_f8", "f4_f6", "f4_f4"] in {2200 let isAsmParserOnly = true in { // Disable ambiguous disassembly.2201 defm _#I#_w32 : VOP3P_Real_WMMA_F8F6F4_gfx1250<op, !cast<VOP3PWMMA_Profile>(WMMAP # "_" # I # "_w32")>;2202 }2203 }2204}2205 2206class VOP3PX2e <bits<8> op, bits<8> LdScaleOp, VOP3PWMMA_Profile P> : Enc128, VOP3Pe_Base {2207 bits<9> scale_src0;2208 bits<9> scale_src1;2209 2210 // Inst{7-0} = unused2211 let Inst{10-8} = {0, matrix_b_scale_fmt{1-0}}; // neg_hi2212 let Inst{11} = matrix_a_scale{0}; // scale_op_sel(0)2213 let Inst{12} = 0; // scale_op_sel(1)2214 let Inst{13} = matrix_a_reuse; // scale_op_sel(2)2215 let Inst{14} = matrix_b_reuse; // scale_op_sel_hi(2)2216 let Inst{15} = 0; // scale_clamp2217 let Inst{31-24} = 0xcc; // Encoding2218 let Inst{23-16} = LdScaleOp;2219 let Inst{40-32} = scale_src0;2220 let Inst{49-41} = scale_src1;2221 let Inst{58-50} = 0x100; // scale src2 = vgpr0 (dummy)2222 let Inst{59} = matrix_b_scale{0}; // scale_op_sel_hi(0)2223 let Inst{60} = 0; // scale_op_sel_hi(1)2224 let Inst{63-61} = {0, matrix_a_scale_fmt{1-0}}; // neg (lo)2225 2226 // The high half of the encoding is the unscaled wmma op.2227 let Inst{71-64} = vdst;2228 2229 let Inst{72} = !if(P.NegHi01, src0_modifiers{1}, 0); // neg_hi src02230 let Inst{73} = !if(P.NegHi01, src1_modifiers{1}, 0); // neg_hi src12231 let Inst{74} = !if(P.NegHi2, src2_modifiers{1}, 0); // neg_hi src22232 2233 let Inst{77-75} = !if(P.HasMatrixFMT, matrix_a_fmt{2-0}, 0); // op_sel2234 2235 let Inst{78,124,123} = !if(P.HasMatrixFMT, matrix_b_fmt{2-0}, 7); // op_sel_hi2236 let Inst{79} = !if(P.HasClamp, clamp{0}, 0);2237 2238 let Inst{87-80} = op;2239 let Inst{95-88} = 0xcc; //encoding2240 let Inst{104-96} = !if(P.HasSrc0, src0, 0);2241 let Inst{113-105} = !if(P.HasSrc1, src1, 0);2242 let Inst{122-114} = !if(P.HasSrc2, src2, 0);2243 2244 // neg_lo2245 let Inst{125} = !if(P.NegLo01, src0_modifiers{0}, 0);2246 let Inst{126} = !if(P.NegLo01, src1_modifiers{0}, 0);2247 let Inst{127} = !if(P.NegLo2, src2_modifiers{0}, 0);2248}2249 2250multiclass VOP3PX2_Real_ScaledWMMA_F4<bits<8> op, bits<8> LdScaleOp, VOP3PWMMA_Profile WMMAP> {2251 defvar PS = !cast<VOP3P_Pseudo>(NAME # "_twoaddr");2252 let SubtargetPredicate = isGFX1250Plus, WaveSizePredicate = isWave32,2253 DecoderNamespace = "GFX1250" in {2254 def _gfx1250 : VOP3P_Real_Gen<PS, GFX1250Gen, PS.Mnemonic>,2255 VOP3PX2e <op, LdScaleOp, WMMAP>;2256 }2257}2258 2259multiclass VOP3PX2_Real_ScaledWMMA<bits<8> op, bits<8> LdScaleOp, VOP3PWMMA_Profile WMMAP> {2260 defvar PS = !cast<VOP3P_Pseudo>(NAME # "_twoaddr");2261 defvar asmName = !substr(PS.Mnemonic, 0, !sub(!size(PS.Mnemonic), !size("_f8_f8_w32")));2262 defvar psName = !substr(NAME, 0, !sub(!size(PS.Mnemonic), !size("_f8_f8_w32")));2263 let SubtargetPredicate = isGFX1250Plus, WaveSizePredicate = isWave32,2264 DecoderNamespace = "GFX1250" in {2265 def _gfx1250 : VOP3P_Real_Gen<PS, GFX1250Gen, asmName>,2266 VOP3PX2e <op, LdScaleOp, WMMAP>,2267 MFMA_F8F6F4_WithSizeTable_Helper<PS, psName # "_f8_f8_w32_gfx1250"> {2268 let AsmString = asmName # PS.AsmOperands;2269 }2270 }2271}2272 2273multiclass VOP3PX2_Real_ScaledWMMA_SrcFormats<bits<8> op, bits<8> LdScaleOp, string WMMAP> {2274 defm _f8_f8_w32 : VOP3PX2_Real_ScaledWMMA<op, LdScaleOp, !cast<VOP3PWMMA_Profile>(WMMAP # "_f8_f8_w32")>;2275 foreach I = ["f8_f6", "f8_f4", "f6_f8", "f6_f6", "f6_f4", "f4_f8", "f4_f6", "f4_f4"] in {2276 let isAsmParserOnly = true in { // Disable ambiguous disassembly.2277 defm _#I#_w32 : VOP3PX2_Real_ScaledWMMA<op, LdScaleOp, !cast<VOP3PWMMA_Profile>(WMMAP # "_" # I # "_w32")>;2278 }2279 }2280}2281 2282defm V_WMMA_F32_16X16X16_F16_w32 : VOP3P_Real_WMMA_gfx12 <0x040, F32_F16_WMMA_w32>;2283defm V_WMMA_F32_16X16X16_BF16_w32 : VOP3P_Real_WMMA_gfx12 <0x041, F32_BF16_WMMA_w32>;2284defm V_WMMA_F16_16X16X16_F16_w32 : VOP3P_Real_WMMA_gfx12 <0x042, F16_F16_WMMA_w32>;2285defm V_WMMA_BF16_16X16X16_BF16_w32 : VOP3P_Real_WMMA_gfx12 <0x043, BF16_BF16_WMMA_w32>;2286defm V_WMMA_I32_16X16X16_IU8_w32 : VOP3P_Real_WMMA_gfx12 <0x044, I32_IU8_WMMA_w32>;2287defm V_WMMA_I32_16X16X16_IU4_w32 : VOP3P_Real_WMMA_gfx12 <0x045, I32_IU4X16_WMMA_w32>;2288defm V_WMMA_F32_16X16X16_FP8_FP8_w32 : VOP3P_Real_WMMA_gfx12 <0x046, F32_FP8BF8_WMMA_w32>;2289defm V_WMMA_F32_16X16X16_FP8_BF8_w32 : VOP3P_Real_WMMA_gfx12 <0x047, F32_FP8BF8_WMMA_w32>;2290defm V_WMMA_F32_16X16X16_BF8_FP8_w32 : VOP3P_Real_WMMA_gfx12 <0x048, F32_FP8BF8_WMMA_w32>;2291defm V_WMMA_F32_16X16X16_BF8_BF8_w32 : VOP3P_Real_WMMA_gfx12 <0x049, F32_FP8BF8_WMMA_w32>;2292defm V_WMMA_I32_16X16X32_IU4_w32 : VOP3P_Real_WMMA_gfx12 <0x04a, I32_IU4X32_WMMA_w32>;2293 2294defm V_WMMA_F32_16X16X16_F16_w64 : VOP3P_Real_WMMA_gfx12w64 <0x040, F32_F16_WMMA_w64>;2295defm V_WMMA_F32_16X16X16_BF16_w64 : VOP3P_Real_WMMA_gfx12w64 <0x041, F32_BF16_WMMA_w64>;2296defm V_WMMA_F16_16X16X16_F16_w64 : VOP3P_Real_WMMA_gfx12w64 <0x042, F16_F16_WMMA_w64>;2297defm V_WMMA_BF16_16X16X16_BF16_w64 : VOP3P_Real_WMMA_gfx12w64 <0x043, BF16_BF16_WMMA_w64>;2298defm V_WMMA_I32_16X16X16_IU8_w64 : VOP3P_Real_WMMA_gfx12w64 <0x044, I32_IU8_WMMA_w64>;2299defm V_WMMA_I32_16X16X16_IU4_w64 : VOP3P_Real_WMMA_gfx12w64 <0x045, I32_IU4X16_WMMA_w64>;2300defm V_WMMA_F32_16X16X16_FP8_FP8_w64 : VOP3P_Real_WMMA_gfx12w64 <0x046, F32_FP8BF8_WMMA_w64>;2301defm V_WMMA_F32_16X16X16_FP8_BF8_w64 : VOP3P_Real_WMMA_gfx12w64 <0x047, F32_FP8BF8_WMMA_w64>;2302defm V_WMMA_F32_16X16X16_BF8_FP8_w64 : VOP3P_Real_WMMA_gfx12w64 <0x048, F32_FP8BF8_WMMA_w64>;2303defm V_WMMA_F32_16X16X16_BF8_BF8_w64 : VOP3P_Real_WMMA_gfx12w64 <0x049, F32_FP8BF8_WMMA_w64>;2304defm V_WMMA_I32_16X16X32_IU4_w64 : VOP3P_Real_WMMA_gfx12w64 <0x04a, I32_IU4X32_WMMA_w64>;2305 2306 2307defm V_SWMMAC_F32_16X16X32_F16_w32 : VOP3P_Real_WMMA_gfx12 <0x050, F32_F16_SWMMAC_w32>;2308defm V_SWMMAC_F32_16X16X32_BF16_w32 : VOP3P_Real_WMMA_gfx12 <0x051, F32_BF16_SWMMAC_w32>;2309defm V_SWMMAC_F16_16X16X32_F16_w32 : VOP3P_Real_WMMA_gfx12 <0x052, F16_F16_SWMMAC_w32>;2310defm V_SWMMAC_BF16_16X16X32_BF16_w32 : VOP3P_Real_WMMA_gfx12 <0x053, BF16_BF16_SWMMAC_w32>;2311defm V_SWMMAC_I32_16X16X32_IU8_w32 : VOP3P_Real_WMMA_gfx12 <0x054, I32_IU8_SWMMAC_w32>;2312defm V_SWMMAC_I32_16X16X32_IU4_w32 : VOP3P_Real_WMMA_gfx12 <0x055, I32_IU4X32_SWMMAC_w32>;2313defm V_SWMMAC_I32_16X16X64_IU4_w32 : VOP3P_Real_WMMA_gfx12 <0x056, I32_IU4X64_SWMMAC_w32>;2314defm V_SWMMAC_F32_16X16X32_FP8_FP8_w32 : VOP3P_Real_WMMA_gfx12 <0x057, F32_FP8BF8_SWMMAC_w32>;2315defm V_SWMMAC_F32_16X16X32_FP8_BF8_w32 : VOP3P_Real_WMMA_gfx12 <0x058, F32_FP8BF8_SWMMAC_w32>;2316defm V_SWMMAC_F32_16X16X32_BF8_FP8_w32 : VOP3P_Real_WMMA_gfx12 <0x059, F32_FP8BF8_SWMMAC_w32>;2317defm V_SWMMAC_F32_16X16X32_BF8_BF8_w32 : VOP3P_Real_WMMA_gfx12 <0x05a, F32_FP8BF8_SWMMAC_w32>;2318 2319defm V_SWMMAC_F32_16X16X32_F16_w64 : VOP3P_Real_WMMA_gfx12w64 <0x050, F32_F16_SWMMAC_w64>;2320defm V_SWMMAC_F32_16X16X32_BF16_w64 : VOP3P_Real_WMMA_gfx12w64 <0x051, F32_BF16_SWMMAC_w64>;2321defm V_SWMMAC_F16_16X16X32_F16_w64 : VOP3P_Real_WMMA_gfx12w64 <0x052, F16_F16_SWMMAC_w64>;2322defm V_SWMMAC_BF16_16X16X32_BF16_w64 : VOP3P_Real_WMMA_gfx12w64 <0x053, BF16_BF16_SWMMAC_w64>;2323defm V_SWMMAC_I32_16X16X32_IU8_w64 : VOP3P_Real_WMMA_gfx12w64 <0x054, I32_IU8_SWMMAC_w64>;2324defm V_SWMMAC_I32_16X16X32_IU4_w64 : VOP3P_Real_WMMA_gfx12w64 <0x055, I32_IU4X32_SWMMAC_w64>;2325defm V_SWMMAC_I32_16X16X64_IU4_w64 : VOP3P_Real_WMMA_gfx12w64 <0x056, I32_IU4X64_SWMMAC_w64>;2326defm V_SWMMAC_F32_16X16X32_FP8_FP8_w64 : VOP3P_Real_WMMA_gfx12w64 <0x057, F32_FP8BF8_SWMMAC_w64>;2327defm V_SWMMAC_F32_16X16X32_FP8_BF8_w64 : VOP3P_Real_WMMA_gfx12w64 <0x058, F32_FP8BF8_SWMMAC_w64>;2328defm V_SWMMAC_F32_16X16X32_BF8_FP8_w64 : VOP3P_Real_WMMA_gfx12w64 <0x059, F32_FP8BF8_SWMMAC_w64>;2329defm V_SWMMAC_F32_16X16X32_BF8_BF8_w64 : VOP3P_Real_WMMA_gfx12w64 <0x05a, F32_FP8BF8_SWMMAC_w64>;2330 2331defm V_WMMA_F32_16X16X4_F32_w32 : VOP3P_Real_WMMA_gfx1250 <0x05d, F32_F32_WMMA_w32>;2332defm V_WMMA_F32_16X16X32_BF16_w32 : VOP3P_Real_WMMA_gfx1250 <0x062, F32_BF16X32_WMMA_w32>;2333defm V_WMMA_F32_16X16X32_F16_w32 : VOP3P_Real_WMMA_gfx1250 <0x060, F32_F16X32_WMMA_w32>;2334defm V_WMMA_F16_16X16X32_F16_w32 : VOP3P_Real_WMMA_gfx1250 <0x061, F16_F16X32_WMMA_w32>;2335defm V_WMMA_BF16_16X16X32_BF16_w32 : VOP3P_Real_WMMA_gfx1250 <0x063, BF16_BF16X32_WMMA_w32>;2336defm V_WMMA_BF16F32_16X16X32_BF16_w32 : VOP3P_Real_WMMA_gfx1250 <0x064, BF16F32_BF16_WMMA_w32>;2337defm V_WMMA_F32_16X16X64_FP8_FP8_w32 : VOP3P_Real_WMMA_gfx1250 <0x06a, F32_FP8BF8X64_WMMA_w32>;2338defm V_WMMA_F32_16X16X64_FP8_BF8_w32 : VOP3P_Real_WMMA_gfx1250 <0x06b, F32_FP8BF8X64_WMMA_w32>;2339defm V_WMMA_F32_16X16X64_BF8_FP8_w32 : VOP3P_Real_WMMA_gfx1250 <0x06c, F32_FP8BF8X64_WMMA_w32>;2340defm V_WMMA_F32_16X16X64_BF8_BF8_w32 : VOP3P_Real_WMMA_gfx1250 <0x06d, F32_FP8BF8X64_WMMA_w32>;2341defm V_WMMA_F16_16X16X64_FP8_FP8_w32 : VOP3P_Real_WMMA_gfx1250 <0x06e, F16_FP8BF8X64_WMMA_w32>;2342defm V_WMMA_F16_16X16X64_FP8_BF8_w32 : VOP3P_Real_WMMA_gfx1250 <0x06f, F16_FP8BF8X64_WMMA_w32>;2343defm V_WMMA_F16_16X16X64_BF8_FP8_w32 : VOP3P_Real_WMMA_gfx1250 <0x070, F16_FP8BF8X64_WMMA_w32>;2344defm V_WMMA_F16_16X16X64_BF8_BF8_w32 : VOP3P_Real_WMMA_gfx1250 <0x071, F16_FP8BF8X64_WMMA_w32>;2345defm V_WMMA_I32_16X16X64_IU8_w32 : VOP3P_Real_WMMA_gfx1250 <0x072, I32_IU8X64_WMMA_w32>;2346defm V_WMMA_F32_16X16X128_FP8_FP8_w32 : VOP3P_Real_WMMA_gfx1250 <0x080, F32_FP8BF8X128_WMMA_w32>;2347defm V_WMMA_F32_16X16X128_FP8_BF8_w32 : VOP3P_Real_WMMA_gfx1250 <0x081, F32_FP8BF8X128_WMMA_w32>;2348defm V_WMMA_F32_16X16X128_BF8_FP8_w32 : VOP3P_Real_WMMA_gfx1250 <0x082, F32_FP8BF8X128_WMMA_w32>;2349defm V_WMMA_F32_16X16X128_BF8_BF8_w32 : VOP3P_Real_WMMA_gfx1250 <0x083, F32_FP8BF8X128_WMMA_w32>;2350defm V_WMMA_F16_16X16X128_FP8_FP8_w32 : VOP3P_Real_WMMA_gfx1250 <0x084, F16_FP8BF8X128_WMMA_w32>;2351defm V_WMMA_F16_16X16X128_FP8_BF8_w32 : VOP3P_Real_WMMA_gfx1250 <0x085, F16_FP8BF8X128_WMMA_w32>;2352defm V_WMMA_F16_16X16X128_BF8_FP8_w32 : VOP3P_Real_WMMA_gfx1250 <0x086, F16_FP8BF8X128_WMMA_w32>;2353defm V_WMMA_F16_16X16X128_BF8_BF8_w32 : VOP3P_Real_WMMA_gfx1250 <0x087, F16_FP8BF8X128_WMMA_w32>;2354defm V_WMMA_F32_32X16X128_F4_w32 : VOP3P_Real_WMMA_gfx1250 <0x088, F32_32X16X128_F4_WMMA_w32>;2355 2356defm V_WMMA_F32_16X16X128_F8F6F4 : VOP3P_Real_WMMA_gfx1250_SrcFormats<0x033, "F32_16X16X128_F8F6F4">;2357defm V_WMMA_SCALE_F32_16X16X128_F8F6F4 : VOP3PX2_Real_ScaledWMMA_SrcFormats<0x033, 0x35, "F32_16X16X128_F8F6F4_SCALE">;2358defm V_WMMA_SCALE16_F32_16X16X128_F8F6F4 : VOP3PX2_Real_ScaledWMMA_SrcFormats<0x033, 0x3a, "F32_16X16X128_F8F6F4_SCALE16">;2359 2360defm V_WMMA_SCALE_F32_32X16X128_F4_w32 : VOP3PX2_Real_ScaledWMMA_F4<0x088, 0x35, F32_32X16X128_F4_SCALE_w32>;2361defm V_WMMA_SCALE16_F32_32X16X128_F4_w32 : VOP3PX2_Real_ScaledWMMA_F4<0x088, 0x3a, F32_32X16X128_F4_SCALE16_w32>;2362 2363defm V_SWMMAC_F32_16X16X64_F16_w32 : VOP3P_Real_WMMA_gfx1250 <0x065, F32_F16X64_SWMMAC_w32>;2364defm V_SWMMAC_F32_16X16X64_BF16_w32 : VOP3P_Real_WMMA_gfx1250 <0x066, F32_BF16X64_SWMMAC_w32>;2365defm V_SWMMAC_F16_16X16X64_F16_w32 : VOP3P_Real_WMMA_gfx1250 <0x067, F16_F16X64_SWMMAC_w32>;2366defm V_SWMMAC_BF16_16X16X64_BF16_w32 : VOP3P_Real_WMMA_gfx1250 <0x068, BF16_BF16X64_SWMMAC_w32>;2367defm V_SWMMAC_BF16F32_16X16X64_BF16_w32 : VOP3P_Real_WMMA_gfx1250 <0x069, F32_BF16X64_SWMMAC_w32>;2368defm V_SWMMAC_F32_16X16X128_FP8_FP8_w32 : VOP3P_Real_WMMA_gfx1250 <0x073, F32_FP8BF8X128_SWMMAC_w32>;2369defm V_SWMMAC_F32_16X16X128_FP8_BF8_w32 : VOP3P_Real_WMMA_gfx1250 <0x074, F32_FP8BF8X128_SWMMAC_w32>;2370defm V_SWMMAC_F32_16X16X128_BF8_FP8_w32 : VOP3P_Real_WMMA_gfx1250 <0x075, F32_FP8BF8X128_SWMMAC_w32>;2371defm V_SWMMAC_F32_16X16X128_BF8_BF8_w32 : VOP3P_Real_WMMA_gfx1250 <0x076, F32_FP8BF8X128_SWMMAC_w32>;2372defm V_SWMMAC_F16_16X16X128_FP8_FP8_w32 : VOP3P_Real_WMMA_gfx1250 <0x077, F16_FP8BF8X128_SWMMAC_w32>;2373defm V_SWMMAC_F16_16X16X128_FP8_BF8_w32 : VOP3P_Real_WMMA_gfx1250 <0x078, F16_FP8BF8X128_SWMMAC_w32>;2374defm V_SWMMAC_F16_16X16X128_BF8_FP8_w32 : VOP3P_Real_WMMA_gfx1250 <0x079, F16_FP8BF8X128_SWMMAC_w32>;2375defm V_SWMMAC_F16_16X16X128_BF8_BF8_w32 : VOP3P_Real_WMMA_gfx1250 <0x07a, F16_FP8BF8X128_SWMMAC_w32>;2376defm V_SWMMAC_I32_16X16X128_IU8_w32 : VOP3P_Real_WMMA_gfx1250 <0x07b, I32_IU8X128_SWMMAC_w32>;2377 2378multiclass VOP3P_Real_with_name<GFXGen Gen, bits<8> op,2379 string backing_ps_name = NAME,2380 string asmName = !cast<VOP3P_Pseudo>(NAME).Mnemonic> {2381 defvar ps = !cast<VOP3P_Pseudo>(backing_ps_name);2382 let AsmString = asmName # ps.AsmOperands in2383 def Gen.Suffix :2384 VOP3P_Real_Gen<!cast<VOP3P_Pseudo>(backing_ps_name), Gen, asmName>,2385 VOP3Pe_gfx11_gfx12<op, !cast<VOP3P_Pseudo>(backing_ps_name).Pfl>;2386 2387 def : AMDGPUMnemonicAlias<ps.Mnemonic, asmName> {2388 let AssemblerPredicate = Gen.AssemblerPredicate;2389 }2390}2391 2392multiclass VOP3P_Real_dpp<GFXGen Gen, bits<8> op, string backing_ps_name = NAME,2393 string asmName = !cast<VOP3P_Pseudo>(NAME).Mnemonic> {2394 defvar ps = !cast<VOP3P_Pseudo>(backing_ps_name);2395 def _dpp#Gen.Suffix2396 : VOP3P_DPP16<op, !cast<VOP_DPP_Pseudo>(backing_ps_name #"_dpp"),2397 Gen.Subtarget> {2398 let AsmString = asmName #ps.Pfl.AsmVOP3DPP16;2399 let DecoderNamespace = Gen.DecoderNamespace;2400 let AssemblerPredicate = Gen.AssemblerPredicate;2401 }2402}2403 2404multiclass VOP3P_Real_dpp8<GFXGen Gen, bits<8> op, string backing_ps_name = NAME,2405 string asmName = !cast<VOP3P_Pseudo>(NAME).Mnemonic> {2406 defvar ps = !cast<VOP3P_Pseudo>(backing_ps_name);2407 def _dpp8#Gen.Suffix : VOP3P_DPP8_Base<op, ps> {2408 let AsmString = asmName #ps.Pfl.AsmVOP3DPP8;2409 let DecoderNamespace = Gen.DecoderNamespace;2410 let AssemblerPredicate = Gen.AssemblerPredicate;2411 }2412}2413 2414multiclass VOP3P_Realtriple<GFXGen Gen, bits<8> op, string backing_ps_name = NAME,2415 string asmName = !cast<VOP3P_Pseudo>(NAME).Mnemonic>2416 : VOP3P_Real_Base<Gen, op, backing_ps_name, asmName>,2417 VOP3P_Real_dpp<Gen, op, backing_ps_name, asmName>,2418 VOP3P_Real_dpp8<Gen, op, backing_ps_name, asmName>;2419 2420multiclass VOP3P_Realtriple_gfx11_gfx12<bits<8> op>2421 : VOP3P_Realtriple<GFX11Gen, op>, VOP3P_Realtriple<GFX12Gen, op>;2422 2423//===----------------------------------------------------------------------===//2424// GFX122425//===----------------------------------------------------------------------===//2426 2427multiclass VOP3P_Real_gfx12<bits<8> op> : VOP3P_Real_Base<GFX12Gen, op>;2428 2429multiclass VOP3P_Real_gfx1250<bits<8> op> : VOP3P_Real_Base<GFX1250Gen, op>;2430 2431multiclass VOP3P_Real_with_name_gfx12<bits<8> op,2432 string backing_ps_name = NAME,2433 string asmName = !cast<VOP3P_Pseudo>(NAME).Mnemonic> :2434 VOP3P_Real_with_name<GFX12Gen, op, backing_ps_name, asmName>;2435 2436multiclass VOP3P_Real_LD_SCALE_gfx1250<bits<8> op> {2437 defvar ps = !cast<VOP3P_Pseudo>(NAME);2438 def _gfx1250 :2439 VOP3P_Real_Gen<ps, GFX1250Gen, ps.Mnemonic>,2440 VOP3Pe_gfx11_gfx12<op, ps.Pfl> {2441 let Inst{58-50} = 0x100; // scale src2 = vgpr0 (dummy)2442 }2443}2444 2445defm V_PK_MIN_NUM_F16 : VOP3P_Real_with_name_gfx12<0x1b, "V_PK_MIN_F16", "v_pk_min_num_f16">;2446defm V_PK_MAX_NUM_F16 : VOP3P_Real_with_name_gfx12<0x1c, "V_PK_MAX_F16", "v_pk_max_num_f16">;2447 2448defm V_PK_FMA_F32 : VOP3P_Real_gfx12<0x1f>;2449defm V_PK_MUL_F32 : VOP3P_Real_gfx12<0x28>;2450defm V_PK_ADD_F32 : VOP3P_Real_gfx12<0x29>;2451 2452defm V_PK_ADD_MAX_I16 : VOP3P_Real_gfx1250<0x14>;2453defm V_PK_ADD_MAX_U16 : VOP3P_Real_gfx1250<0x15>;2454defm V_PK_ADD_MIN_I16 : VOP3P_Real_gfx1250<0x2d>;2455defm V_PK_ADD_MIN_U16 : VOP3P_Real_gfx1250<0x2e>;2456defm V_PK_MAX3_I16 : VOP3P_Real_gfx1250<0x2f>;2457defm V_PK_MAX3_U16 : VOP3P_Real_gfx1250<0x30>;2458defm V_PK_MIN3_I16 : VOP3P_Real_gfx1250<0x31>;2459defm V_PK_MIN3_U16 : VOP3P_Real_gfx1250<0x32>;2460defm V_PK_FMA_BF16 : VOP3P_Real_gfx1250<0x11>;2461defm V_PK_ADD_BF16 : VOP3P_Real_gfx1250<0x23>;2462defm V_PK_MUL_BF16 : VOP3P_Real_gfx1250<0x2a>;2463defm V_PK_MIN_NUM_BF16 : VOP3P_Real_gfx1250<0x2b>;2464defm V_PK_MAX_NUM_BF16 : VOP3P_Real_gfx1250<0x2c>;2465defm V_PK_MINIMUM3_F16 : VOP3P_Real_gfx1250<0x36>;2466defm V_PK_MAXIMUM3_F16 : VOP3P_Real_gfx1250<0x37>;2467defm V_PK_MIN3_NUM_F16 : VOP3P_Real_gfx1250<0x38>;2468defm V_PK_MAX3_NUM_F16 : VOP3P_Real_gfx1250<0x39>;2469 2470defm V_FMA_MIX_F32_BF16 : VOP3P_Realtriple<GFX1250Gen, 0x3d>;2471defm V_FMA_MIXLO_BF16 : VOP3P_Realtriple<GFX1250Gen, 0x3e>;2472defm V_FMA_MIXHI_BF16 : VOP3P_Realtriple<GFX1250Gen, 0x3f>;2473 2474defm V_WMMA_LD_SCALE_PAIRED_B32 : VOP3P_Real_LD_SCALE_gfx1250<0x35>;2475defm V_WMMA_LD_SCALE16_PAIRED_B64 : VOP3P_Real_LD_SCALE_gfx1250<0x3a>;2476 2477let AssemblerPredicate = isGFX1250Plus in2478def : AMDGPUMnemonicAlias<"v_fma_mix_f32_f16", "v_fma_mix_f32">;2479 2480defm V_PK_MINIMUM_F16 : VOP3P_Real_gfx12<0x1d>;2481defm V_PK_MAXIMUM_F16 : VOP3P_Real_gfx12<0x1e>;2482 2483defm V_DOT4_F32_FP8_BF8 : VOP3P_Realtriple<GFX12Gen, 0x24>;2484defm V_DOT4_F32_BF8_FP8 : VOP3P_Realtriple<GFX12Gen, 0x25>;2485defm V_DOT4_F32_FP8_FP8 : VOP3P_Realtriple<GFX12Gen, 0x26>;2486defm V_DOT4_F32_BF8_BF8 : VOP3P_Realtriple<GFX12Gen, 0x27>;2487 2488//===----------------------------------------------------------------------===//2489// GFX112490//===----------------------------------------------------------------------===//2491 2492multiclass VOP3P_Real_gfx11_gfx12<bits<8> op> :2493 VOP3P_Real_Base<GFX11Gen, op>, VOP3P_Real_Base<GFX12Gen, op>;2494 2495defm V_DOT4_I32_IU8 : VOP3P_Real_gfx11_gfx12<0x16>;2496defm V_DOT8_I32_IU4 : VOP3P_Real_gfx11_gfx12<0x18>;2497defm V_DOT2_F32_BF16 : VOP3P_Realtriple_gfx11_gfx12<0x1a>;2498 2499let AssemblerPredicate = isGFX11Plus in {2500 def : AMDGPUMnemonicAlias<"v_dot4_i32_i8", "v_dot4_i32_iu8">;2501 def : AMDGPUMnemonicAlias<"v_dot8_i32_i4", "v_dot8_i32_iu4">;2502}2503 2504multiclass VOP3P_Real_WMMA <bits<8> op> {2505 let WaveSizePredicate = isWave32, DecoderNamespace = "GFX11" in {2506 defm _twoaddr_w32 : VOP3P_Real_Base <GFX11Gen, op>;2507 }2508 let WaveSizePredicate = isWave64, DecoderNamespace = "GFX11W64" in {2509 defm _twoaddr_w64 : VOP3P_Real_Base <GFX11Gen, op>;2510 }2511}2512 2513defm V_WMMA_F32_16X16X16_F16 : VOP3P_Real_WMMA <0x040>;2514defm V_WMMA_F32_16X16X16_BF16 : VOP3P_Real_WMMA <0x041>;2515defm V_WMMA_F16_16X16X16_F16 : VOP3P_Real_WMMA <0x042>;2516defm V_WMMA_BF16_16X16X16_BF16 : VOP3P_Real_WMMA <0x043>;2517defm V_WMMA_I32_16X16X16_IU8 : VOP3P_Real_WMMA <0x044>;2518defm V_WMMA_I32_16X16X16_IU4 : VOP3P_Real_WMMA <0x045>;2519 2520//===----------------------------------------------------------------------===//2521// GFX8 (VI)2522//===----------------------------------------------------------------------===//2523 2524multiclass VOP3P_Real_vi<bits<7> op> {2525 def _vi : VOP3P_Real<!cast<VOP3_Pseudo>(NAME), SIEncodingFamily.VI>,2526 VOP3Pe_vi <op, !cast<VOP3_Pseudo>(NAME).Pfl> {2527 let AssemblerPredicate = HasVOP3PInsts;2528 let DecoderNamespace = "GFX8";2529 let VOP3P = 1;2530 }2531}2532 2533multiclass VOP3P_Real_MAI<bits<7> op> {2534 def _vi : VOP3P_Real<!cast<VOP3_Pseudo>(NAME#"_e64"), SIEncodingFamily.VI>,2535 VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME#"_e64").Pfl, ?> {2536 let AssemblerPredicate = HasMAIInsts;2537 let DecoderNamespace = "GFX8";2538 let Inst{14} = ?; // op_sel_hi(2)2539 let Inst{59} = ?; // op_sel_hi(0)2540 let Inst{60} = ?; // op_sel_hi(1)2541 }2542}2543 2544let Constraints = "" in {2545multiclass VOP3P_Real_MFMA_gfx90a<bits<7> op> {2546 let SubtargetPredicate = isGFX90AOnly,2547 AssemblerPredicate = isGFX90AOnly, DecoderNamespace = "GFX90A" in {2548 def _gfx90a_acd : VOP3P_Real<!cast<VOP3_Pseudo>(NAME#"_e64"), SIEncodingFamily.GFX90A>,2549 VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME#"_e64").Pfl, 1>;2550 2551 def _gfx90a_vcd : VOP3P_Real<!cast<VOP3_Pseudo>(NAME # "_vgprcd" # "_e64"), SIEncodingFamily.GFX90A>,2552 VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME # "_vgprcd" # "_e64").Pfl, 0>;2553 } // End AssemblerPredicate = isGFX90AOnly, DecoderNamespace = "GFX90A"2554}2555}2556 2557multiclass VOP3P_Real_MFMA_gfx940_aliases<string NameFrom, string NameTo, string Op,2558 VOP3_Pseudo PS_ACD = !cast<VOP3_Pseudo>(Op # "_e64"),2559 VOP3_Pseudo PS_VCD = !cast<VOP3_Pseudo>(Op # "_vgprcd" # "_e64"),2560 VOPProfile Pfl_ACD = PS_ACD.Pfl,2561 VOPProfile Pfl_VCD = PS_VCD.Pfl> {2562 if !ne(NameFrom, NameTo) then {2563 let SubtargetPredicate = PS_ACD.SubtargetPredicate,2564 OtherPredicates = PS_ACD.OtherPredicates in {2565 def : InstAlias <NameTo # " " # PS_ACD.AsmOperands,2566 (!cast<VOP3P_Real>(Op # "_gfx940_acd") Pfl_ACD.DstRC:$vdst,2567 Pfl_ACD.Src0RC64:$src0, Pfl_ACD.Src1RC64:$src1, Pfl_ACD.Src2RC64:$src2,2568 CBSZ:$cbsz, ABID:$abid, blgp:$blgp)>, PredicateControl;2569 def : InstAlias <NameTo # " " # PS_VCD.AsmOperands,2570 (!cast<VOP3P_Real>(Op # "_gfx940_vcd") Pfl_VCD.DstRC:$vdst,2571 Pfl_VCD.Src0RC64:$src0, Pfl_VCD.Src1RC64:$src1, Pfl_VCD.Src2RC64:$src2,2572 CBSZ:$cbsz, ABID:$abid, blgp:$blgp)>, PredicateControl;2573 }2574 }2575}2576 2577multiclass VOP3P_Real_MFMA_gfx940<bits<7> op, string Name = !cast<VOP3_Pseudo>(NAME#"_e64").Mnemonic,2578 VOP3_Pseudo PS_ACD = !cast<VOP3_Pseudo>(NAME # "_e64"),2579 VOP3_Pseudo PS_VCD = !cast<VOP3_Pseudo>(NAME # "_vgprcd" # "_e64")> {2580 let AssemblerPredicate = isGFX940Plus,2581 DecoderNamespace = "GFX940",2582 AsmString = Name # PS_ACD.AsmOperands, Constraints = "" in {2583 def _gfx940_acd : VOP3P_Real<PS_ACD, SIEncodingFamily.GFX940>,2584 VOP3Pe_MAI <op, PS_ACD.Pfl, 1>;2585 2586 def _gfx940_vcd : VOP3P_Real<PS_VCD, SIEncodingFamily.GFX940>,2587 VOP3Pe_MAI <op, PS_VCD.Pfl, 0>;2588 } // End AssemblerPredicate = isGFX940Plus, DecoderNamespace = "GFX940"2589 2590 let SubtargetPredicate = PS_ACD.SubtargetPredicate,2591 OtherPredicates = PS_ACD.OtherPredicates,2592 AssemblerPredicate = isGFX940Plus2593 in {2594 defm : VOP3P_Real_MFMA_gfx940_aliases<Name, PS_ACD.Mnemonic, NAME>;2595 2596 if !ne(!subst("_1k", "", PS_ACD.Mnemonic), PS_ACD.Mnemonic) then2597 defm : VOP3P_Real_MFMA_gfx940_aliases<Name, !subst("_1k", "", PS_ACD.Mnemonic), NAME>;2598 }2599}2600 2601multiclass VOP3P_Real_MFMA_F8F6F4_gfx940<bits<7> op, string Name = !cast<VOP3_Pseudo>(NAME#"_e64").Mnemonic,2602 VOP3_Pseudo PS_ACD = !cast<VOP3_Pseudo>(NAME # "_e64"),2603 VOP3_Pseudo PS_VCD = !cast<VOP3_Pseudo>(NAME # "_vgprcd" # "_e64")> {2604 2605 defvar F8F8Name = !substr(NAME, 0, !sub(!size(NAME), !size("_fN_fM")))#"_f8_f8";2606 2607 let AssemblerPredicate = isGFX940Plus,2608 DecoderNamespace = "GFX940",2609 AsmString = Name # PS_ACD.AsmOperands,2610 Constraints = "" in {2611 def _gfx940_acd : VOP3P_Real<PS_ACD, SIEncodingFamily.GFX940>,2612 VOP3Pe_MAI <op, PS_ACD.Pfl, 1>,2613 MFMA_F8F6F4_WithSizeTable_Helper<PS_ACD, F8F8Name#"_gfx940_acd">;2614 2615 def _gfx940_vcd : VOP3P_Real<PS_VCD, SIEncodingFamily.GFX940>,2616 VOP3Pe_MAI <op, PS_VCD.Pfl, 0>,2617 MFMA_F8F6F4_WithSizeTable_Helper<PS_VCD, F8F8Name#"_gfx940_vcd">;2618 } // End AssemblerPredicate = isGFX940Plus, DecoderNamespace = "GFX940"2619}2620 2621multiclass VOP3P_Real_MFMA_gfx950<bits<7> op, string Name = !cast<VOP3_Pseudo>(NAME#"_e64").Mnemonic,2622 VOP3_Pseudo PS_ACD = !cast<VOP3_Pseudo>(NAME # "_e64"),2623 VOP3_Pseudo PS_VCD = !cast<VOP3_Pseudo>(NAME # "_vgprcd" # "_e64")> {2624 let SubtargetPredicate = HasGFX950Insts,2625 AssemblerPredicate = HasGFX950Insts in {2626 defm "" : VOP3P_Real_MFMA_gfx940<op, Name, PS_ACD, PS_VCD>;2627 }2628}2629 2630 2631multiclass VOP3P_Real_MFMA_F8F6F4_gfx950_mc<bits<7> op, string Name> {2632 defm _f8_f8 : VOP3P_Real_MFMA_F8F6F4_gfx940<op, Name>;2633 2634 let isAsmParserOnly = true in { // Disable ambiguous disassembly.2635 defm _f8_f6 : VOP3P_Real_MFMA_F8F6F4_gfx940<op, Name>;2636 defm _f6_f8 : VOP3P_Real_MFMA_F8F6F4_gfx940<op, Name>;2637 defm _f8_f4 : VOP3P_Real_MFMA_F8F6F4_gfx940<op, Name>;2638 defm _f4_f8 : VOP3P_Real_MFMA_F8F6F4_gfx940<op, Name>;2639 defm _f6_f6 : VOP3P_Real_MFMA_F8F6F4_gfx940<op, Name>;2640 defm _f6_f4 : VOP3P_Real_MFMA_F8F6F4_gfx940<op, Name>;2641 defm _f4_f6 : VOP3P_Real_MFMA_F8F6F4_gfx940<op, Name>;2642 defm _f4_f4 : VOP3P_Real_MFMA_F8F6F4_gfx940<op, Name>;2643 }2644}2645 2646multiclass VOP3PX_Real_ScaledMFMA<bits<7> op> {2647 defvar PS_ACD = !cast<VOP3_Pseudo>(NAME # "_e64");2648 defvar PS_VCD = !cast<VOP3_Pseudo>(NAME # "_vgprcd" # "_e64");2649 defvar Name = PS_ACD.Mnemonic;2650 defvar F8F8Name = !substr(NAME, 0, !sub(!size(NAME), !size("_fN_fM")))#"_f8_f8";2651 let SubtargetPredicate = HasGFX950Insts,2652 DecoderNamespace = "GFX940",2653 AsmString = Name # PS_ACD.AsmOperands, Constraints = "" in {2654 def _gfx940_acd : VOP3P_Real<PS_ACD, SIEncodingFamily.GFX940>,2655 VOP3PXe <op, PS_ACD.Pfl, /*acc_cd=*/1>,2656 MFMA_F8F6F4_WithSizeTable_Helper<PS_ACD, F8F8Name#"_gfx940_acd">;2657 2658 def _gfx940_vcd : VOP3P_Real<PS_VCD, SIEncodingFamily.GFX940>,2659 VOP3PXe <op, PS_VCD.Pfl, /*acc_cd=*/0>,2660 MFMA_F8F6F4_WithSizeTable_Helper<PS_VCD, F8F8Name#"_gfx940_vcd">;2661 }2662}2663 2664multiclass VOP3PX_Real_ScaledMFMA_F8F6F4_mc<bits<7> op> {2665 defm _f8_f8 : VOP3PX_Real_ScaledMFMA<op>;2666 2667 let isAsmParserOnly = 1 in { // Disable ambiguous disassembly.2668 defm _f8_f6 : VOP3PX_Real_ScaledMFMA<op>;2669 defm _f6_f8 : VOP3PX_Real_ScaledMFMA<op>;2670 defm _f8_f4 : VOP3PX_Real_ScaledMFMA<op>;2671 defm _f4_f8 : VOP3PX_Real_ScaledMFMA<op>;2672 defm _f6_f6 : VOP3PX_Real_ScaledMFMA<op>;2673 defm _f6_f4 : VOP3PX_Real_ScaledMFMA<op>;2674 defm _f4_f6 : VOP3PX_Real_ScaledMFMA<op>;2675 defm _f4_f4 : VOP3PX_Real_ScaledMFMA<op>;2676 }2677}2678 2679multiclass VOP3P_Real_MFMA_vi<bits<7> op> {2680 def _vi : VOP3P_Real<!cast<VOP3_Pseudo>(NAME#"_e64"), SIEncodingFamily.VI>,2681 VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME#"_e64").Pfl, ?> {2682 let SubtargetPredicate = isGFX8GFX9NotGFX90A;2683 let AssemblerPredicate = HasMAIInsts;2684 let DecoderNamespace = "GFX8";2685 let Constraints = "";2686 }2687}2688 2689multiclass VOP3P_Real_MFMA_vi_gfx90a<bits<7> op> :2690 VOP3P_Real_MFMA_gfx90a <op>,2691 VOP3P_Real_MFMA_vi <op>;2692 2693multiclass VOP3P_Real_MFMA<bits<7> op, string GFX940Name = !cast<VOP3_Pseudo>(NAME#"_e64").Mnemonic> :2694 VOP3P_Real_MFMA_vi_gfx90a <op>,2695 VOP3P_Real_MFMA_gfx940 <op, GFX940Name>;2696 2697multiclass VOP3P_Real_SMFMAC<bits<7> op, string alias> {2698 def _gfx940 : VOP3P_Real<!cast<VOP3_Pseudo>(NAME#"_e64"), SIEncodingFamily.VI>,2699 VOP3Pe_SMFMAC <op> {2700 let AssemblerPredicate = isGFX940Plus;2701 let DecoderNamespace = "GFX8";2702 }2703 def : AMDGPUMnemonicAlias<alias, !cast<VOP3_Pseudo>(NAME#"_e64").Mnemonic> {2704 let AssemblerPredicate = isGFX940Plus;2705 }2706}2707 2708defm V_PK_MAD_I16 : VOP3P_Real_vi <0x00>;2709defm V_PK_MUL_LO_U16 : VOP3P_Real_vi <0x01>;2710defm V_PK_ADD_I16 : VOP3P_Real_vi <0x02>;2711defm V_PK_SUB_I16 : VOP3P_Real_vi <0x03>;2712defm V_PK_LSHLREV_B16 : VOP3P_Real_vi <0x04>;2713defm V_PK_LSHRREV_B16 : VOP3P_Real_vi <0x05>;2714defm V_PK_ASHRREV_I16 : VOP3P_Real_vi <0x06>;2715defm V_PK_MAX_I16 : VOP3P_Real_vi <0x07>;2716defm V_PK_MIN_I16 : VOP3P_Real_vi <0x08>;2717defm V_PK_MAD_U16 : VOP3P_Real_vi <0x09>;2718 2719defm V_PK_ADD_U16 : VOP3P_Real_vi <0x0a>;2720defm V_PK_SUB_U16 : VOP3P_Real_vi <0x0b>;2721defm V_PK_MAX_U16 : VOP3P_Real_vi <0x0c>;2722defm V_PK_MIN_U16 : VOP3P_Real_vi <0x0d>;2723defm V_PK_FMA_F16 : VOP3P_Real_vi <0x0e>;2724defm V_PK_ADD_F16 : VOP3P_Real_vi <0x0f>;2725defm V_PK_MUL_F16 : VOP3P_Real_vi <0x10>;2726defm V_PK_MIN_F16 : VOP3P_Real_vi <0x11>;2727defm V_PK_MAX_F16 : VOP3P_Real_vi <0x12>;2728 2729defm V_DOT2_F32_BF16 : VOP3P_Real_vi<0x1a>;2730defm V_PK_MINIMUM3_F16 : VOP3P_Real_vi <0x1b>;2731defm V_PK_MAXIMUM3_F16 : VOP3P_Real_vi <0x1c>;2732 2733defm V_MAD_MIX_F32 : VOP3P_Real_vi <0x20>;2734defm V_MAD_MIXLO_F16 : VOP3P_Real_vi <0x21>;2735defm V_MAD_MIXHI_F16 : VOP3P_Real_vi <0x22>;2736 2737let OtherPredicates = [HasFmaMixInsts],2738 DecoderNamespace = "GFX9_DL" in {2739// The mad_mix instructions were renamed and their behaviors changed,2740// but the opcode stayed the same so we need to put these in a2741// different DecoderNamespace to avoid the ambiguity.2742defm V_FMA_MIX_F32 : VOP3P_Real_vi <0x20>;2743defm V_FMA_MIXLO_F16 : VOP3P_Real_vi <0x21>;2744defm V_FMA_MIXHI_F16 : VOP3P_Real_vi <0x22>;2745}2746 2747defm V_DOT2_I32_I16 : VOP3P_Real_vi <0x26>;2748defm V_DOT2_U32_U16 : VOP3P_Real_vi <0x27>;2749 2750defm V_DOT2_F32_F16 : VOP3P_Real_vi <0x23>;2751defm V_DOT4_U32_U8 : VOP3P_Real_vi <0x29>;2752defm V_DOT8_U32_U4 : VOP3P_Real_vi <0x2b>;2753 2754defm V_DOT4_I32_I8 : VOP3P_Real_vi <0x28>;2755defm V_DOT8_I32_I4 : VOP3P_Real_vi <0x2a>;2756 2757defm V_ACCVGPR_READ_B32 : VOP3P_Real_MAI <0x58>;2758defm V_ACCVGPR_WRITE_B32 : VOP3P_Real_MAI <0x59>;2759defm V_MFMA_F32_32X32X1F32 : VOP3P_Real_MFMA <0x40, "v_mfma_f32_32x32x1_2b_f32">;2760defm V_MFMA_F32_16X16X1F32 : VOP3P_Real_MFMA <0x41, "v_mfma_f32_16x16x1_4b_f32">;2761defm V_MFMA_F32_4X4X1F32 : VOP3P_Real_MFMA <0x42, "v_mfma_f32_4x4x1_16b_f32">;2762defm V_MFMA_F32_32X32X2F32 : VOP3P_Real_MFMA <0x44, "v_mfma_f32_32x32x2_f32">;2763defm V_MFMA_F32_16X16X4F32 : VOP3P_Real_MFMA <0x45, "v_mfma_f32_16x16x4_f32">;2764defm V_MFMA_F32_32X32X4F16 : VOP3P_Real_MFMA <0x48, "v_mfma_f32_32x32x4_2b_f16">;2765defm V_MFMA_F32_16X16X4F16 : VOP3P_Real_MFMA <0x49, "v_mfma_f32_16x16x4_4b_f16">;2766defm V_MFMA_F32_4X4X4F16 : VOP3P_Real_MFMA <0x4a, "v_mfma_f32_4x4x4_16b_f16">;2767defm V_MFMA_F32_32X32X8F16 : VOP3P_Real_MFMA <0x4c, "v_mfma_f32_32x32x8_f16">;2768defm V_MFMA_F32_16X16X16F16 : VOP3P_Real_MFMA <0x4d, "v_mfma_f32_16x16x16_f16">;2769defm V_MFMA_I32_32X32X4I8 : VOP3P_Real_MFMA <0x50, "v_mfma_i32_32x32x4_2b_i8">;2770defm V_MFMA_I32_16X16X4I8 : VOP3P_Real_MFMA <0x51, "v_mfma_i32_16x16x4_4b_i8">;2771defm V_MFMA_I32_4X4X4I8 : VOP3P_Real_MFMA <0x52, "v_mfma_i32_4x4x4_16b_i8">;2772 2773defm V_MFMA_I32_16X16X16I8 : VOP3P_Real_MFMA_vi_gfx90a <0x55>;2774defm V_MFMA_I32_32X32X8I8 : VOP3P_Real_MFMA_vi_gfx90a <0x54>;2775defm V_MFMA_F32_32X32X2BF16 : VOP3P_Real_MFMA_vi_gfx90a <0x68>;2776defm V_MFMA_F32_16X16X2BF16 : VOP3P_Real_MFMA_vi_gfx90a <0x69>;2777defm V_MFMA_F32_4X4X2BF16 : VOP3P_Real_MFMA_vi_gfx90a <0x6b>;2778defm V_MFMA_F32_32X32X4BF16 : VOP3P_Real_MFMA_vi_gfx90a <0x6c>;2779defm V_MFMA_F32_16X16X8BF16 : VOP3P_Real_MFMA_vi_gfx90a <0x6d>;2780 2781defm V_MFMA_F32_32X32X4BF16_1K : VOP3P_Real_MFMA_gfx90a <0x63>;2782defm V_MFMA_F32_16X16X4BF16_1K : VOP3P_Real_MFMA_gfx90a <0x64>;2783defm V_MFMA_F32_4X4X4BF16_1K : VOP3P_Real_MFMA_gfx90a <0x65>;2784defm V_MFMA_F32_32X32X8BF16_1K : VOP3P_Real_MFMA_gfx90a <0x66>;2785defm V_MFMA_F32_16X16X16BF16_1K : VOP3P_Real_MFMA_gfx90a <0x67>;2786defm V_MFMA_F64_16X16X4F64 : VOP3P_Real_MFMA_gfx90a <0x6e>;2787defm V_MFMA_F64_4X4X4F64 : VOP3P_Real_MFMA_gfx90a <0x6f>;2788 2789defm V_MFMA_F32_16X16X32_F16 : VOP3P_Real_MFMA_gfx950 <0x54, "v_mfma_f32_16x16x32_f16">;2790defm V_MFMA_F32_32X32X16_F16 : VOP3P_Real_MFMA_gfx950 <0x55, "v_mfma_f32_32x32x16_f16">;2791defm V_MFMA_F32_16X16X32_BF16 : VOP3P_Real_MFMA_gfx950 <0x35, "v_mfma_f32_16x16x32_bf16">;2792defm V_MFMA_I32_16X16X64_I8 : VOP3P_Real_MFMA_gfx950 <0x36, "v_mfma_i32_16x16x64_i8">;2793defm V_MFMA_F32_32X32X16_BF16 : VOP3P_Real_MFMA_gfx950 <0x37, "v_mfma_f32_32x32x16_bf16">;2794defm V_MFMA_I32_32X32X32_I8 : VOP3P_Real_MFMA_gfx950 <0x38, "v_mfma_i32_32x32x32_i8">;2795 2796defm V_MFMA_LD_SCALE_B32 : VOP3P_Real_vi <0x2c>;2797defm V_MFMA_F32_16X16X128_F8F6F4 : VOP3P_Real_MFMA_F8F6F4_gfx950_mc <0x2d, "v_mfma_f32_16x16x128_f8f6f4">;2798defm V_MFMA_SCALE_F32_16X16X128_F8F6F4 : VOP3PX_Real_ScaledMFMA_F8F6F4_mc <0x2d>;2799defm V_MFMA_F32_32X32X64_F8F6F4 : VOP3P_Real_MFMA_F8F6F4_gfx950_mc <0x2e, "v_mfma_f32_32x32x64_f8f6f4">;2800defm V_MFMA_SCALE_F32_32X32X64_F8F6F4 : VOP3PX_Real_ScaledMFMA_F8F6F4_mc <0x2e>;2801 2802defm V_MFMA_I32_32X32X16I8 : VOP3P_Real_MFMA_gfx940 <0x56, "v_mfma_i32_32x32x16_i8">;2803defm V_MFMA_I32_16X16X32I8 : VOP3P_Real_MFMA_gfx940 <0x57, "v_mfma_i32_16x16x32_i8">;2804defm V_MFMA_F32_16X16X8XF32 : VOP3P_Real_MFMA_gfx940 <0x3e, "v_mfma_f32_16x16x8_xf32">;2805defm V_MFMA_F32_32X32X4XF32 : VOP3P_Real_MFMA_gfx940 <0x3f, "v_mfma_f32_32x32x4_xf32">;2806 2807defm V_MFMA_F32_16X16X32_BF8_BF8 : VOP3P_Real_MFMA_gfx940 <0x70>;2808defm V_MFMA_F32_16X16X32_BF8_FP8 : VOP3P_Real_MFMA_gfx940 <0x71>;2809defm V_MFMA_F32_16X16X32_FP8_BF8 : VOP3P_Real_MFMA_gfx940 <0x72>;2810defm V_MFMA_F32_16X16X32_FP8_FP8 : VOP3P_Real_MFMA_gfx940 <0x73>;2811defm V_MFMA_F32_32X32X16_BF8_BF8 : VOP3P_Real_MFMA_gfx940 <0x74>;2812defm V_MFMA_F32_32X32X16_BF8_FP8 : VOP3P_Real_MFMA_gfx940 <0x75>;2813defm V_MFMA_F32_32X32X16_FP8_BF8 : VOP3P_Real_MFMA_gfx940 <0x76>;2814defm V_MFMA_F32_32X32X16_FP8_FP8 : VOP3P_Real_MFMA_gfx940 <0x77>;2815 2816defm V_MFMA_F32_32X32X4BF16_1K : VOP3P_Real_MFMA_gfx940 <0x5d, "v_mfma_f32_32x32x4_2b_bf16">;2817defm V_MFMA_F32_16X16X4BF16_1K : VOP3P_Real_MFMA_gfx940 <0x5e, "v_mfma_f32_16x16x4_4b_bf16">;2818defm V_MFMA_F32_4X4X4BF16_1K : VOP3P_Real_MFMA_gfx940 <0x5f, "v_mfma_f32_4x4x4_16b_bf16">;2819defm V_MFMA_F32_32X32X8BF16_1K : VOP3P_Real_MFMA_gfx940 <0x60, "v_mfma_f32_32x32x8_bf16">;2820defm V_MFMA_F32_16X16X16BF16_1K : VOP3P_Real_MFMA_gfx940 <0x61, "v_mfma_f32_16x16x16_bf16">;2821 2822defm V_MFMA_F64_16X16X4F64 : VOP3P_Real_MFMA_gfx940 <0x6e, "v_mfma_f64_16x16x4_f64">;2823defm V_MFMA_F64_4X4X4F64 : VOP3P_Real_MFMA_gfx940 <0x6f, "v_mfma_f64_4x4x4_4b_f64">;2824 2825defm V_SMFMAC_F32_16X16X32_F16 : VOP3P_Real_SMFMAC <0x62, "v_smfmac_f32_16x16x32f16">;2826defm V_SMFMAC_F32_32X32X16_F16 : VOP3P_Real_SMFMAC <0x64, "v_smfmac_f32_32x32x16f16">;2827defm V_SMFMAC_F32_16X16X32_BF16 : VOP3P_Real_SMFMAC <0x66, "v_smfmac_f32_16x16x32bf16">;2828defm V_SMFMAC_F32_32X32X16_BF16 : VOP3P_Real_SMFMAC <0x68, "v_smfmac_f32_32x32x16bf16">;2829defm V_SMFMAC_I32_16X16X64_I8 : VOP3P_Real_SMFMAC <0x6a, "v_smfmac_i32_16x16x64i8">;2830defm V_SMFMAC_I32_32X32X32_I8 : VOP3P_Real_SMFMAC <0x6c, "v_smfmac_i32_32x32x32i8">;2831defm V_SMFMAC_F32_16X16X64_BF8_BF8 : VOP3P_Real_SMFMAC <0x78, "v_smfmac_f32_16x16x64bf8bf8">;2832defm V_SMFMAC_F32_16X16X64_BF8_FP8 : VOP3P_Real_SMFMAC <0x79, "v_smfmac_f32_16x16x64bf8fp8">;2833defm V_SMFMAC_F32_16X16X64_FP8_BF8 : VOP3P_Real_SMFMAC <0x7a, "v_smfmac_f32_16x16x64fp8bf8">;2834defm V_SMFMAC_F32_16X16X64_FP8_FP8 : VOP3P_Real_SMFMAC <0x7b, "v_smfmac_f32_16x16x64fp8fp8">;2835defm V_SMFMAC_F32_32X32X32_BF8_BF8 : VOP3P_Real_SMFMAC <0x7c, "v_smfmac_f32_32x32x32bf8bf8">;2836defm V_SMFMAC_F32_32X32X32_BF8_FP8 : VOP3P_Real_SMFMAC <0x7d, "v_smfmac_f32_32x32x32bf8fp8">;2837defm V_SMFMAC_F32_32X32X32_FP8_BF8 : VOP3P_Real_SMFMAC <0x7e, "v_smfmac_f32_32x32x32fp8bf8">;2838defm V_SMFMAC_F32_32X32X32_FP8_FP8 : VOP3P_Real_SMFMAC <0x7f, "v_smfmac_f32_32x32x32fp8fp8">;2839 2840defm V_SMFMAC_F32_16X16X64_F16 : VOP3P_Real_SMFMAC <0x5a, "v_smfmac_f32_16x16x64f16">;2841defm V_SMFMAC_F32_32X32X32_F16 : VOP3P_Real_SMFMAC <0x5b, "v_smfmac_f32_32x32x32f16">;2842defm V_SMFMAC_F32_16X16X64_BF16 : VOP3P_Real_SMFMAC <0x39, "v_smfmac_f32_16x16x64bf16">;2843defm V_SMFMAC_F32_32X32X32_BF16 : VOP3P_Real_SMFMAC <0x46, "v_smfmac_f32_32x32x32bf16">;2844defm V_SMFMAC_I32_16X16X128_I8 : VOP3P_Real_SMFMAC <0x3a, "v_smfmac_i32_16x16x128i8">;2845defm V_SMFMAC_I32_32X32X64_I8 : VOP3P_Real_SMFMAC <0x47, "v_smfmac_i32_32x32x64i8">;2846 2847defm V_SMFMAC_F32_16X16X128_BF8_BF8 : VOP3P_Real_SMFMAC <0x3b, "v_smfmac_f32_16x16x128bf8bf8">;2848defm V_SMFMAC_F32_16X16X128_BF8_FP8 : VOP3P_Real_SMFMAC <0x3c, "v_smfmac_f32_16x16x128bf8fp8">;2849defm V_SMFMAC_F32_16X16X128_FP8_BF8 : VOP3P_Real_SMFMAC <0x3d, "v_smfmac_f32_16x16x128fp8bf8">;2850defm V_SMFMAC_F32_16X16X128_FP8_FP8 : VOP3P_Real_SMFMAC <0x43, "v_smfmac_f32_16x16x128fp8fp8">;2851defm V_SMFMAC_F32_32X32X64_BF8_BF8 : VOP3P_Real_SMFMAC <0x4b, "v_smfmac_f32_32x32x64bf8bf8">;2852defm V_SMFMAC_F32_32X32X64_BF8_FP8 : VOP3P_Real_SMFMAC <0x4e, "v_smfmac_f32_32x32x64bf8fp8">;2853defm V_SMFMAC_F32_32X32X64_FP8_BF8 : VOP3P_Real_SMFMAC <0x4f, "v_smfmac_f32_32x32x64fp8bf8">;2854defm V_SMFMAC_F32_32X32X64_FP8_FP8 : VOP3P_Real_SMFMAC <0x53, "v_smfmac_f32_32x32x64fp8fp8">;2855 2856defm V_PK_FMA_F32 : VOP3P_Real_vi <0x30>;2857defm V_PK_MUL_F32 : VOP3P_Real_vi <0x31>;2858defm V_PK_ADD_F32 : VOP3P_Real_vi <0x32>;2859defm V_PK_MOV_B32 : VOP3P_Real_vi <0x33>;2860 2861//===----------------------------------------------------------------------===//2862// GFX10.2863//===----------------------------------------------------------------------===//2864 2865let AssemblerPredicate = isGFX10Only, DecoderNamespace = "GFX10", VOP3P = 1 in {2866 multiclass VOP3P_Real_gfx10<bits<8> op> {2867 def _gfx10 : VOP3P_Real<!cast<VOP3P_Pseudo>(NAME), SIEncodingFamily.GFX10>,2868 VOP3Pe_gfx10 <op, !cast<VOP3P_Pseudo>(NAME).Pfl>;2869 }2870} // End AssemblerPredicate = isGFX10Only, DecoderNamespace = "GFX10", VOP3P = 12871 2872multiclass VOP3P_Real_gfx10_gfx11<bits<8> op> :2873 VOP3P_Real_gfx10<op>, VOP3P_Real_Base<GFX11Gen, op>;2874 2875multiclass VOP3P_Real_gfx10_gfx11_gfx12<bits<8> op> :2876 VOP3P_Real_gfx10_gfx11<op>, VOP3P_Real_Base<GFX12Gen, op>;2877 2878multiclass VOP3P_Real_gfx10_gfx11_gfx12_Triple<bits<8> op> :2879 VOP3P_Real_gfx10<op>, VOP3P_Realtriple<GFX11Gen, op>,2880 VOP3P_Realtriple<GFX12Gen, op>;2881 2882defm V_PK_MAD_I16 : VOP3P_Real_gfx10_gfx11_gfx12<0x00>;2883defm V_PK_MUL_LO_U16 : VOP3P_Real_gfx10_gfx11_gfx12<0x01>;2884defm V_PK_ADD_I16 : VOP3P_Real_gfx10_gfx11_gfx12<0x02>;2885defm V_PK_SUB_I16 : VOP3P_Real_gfx10_gfx11_gfx12<0x03>;2886defm V_PK_LSHLREV_B16 : VOP3P_Real_gfx10_gfx11_gfx12<0x04>;2887defm V_PK_LSHRREV_B16 : VOP3P_Real_gfx10_gfx11_gfx12<0x05>;2888defm V_PK_ASHRREV_I16 : VOP3P_Real_gfx10_gfx11_gfx12<0x06>;2889defm V_PK_MAX_I16 : VOP3P_Real_gfx10_gfx11_gfx12<0x07>;2890defm V_PK_MIN_I16 : VOP3P_Real_gfx10_gfx11_gfx12<0x08>;2891defm V_PK_MAD_U16 : VOP3P_Real_gfx10_gfx11_gfx12<0x09>;2892defm V_PK_ADD_U16 : VOP3P_Real_gfx10_gfx11_gfx12<0x0a>;2893defm V_PK_SUB_U16 : VOP3P_Real_gfx10_gfx11_gfx12<0x0b>;2894defm V_PK_MAX_U16 : VOP3P_Real_gfx10_gfx11_gfx12<0x0c>;2895defm V_PK_MIN_U16 : VOP3P_Real_gfx10_gfx11_gfx12<0x0d>;2896defm V_PK_FMA_F16 : VOP3P_Real_gfx10_gfx11_gfx12<0x0e>;2897defm V_PK_ADD_F16 : VOP3P_Real_gfx10_gfx11_gfx12<0x0f>;2898defm V_PK_MUL_F16 : VOP3P_Real_gfx10_gfx11_gfx12<0x10>;2899defm V_PK_MIN_F16 : VOP3P_Real_gfx10_gfx11<0x11>;2900defm V_PK_MAX_F16 : VOP3P_Real_gfx10_gfx11<0x12>;2901defm V_FMA_MIX_F32 : VOP3P_Real_gfx10_gfx11_gfx12_Triple<0x20>;2902defm V_FMA_MIXLO_F16 : VOP3P_Real_gfx10_gfx11_gfx12_Triple<0x21>;2903defm V_FMA_MIXHI_F16 : VOP3P_Real_gfx10_gfx11_gfx12_Triple<0x22>;2904 2905defm V_DOT2_I32_I16 : VOP3P_Real_gfx10 <0x14>;2906defm V_DOT2_U32_U16 : VOP3P_Real_gfx10 <0x15>;2907 2908defm V_DOT2_F32_F16 : VOP3P_Real_gfx10_gfx11_gfx12_Triple<0x13>;2909defm V_DOT4_U32_U8 : VOP3P_Real_gfx10_gfx11_gfx12<0x17>;2910defm V_DOT8_U32_U4 : VOP3P_Real_gfx10_gfx11_gfx12<0x19>;2911 2912defm V_DOT4_I32_I8 : VOP3P_Real_gfx10 <0x16>;2913defm V_DOT8_I32_I4 : VOP3P_Real_gfx10 <0x18>;2914