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-exception 6// 7//===----------------------------------------------------------------------===// 8 9//===----------------------------------------------------------------------===// 10// VOP3P Classes 11//===----------------------------------------------------------------------===// 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 are 18 // overloaded in meaning and the logic in printOperandAndFPInputMods is 19 // wrong for vop3p 20 let AsmVOP3DPPBase = AsmVOP3P; 21} 22 23// Used for FMA_MIX* and MAD_MIX* insts 24// Their operands are only sort of f16 operands. Depending on 25// op_sel_hi, these may be interpreted as f32. The inline immediate 26// values are really f16 converted to f32, so we treat these as f16 27// operands. 28class VOP3P_Mix_Profile<VOPProfile P, VOP3Features Features = VOP3_REGULAR, 29 bit useTiedOutput = 0> : VOP3P_Profile<P, Features, 1> { 30 bit UseTiedOutput = useTiedOutput; 31 32 dag srcs = 33 (ins FP16InputMods:$src0_modifiers, VCSrc_f16:$src0, 34 FP16InputMods:$src1_modifiers, VCSrc_f16:$src1, 35 FP16InputMods:$src2_modifiers, VCSrc_f16:$src2); 36 dag dpp_srcs = 37 (ins FPVRegInputMods:$src0_modifiers, VGPRSrc_32:$src0, 38 FP16InputMods:$src1_modifiers, VCSrc_f16:$src1, 39 FP16InputMods:$src2_modifiers, VCSrc_f16:$src2); 40 41 // FIXME: clampmod0 misbehaves with the non-default vdst_in 42 // following it. For now workaround this by requiring clamp 43 // in tied patterns. This should use undef_tied_input, but it 44 // seems underdeveloped and doesn't apply the right register 45 // class constraints. 46 dag mods = !con(!if(UseTiedOutput, (ins clampmod:$clamp, VGPR_32:$vdst_in), 47 (ins clampmod0:$clamp)), 48 (ins op_sel0:$op_sel, op_sel_hi0:$op_sel_hi)); 49 // We use Ins64 because that is the one which populates InOperandList 50 // due to the logic in class VOP3_Pseudo 51 let Ins64 = !con(srcs, mods); 52 let InsVOP3Base = !con(dpp_srcs, mods); 53 let Asm64 = 54 "$vdst, $src0_modifiers, $src1_modifiers, $src2_modifiers$op_sel$op_sel_hi$clamp"; 55 let AsmVOP3DPPBase = Asm64; 56} 57 58multiclass VOP3PInst<string OpName, VOPProfile P, 59 SDPatternOperator node = null_frag, bit IsDOT = 0> { 60 def NAME : VOP3P_Pseudo<OpName, P, 61 !if (P.HasModifiers, 62 getVOP3PModPat<P, node, IsDOT, IsDOT>.ret, 63 getVOP3Pat<P, node>.ret)>; 64 let SubtargetPredicate = isGFX11Plus in { 65 if P.HasExtVOP3DPP then 66 def _dpp : VOP3_DPP_Pseudo<OpName, P> { 67 let VOP3P = 1; 68 let PseudoInstr = OpName #"_dpp"; 69 } 70 } // end SubtargetPredicate = isGFX11Plus 71} 72 73// Non-packed instructions that use the VOP3P encoding. 74// VOP3 neg/abs and VOP3P opsel/opsel_hi modifiers are allowed. 75multiclass VOP3_VOP3PInst<string OpName, VOP3P_Mix_Profile P> { 76 def NAME : VOP3P_Pseudo<OpName, P> { 77 let Constraints = !if(P.UseTiedOutput, "$vdst = $vdst_in", ""); 78 let DisableEncoding = !if(P.UseTiedOutput, "$vdst_in", ""); 79 } 80 let SubtargetPredicate = isGFX11Plus in { 81 if P.HasExtVOP3DPP then 82 def _dpp : VOP3_DPP_Pseudo<OpName, P> { 83 let VOP3P = 1; 84 let PseudoInstr = OpName#"_dpp"; 85 let Constraints = !if(P.UseTiedOutput, "$vdst = $vdst_in", ""); 86 let DisableEncoding = !if(P.UseTiedOutput, "$vdst_in", ""); 87 } 88 } // end SubtargetPredicate = isGFX11Plus 89} 90 91let isCommutable = 1 in { 92defm V_PK_MAD_I16 : VOP3PInst<"v_pk_mad_i16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16_V2I16>>; 93defm V_PK_MAD_U16 : VOP3PInst<"v_pk_mad_u16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16_V2I16>>; 94 95let FPDPRounding = 1 in { 96defm V_PK_FMA_F16 : VOP3PInst<"v_pk_fma_f16", VOP3P_Profile<VOP_V2F16_V2F16_V2F16_V2F16>, any_fma>; 97defm V_PK_ADD_F16 : VOP3PInst<"v_pk_add_f16", VOP3P_Profile<VOP_V2F16_V2F16_V2F16>, any_fadd>; 98defm V_PK_MUL_F16 : VOP3PInst<"v_pk_mul_f16", VOP3P_Profile<VOP_V2F16_V2F16_V2F16>, any_fmul>; 99} // End FPDPRounding = 1 100defm V_PK_MAX_F16 : VOP3PInst<"v_pk_max_f16", VOP3P_Profile<VOP_V2F16_V2F16_V2F16>, fmaxnum_like>; 101defm V_PK_MIN_F16 : VOP3PInst<"v_pk_min_f16", VOP3P_Profile<VOP_V2F16_V2F16_V2F16>, fminnum_like>; 102 103defm V_PK_ADD_U16 : VOP3PInst<"v_pk_add_u16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, add>; 104defm V_PK_ADD_I16 : VOP3PInst<"v_pk_add_i16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>>; 105defm V_PK_MUL_LO_U16 : VOP3PInst<"v_pk_mul_lo_u16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, mul>; 106 107defm V_PK_MIN_I16 : VOP3PInst<"v_pk_min_i16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, smin>; 108defm V_PK_MIN_U16 : VOP3PInst<"v_pk_min_u16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, umin>; 109defm V_PK_MAX_I16 : VOP3PInst<"v_pk_max_i16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, smax>; 110defm V_PK_MAX_U16 : VOP3PInst<"v_pk_max_u16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, umax>; 111} 112 113defm V_PK_SUB_U16 : VOP3PInst<"v_pk_sub_u16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>>; 114defm V_PK_SUB_I16 : VOP3PInst<"v_pk_sub_i16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, sub>; 115 116defm V_PK_LSHLREV_B16 : VOP3PInst<"v_pk_lshlrev_b16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, clshl_rev_16>; 117defm V_PK_ASHRREV_I16 : VOP3PInst<"v_pk_ashrrev_i16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, cashr_rev_16>; 118defm V_PK_LSHRREV_B16 : VOP3PInst<"v_pk_lshrrev_b16", VOP3P_Profile<VOP_V2I16_V2I16_V2I16>, clshr_rev_16>; 119 120 121let SubtargetPredicate = HasVOP3PInsts in { 122 123// Undo sub x, c -> add x, -c canonicalization since c is more likely 124// an inline immediate than -c. 125// The constant will be emitted as a mov, and folded later. 126// TODO: We could directly encode the immediate now 127def : GCNPat< 128 (add (v2i16 (VOP3PMods v2i16:$src0, i32:$src0_modifiers)), NegSubInlineConstV216:$src1), 129 (V_PK_SUB_U16 $src0_modifiers, $src0, SRCMODS.OP_SEL_1, NegSubInlineConstV216:$src1) 130>; 131 132// Integer operations with clamp bit set. 133class VOP3PSatPat<SDPatternOperator pat, Instruction inst> : GCNPat< 134 (pat (v2i16 (VOP3PMods v2i16:$src0, i32:$src0_modifiers)), 135 (v2i16 (VOP3PMods v2i16:$src1, i32:$src1_modifiers))), 136 (inst $src0_modifiers, $src0, $src1_modifiers, $src1, DSTCLAMP.ENABLE) 137>; 138 139def : VOP3PSatPat<uaddsat, V_PK_ADD_U16>; 140def : VOP3PSatPat<saddsat, V_PK_ADD_I16>; 141def : VOP3PSatPat<usubsat, V_PK_SUB_U16>; 142def : VOP3PSatPat<ssubsat, V_PK_SUB_I16>; 143} // End SubtargetPredicate = HasVOP3PInsts 144 145multiclass MadFmaMixPats<SDPatternOperator fma_like, 146 Instruction mixlo_inst, 147 Instruction mixhi_inst> { 148 def : GCNPat < 149 (f16 (fpround (fma_like (f32 (VOP3PMadMixMods f16:$src0, i32:$src0_modifiers)), 150 (f32 (VOP3PMadMixMods f16:$src1, i32:$src1_modifiers)), 151 (f32 (VOP3PMadMixMods f16:$src2, i32:$src2_modifiers))))), 152 (mixlo_inst $src0_modifiers, $src0, 153 $src1_modifiers, $src1, 154 $src2_modifiers, $src2, 155 DSTCLAMP.NONE, 156 (i32 (IMPLICIT_DEF))) 157 >; 158 159 // FIXME: Special case handling for maxhi (especially for clamp) 160 // because dealing with the write to high half of the register is 161 // difficult. 162 def : GCNPat < 163 (build_vector f16:$elt0, (fpround (fma_like (f32 (VOP3PMadMixMods f16:$src0, i32:$src0_modifiers)), 164 (f32 (VOP3PMadMixMods f16:$src1, i32:$src1_modifiers)), 165 (f32 (VOP3PMadMixMods f16:$src2, i32:$src2_modifiers))))), 166 (v2f16 (mixhi_inst $src0_modifiers, $src0, 167 $src1_modifiers, $src1, 168 $src2_modifiers, $src2, 169 DSTCLAMP.NONE, 170 $elt0)) 171 >; 172 173 def : GCNPat < 174 (build_vector 175 f16:$elt0, 176 (AMDGPUclamp (fpround (fma_like (f32 (VOP3PMadMixMods f16:$src0, i32:$src0_modifiers)), 177 (f32 (VOP3PMadMixMods f16:$src1, i32:$src1_modifiers)), 178 (f32 (VOP3PMadMixMods f16:$src2, i32:$src2_modifiers)))))), 179 (v2f16 (mixhi_inst $src0_modifiers, $src0, 180 $src1_modifiers, $src1, 181 $src2_modifiers, $src2, 182 DSTCLAMP.ENABLE, 183 $elt0)) 184 >; 185 186 def : GCNPat < 187 (AMDGPUclamp (build_vector 188 (fpround (fma_like (f32 (VOP3PMadMixMods f16:$lo_src0, i32:$lo_src0_modifiers)), 189 (f32 (VOP3PMadMixMods f16:$lo_src1, i32:$lo_src1_modifiers)), 190 (f32 (VOP3PMadMixMods f16:$lo_src2, i32:$lo_src2_modifiers)))), 191 (fpround (fma_like (f32 (VOP3PMadMixMods f16:$hi_src0, i32:$hi_src0_modifiers)), 192 (f32 (VOP3PMadMixMods f16:$hi_src1, i32:$hi_src1_modifiers)), 193 (f32 (VOP3PMadMixMods f16:$hi_src2, i32:$hi_src2_modifiers)))))), 194 (v2f16 (mixhi_inst $hi_src0_modifiers, $hi_src0, 195 $hi_src1_modifiers, $hi_src1, 196 $hi_src2_modifiers, $hi_src2, 197 DSTCLAMP.ENABLE, 198 (mixlo_inst $lo_src0_modifiers, $lo_src0, 199 $lo_src1_modifiers, $lo_src1, 200 $lo_src2_modifiers, $lo_src2, 201 DSTCLAMP.ENABLE, 202 (i32 (IMPLICIT_DEF))))) 203 >; 204} 205 206let SubtargetPredicate = HasMadMixInsts in { 207 208// These are VOP3a-like opcodes which accept no omod. 209// Size of src arguments (16/32) is controlled by op_sel. 210// For 16-bit src arguments their location (hi/lo) are controlled by op_sel_hi. 211let isCommutable = 1, mayRaiseFPException = 0 in { 212defm V_MAD_MIX_F32 : VOP3_VOP3PInst<"v_mad_mix_f32", VOP3P_Mix_Profile<VOP_F32_F16_F16_F16, VOP3_OPSEL>>; 213 214let FPDPRounding = 1 in { 215// Clamp modifier is applied after conversion to f16. 216defm V_MAD_MIXLO_F16 : VOP3_VOP3PInst<"v_mad_mixlo_f16", VOP3P_Mix_Profile<VOP_F16_F16_F16_F16, VOP3_OPSEL, 1>>; 217 218let ClampLo = 0, ClampHi = 1 in { 219defm V_MAD_MIXHI_F16 : VOP3_VOP3PInst<"v_mad_mixhi_f16", VOP3P_Mix_Profile<VOP_F16_F16_F16_F16, VOP3_OPSEL, 1>>; 220} 221} // End FPDPRounding = 1 222} 223 224defm : MadFmaMixPats<fmad, V_MAD_MIXLO_F16, V_MAD_MIXHI_F16>; 225} // End SubtargetPredicate = HasMadMixInsts 226 227 228// Essentially the same as the mad_mix versions 229let SubtargetPredicate = HasFmaMixInsts in { 230let isCommutable = 1 in { 231defm V_FMA_MIX_F32 : VOP3_VOP3PInst<"v_fma_mix_f32", VOP3P_Mix_Profile<VOP_F32_F16_F16_F16, VOP3_OPSEL>>; 232 233let FPDPRounding = 1 in { 234// Clamp modifier is applied after conversion to f16. 235defm V_FMA_MIXLO_F16 : VOP3_VOP3PInst<"v_fma_mixlo_f16", VOP3P_Mix_Profile<VOP_F16_F16_F16_F16, VOP3_OPSEL, 1>>; 236 237let ClampLo = 0, ClampHi = 1 in { 238defm V_FMA_MIXHI_F16 : VOP3_VOP3PInst<"v_fma_mixhi_f16", VOP3P_Mix_Profile<VOP_F16_F16_F16_F16, VOP3_OPSEL, 1>>; 239} 240} // End FPDPRounding = 1 241} 242 243defm : MadFmaMixPats<fma, V_FMA_MIXLO_F16, V_FMA_MIXHI_F16>; 244} 245 246// Defines patterns that extract signed 4bit from each Idx[0]. 247foreach Idx = [[0,28],[4,24],[8,20],[12,16],[16,12],[20,8],[24,4]] in 248 def ExtractSigned4bit_#Idx[0] : PatFrag<(ops node:$src), 249 (sra (shl node:$src, (i32 Idx[1])), (i32 28))>; 250 251// Defines code pattern that extracts U(unsigned/signed) 4/8bit from FromBitIndex. 252class Extract<int FromBitIndex, int BitMask, bit U>: PatFrag< 253 (ops node:$src), 254 !if (!or (!and (!eq (BitMask, 255), !eq (FromBitIndex, 24)), !eq (FromBitIndex, 28)), // last element 255 !if (U, (srl node:$src, (i32 FromBitIndex)), (sra node:$src, (i32 FromBitIndex))), 256 !if (!eq (FromBitIndex, 0), // first element 257 !if (U, (and node:$src, (i32 BitMask)), 258 !if (!eq (BitMask, 15), (!cast<PatFrag>("ExtractSigned4bit_"#FromBitIndex) node:$src), 259 (sext_inreg node:$src, i8))), 260 !if (U, (and (srl node:$src, (i32 FromBitIndex)), (i32 BitMask)), 261 !if (!eq (BitMask, 15), (!cast<PatFrag>("ExtractSigned4bit_"#FromBitIndex) node:$src), 262 (sext_inreg (srl node:$src, (i32 FromBitIndex)), i8)))))>; 263 264 265foreach Type = ["I", "U"] in 266 foreach Index = 0-3 in { 267 // Defines patterns that extract each Index'ed 8bit from an unsigned 268 // 32bit scalar value; 269 def Type#Index#"_8bit" : Extract<!shl(Index, 3), 255, !eq (Type, "U")>; 270 271 // Defines multiplication patterns where the multiplication is happening on each 272 // Index'ed 8bit of a 32bit scalar value. 273 274 def Mul#Type#_Elt#Index : PatFrag< 275 (ops node:$src0, node:$src1), 276 (!cast<HasOneUseBinOp>(!if (!eq (Type, "I"), AMDGPUmul_i24_oneuse, AMDGPUmul_u24_oneuse)) 277 (!cast<Extract>(Type#Index#"_8bit") node:$src0), 278 (!cast<Extract>(Type#Index#"_8bit") node:$src1))>; 279 } 280 281// Different variants of dot8 patterns cause a huge increase in the compile time. 282// Define non-associative/commutative add/mul to prevent permutation in the dot8 283// pattern. 284def NonACAdd : SDNode<"ISD::ADD" , SDTIntBinOp>; 285def NonACAdd_oneuse : HasOneUseBinOp<NonACAdd>; 286 287def NonACAMDGPUmul_u24 : SDNode<"AMDGPUISD::MUL_U24" , SDTIntBinOp>; 288def NonACAMDGPUmul_u24_oneuse : HasOneUseBinOp<NonACAMDGPUmul_u24>; 289 290def NonACAMDGPUmul_i24 : SDNode<"AMDGPUISD::MUL_I24" , SDTIntBinOp>; 291def NonACAMDGPUmul_i24_oneuse : HasOneUseBinOp<NonACAMDGPUmul_i24>; 292 293foreach Type = ["I", "U"] in 294 foreach Index = 0-7 in { 295 // Defines patterns that extract each Index'ed 4bit from an unsigned 296 // 32bit scalar value; 297 def Type#Index#"_4bit" : Extract<!shl(Index, 2), 15, !eq (Type, "U")>; 298 299 // Defines multiplication patterns where the multiplication is happening on each 300 // Index'ed 8bit of a 32bit scalar value. 301 def Mul#Type#Index#"_4bit" : PatFrag< 302 (ops node:$src0, node:$src1), 303 (!cast<HasOneUseBinOp>(!if (!eq (Type, "I"), NonACAMDGPUmul_i24_oneuse, NonACAMDGPUmul_u24_oneuse)) 304 (!cast<Extract>(Type#Index#"_4bit") node:$src0), 305 (!cast<Extract>(Type#Index#"_4bit") node:$src1))>; 306 } 307 308class UDot2Pat<Instruction Inst> : GCNPat < 309 (add (add_oneuse (AMDGPUmul_u24_oneuse (srl i32:$src0, (i32 16)), 310 (srl i32:$src1, (i32 16))), i32:$src2), 311 (AMDGPUmul_u24_oneuse (and i32:$src0, (i32 65535)), 312 (and i32:$src1, (i32 65535))) 313 ), 314 (Inst (i32 8), $src0, (i32 8), $src1, (i32 8), $src2, (i1 0))> { 315 let SubtargetPredicate = !cast<VOP_Pseudo>(Inst).SubtargetPredicate; 316} 317 318class SDot2Pat<Instruction Inst> : GCNPat < 319 (add (add_oneuse (AMDGPUmul_i24_oneuse (sra i32:$src0, (i32 16)), 320 (sra i32:$src1, (i32 16))), i32:$src2), 321 (AMDGPUmul_i24_oneuse (sext_inreg i32:$src0, i16), 322 (sext_inreg i32:$src1, i16))), 323 (Inst (i32 8), $src0, (i32 8), $src1, (i32 8), $src2, (i1 0))> { 324 let SubtargetPredicate = !cast<VOP_Pseudo>(Inst).SubtargetPredicate; 325} 326 327let IsDOT = 1 in { 328let SubtargetPredicate = HasDot2Insts in { 329 330defm V_DOT2_I32_I16 : VOP3PInst<"v_dot2_i32_i16", 331 VOP3P_Profile<VOP_I32_V2I16_V2I16_I32>, int_amdgcn_sdot2, 1>; 332defm V_DOT2_U32_U16 : VOP3PInst<"v_dot2_u32_u16", 333 VOP3P_Profile<VOP_I32_V2I16_V2I16_I32>, int_amdgcn_udot2, 1>; 334 335} // End SubtargetPredicate = HasDot2Insts 336 337let SubtargetPredicate = HasDot7Insts in { 338 339defm V_DOT2_F32_F16 : VOP3PInst<"v_dot2_f32_f16", 340 VOP3P_Profile<VOP_F32_V2F16_V2F16_F32, VOP3_REGULAR, /*HasDPP*/ 1>, 341 AMDGPUfdot2, 1/*ExplicitClamp*/>; 342defm V_DOT4_U32_U8 : VOP3PInst<"v_dot4_u32_u8", 343 VOP3P_Profile<VOP_I32_I32_I32_I32, VOP3_PACKED>, int_amdgcn_udot4, 1>; 344defm V_DOT8_U32_U4 : VOP3PInst<"v_dot8_u32_u4", 345 VOP3P_Profile<VOP_I32_I32_I32_I32, VOP3_PACKED>, int_amdgcn_udot8, 1>; 346 347} // End SubtargetPredicate = HasDot7Insts 348 349let SubtargetPredicate = HasDot1Insts in { 350 351defm V_DOT4_I32_I8 : VOP3PInst<"v_dot4_i32_i8", 352 VOP3P_Profile<VOP_I32_I32_I32_I32, VOP3_PACKED>, int_amdgcn_sdot4, 1>; 353defm V_DOT8_I32_I4 : VOP3PInst<"v_dot8_i32_i4", 354 VOP3P_Profile<VOP_I32_I32_I32_I32, VOP3_PACKED>, int_amdgcn_sdot8, 1>; 355 356} // End SubtargetPredicate = HasDot1Insts 357 358let SubtargetPredicate = HasDot8Insts in { 359 360defm V_DOT2_F32_BF16 : VOP3PInst<"v_dot2_f32_bf16", 361 VOP3P_Profile<VOP_F32_V2I16_V2I16_F32, VOP3_REGULAR, /*HasDPP*/ 1>, 362 int_amdgcn_fdot2_f32_bf16, 1>; 363 364} // End SubtargetPredicate = HasDot8Insts 365 366} // End let IsDOT = 1 367 368multiclass VOP3PDOTIUInst <string OpName, SDPatternOperator intrinsic_node> { 369 let IsDOT = 1 in 370 defm NAME : VOP3PInst<OpName, VOP3P_Profile<VOP_I32_I32_I32_I32, VOP3_PACKED>, 371 null_frag, 1>; 372 // Dot-iu instructions consider input as signed if imod neg bits are set. Thus 373 // Dot-iu Intrinsics have extra operands and require separate codegen pattern. 374 def : GCNPat < (intrinsic_node (DotIUVOP3PMods i32:$src0_mods), i32:$src0, 375 (DotIUVOP3PMods i32:$src1_mods), i32:$src1, 376 i32:$src2, (i1 timm:$clamp)), 377 (!cast<Instruction>(NAME) $src0_mods, i32:$src0, 378 $src1_mods, i32:$src1, 379 (i32 8), i32:$src2, i1:$clamp) 380 >; 381} 382 383let SubtargetPredicate = HasDot8Insts in { 384defm V_DOT4_I32_IU8 : VOP3PDOTIUInst<"v_dot4_i32_iu8", int_amdgcn_sudot4>; 385defm V_DOT8_I32_IU4 : VOP3PDOTIUInst<"v_dot8_i32_iu4", int_amdgcn_sudot8>; 386} // End SubtargetPredicate = HasDot8Insts 387 388def : UDot2Pat<V_DOT2_U32_U16>; 389def : SDot2Pat<V_DOT2_I32_I16>; 390 391foreach Type = ["U", "I"] in 392 let SubtargetPredicate = !cast<VOP_Pseudo>("V_DOT4_"#Type#"32_"#Type#8).SubtargetPredicate in 393 def : GCNPat < 394 !cast<dag>(!foldl((i32 i32:$src2), [0, 1, 2, 3], lhs, y, 395 (add_oneuse lhs, (!cast<PatFrag>("Mul"#Type#"_Elt"#y) i32:$src0, i32:$src1)))), 396 (!cast<VOP3P_Pseudo>("V_DOT4_"#Type#"32_"#Type#8) (i32 8), $src0, (i32 8), $src1, (i32 8), $src2, (i1 0))>; 397 398foreach Type = ["U", "I"] in 399 let SubtargetPredicate = !cast<VOP_Pseudo>("V_DOT8_"#Type#"32_"#Type#4).SubtargetPredicate in 400 def : GCNPat < 401 !cast<dag>(!foldl((add_oneuse i32:$src2, (!cast<PatFrag>("Mul"#Type#"0_4bit") i32:$src0, i32:$src1)), 402 [1, 2, 3, 4, 5, 6, 7], lhs, y, 403 (NonACAdd_oneuse lhs, (!cast<PatFrag>("Mul"#Type#y#"_4bit") i32:$src0, i32:$src1)))), 404 (!cast<VOP3P_Pseudo>("V_DOT8_"#Type#"32_"#Type#4) (i32 8), $src0, (i32 8), $src1, (i32 8), $src2, (i1 0))>; 405 406// Different variants of dot8 code-gen dag patterns are not generated through table-gen due to a huge increase 407// in the compile time. Directly handle the pattern generated by the FE here. 408foreach Type = ["U", "I"] in 409 let SubtargetPredicate = !cast<VOP_Pseudo>("V_DOT8_"#Type#"32_"#Type#4).SubtargetPredicate in 410 def : GCNPat < 411 !cast<dag>(!foldl((add_oneuse i32:$src2, (!cast<PatFrag>("Mul"#Type#"0_4bit") i32:$src0, i32:$src1)), 412 [7, 1, 2, 3, 4, 5, 6], lhs, y, 413 (NonACAdd_oneuse lhs, (!cast<PatFrag>("Mul"#Type#y#"_4bit") i32:$src0, i32:$src1)))), 414 (!cast<VOP3P_Pseudo>("V_DOT8_"#Type#"32_"#Type#4) (i32 8), $src0, (i32 8), $src1, (i32 8), $src2, (i1 0))>; 415 416def ADst_32 : VOPDstOperand<AGPR_32>; 417def ADst_64 : VOPDstOperand<AReg_64>; 418def ADst_128 : VOPDstOperand<AReg_128>; 419def ADst_256 : VOPDstOperand<AReg_256>; 420def ADst_512 : VOPDstOperand<AReg_512>; 421def ADst_1024 : VOPDstOperand<AReg_1024>; 422def VDst_64 : VOPDstOperand<VReg_64>; 423def VDst_128 : VOPDstOperand<VReg_128>; 424def VDst_256 : VOPDstOperand<VReg_256>; 425def VDst_512 : VOPDstOperand<VReg_512>; 426def VDst_1024 : VOPDstOperand<VReg_1024>; 427 428def VOPProfileAccRead : VOP3P_Profile<VOP_I32_I32, VOP3_MAI> { 429 let Src0RC64 = ARegSrc_32; 430} 431 432def VOPProfileAccWrite : VOP3P_Profile<VOP_I32_I32, VOP3_MAI> { 433 let DstRC = ADst_32; 434 let Src0RC64 = VCSrc_b32; 435} 436 437class VOPProfileMAI<VOPProfile P, RegisterOperand _SrcRC, RegisterOperand _DstRC, 438 RegisterOperand SrcABRC = AVSrc_32> 439 : VOP3P_Profile<P, VOP3_MAI> { 440 let DstRC = _DstRC; 441 let Src0RC64 = SrcABRC; 442 let Src1RC64 = SrcABRC; 443 let Src2RC64 = _SrcRC; 444 let HasOpSel = 0; 445 let HasClamp = 0; 446 let HasIntClamp = 0; 447 let HasOMod = 0; 448 let HasModifiers = 0; 449 let Asm64 = "$vdst, $src0, $src1, $src2$cbsz$abid$blgp"; 450 let AsmVOP3DPPBase = Asm64; 451 let Ins64 = (ins Src0RC64:$src0, Src1RC64:$src1, Src2RC64:$src2, cbsz:$cbsz, abid:$abid, blgp:$blgp); 452 let InsVOP3Base = Ins64; 453 // Dst and SrcC cannot partially overlap if SrcC/Dst is bigger than 4 VGPRs. 454 // We then create two versions of the instruction: with tied dst and src2 455 // and with the earlyclobber flag on the dst. This is stricter than the 456 // actual HW restriction. In particular earlyclobber also affects src0 and 457 // src1 allocation which is not required. 458 bit NoDstOverlap = !gt(DstVT.Size, 128); 459} 460 461class VOPProfileSMFMAC<VOPProfile P, RegisterOperand _DstRC, 462 RegisterOperand _SrcARC, RegisterOperand _SrcBRC> 463 : VOPProfileMAI<P, _DstRC, _DstRC, _SrcARC> { 464 let Src1RC64 = _SrcBRC; 465 let Src2VT = DstVT; 466 let Asm64 = " $vdst, $src0, $src1, $idx$cbsz$abid"; 467 let Outs64 = (outs DstRC:$vdst); 468 let Ins64 = (ins Src0RC64:$src0, Src1RC64:$src1, VRegSrc_32:$idx, cbsz:$cbsz, abid:$abid, Src2RC64:$src2); 469} 470 471def VOPProfileMAI_F32_F32_X4 : VOPProfileMAI<VOP_V4F32_F32_F32_V4F32, AISrc_128_f32, ADst_128>; 472def VOPProfileMAI_F32_F32_X16 : VOPProfileMAI<VOP_V16F32_F32_F32_V16F32, AISrc_512_f32, ADst_512>; 473def VOPProfileMAI_F32_F32_X32 : VOPProfileMAI<VOP_V32F32_F32_F32_V32F32, AISrc_1024_f32, ADst_1024>; 474def VOPProfileMAI_I32_I32_X4 : VOPProfileMAI<VOP_V4I32_I32_I32_V4I32, AISrc_128_b32, ADst_128>; 475def VOPProfileMAI_I32_I32_X16 : VOPProfileMAI<VOP_V16I32_I32_I32_V16I32, AISrc_512_b32, ADst_512>; 476def VOPProfileMAI_I32_I32_X32 : VOPProfileMAI<VOP_V32I32_I32_I32_V32I32, AISrc_1024_b32, ADst_1024>; 477def VOPProfileMAI_F32_V2I16_X4 : VOPProfileMAI<VOP_V4F32_V2I16_V2I16_V4F32, AISrc_128_b32, ADst_128>; 478def VOPProfileMAI_F32_V2I16_X16 : VOPProfileMAI<VOP_V16F32_V2I16_V2I16_V16F32, AISrc_512_b32, ADst_512>; 479def VOPProfileMAI_F32_V2I16_X32 : VOPProfileMAI<VOP_V32F32_V2I16_V2I16_V32F32, AISrc_1024_b32, ADst_1024>; 480def VOPProfileMAI_F32_V4F16_X4 : VOPProfileMAI<VOP_V4F32_V4F16_V4F16_V4F32, AISrc_128_b32, ADst_128, AVSrc_64>; 481def VOPProfileMAI_F32_V4F16_X16 : VOPProfileMAI<VOP_V16F32_V4F16_V4F16_V16F32, AISrc_512_b32, ADst_512, AVSrc_64>; 482def VOPProfileMAI_F32_V4F16_X32 : VOPProfileMAI<VOP_V32F32_V4F16_V4F16_V32F32, AISrc_1024_b32, ADst_1024, AVSrc_64>; 483def VOPProfileMAI_F32_V4I16_X4 : VOPProfileMAI<VOP_V4F32_V4I16_V4I16_V4F32, AISrc_128_b32, ADst_128, AVSrc_64>; 484def VOPProfileMAI_F32_V4I16_X16 : VOPProfileMAI<VOP_V16F32_V4I16_V4I16_V16F32, AISrc_512_b32, ADst_512, AVSrc_64>; 485def VOPProfileMAI_F32_V4I16_X32 : VOPProfileMAI<VOP_V32F32_V4I16_V4I16_V32F32, AISrc_1024_b32, ADst_1024, AVSrc_64>; 486def VOPProfileMAI_F64_16X16X4F64 : VOPProfileMAI<VOP_V4F64_F64_F64_V4F64, AISrc_256_f64, ADst_256, AVSrc_64>; 487def VOPProfileMAI_F64_4X4X4F64 : VOPProfileMAI<VOP_F64_F64_F64_F64, AISrc_64_f64, ADst_64, AVSrc_64>; 488def VOPProfileMAI_I32_I64_X16 : VOPProfileMAI<VOP_V4I32_I64_I64_V4I32, AISrc_128_b32, ADst_128, AVSrc_64>; 489def VOPProfileMAI_I32_I64_X32 : VOPProfileMAI<VOP_V16I32_I64_I64_V16I32, AISrc_512_b32, ADst_512, AVSrc_64>; 490def VOPProfileMAI_F32_V2F32_X16 : VOPProfileMAI<VOP_V4F32_V2F32_V2F32_V4F32, AISrc_128_b32, ADst_128, AVSrc_64>; 491def VOPProfileMAI_F32_V2F32_X32 : VOPProfileMAI<VOP_V16F32_V2F32_V2F32_V16F32, AISrc_512_b32, ADst_512, AVSrc_64>; 492 493def VOPProfileMAI_F32_F32_X4_VCD : VOPProfileMAI<VOP_V4F32_F32_F32_V4F32, VISrc_128_f32, VDst_128>; 494def VOPProfileMAI_F32_F32_X16_VCD : VOPProfileMAI<VOP_V16F32_F32_F32_V16F32, VISrc_512_f32, VDst_512>; 495def VOPProfileMAI_F32_F32_X32_VCD : VOPProfileMAI<VOP_V32F32_F32_F32_V32F32, VISrc_1024_f32, VDst_1024>; 496def VOPProfileMAI_I32_I32_X4_VCD : VOPProfileMAI<VOP_V4I32_I32_I32_V4I32, VISrc_128_b32, VDst_128>; 497def VOPProfileMAI_I32_I32_X16_VCD : VOPProfileMAI<VOP_V16I32_I32_I32_V16I32, VISrc_512_b32, VDst_512>; 498def VOPProfileMAI_I32_I32_X32_VCD : VOPProfileMAI<VOP_V32I32_I32_I32_V32I32, VISrc_1024_b32, VDst_1024>; 499def VOPProfileMAI_F32_V2I16_X4_VCD : VOPProfileMAI<VOP_V4F32_V2I16_V2I16_V4F32, VISrc_128_b32, VDst_128>; 500def VOPProfileMAI_F32_V2I16_X16_VCD : VOPProfileMAI<VOP_V16F32_V2I16_V2I16_V16F32, VISrc_512_b32, VDst_512>; 501def VOPProfileMAI_F32_V2I16_X32_VCD : VOPProfileMAI<VOP_V32F32_V2I16_V2I16_V32F32, VISrc_1024_b32, VDst_1024>; 502def VOPProfileMAI_F32_V4F16_X4_VCD : VOPProfileMAI<VOP_V4F32_V4F16_V4F16_V4F32, VISrc_128_b32, VDst_128, AVSrc_64>; 503def VOPProfileMAI_F32_V4F16_X16_VCD : VOPProfileMAI<VOP_V16F32_V4F16_V4F16_V16F32, VISrc_512_b32, VDst_512, AVSrc_64>; 504def VOPProfileMAI_F32_V4F16_X32_VCD : VOPProfileMAI<VOP_V32F32_V4F16_V4F16_V32F32, VISrc_1024_b32, VDst_1024, AVSrc_64>; 505def VOPProfileMAI_F32_V4I16_X4_VCD : VOPProfileMAI<VOP_V4F32_V4I16_V4I16_V4F32, VISrc_128_b32, VDst_128, AVSrc_64>; 506def VOPProfileMAI_F32_V4I16_X16_VCD : VOPProfileMAI<VOP_V16F32_V4I16_V4I16_V16F32, VISrc_512_b32, VDst_512, AVSrc_64>; 507def VOPProfileMAI_F32_V4I16_X32_VCD : VOPProfileMAI<VOP_V32F32_V4I16_V4I16_V32F32, VISrc_1024_b32, VDst_1024, AVSrc_64>; 508def VOPProfileMAI_F64_16X16X4F64_VCD : VOPProfileMAI<VOP_V4F64_F64_F64_V4F64, VISrc_256_f64, VDst_256, AVSrc_64>; 509def VOPProfileMAI_F64_4X4X4F64_VCD : VOPProfileMAI<VOP_F64_F64_F64_F64, VISrc_64_f64, VDst_64, AVSrc_64>; 510def VOPProfileMAI_I32_I64_X16_VCD : VOPProfileMAI<VOP_V4I32_I64_I64_V4I32, VISrc_128_b32, VDst_128, AVSrc_64>; 511def VOPProfileMAI_I32_I64_X32_VCD : VOPProfileMAI<VOP_V16I32_I64_I64_V16I32, VISrc_512_b32, VDst_512, AVSrc_64>; 512def VOPProfileMAI_F32_V2F32_X16_VCD : VOPProfileMAI<VOP_V4F32_V2F32_V2F32_V4F32, VISrc_128_b32, VDst_128, AVSrc_64>; 513def VOPProfileMAI_F32_V2F32_X32_VCD : VOPProfileMAI<VOP_V16F32_V2F32_V2F32_V16F32, VISrc_512_b32, VDst_512, AVSrc_64>; 514 515def VOPProfileSMFMAC_F32_16X16X32_F16 : VOPProfileSMFMAC<VOP_V4F32_V4F16_V8F16_I32, AVDst_128, AVSrc_64, AVSrc_128>; 516def VOPProfileSMFMAC_F32_32X32X16_F16 : VOPProfileSMFMAC<VOP_V16F32_V4F16_V8F16_I32, AVDst_512, AVSrc_64, AVSrc_128>; 517def VOPProfileSMFMAC_F32_16X16X32_I16 : VOPProfileSMFMAC<VOP_V4F32_V4I16_V8I16_I32, AVDst_128, AVSrc_64, AVSrc_128>; 518def VOPProfileSMFMAC_F32_32X32X16_I16 : VOPProfileSMFMAC<VOP_V16F32_V4I16_V8I16_I32, AVDst_512, AVSrc_64, AVSrc_128>; 519def VOPProfileSMFMAC_I32_16X16X64_I8 : VOPProfileSMFMAC<VOP_V4I32_V2I32_V4I32_I32, AVDst_128, AVSrc_64, AVSrc_128>; 520def VOPProfileSMFMAC_I32_32X32X32_I8 : VOPProfileSMFMAC<VOP_V16I32_V2I32_V4I32_I32, AVDst_512, AVSrc_64, AVSrc_128>; 521 522class MFMATable <bit is_mac, string Name> { 523 bit IsMac = is_mac; 524 string FMAOp = Name; 525} 526 527class MAIFrag<SDPatternOperator Op, code pred> : PatFrag < 528 (ops node:$src0, node:$src1, node:$src2, node:$cbsz, node:$abid, node:$blgp), 529 (Op $src0, $src1, $src2, $cbsz, $abid, $blgp), 530 pred 531>; 532 533let GISelPredicateCode = [{ return MF.getInfo<SIMachineFunctionInfo>()->mayNeedAGPRs(); }] in 534class AgprMAIFrag<SDPatternOperator Op> : 535 MAIFrag<Op, [{ return MF->getInfo<SIMachineFunctionInfo>()->mayNeedAGPRs(); }]>; 536 537let GISelPredicateCode = [{ return !MF.getInfo<SIMachineFunctionInfo>()->mayNeedAGPRs(); }] in 538class VgprMAIFrag<SDPatternOperator Op> : 539 MAIFrag<Op, [{ return !MF->getInfo<SIMachineFunctionInfo>()->mayNeedAGPRs(); }]>; 540 541let Predicates = [HasMAIInsts] in { 542 543let isAsCheapAsAMove = 1, isReMaterializable = 1 in { 544 defm V_ACCVGPR_READ_B32 : VOP3Inst<"v_accvgpr_read_b32", VOPProfileAccRead>; 545 let isMoveImm = 1 in { 546 defm V_ACCVGPR_WRITE_B32 : VOP3Inst<"v_accvgpr_write_b32", VOPProfileAccWrite>; 547 } // End isMoveImm = 1 548} // End isAsCheapAsAMove = 1, isReMaterializable = 1 549 550class MAIInst<string OpName, VOPProfile P, SDPatternOperator node> 551 : VOP3InstBase<OpName, P, node> { 552 Instruction Opcode = !cast<Instruction>(NAME); 553 bit is_dgemm = 0; 554 bit is_gfx940_xdl = 0; 555} 556 557multiclass MAIInst<string OpName, string P, SDPatternOperator node, 558 bit NoDstOverlap = !cast<VOPProfileMAI>("VOPProfileMAI_" # P).NoDstOverlap> { 559 let isConvergent = 1, mayRaiseFPException = 0, ReadsModeReg = 1 in { 560 // FP32 denorm mode is respected, rounding mode is not. Exceptions are not supported. 561 let Constraints = !if(NoDstOverlap, "@earlyclobber $vdst", "") in { 562 def _e64 : MAIInst<OpName, !cast<VOPProfileMAI>("VOPProfileMAI_" # P), 563 !if(NoDstOverlap, null_frag, AgprMAIFrag<node>)>, 564 MFMATable<0, NAME # "_e64">; 565 566 let SubtargetPredicate = isGFX90APlus, Mnemonic = OpName in 567 def _vgprcd_e64 : MAIInst<OpName # "_vgprcd", !cast<VOPProfileMAI>("VOPProfileMAI_" # P # "_VCD"), 568 !if(NoDstOverlap, null_frag, VgprMAIFrag<node>)>, 569 MFMATable<0, NAME # "_vgprcd_e64">; 570 } 571 572 foreach _ = BoolToList<NoDstOverlap>.ret in { 573 let Constraints = !if(NoDstOverlap, "$vdst = $src2", ""), 574 isConvertibleToThreeAddress = NoDstOverlap, 575 Mnemonic = OpName in { 576 def "_mac_e64" : MAIInst<OpName # "_mac", !cast<VOPProfileMAI>("VOPProfileMAI_" # P), AgprMAIFrag<node>>, 577 MFMATable<1, NAME # "_e64">; 578 579 let SubtargetPredicate = isGFX90APlus in 580 def _mac_vgprcd_e64 : MAIInst<OpName # "_mac_vgprcd", !cast<VOPProfileMAI>("VOPProfileMAI_" # P # "_VCD"), 581 VgprMAIFrag<node>>, 582 MFMATable<1, NAME # "_vgprcd_e64">; 583 } 584 } 585 } // End isConvergent = 1, mayRaiseFPException = 0, ReadsModeReg = 1 586} 587 588defm V_MFMA_F32_4X4X1F32 : MAIInst<"v_mfma_f32_4x4x1f32", "F32_F32_X4", int_amdgcn_mfma_f32_4x4x1f32>; 589defm V_MFMA_F32_16X16X1F32 : MAIInst<"v_mfma_f32_16x16x1f32", "F32_F32_X16", int_amdgcn_mfma_f32_16x16x1f32>; 590defm V_MFMA_F32_16X16X4F32 : MAIInst<"v_mfma_f32_16x16x4f32", "F32_F32_X4", int_amdgcn_mfma_f32_16x16x4f32>; 591defm V_MFMA_F32_32X32X1F32 : MAIInst<"v_mfma_f32_32x32x1f32", "F32_F32_X32", int_amdgcn_mfma_f32_32x32x1f32>; 592defm V_MFMA_F32_32X32X2F32 : MAIInst<"v_mfma_f32_32x32x2f32", "F32_F32_X16", int_amdgcn_mfma_f32_32x32x2f32>; 593 594let is_gfx940_xdl = 1 in { 595defm V_MFMA_F32_4X4X4F16 : MAIInst<"v_mfma_f32_4x4x4f16", "F32_V4F16_X4", int_amdgcn_mfma_f32_4x4x4f16>; 596defm V_MFMA_I32_4X4X4I8 : MAIInst<"v_mfma_i32_4x4x4i8", "I32_I32_X4", int_amdgcn_mfma_i32_4x4x4i8>; 597defm V_MFMA_F32_16X16X4F16 : MAIInst<"v_mfma_f32_16x16x4f16", "F32_V4F16_X16", int_amdgcn_mfma_f32_16x16x4f16>; 598defm V_MFMA_F32_16X16X16F16 : MAIInst<"v_mfma_f32_16x16x16f16", "F32_V4F16_X4", int_amdgcn_mfma_f32_16x16x16f16>; 599defm V_MFMA_I32_16X16X4I8 : MAIInst<"v_mfma_i32_16x16x4i8", "I32_I32_X16", int_amdgcn_mfma_i32_16x16x4i8>; 600defm V_MFMA_F32_32X32X4F16 : MAIInst<"v_mfma_f32_32x32x4f16", "F32_V4F16_X32", int_amdgcn_mfma_f32_32x32x4f16>; 601defm V_MFMA_F32_32X32X8F16 : MAIInst<"v_mfma_f32_32x32x8f16", "F32_V4F16_X16", int_amdgcn_mfma_f32_32x32x8f16>; 602defm V_MFMA_I32_32X32X4I8 : MAIInst<"v_mfma_i32_32x32x4i8", "I32_I32_X32", int_amdgcn_mfma_i32_32x32x4i8>; 603} 604 605let Predicates = [isGFX908orGFX90A] in { 606defm V_MFMA_I32_16X16X16I8 : MAIInst<"v_mfma_i32_16x16x16i8", "I32_I32_X4", int_amdgcn_mfma_i32_16x16x16i8>; 607defm V_MFMA_I32_32X32X8I8 : MAIInst<"v_mfma_i32_32x32x8i8", "I32_I32_X16", int_amdgcn_mfma_i32_32x32x8i8>; 608defm V_MFMA_F32_4X4X2BF16 : MAIInst<"v_mfma_f32_4x4x2bf16", "F32_V2I16_X4", int_amdgcn_mfma_f32_4x4x2bf16>; 609defm V_MFMA_F32_16X16X2BF16 : MAIInst<"v_mfma_f32_16x16x2bf16", "F32_V2I16_X16", int_amdgcn_mfma_f32_16x16x2bf16>; 610defm V_MFMA_F32_16X16X8BF16 : MAIInst<"v_mfma_f32_16x16x8bf16", "F32_V2I16_X4", int_amdgcn_mfma_f32_16x16x8bf16>; 611defm V_MFMA_F32_32X32X2BF16 : MAIInst<"v_mfma_f32_32x32x2bf16", "F32_V2I16_X32", int_amdgcn_mfma_f32_32x32x2bf16>; 612defm V_MFMA_F32_32X32X4BF16 : MAIInst<"v_mfma_f32_32x32x4bf16", "F32_V2I16_X16", int_amdgcn_mfma_f32_32x32x4bf16>; 613} 614 615} // End SubtargetPredicate = HasMAIInsts 616 617let Predicates = [isGFX90APlus] in { 618 let is_gfx940_xdl = 1 in { 619 defm V_MFMA_F32_32X32X4BF16_1K : MAIInst<"v_mfma_f32_32x32x4bf16_1k", "F32_V4I16_X32", int_amdgcn_mfma_f32_32x32x4bf16_1k>; 620 defm V_MFMA_F32_16X16X4BF16_1K : MAIInst<"v_mfma_f32_16x16x4bf16_1k", "F32_V4I16_X16", int_amdgcn_mfma_f32_16x16x4bf16_1k>; 621 defm V_MFMA_F32_4X4X4BF16_1K : MAIInst<"v_mfma_f32_4x4x4bf16_1k", "F32_V4I16_X4", int_amdgcn_mfma_f32_4x4x4bf16_1k>; 622 defm V_MFMA_F32_32X32X8BF16_1K : MAIInst<"v_mfma_f32_32x32x8bf16_1k", "F32_V4I16_X16", int_amdgcn_mfma_f32_32x32x8bf16_1k>; 623 defm V_MFMA_F32_16X16X16BF16_1K : MAIInst<"v_mfma_f32_16x16x16bf16_1k", "F32_V4I16_X4", int_amdgcn_mfma_f32_16x16x16bf16_1k>; 624 } 625 626 let is_dgemm = 1 in { 627 defm V_MFMA_F64_16X16X4F64 : MAIInst<"v_mfma_f64_16x16x4f64", "F64_16X16X4F64", int_amdgcn_mfma_f64_16x16x4f64>; 628 defm V_MFMA_F64_4X4X4F64 : MAIInst<"v_mfma_f64_4x4x4f64", "F64_4X4X4F64", int_amdgcn_mfma_f64_4x4x4f64>; 629 } 630} // End Predicates = [isGFX90APlus] 631 632let Predicates = [isGFX940Plus], is_gfx940_xdl = 1 in { 633 defm V_MFMA_I32_32X32X16I8 : MAIInst<"v_mfma_i32_32x32x16i8", "I32_I64_X32", int_amdgcn_mfma_i32_32x32x16_i8>; 634 defm V_MFMA_I32_16X16X32I8 : MAIInst<"v_mfma_i32_16x16x32i8", "I32_I64_X16", int_amdgcn_mfma_i32_16x16x32_i8>; 635 defm V_MFMA_F32_16X16X8XF32 : MAIInst<"v_mfma_f32_16x16x8xf32", "F32_V2F32_X16", int_amdgcn_mfma_f32_16x16x8_xf32>; 636 defm V_MFMA_F32_32X32X4XF32 : MAIInst<"v_mfma_f32_32x32x4xf32", "F32_V2F32_X32", int_amdgcn_mfma_f32_32x32x4_xf32>; 637} // End Predicates = [isGFX940Plus], is_gfx940_xdl = 1 638 639multiclass SMFMACInst<string OpName, string P, SDPatternOperator node> { 640 let Constraints = "$vdst = $src2", DisableEncoding = "$src2", 641 isConvergent = 1, mayRaiseFPException = 0, ReadsModeReg = 1, is_gfx940_xdl = 1 in { 642 def _e64 : MAIInst<OpName, !cast<VOPProfileSMFMAC>("VOPProfileSMFMAC_" # P), node>; 643 } 644} 645 646let SubtargetPredicate = isGFX940Plus in { 647defm V_SMFMAC_F32_16X16X32_F16 : SMFMACInst<"v_smfmac_f32_16x16x32_f16", "F32_16X16X32_F16", int_amdgcn_smfmac_f32_16x16x32_f16>; 648defm V_SMFMAC_F32_32X32X16_F16 : SMFMACInst<"v_smfmac_f32_32x32x16_f16", "F32_32X32X16_F16", int_amdgcn_smfmac_f32_32x32x16_f16>; 649defm V_SMFMAC_F32_16X16X32_BF16 : SMFMACInst<"v_smfmac_f32_16x16x32_bf16", "F32_16X16X32_I16", int_amdgcn_smfmac_f32_16x16x32_bf16>; 650defm V_SMFMAC_F32_32X32X16_BF16 : SMFMACInst<"v_smfmac_f32_32x32x16_bf16", "F32_32X32X16_I16", int_amdgcn_smfmac_f32_32x32x16_bf16>; 651defm V_SMFMAC_I32_16X16X64_I8 : SMFMACInst<"v_smfmac_i32_16x16x64_i8", "I32_16X16X64_I8", int_amdgcn_smfmac_i32_16x16x64_i8>; 652defm V_SMFMAC_I32_32X32X32_I8 : SMFMACInst<"v_smfmac_i32_32x32x32_i8", "I32_32X32X32_I8", int_amdgcn_smfmac_i32_32x32x32_i8>; 653} 654 655def MAIInstInfoTable : GenericTable { 656 let FilterClass = "MAIInst"; 657 let CppTypeName = "MAIInstInfo"; 658 let Fields = [ 659 "Opcode", "is_dgemm", "is_gfx940_xdl" 660 ]; 661 662 let PrimaryKey = ["Opcode"]; 663 let PrimaryKeyName = "getMAIInstInfoHelper"; 664} 665 666let SubtargetPredicate = HasPackedFP32Ops, isCommutable = 1 in { 667 defm V_PK_FMA_F32 : VOP3PInst<"v_pk_fma_f32", VOP3P_Profile<VOP_V2F32_V2F32_V2F32_V2F32, VOP3_PACKED>, any_fma>; 668 defm V_PK_MUL_F32 : VOP3PInst<"v_pk_mul_f32", VOP3P_Profile<VOP_V2F32_V2F32_V2F32, VOP3_PACKED>, any_fmul>; 669 defm V_PK_ADD_F32 : VOP3PInst<"v_pk_add_f32", VOP3P_Profile<VOP_V2F32_V2F32_V2F32, VOP3_PACKED>, any_fadd>; 670 defm V_PK_MOV_B32 : VOP3PInst<"v_pk_mov_b32", VOP3P_Profile<VOP_V2I32_V2I32_V2I32, VOP3_PACKED>>; 671} // End SubtargetPredicate = HasPackedFP32Ops, isCommutable = 1 672 673def : MnemonicAlias<"v_accvgpr_read", "v_accvgpr_read_b32">; 674def : MnemonicAlias<"v_accvgpr_write", "v_accvgpr_write_b32">; 675 676//===----------------------------------------------------------------------===// 677// Begin Real Encodings 678//===----------------------------------------------------------------------===// 679 680class VOP3P_DPP16<bits<7> op, VOP_DPP_Pseudo ps, int subtarget, 681 string opName = ps.OpName> 682 : VOP3P_DPP<op, opName, ps.Pfl, 1>, SIMCInstr<ps.PseudoInstr, subtarget> { 683 let hasSideEffects = ps.hasSideEffects; 684 let Defs = ps.Defs; 685 let SchedRW = ps.SchedRW; 686 let Uses = ps.Uses; 687 let AssemblerPredicate = HasDPP16; 688 let SubtargetPredicate = HasDPP16; 689 let OtherPredicates = ps.OtherPredicates; 690} 691 692class VOP3P_DPP8_Base<bits<7> op, VOP_Pseudo ps, string opName = ps.OpName> 693 : VOP3P_DPP8<op, opName, ps.Pfl> { 694 let hasSideEffects = ps.hasSideEffects; 695 let Defs = ps.Defs; 696 let SchedRW = ps.SchedRW; 697 let Uses = ps.Uses; 698 let OtherPredicates = ps.OtherPredicates; 699} 700 701//===----------------------------------------------------------------------===// 702// GFX11. 703//===----------------------------------------------------------------------===// 704 705let AssemblerPredicate = isGFX11Plus, 706 DecoderNamespace = "GFX11" in { 707 708 multiclass VOP3P_Real_gfx11<bits<7> op, string backing_ps_name = NAME, 709 string asmName = !cast<VOP3P_Pseudo>(NAME).Mnemonic> { 710 def _gfx11 : VOP3P_Real<!cast<VOP3P_Pseudo>(backing_ps_name), 711 SIEncodingFamily.GFX11, asmName>, 712 VOP3Pe_gfx11<op, !cast<VOP3P_Pseudo>(backing_ps_name).Pfl>; 713 } 714 715 multiclass VOP3P_Real_dpp_gfx11<bits<7> op, string backing_ps_name = NAME, 716 string asmName = !cast<VOP3P_Pseudo>(NAME).Mnemonic> { 717 defvar ps = !cast<VOP3P_Pseudo>(backing_ps_name); 718 def _dpp_gfx11 719 : VOP3P_DPP16<op, !cast<VOP_DPP_Pseudo>(backing_ps_name #"_dpp"), 720 SIEncodingFamily.GFX11> { 721 let AsmString = asmName #ps.Pfl.AsmVOP3DPP16; 722 let DecoderNamespace = "DPPGFX11"; 723 } 724 } 725 726 multiclass VOP3P_Real_dpp8_gfx11<bits<7> op, string backing_ps_name = NAME, 727 string asmName = !cast<VOP3P_Pseudo>(NAME).Mnemonic> { 728 defvar ps = !cast<VOP3P_Pseudo>(backing_ps_name); 729 def _dpp8_gfx11 : VOP3P_DPP8_Base<op, ps> { 730 let AsmString = asmName #ps.Pfl.AsmVOP3DPP8; 731 let DecoderNamespace = "DPP8GFX11"; 732 } 733 } 734 735 multiclass VOP3P_Realtriple_gfx11<bits<7> op, string backing_ps_name = NAME, 736 string asmName = !cast<VOP3P_Pseudo>(NAME).Mnemonic> 737 : VOP3P_Real_gfx11<op, backing_ps_name, asmName>, 738 VOP3P_Real_dpp_gfx11<op, backing_ps_name, asmName>, 739 VOP3P_Real_dpp8_gfx11<op, backing_ps_name, asmName>; 740} // End AssemblerPredicate = isGFX11Plus, DecoderNamespace = "GFX11" 741 742defm V_DOT4_I32_IU8 : VOP3P_Real_gfx11 <0x16>; 743defm V_DOT8_I32_IU4 : VOP3P_Real_gfx11 <0x18>; 744defm V_DOT2_F32_BF16 : VOP3P_Real_gfx11 <0x1a>; 745 746//===----------------------------------------------------------------------===// 747// GFX8 (VI) 748//===----------------------------------------------------------------------===// 749 750multiclass VOP3P_Real_vi<bits<7> op> { 751 def _vi : VOP3P_Real<!cast<VOP3_Pseudo>(NAME), SIEncodingFamily.VI>, 752 VOP3Pe <op, !cast<VOP3_Pseudo>(NAME).Pfl> { 753 let AssemblerPredicate = HasVOP3PInsts; 754 let DecoderNamespace = "GFX8"; 755 let VOP3P = 1; 756 } 757} 758 759multiclass VOP3P_Real_MAI<bits<7> op> { 760 def _vi : VOP3P_Real<!cast<VOP3_Pseudo>(NAME#"_e64"), SIEncodingFamily.VI>, 761 VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME#"_e64").Pfl, ?> { 762 let AssemblerPredicate = HasMAIInsts; 763 let DecoderNamespace = "GFX8"; 764 let Inst{14} = ?; // op_sel_hi(2) 765 let Inst{59} = ?; // op_sel_hi(0) 766 let Inst{60} = ?; // op_sel_hi(1) 767 } 768} 769 770let Constraints = "" in { 771multiclass VOP3P_Real_MFMA_gfx90a<bits<7> op> { 772 let SubtargetPredicate = isGFX90AOnly, 773 AssemblerPredicate = isGFX90AOnly, DecoderNamespace = "GFX90A" in { 774 def _gfx90a_acd : VOP3P_Real<!cast<VOP3_Pseudo>(NAME#"_e64"), SIEncodingFamily.GFX90A>, 775 VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME#"_e64").Pfl, 1>; 776 777 def _gfx90a_vcd : VOP3P_Real<!cast<VOP3_Pseudo>(NAME # "_vgprcd" # "_e64"), SIEncodingFamily.GFX90A>, 778 VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME # "_vgprcd" # "_e64").Pfl, 0>; 779 } // End AssemblerPredicate = isGFX90AOnly, DecoderNamespace = "GFX90A" 780} 781} 782 783multiclass VOP3P_Real_MFMA_gfx940_aliases<string NameFrom, string NameTo, string Op, 784 VOP3_Pseudo PS_ACD = !cast<VOP3_Pseudo>(Op # "_e64"), 785 VOP3_Pseudo PS_VCD = !cast<VOP3_Pseudo>(Op # "_vgprcd" # "_e64"), 786 VOPProfile Pfl_ACD = PS_ACD.Pfl, 787 VOPProfile Pfl_VCD = PS_VCD.Pfl> { 788 let Predicates = [isGFX940Plus] in { 789 foreach _ = BoolToList<!ne(NameFrom, NameTo)>.ret in { 790 def : InstAlias <NameTo # " " # PS_ACD.AsmOperands, 791 (!cast<VOP3P_Real>(Op # "_gfx940_acd") Pfl_ACD.DstRC:$vdst, 792 Pfl_ACD.Src0RC64:$src0, Pfl_ACD.Src1RC64:$src1, Pfl_ACD.Src2RC64:$src2, 793 cbsz:$cbsz, abid:$abid, blgp:$blgp)>, PredicateControl; 794 def : InstAlias <NameTo # " " # PS_VCD.AsmOperands, 795 (!cast<VOP3P_Real>(Op # "_gfx940_vcd") Pfl_VCD.DstRC:$vdst, 796 Pfl_VCD.Src0RC64:$src0, Pfl_VCD.Src1RC64:$src1, Pfl_VCD.Src2RC64:$src2, 797 cbsz:$cbsz, abid:$abid, blgp:$blgp)>, PredicateControl; 798 } 799 } // End Predicates = [isGFX940Plus] 800} 801 802multiclass VOP3P_Real_MFMA_gfx940<bits<7> op, string Name = !cast<VOP3_Pseudo>(NAME#"_e64").Mnemonic, 803 VOP3_Pseudo PS_ACD = !cast<VOP3_Pseudo>(NAME # "_e64"), 804 VOP3_Pseudo PS_VCD = !cast<VOP3_Pseudo>(NAME # "_vgprcd" # "_e64")> { 805 let SubtargetPredicate = isGFX940Plus, 806 AssemblerPredicate = isGFX940Plus, DecoderNamespace = "GFX9", 807 AsmString = Name # PS_ACD.AsmOperands, Constraints = "" in { 808 def _gfx940_acd : VOP3P_Real<PS_ACD, SIEncodingFamily.GFX940>, 809 VOP3Pe_MAI <op, PS_ACD.Pfl, 1>; 810 811 def _gfx940_vcd : VOP3P_Real<PS_VCD, SIEncodingFamily.GFX940>, 812 VOP3Pe_MAI <op, PS_VCD.Pfl, 0>; 813 } // End AssemblerPredicate = isGFX940Plus, DecoderNamespace = "GFX9" 814 815 defm : VOP3P_Real_MFMA_gfx940_aliases<Name, PS_ACD.Mnemonic, NAME>; 816 817 foreach _ = BoolToList<!ne(!subst("_1k", "", PS_ACD.Mnemonic), PS_ACD.Mnemonic)>.ret in 818 defm : VOP3P_Real_MFMA_gfx940_aliases<Name, !subst("_1k", "", PS_ACD.Mnemonic), NAME>; 819} 820 821multiclass VOP3P_Real_MFMA<bits<7> op, string GFX940Name = !cast<VOP3_Pseudo>(NAME#"_e64").Mnemonic> : 822 VOP3P_Real_MFMA_gfx90a <op>, 823 VOP3P_Real_MFMA_gfx940 <op, GFX940Name> { 824 def _vi : VOP3P_Real<!cast<VOP3_Pseudo>(NAME#"_e64"), SIEncodingFamily.VI>, 825 VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME#"_e64").Pfl, ?> { 826 let AssemblerPredicate = HasMAIInsts; 827 let DecoderNamespace = "GFX8"; 828 let Constraints = ""; 829 } 830} 831 832multiclass VOP3P_Real_SMFMAC<bits<7> op, string alias> { 833 def _gfx940 : VOP3P_Real<!cast<VOP3_Pseudo>(NAME#"_e64"), SIEncodingFamily.VI>, 834 VOP3Pe_SMFMAC <op> { 835 let AssemblerPredicate = isGFX940Plus; 836 let DecoderNamespace = "GFX8"; 837 } 838 def : MnemonicAlias<alias, !cast<VOP3_Pseudo>(NAME#"_e64").Mnemonic>; 839} 840 841defm V_PK_MAD_I16 : VOP3P_Real_vi <0x00>; 842defm V_PK_MUL_LO_U16 : VOP3P_Real_vi <0x01>; 843defm V_PK_ADD_I16 : VOP3P_Real_vi <0x02>; 844defm V_PK_SUB_I16 : VOP3P_Real_vi <0x03>; 845defm V_PK_LSHLREV_B16 : VOP3P_Real_vi <0x04>; 846defm V_PK_LSHRREV_B16 : VOP3P_Real_vi <0x05>; 847defm V_PK_ASHRREV_I16 : VOP3P_Real_vi <0x06>; 848defm V_PK_MAX_I16 : VOP3P_Real_vi <0x07>; 849defm V_PK_MIN_I16 : VOP3P_Real_vi <0x08>; 850defm V_PK_MAD_U16 : VOP3P_Real_vi <0x09>; 851 852defm V_PK_ADD_U16 : VOP3P_Real_vi <0x0a>; 853defm V_PK_SUB_U16 : VOP3P_Real_vi <0x0b>; 854defm V_PK_MAX_U16 : VOP3P_Real_vi <0x0c>; 855defm V_PK_MIN_U16 : VOP3P_Real_vi <0x0d>; 856defm V_PK_FMA_F16 : VOP3P_Real_vi <0x0e>; 857defm V_PK_ADD_F16 : VOP3P_Real_vi <0x0f>; 858defm V_PK_MUL_F16 : VOP3P_Real_vi <0x10>; 859defm V_PK_MIN_F16 : VOP3P_Real_vi <0x11>; 860defm V_PK_MAX_F16 : VOP3P_Real_vi <0x12>; 861 862 863let SubtargetPredicate = HasMadMixInsts in { 864defm V_MAD_MIX_F32 : VOP3P_Real_vi <0x20>; 865defm V_MAD_MIXLO_F16 : VOP3P_Real_vi <0x21>; 866defm V_MAD_MIXHI_F16 : VOP3P_Real_vi <0x22>; 867} 868 869let SubtargetPredicate = HasFmaMixInsts in { 870let DecoderNamespace = "GFX9_DL" in { 871// The mad_mix instructions were renamed and their behaviors changed, 872// but the opcode stayed the same so we need to put these in a 873// different DecoderNamespace to avoid the ambiguity. 874defm V_FMA_MIX_F32 : VOP3P_Real_vi <0x20>; 875defm V_FMA_MIXLO_F16 : VOP3P_Real_vi <0x21>; 876defm V_FMA_MIXHI_F16 : VOP3P_Real_vi <0x22>; 877} 878} 879 880 881let SubtargetPredicate = HasDot2Insts in { 882 883defm V_DOT2_I32_I16 : VOP3P_Real_vi <0x26>; 884defm V_DOT2_U32_U16 : VOP3P_Real_vi <0x27>; 885 886} // End SubtargetPredicate = HasDot2Insts 887 888let SubtargetPredicate = HasDot7Insts in { 889 890defm V_DOT2_F32_F16 : VOP3P_Real_vi <0x23>; 891defm V_DOT4_U32_U8 : VOP3P_Real_vi <0x29>; 892defm V_DOT8_U32_U4 : VOP3P_Real_vi <0x2b>; 893 894} // End SubtargetPredicate = HasDot7Insts 895 896let SubtargetPredicate = HasDot1Insts in { 897 898defm V_DOT4_I32_I8 : VOP3P_Real_vi <0x28>; 899defm V_DOT8_I32_I4 : VOP3P_Real_vi <0x2a>; 900 901} // End SubtargetPredicate = HasDot1Insts 902 903let SubtargetPredicate = HasMAIInsts in { 904 905defm V_ACCVGPR_READ_B32 : VOP3P_Real_MAI <0x58>; 906defm V_ACCVGPR_WRITE_B32 : VOP3P_Real_MAI <0x59>; 907defm V_MFMA_F32_32X32X1F32 : VOP3P_Real_MFMA <0x40, "v_mfma_f32_32x32x1_2b_f32">; 908defm V_MFMA_F32_16X16X1F32 : VOP3P_Real_MFMA <0x41, "v_mfma_f32_16x16x1_4b_f32">; 909defm V_MFMA_F32_4X4X1F32 : VOP3P_Real_MFMA <0x42, "v_mfma_f32_4x4x1_16b_f32">; 910defm V_MFMA_F32_32X32X2F32 : VOP3P_Real_MFMA <0x44, "v_mfma_f32_32x32x2_f32">; 911defm V_MFMA_F32_16X16X4F32 : VOP3P_Real_MFMA <0x45, "v_mfma_f32_16x16x4_f32">; 912defm V_MFMA_F32_32X32X4F16 : VOP3P_Real_MFMA <0x48, "v_mfma_f32_32x32x4_2b_f16">; 913defm V_MFMA_F32_16X16X4F16 : VOP3P_Real_MFMA <0x49, "v_mfma_f32_16x16x4_4b_f16">; 914defm V_MFMA_F32_4X4X4F16 : VOP3P_Real_MFMA <0x4a, "v_mfma_f32_4x4x4_16b_f16">; 915defm V_MFMA_F32_32X32X8F16 : VOP3P_Real_MFMA <0x4c, "v_mfma_f32_32x32x8_f16">; 916defm V_MFMA_F32_16X16X16F16 : VOP3P_Real_MFMA <0x4d, "v_mfma_f32_16x16x16_f16">; 917defm V_MFMA_I32_32X32X4I8 : VOP3P_Real_MFMA <0x50, "v_mfma_i32_32x32x4_2b_i8">; 918defm V_MFMA_I32_16X16X4I8 : VOP3P_Real_MFMA <0x51, "v_mfma_i32_16x16x4_4b_i8">; 919defm V_MFMA_I32_4X4X4I8 : VOP3P_Real_MFMA <0x52, "v_mfma_i32_4x4x4_16b_i8">; 920 921let SubtargetPredicate = isGFX908orGFX90A in { 922defm V_MFMA_I32_16X16X16I8 : VOP3P_Real_MFMA <0x55>; 923defm V_MFMA_I32_32X32X8I8 : VOP3P_Real_MFMA <0x54>; 924defm V_MFMA_F32_32X32X2BF16 : VOP3P_Real_MFMA <0x68>; 925defm V_MFMA_F32_16X16X2BF16 : VOP3P_Real_MFMA <0x69>; 926defm V_MFMA_F32_4X4X2BF16 : VOP3P_Real_MFMA <0x6b>; 927defm V_MFMA_F32_32X32X4BF16 : VOP3P_Real_MFMA <0x6c>; 928defm V_MFMA_F32_16X16X8BF16 : VOP3P_Real_MFMA <0x6d>; 929} 930 931} // End SubtargetPredicate = HasMAIInsts 932 933defm V_MFMA_F32_32X32X4BF16_1K : VOP3P_Real_MFMA_gfx90a <0x63>; 934defm V_MFMA_F32_16X16X4BF16_1K : VOP3P_Real_MFMA_gfx90a <0x64>; 935defm V_MFMA_F32_4X4X4BF16_1K : VOP3P_Real_MFMA_gfx90a <0x65>; 936defm V_MFMA_F32_32X32X8BF16_1K : VOP3P_Real_MFMA_gfx90a <0x66>; 937defm V_MFMA_F32_16X16X16BF16_1K : VOP3P_Real_MFMA_gfx90a <0x67>; 938defm V_MFMA_F64_16X16X4F64 : VOP3P_Real_MFMA_gfx90a <0x6e>; 939defm V_MFMA_F64_4X4X4F64 : VOP3P_Real_MFMA_gfx90a <0x6f>; 940 941defm V_MFMA_I32_32X32X16I8 : VOP3P_Real_MFMA_gfx940 <0x56, "v_mfma_i32_32x32x16_i8">; 942defm V_MFMA_I32_16X16X32I8 : VOP3P_Real_MFMA_gfx940 <0x57, "v_mfma_i32_16x16x32_i8">; 943defm V_MFMA_F32_16X16X8XF32 : VOP3P_Real_MFMA_gfx940 <0x3e, "v_mfma_f32_16x16x8_xf32">; 944defm V_MFMA_F32_32X32X4XF32 : VOP3P_Real_MFMA_gfx940 <0x3f, "v_mfma_f32_32x32x4_xf32">; 945 946defm V_MFMA_F32_32X32X4BF16_1K : VOP3P_Real_MFMA_gfx940 <0x5d, "v_mfma_f32_32x32x4_2b_bf16">; 947defm V_MFMA_F32_16X16X4BF16_1K : VOP3P_Real_MFMA_gfx940 <0x5e, "v_mfma_f32_16x16x4_4b_bf16">; 948defm V_MFMA_F32_4X4X4BF16_1K : VOP3P_Real_MFMA_gfx940 <0x5f, "v_mfma_f32_4x4x4_16b_bf16">; 949defm V_MFMA_F32_32X32X8BF16_1K : VOP3P_Real_MFMA_gfx940 <0x60, "v_mfma_f32_32x32x8_bf16">; 950defm V_MFMA_F32_16X16X16BF16_1K : VOP3P_Real_MFMA_gfx940 <0x61, "v_mfma_f32_16x16x16_bf16">; 951 952defm V_MFMA_F64_16X16X4F64 : VOP3P_Real_MFMA_gfx940 <0x6e, "v_mfma_f64_16x16x4_f64">; 953defm V_MFMA_F64_4X4X4F64 : VOP3P_Real_MFMA_gfx940 <0x6f, "v_mfma_f64_4x4x4_4b_f64">; 954 955defm V_SMFMAC_F32_16X16X32_F16 : VOP3P_Real_SMFMAC <0x62, "v_smfmac_f32_16x16x32f16">; 956defm V_SMFMAC_F32_32X32X16_F16 : VOP3P_Real_SMFMAC <0x64, "v_smfmac_f32_32x32x16f16">; 957defm V_SMFMAC_F32_16X16X32_BF16 : VOP3P_Real_SMFMAC <0x66, "v_smfmac_f32_16x16x32bf16">; 958defm V_SMFMAC_F32_32X32X16_BF16 : VOP3P_Real_SMFMAC <0x68, "v_smfmac_f32_32x32x16bf16">; 959defm V_SMFMAC_I32_16X16X64_I8 : VOP3P_Real_SMFMAC <0x6a, "v_smfmac_i32_16x16x64i8">; 960defm V_SMFMAC_I32_32X32X32_I8 : VOP3P_Real_SMFMAC <0x6c, "v_smfmac_i32_32x32x32i8">; 961 962let SubtargetPredicate = HasPackedFP32Ops in { 963 defm V_PK_FMA_F32 : VOP3P_Real_vi <0x30>; 964 defm V_PK_MUL_F32 : VOP3P_Real_vi <0x31>; 965 defm V_PK_ADD_F32 : VOP3P_Real_vi <0x32>; 966 defm V_PK_MOV_B32 : VOP3P_Real_vi <0x33>; 967} // End SubtargetPredicate = HasPackedFP32Ops 968 969//===----------------------------------------------------------------------===// 970// GFX10. 971//===----------------------------------------------------------------------===// 972 973let AssemblerPredicate = isGFX10Only, DecoderNamespace = "GFX10", VOP3P = 1 in { 974 multiclass VOP3P_Real_gfx10<bits<7> op> { 975 def _gfx10 : VOP3P_Real<!cast<VOP3P_Pseudo>(NAME), SIEncodingFamily.GFX10>, 976 VOP3Pe_gfx10 <op, !cast<VOP3P_Pseudo>(NAME).Pfl>; 977 } 978} // End AssemblerPredicate = isGFX10Only, DecoderNamespace = "GFX10", VOP3P = 1 979 980multiclass VOP3P_Real_gfx10_gfx11<bits<7> op> 981 : VOP3P_Real_gfx10<op>, VOP3P_Real_gfx11<op>; 982 983multiclass VOP3P_Real_gfx10_gfx11_Triple<bits<7> op> 984 : VOP3P_Real_gfx10<op>, VOP3P_Realtriple_gfx11<op>; 985 986defm V_PK_MAD_I16 : VOP3P_Real_gfx10_gfx11<0x00>; 987defm V_PK_MUL_LO_U16 : VOP3P_Real_gfx10_gfx11<0x01>; 988defm V_PK_ADD_I16 : VOP3P_Real_gfx10_gfx11<0x02>; 989defm V_PK_SUB_I16 : VOP3P_Real_gfx10_gfx11<0x03>; 990defm V_PK_LSHLREV_B16 : VOP3P_Real_gfx10_gfx11<0x04>; 991defm V_PK_LSHRREV_B16 : VOP3P_Real_gfx10_gfx11<0x05>; 992defm V_PK_ASHRREV_I16 : VOP3P_Real_gfx10_gfx11<0x06>; 993defm V_PK_MAX_I16 : VOP3P_Real_gfx10_gfx11<0x07>; 994defm V_PK_MIN_I16 : VOP3P_Real_gfx10_gfx11<0x08>; 995defm V_PK_MAD_U16 : VOP3P_Real_gfx10_gfx11<0x09>; 996defm V_PK_ADD_U16 : VOP3P_Real_gfx10_gfx11<0x0a>; 997defm V_PK_SUB_U16 : VOP3P_Real_gfx10_gfx11<0x0b>; 998defm V_PK_MAX_U16 : VOP3P_Real_gfx10_gfx11<0x0c>; 999defm V_PK_MIN_U16 : VOP3P_Real_gfx10_gfx11<0x0d>; 1000defm V_PK_FMA_F16 : VOP3P_Real_gfx10_gfx11<0x0e>; 1001defm V_PK_ADD_F16 : VOP3P_Real_gfx10_gfx11<0x0f>; 1002defm V_PK_MUL_F16 : VOP3P_Real_gfx10_gfx11<0x10>; 1003defm V_PK_MIN_F16 : VOP3P_Real_gfx10_gfx11<0x11>; 1004defm V_PK_MAX_F16 : VOP3P_Real_gfx10_gfx11<0x12>; 1005defm V_FMA_MIX_F32 : VOP3P_Real_gfx10_gfx11_Triple <0x20>; 1006defm V_FMA_MIXLO_F16 : VOP3P_Real_gfx10_gfx11_Triple <0x21>; 1007defm V_FMA_MIXHI_F16 : VOP3P_Real_gfx10_gfx11_Triple <0x22>; 1008 1009let SubtargetPredicate = HasDot2Insts in { 1010 1011defm V_DOT2_I32_I16 : VOP3P_Real_gfx10 <0x14>; 1012defm V_DOT2_U32_U16 : VOP3P_Real_gfx10 <0x15>; 1013 1014} // End SubtargetPredicate = HasDot2Insts 1015 1016let SubtargetPredicate = HasDot7Insts in { 1017 1018defm V_DOT2_F32_F16 : VOP3P_Real_gfx10_gfx11_Triple <0x13>; 1019defm V_DOT4_U32_U8 : VOP3P_Real_gfx10_gfx11 <0x17>; 1020defm V_DOT8_U32_U4 : VOP3P_Real_gfx10_gfx11 <0x19>; 1021 1022} // End SubtargetPredicate = HasDot7Insts 1023 1024let SubtargetPredicate = HasDot1Insts in { 1025 1026defm V_DOT4_I32_I8 : VOP3P_Real_gfx10 <0x16>; 1027defm V_DOT8_I32_I4 : VOP3P_Real_gfx10 <0x18>; 1028 1029} // End SubtargetPredicate = HasDot1Insts 1030