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 13// Used for FMA_MIX* and MAD_MIX* insts 14// Their operands are only sort of f16 operands. Depending on 15// op_sel_hi, these may be interpreted as f32. The inline immediate 16// values are really f16 converted to f32, so we treat these as f16 17// operands. 18class VOP3P_Mix_Profile<VOPProfile P, VOP3Features Features = VOP3_REGULAR, 19 bit useTiedOutput = 0> : VOP3_Profile<P, Features> { 20 bit UseTiedOutput = useTiedOutput; 21 22 dag srcs = 23 (ins FP16InputMods:$src0_modifiers, VCSrc_f16:$src0, 24 FP16InputMods:$src1_modifiers, VCSrc_f16:$src1, 25 FP16InputMods:$src2_modifiers, VCSrc_f16:$src2); 26 27 // FIXME: clampmod0 misbehaves with the non-default vdst_in 28 // following it. For now workaround this by requiring clamp 29 // in tied patterns. This should use undef_tied_input, but it 30 // seems underdeveloped and doesn't apply the right register 31 // class constraints. 32 dag mods = !con(!if(UseTiedOutput, (ins clampmod:$clamp, VGPR_32:$vdst_in), 33 (ins clampmod0:$clamp)), 34 (ins op_sel0:$op_sel, op_sel_hi0:$op_sel_hi)); 35 // We use Ins64 because that is the one which populates InOperandList 36 // due to the logic in class VOP3_Pseudo 37 let Ins64 = !con(srcs, mods); 38 let Asm64 = 39 "$vdst, $src0_modifiers, $src1_modifiers, $src2_modifiers$op_sel$op_sel_hi$clamp"; 40} 41 42multiclass VOP3PInst<string OpName, VOPProfile P, 43 SDPatternOperator node = null_frag, bit IsDOT = 0> { 44 def NAME : VOP3P_Pseudo<OpName, P, 45 !if (P.HasModifiers, 46 getVOP3PModPat<P, node, IsDOT, IsDOT>.ret, 47 getVOP3Pat<P, node>.ret)>; 48} 49 50// Non-packed instructions that use the VOP3P encoding. 51// VOP3 neg/abs and VOP3P opsel/opsel_hi modifiers are allowed. 52multiclass VOP3_VOP3PInst<string OpName, VOP3P_Mix_Profile P> { 53 def NAME : VOP3P_Pseudo<OpName, P> { 54 let Constraints = !if(P.UseTiedOutput, "$vdst = $vdst_in", ""); 55 let DisableEncoding = !if(P.UseTiedOutput, "$vdst_in", ""); 56 } 57} 58 59let isCommutable = 1 in { 60defm V_PK_MAD_I16 : VOP3PInst<"v_pk_mad_i16", VOP3_Profile<VOP_V2I16_V2I16_V2I16_V2I16>>; 61defm V_PK_MAD_U16 : VOP3PInst<"v_pk_mad_u16", VOP3_Profile<VOP_V2I16_V2I16_V2I16_V2I16>>; 62 63let FPDPRounding = 1 in { 64defm V_PK_FMA_F16 : VOP3PInst<"v_pk_fma_f16", VOP3_Profile<VOP_V2F16_V2F16_V2F16_V2F16>, any_fma>; 65defm V_PK_ADD_F16 : VOP3PInst<"v_pk_add_f16", VOP3_Profile<VOP_V2F16_V2F16_V2F16>, any_fadd>; 66defm V_PK_MUL_F16 : VOP3PInst<"v_pk_mul_f16", VOP3_Profile<VOP_V2F16_V2F16_V2F16>, any_fmul>; 67} // End FPDPRounding = 1 68defm V_PK_MAX_F16 : VOP3PInst<"v_pk_max_f16", VOP3_Profile<VOP_V2F16_V2F16_V2F16>, fmaxnum_like>; 69defm V_PK_MIN_F16 : VOP3PInst<"v_pk_min_f16", VOP3_Profile<VOP_V2F16_V2F16_V2F16>, fminnum_like>; 70 71defm V_PK_ADD_U16 : VOP3PInst<"v_pk_add_u16", VOP3_Profile<VOP_V2I16_V2I16_V2I16>, add>; 72defm V_PK_ADD_I16 : VOP3PInst<"v_pk_add_i16", VOP3_Profile<VOP_V2I16_V2I16_V2I16>>; 73defm V_PK_MUL_LO_U16 : VOP3PInst<"v_pk_mul_lo_u16", VOP3_Profile<VOP_V2I16_V2I16_V2I16>, mul>; 74 75defm V_PK_MIN_I16 : VOP3PInst<"v_pk_min_i16", VOP3_Profile<VOP_V2I16_V2I16_V2I16>, smin>; 76defm V_PK_MIN_U16 : VOP3PInst<"v_pk_min_u16", VOP3_Profile<VOP_V2I16_V2I16_V2I16>, umin>; 77defm V_PK_MAX_I16 : VOP3PInst<"v_pk_max_i16", VOP3_Profile<VOP_V2I16_V2I16_V2I16>, smax>; 78defm V_PK_MAX_U16 : VOP3PInst<"v_pk_max_u16", VOP3_Profile<VOP_V2I16_V2I16_V2I16>, umax>; 79} 80 81defm V_PK_SUB_U16 : VOP3PInst<"v_pk_sub_u16", VOP3_Profile<VOP_V2I16_V2I16_V2I16>>; 82defm V_PK_SUB_I16 : VOP3PInst<"v_pk_sub_i16", VOP3_Profile<VOP_V2I16_V2I16_V2I16>, sub>; 83 84defm V_PK_LSHLREV_B16 : VOP3PInst<"v_pk_lshlrev_b16", VOP3_Profile<VOP_V2I16_V2I16_V2I16>, clshl_rev_16>; 85defm V_PK_ASHRREV_I16 : VOP3PInst<"v_pk_ashrrev_i16", VOP3_Profile<VOP_V2I16_V2I16_V2I16>, cashr_rev_16>; 86defm V_PK_LSHRREV_B16 : VOP3PInst<"v_pk_lshrrev_b16", VOP3_Profile<VOP_V2I16_V2I16_V2I16>, clshr_rev_16>; 87 88 89let SubtargetPredicate = HasVOP3PInsts in { 90 91// Undo sub x, c -> add x, -c canonicalization since c is more likely 92// an inline immediate than -c. 93// The constant will be emitted as a mov, and folded later. 94// TODO: We could directly encode the immediate now 95def : GCNPat< 96 (add (v2i16 (VOP3PMods v2i16:$src0, i32:$src0_modifiers)), NegSubInlineConstV216:$src1), 97 (V_PK_SUB_U16 $src0_modifiers, $src0, SRCMODS.OP_SEL_1, NegSubInlineConstV216:$src1) 98>; 99 100// Integer operations with clamp bit set. 101class VOP3PSatPat<SDPatternOperator pat, Instruction inst> : GCNPat< 102 (pat (v2i16 (VOP3PMods v2i16:$src0, i32:$src0_modifiers)), 103 (v2i16 (VOP3PMods v2i16:$src1, i32:$src1_modifiers))), 104 (inst $src0_modifiers, $src0, $src1_modifiers, $src1, DSTCLAMP.ENABLE) 105>; 106 107def : VOP3PSatPat<uaddsat, V_PK_ADD_U16>; 108def : VOP3PSatPat<saddsat, V_PK_ADD_I16>; 109def : VOP3PSatPat<usubsat, V_PK_SUB_U16>; 110def : VOP3PSatPat<ssubsat, V_PK_SUB_I16>; 111} // End SubtargetPredicate = HasVOP3PInsts 112 113multiclass MadFmaMixPats<SDPatternOperator fma_like, 114 Instruction mixlo_inst, 115 Instruction mixhi_inst> { 116 def : GCNPat < 117 (f16 (fpround (fma_like (f32 (VOP3PMadMixMods f16:$src0, i32:$src0_modifiers)), 118 (f32 (VOP3PMadMixMods f16:$src1, i32:$src1_modifiers)), 119 (f32 (VOP3PMadMixMods f16:$src2, i32:$src2_modifiers))))), 120 (mixlo_inst $src0_modifiers, $src0, 121 $src1_modifiers, $src1, 122 $src2_modifiers, $src2, 123 DSTCLAMP.NONE, 124 (i32 (IMPLICIT_DEF))) 125 >; 126 127 // FIXME: Special case handling for maxhi (especially for clamp) 128 // because dealing with the write to high half of the register is 129 // difficult. 130 def : GCNPat < 131 (build_vector f16:$elt0, (fpround (fma_like (f32 (VOP3PMadMixMods f16:$src0, i32:$src0_modifiers)), 132 (f32 (VOP3PMadMixMods f16:$src1, i32:$src1_modifiers)), 133 (f32 (VOP3PMadMixMods f16:$src2, i32:$src2_modifiers))))), 134 (v2f16 (mixhi_inst $src0_modifiers, $src0, 135 $src1_modifiers, $src1, 136 $src2_modifiers, $src2, 137 DSTCLAMP.NONE, 138 $elt0)) 139 >; 140 141 def : GCNPat < 142 (build_vector 143 f16:$elt0, 144 (AMDGPUclamp (fpround (fma_like (f32 (VOP3PMadMixMods f16:$src0, i32:$src0_modifiers)), 145 (f32 (VOP3PMadMixMods f16:$src1, i32:$src1_modifiers)), 146 (f32 (VOP3PMadMixMods f16:$src2, i32:$src2_modifiers)))))), 147 (v2f16 (mixhi_inst $src0_modifiers, $src0, 148 $src1_modifiers, $src1, 149 $src2_modifiers, $src2, 150 DSTCLAMP.ENABLE, 151 $elt0)) 152 >; 153 154 def : GCNPat < 155 (AMDGPUclamp (build_vector 156 (fpround (fma_like (f32 (VOP3PMadMixMods f16:$lo_src0, i32:$lo_src0_modifiers)), 157 (f32 (VOP3PMadMixMods f16:$lo_src1, i32:$lo_src1_modifiers)), 158 (f32 (VOP3PMadMixMods f16:$lo_src2, i32:$lo_src2_modifiers)))), 159 (fpround (fma_like (f32 (VOP3PMadMixMods f16:$hi_src0, i32:$hi_src0_modifiers)), 160 (f32 (VOP3PMadMixMods f16:$hi_src1, i32:$hi_src1_modifiers)), 161 (f32 (VOP3PMadMixMods f16:$hi_src2, i32:$hi_src2_modifiers)))))), 162 (v2f16 (mixhi_inst $hi_src0_modifiers, $hi_src0, 163 $hi_src1_modifiers, $hi_src1, 164 $hi_src2_modifiers, $hi_src2, 165 DSTCLAMP.ENABLE, 166 (mixlo_inst $lo_src0_modifiers, $lo_src0, 167 $lo_src1_modifiers, $lo_src1, 168 $lo_src2_modifiers, $lo_src2, 169 DSTCLAMP.ENABLE, 170 (i32 (IMPLICIT_DEF))))) 171 >; 172} 173 174let SubtargetPredicate = HasMadMixInsts in { 175 176// These are VOP3a-like opcodes which accept no omod. 177// Size of src arguments (16/32) is controlled by op_sel. 178// For 16-bit src arguments their location (hi/lo) are controlled by op_sel_hi. 179let isCommutable = 1, mayRaiseFPException = 0 in { 180defm V_MAD_MIX_F32 : VOP3_VOP3PInst<"v_mad_mix_f32", VOP3P_Mix_Profile<VOP_F32_F16_F16_F16, VOP3_OPSEL>>; 181 182let FPDPRounding = 1 in { 183// Clamp modifier is applied after conversion to f16. 184defm V_MAD_MIXLO_F16 : VOP3_VOP3PInst<"v_mad_mixlo_f16", VOP3P_Mix_Profile<VOP_F16_F16_F16_F16, VOP3_OPSEL, 1>>; 185 186let ClampLo = 0, ClampHi = 1 in { 187defm V_MAD_MIXHI_F16 : VOP3_VOP3PInst<"v_mad_mixhi_f16", VOP3P_Mix_Profile<VOP_F16_F16_F16_F16, VOP3_OPSEL, 1>>; 188} 189} // End FPDPRounding = 1 190} 191 192defm : MadFmaMixPats<fmad, V_MAD_MIXLO_F16, V_MAD_MIXHI_F16>; 193} // End SubtargetPredicate = HasMadMixInsts 194 195 196// Essentially the same as the mad_mix versions 197let SubtargetPredicate = HasFmaMixInsts in { 198let isCommutable = 1 in { 199defm V_FMA_MIX_F32 : VOP3_VOP3PInst<"v_fma_mix_f32", VOP3P_Mix_Profile<VOP_F32_F16_F16_F16, VOP3_OPSEL>>; 200 201let FPDPRounding = 1 in { 202// Clamp modifier is applied after conversion to f16. 203defm V_FMA_MIXLO_F16 : VOP3_VOP3PInst<"v_fma_mixlo_f16", VOP3P_Mix_Profile<VOP_F16_F16_F16_F16, VOP3_OPSEL, 1>>; 204 205let ClampLo = 0, ClampHi = 1 in { 206defm V_FMA_MIXHI_F16 : VOP3_VOP3PInst<"v_fma_mixhi_f16", VOP3P_Mix_Profile<VOP_F16_F16_F16_F16, VOP3_OPSEL, 1>>; 207} 208} // End FPDPRounding = 1 209} 210 211defm : MadFmaMixPats<fma, V_FMA_MIXLO_F16, V_FMA_MIXHI_F16>; 212} 213 214// Defines patterns that extract signed 4bit from each Idx[0]. 215foreach Idx = [[0,28],[4,24],[8,20],[12,16],[16,12],[20,8],[24,4]] in 216 def ExtractSigned4bit_#Idx[0] : PatFrag<(ops node:$src), 217 (sra (shl node:$src, (i32 Idx[1])), (i32 28))>; 218 219// Defines code pattern that extracts U(unsigned/signed) 4/8bit from FromBitIndex. 220class Extract<int FromBitIndex, int BitMask, bit U>: PatFrag< 221 (ops node:$src), 222 !if (!or (!and (!eq (BitMask, 255), !eq (FromBitIndex, 24)), !eq (FromBitIndex, 28)), // last element 223 !if (U, (srl node:$src, (i32 FromBitIndex)), (sra node:$src, (i32 FromBitIndex))), 224 !if (!eq (FromBitIndex, 0), // first element 225 !if (U, (and node:$src, (i32 BitMask)), 226 !if (!eq (BitMask, 15), (!cast<PatFrag>("ExtractSigned4bit_"#FromBitIndex) node:$src), 227 (sext_inreg node:$src, i8))), 228 !if (U, (and (srl node:$src, (i32 FromBitIndex)), (i32 BitMask)), 229 !if (!eq (BitMask, 15), (!cast<PatFrag>("ExtractSigned4bit_"#FromBitIndex) node:$src), 230 (sext_inreg (srl node:$src, (i32 FromBitIndex)), i8)))))>; 231 232 233foreach Type = ["I", "U"] in 234 foreach Index = 0-3 in { 235 // Defines patterns that extract each Index'ed 8bit from an unsigned 236 // 32bit scalar value; 237 def Type#Index#"_8bit" : Extract<!shl(Index, 3), 255, !eq (Type, "U")>; 238 239 // Defines multiplication patterns where the multiplication is happening on each 240 // Index'ed 8bit of a 32bit scalar value. 241 242 def Mul#Type#_Elt#Index : PatFrag< 243 (ops node:$src0, node:$src1), 244 (!cast<HasOneUseBinOp>(!if (!eq (Type, "I"), AMDGPUmul_i24_oneuse, AMDGPUmul_u24_oneuse)) 245 (!cast<Extract>(Type#Index#"_8bit") node:$src0), 246 (!cast<Extract>(Type#Index#"_8bit") node:$src1))>; 247 } 248 249// Different variants of dot8 patterns cause a huge increase in the compile time. 250// Define non-associative/commutative add/mul to prevent permutation in the dot8 251// pattern. 252def NonACAdd : SDNode<"ISD::ADD" , SDTIntBinOp>; 253def NonACAdd_oneuse : HasOneUseBinOp<NonACAdd>; 254 255def NonACAMDGPUmul_u24 : SDNode<"AMDGPUISD::MUL_U24" , SDTIntBinOp>; 256def NonACAMDGPUmul_u24_oneuse : HasOneUseBinOp<NonACAMDGPUmul_u24>; 257 258def NonACAMDGPUmul_i24 : SDNode<"AMDGPUISD::MUL_I24" , SDTIntBinOp>; 259def NonACAMDGPUmul_i24_oneuse : HasOneUseBinOp<NonACAMDGPUmul_i24>; 260 261foreach Type = ["I", "U"] in 262 foreach Index = 0-7 in { 263 // Defines patterns that extract each Index'ed 4bit from an unsigned 264 // 32bit scalar value; 265 def Type#Index#"_4bit" : Extract<!shl(Index, 2), 15, !eq (Type, "U")>; 266 267 // Defines multiplication patterns where the multiplication is happening on each 268 // Index'ed 8bit of a 32bit scalar value. 269 def Mul#Type#Index#"_4bit" : PatFrag< 270 (ops node:$src0, node:$src1), 271 (!cast<HasOneUseBinOp>(!if (!eq (Type, "I"), NonACAMDGPUmul_i24_oneuse, NonACAMDGPUmul_u24_oneuse)) 272 (!cast<Extract>(Type#Index#"_4bit") node:$src0), 273 (!cast<Extract>(Type#Index#"_4bit") node:$src1))>; 274 } 275 276class UDot2Pat<Instruction Inst> : GCNPat < 277 (add (add_oneuse (AMDGPUmul_u24_oneuse (srl i32:$src0, (i32 16)), 278 (srl i32:$src1, (i32 16))), i32:$src2), 279 (AMDGPUmul_u24_oneuse (and i32:$src0, (i32 65535)), 280 (and i32:$src1, (i32 65535))) 281 ), 282 (Inst (i32 8), $src0, (i32 8), $src1, (i32 8), $src2, (i1 0))> { 283 let SubtargetPredicate = !cast<VOP_Pseudo>(Inst).SubtargetPredicate; 284} 285 286class SDot2Pat<Instruction Inst> : GCNPat < 287 (add (add_oneuse (AMDGPUmul_i24_oneuse (sra i32:$src0, (i32 16)), 288 (sra i32:$src1, (i32 16))), i32:$src2), 289 (AMDGPUmul_i24_oneuse (sext_inreg i32:$src0, i16), 290 (sext_inreg i32:$src1, i16))), 291 (Inst (i32 8), $src0, (i32 8), $src1, (i32 8), $src2, (i1 0))> { 292 let SubtargetPredicate = !cast<VOP_Pseudo>(Inst).SubtargetPredicate; 293} 294 295let IsDOT = 1 in { 296let SubtargetPredicate = HasDot2Insts in { 297 298defm V_DOT2_I32_I16 : VOP3PInst<"v_dot2_i32_i16", 299 VOP3_Profile<VOP_I32_V2I16_V2I16_I32>, int_amdgcn_sdot2, 1>; 300defm V_DOT2_U32_U16 : VOP3PInst<"v_dot2_u32_u16", 301 VOP3_Profile<VOP_I32_V2I16_V2I16_I32>, int_amdgcn_udot2, 1>; 302 303} // End SubtargetPredicate = HasDot2Insts 304 305let SubtargetPredicate = HasDot7Insts in { 306 307defm V_DOT2_F32_F16 : VOP3PInst<"v_dot2_f32_f16", 308 VOP3_Profile<VOP_F32_V2F16_V2F16_F32>, 309 AMDGPUfdot2, 1/*ExplicitClamp*/>; 310defm V_DOT4_U32_U8 : VOP3PInst<"v_dot4_u32_u8", 311 VOP3_Profile<VOP_I32_I32_I32_I32, VOP3_PACKED>, int_amdgcn_udot4, 1>; 312defm V_DOT8_U32_U4 : VOP3PInst<"v_dot8_u32_u4", 313 VOP3_Profile<VOP_I32_I32_I32_I32, VOP3_PACKED>, int_amdgcn_udot8, 1>; 314 315} // End SubtargetPredicate = HasDot7Insts 316 317let SubtargetPredicate = HasDot1Insts in { 318 319defm V_DOT4_I32_I8 : VOP3PInst<"v_dot4_i32_i8", 320 VOP3_Profile<VOP_I32_I32_I32_I32, VOP3_PACKED>, int_amdgcn_sdot4, 1>; 321defm V_DOT8_I32_I4 : VOP3PInst<"v_dot8_i32_i4", 322 VOP3_Profile<VOP_I32_I32_I32_I32, VOP3_PACKED>, int_amdgcn_sdot8, 1>; 323 324} // End SubtargetPredicate = HasDot1Insts 325} // End let IsDOT = 1 326 327def : UDot2Pat<V_DOT2_U32_U16>; 328def : SDot2Pat<V_DOT2_I32_I16>; 329 330foreach Type = ["U", "I"] in 331 let SubtargetPredicate = !cast<VOP_Pseudo>("V_DOT4_"#Type#"32_"#Type#8).SubtargetPredicate in 332 def : GCNPat < 333 !cast<dag>(!foldl((i32 i32:$src2), [0, 1, 2, 3], lhs, y, 334 (add_oneuse lhs, (!cast<PatFrag>("Mul"#Type#"_Elt"#y) i32:$src0, i32:$src1)))), 335 (!cast<VOP3P_Pseudo>("V_DOT4_"#Type#"32_"#Type#8) (i32 8), $src0, (i32 8), $src1, (i32 8), $src2, (i1 0))>; 336 337foreach Type = ["U", "I"] in 338 let SubtargetPredicate = !cast<VOP_Pseudo>("V_DOT8_"#Type#"32_"#Type#4).SubtargetPredicate in 339 def : GCNPat < 340 !cast<dag>(!foldl((add_oneuse i32:$src2, (!cast<PatFrag>("Mul"#Type#"0_4bit") i32:$src0, i32:$src1)), 341 [1, 2, 3, 4, 5, 6, 7], lhs, y, 342 (NonACAdd_oneuse lhs, (!cast<PatFrag>("Mul"#Type#y#"_4bit") i32:$src0, i32:$src1)))), 343 (!cast<VOP3P_Pseudo>("V_DOT8_"#Type#"32_"#Type#4) (i32 8), $src0, (i32 8), $src1, (i32 8), $src2, (i1 0))>; 344 345// Different variants of dot8 code-gen dag patterns are not generated through table-gen due to a huge increase 346// in the compile time. Directly handle the pattern generated by the FE here. 347foreach Type = ["U", "I"] in 348 let SubtargetPredicate = !cast<VOP_Pseudo>("V_DOT8_"#Type#"32_"#Type#4).SubtargetPredicate in 349 def : GCNPat < 350 !cast<dag>(!foldl((add_oneuse i32:$src2, (!cast<PatFrag>("Mul"#Type#"0_4bit") i32:$src0, i32:$src1)), 351 [7, 1, 2, 3, 4, 5, 6], lhs, y, 352 (NonACAdd_oneuse lhs, (!cast<PatFrag>("Mul"#Type#y#"_4bit") i32:$src0, i32:$src1)))), 353 (!cast<VOP3P_Pseudo>("V_DOT8_"#Type#"32_"#Type#4) (i32 8), $src0, (i32 8), $src1, (i32 8), $src2, (i1 0))>; 354 355def ADst_32 : VOPDstOperand<AGPR_32>; 356def ADst_64 : VOPDstOperand<AReg_64>; 357def ADst_128 : VOPDstOperand<AReg_128>; 358def ADst_256 : VOPDstOperand<AReg_256>; 359def ADst_512 : VOPDstOperand<AReg_512>; 360def ADst_1024 : VOPDstOperand<AReg_1024>; 361def VDst_64 : VOPDstOperand<VReg_64>; 362def VDst_128 : VOPDstOperand<VReg_128>; 363def VDst_256 : VOPDstOperand<VReg_256>; 364def VDst_512 : VOPDstOperand<VReg_512>; 365def VDst_1024 : VOPDstOperand<VReg_1024>; 366 367def VOPProfileAccRead : VOP3_Profile<VOP_I32_I32, VOP3_MAI> { 368 let Src0RC64 = ARegSrc_32; 369} 370 371def VOPProfileAccWrite : VOP3_Profile<VOP_I32_I32, VOP3_MAI> { 372 let DstRC = ADst_32; 373 let Src0RC64 = VCSrc_b32; 374} 375 376class VOPProfileMAI<VOPProfile P, RegisterOperand _SrcRC, RegisterOperand _DstRC, 377 RegisterOperand SrcABRC = AVSrc_32> 378 : VOP3_Profile<P, VOP3_MAI> { 379 let DstRC = _DstRC; 380 let Src0RC64 = SrcABRC; 381 let Src1RC64 = SrcABRC; 382 let Src2RC64 = _SrcRC; 383 let HasOpSel = 0; 384 let HasClamp = 0; 385 let HasIntClamp = 0; 386 let HasOMod = 0; 387 let HasModifiers = 0; 388 let Asm64 = "$vdst, $src0, $src1, $src2$cbsz$abid$blgp"; 389 let Ins64 = (ins Src0RC64:$src0, Src1RC64:$src1, Src2RC64:$src2, cbsz:$cbsz, abid:$abid, blgp:$blgp); 390 // Dst and SrcC cannot partially overlap if SrcC/Dst is bigger than 4 VGPRs. 391 // We then create two versions of the instruction: with tied dst and src2 392 // and with the earlyclobber flag on the dst. This is stricter than the 393 // actual HW restriction. In particular earlyclobber also affects src0 and 394 // src1 allocation which is not required. 395 bit NoDstOverlap = !gt(DstVT.Size, 128); 396} 397 398def VOPProfileMAI_F32_F32_X4 : VOPProfileMAI<VOP_V4F32_F32_F32_V4F32, AISrc_128_f32, ADst_128>; 399def VOPProfileMAI_F32_F32_X16 : VOPProfileMAI<VOP_V16F32_F32_F32_V16F32, AISrc_512_f32, ADst_512>; 400def VOPProfileMAI_F32_F32_X32 : VOPProfileMAI<VOP_V32F32_F32_F32_V32F32, AISrc_1024_f32, ADst_1024>; 401def VOPProfileMAI_I32_I32_X4 : VOPProfileMAI<VOP_V4I32_I32_I32_V4I32, AISrc_128_b32, ADst_128>; 402def VOPProfileMAI_I32_I32_X16 : VOPProfileMAI<VOP_V16I32_I32_I32_V16I32, AISrc_512_b32, ADst_512>; 403def VOPProfileMAI_I32_I32_X32 : VOPProfileMAI<VOP_V32I32_I32_I32_V32I32, AISrc_1024_b32, ADst_1024>; 404def VOPProfileMAI_F32_V2I16_X4 : VOPProfileMAI<VOP_V4F32_V2I16_V2I16_V4F32, AISrc_128_b32, ADst_128>; 405def VOPProfileMAI_F32_V2I16_X16 : VOPProfileMAI<VOP_V16F32_V2I16_V2I16_V16F32, AISrc_512_b32, ADst_512>; 406def VOPProfileMAI_F32_V2I16_X32 : VOPProfileMAI<VOP_V32F32_V2I16_V2I16_V32F32, AISrc_1024_b32, ADst_1024>; 407def VOPProfileMAI_F32_V4F16_X4 : VOPProfileMAI<VOP_V4F32_V4F16_V4F16_V4F32, AISrc_128_b32, ADst_128, AVSrc_64>; 408def VOPProfileMAI_F32_V4F16_X16 : VOPProfileMAI<VOP_V16F32_V4F16_V4F16_V16F32, AISrc_512_b32, ADst_512, AVSrc_64>; 409def VOPProfileMAI_F32_V4F16_X32 : VOPProfileMAI<VOP_V32F32_V4F16_V4F16_V32F32, AISrc_1024_b32, ADst_1024, AVSrc_64>; 410def VOPProfileMAI_F32_V4I16_X4 : VOPProfileMAI<VOP_V4F32_V4I16_V4I16_V4F32, AISrc_128_b32, ADst_128, AVSrc_64>; 411def VOPProfileMAI_F32_V4I16_X16 : VOPProfileMAI<VOP_V16F32_V4I16_V4I16_V16F32, AISrc_512_b32, ADst_512, AVSrc_64>; 412def VOPProfileMAI_F32_V4I16_X32 : VOPProfileMAI<VOP_V32F32_V4I16_V4I16_V32F32, AISrc_1024_b32, ADst_1024, AVSrc_64>; 413def VOPProfileMAI_F64_16X16X4F64 : VOPProfileMAI<VOP_V4F64_F64_F64_V4F64, AISrc_256_f64, ADst_256, AVSrc_64>; 414def VOPProfileMAI_F64_4X4X4F64 : VOPProfileMAI<VOP_F64_F64_F64_F64, AISrc_64_f64, ADst_64, AVSrc_64>; 415 416def VOPProfileMAI_F32_F32_X4_VCD : VOPProfileMAI<VOP_V4F32_F32_F32_V4F32, VISrc_128_f32, VDst_128>; 417def VOPProfileMAI_F32_F32_X16_VCD : VOPProfileMAI<VOP_V16F32_F32_F32_V16F32, VISrc_512_f32, VDst_512>; 418def VOPProfileMAI_F32_F32_X32_VCD : VOPProfileMAI<VOP_V32F32_F32_F32_V32F32, VISrc_1024_f32, VDst_1024>; 419def VOPProfileMAI_I32_I32_X4_VCD : VOPProfileMAI<VOP_V4I32_I32_I32_V4I32, VISrc_128_b32, VDst_128>; 420def VOPProfileMAI_I32_I32_X16_VCD : VOPProfileMAI<VOP_V16I32_I32_I32_V16I32, VISrc_512_b32, VDst_512>; 421def VOPProfileMAI_I32_I32_X32_VCD : VOPProfileMAI<VOP_V32I32_I32_I32_V32I32, VISrc_1024_b32, VDst_1024>; 422def VOPProfileMAI_F32_V2I16_X4_VCD : VOPProfileMAI<VOP_V4F32_V2I16_V2I16_V4F32, VISrc_128_b32, VDst_128>; 423def VOPProfileMAI_F32_V2I16_X16_VCD : VOPProfileMAI<VOP_V16F32_V2I16_V2I16_V16F32, VISrc_512_b32, VDst_512>; 424def VOPProfileMAI_F32_V2I16_X32_VCD : VOPProfileMAI<VOP_V32F32_V2I16_V2I16_V32F32, VISrc_1024_b32, VDst_1024>; 425def VOPProfileMAI_F32_V4F16_X4_VCD : VOPProfileMAI<VOP_V4F32_V4F16_V4F16_V4F32, VISrc_128_b32, VDst_128, AVSrc_64>; 426def VOPProfileMAI_F32_V4F16_X16_VCD : VOPProfileMAI<VOP_V16F32_V4F16_V4F16_V16F32, VISrc_512_b32, VDst_512, AVSrc_64>; 427def VOPProfileMAI_F32_V4F16_X32_VCD : VOPProfileMAI<VOP_V32F32_V4F16_V4F16_V32F32, VISrc_1024_b32, VDst_1024, AVSrc_64>; 428def VOPProfileMAI_F32_V4I16_X4_VCD : VOPProfileMAI<VOP_V4F32_V4I16_V4I16_V4F32, VISrc_128_b32, VDst_128, AVSrc_64>; 429def VOPProfileMAI_F32_V4I16_X16_VCD : VOPProfileMAI<VOP_V16F32_V4I16_V4I16_V16F32, VISrc_512_b32, VDst_512, AVSrc_64>; 430def VOPProfileMAI_F32_V4I16_X32_VCD : VOPProfileMAI<VOP_V32F32_V4I16_V4I16_V32F32, VISrc_1024_b32, VDst_1024, AVSrc_64>; 431def VOPProfileMAI_F64_16X16X4F64_VCD : VOPProfileMAI<VOP_V4F64_F64_F64_V4F64, VISrc_256_f64, VDst_256, AVSrc_64>; 432def VOPProfileMAI_F64_4X4X4F64_VCD : VOPProfileMAI<VOP_F64_F64_F64_F64, VISrc_64_f64, VDst_64, AVSrc_64>; 433 434class MFMATable <bit is_mac, string Name> { 435 bit IsMac = is_mac; 436 string FMAOp = Name; 437} 438 439class MAIFrag<SDPatternOperator Op, code pred> : PatFrag < 440 (ops node:$src0, node:$src1, node:$src2, node:$cbsz, node:$abid, node:$blgp), 441 (Op $src0, $src1, $src2, $cbsz, $abid, $blgp), 442 pred 443>; 444 445let GISelPredicateCode = [{ return MF.getInfo<SIMachineFunctionInfo>()->mayNeedAGPRs(); }] in 446class AgprMAIFrag<SDPatternOperator Op> : 447 MAIFrag<Op, [{ return MF->getInfo<SIMachineFunctionInfo>()->mayNeedAGPRs(); }]>; 448 449let GISelPredicateCode = [{ return !MF.getInfo<SIMachineFunctionInfo>()->mayNeedAGPRs(); }] in 450class VgprMAIFrag<SDPatternOperator Op> : 451 MAIFrag<Op, [{ return !MF->getInfo<SIMachineFunctionInfo>()->mayNeedAGPRs(); }]>; 452 453let Predicates = [HasMAIInsts] in { 454 455let isAsCheapAsAMove = 1, isReMaterializable = 1 in { 456 defm V_ACCVGPR_READ_B32 : VOP3Inst<"v_accvgpr_read_b32", VOPProfileAccRead>; 457 let isMoveImm = 1 in { 458 defm V_ACCVGPR_WRITE_B32 : VOP3Inst<"v_accvgpr_write_b32", VOPProfileAccWrite>; 459 } // End isMoveImm = 1 460} // End isAsCheapAsAMove = 1, isReMaterializable = 1 461 462multiclass MAIInst<string OpName, string P, SDPatternOperator node, 463 bit NoDstOverlap = !cast<VOPProfileMAI>("VOPProfileMAI_" # P).NoDstOverlap> { 464 let isConvergent = 1, mayRaiseFPException = 0, ReadsModeReg = 1 in { 465 // FP32 denorm mode is respected, rounding mode is not. Exceptions are not supported. 466 let Constraints = !if(NoDstOverlap, "@earlyclobber $vdst", "") in { 467 defm "" : VOP3Inst<OpName, !cast<VOPProfileMAI>("VOPProfileMAI_" # P), 468 !if(NoDstOverlap, null_frag, AgprMAIFrag<node>)>, 469 MFMATable<0, NAME # "_e64">; 470 471 let SubtargetPredicate = isGFX90APlus, Mnemonic = OpName in 472 defm _vgprcd : VOP3Inst<OpName # "_vgprcd", !cast<VOPProfileMAI>("VOPProfileMAI_" # P # "_VCD"), 473 !if(NoDstOverlap, null_frag, VgprMAIFrag<node>)>, 474 MFMATable<0, NAME # "_vgprcd_e64">; 475 } 476 477 foreach _ = BoolToList<NoDstOverlap>.ret in { 478 let Constraints = !if(NoDstOverlap, "$vdst = $src2", ""), 479 isConvertibleToThreeAddress = NoDstOverlap, 480 Mnemonic = OpName in { 481 defm "_mac" : VOP3Inst<OpName # "_mac", !cast<VOPProfileMAI>("VOPProfileMAI_" # P), AgprMAIFrag<node>>, 482 MFMATable<1, NAME # "_e64">; 483 484 let SubtargetPredicate = isGFX90APlus in 485 defm _mac_vgprcd : VOP3Inst<OpName # "_mac_vgprcd", !cast<VOPProfileMAI>("VOPProfileMAI_" # P # "_VCD"), 486 VgprMAIFrag<node>>, 487 MFMATable<1, NAME # "_vgprcd_e64">; 488 } 489 } 490 } // End isConvergent = 1, mayRaiseFPException = 0, ReadsModeReg = 1 491} 492 493defm V_MFMA_F32_4X4X1F32 : MAIInst<"v_mfma_f32_4x4x1f32", "F32_F32_X4", int_amdgcn_mfma_f32_4x4x1f32>; 494defm V_MFMA_F32_4X4X4F16 : MAIInst<"v_mfma_f32_4x4x4f16", "F32_V4F16_X4", int_amdgcn_mfma_f32_4x4x4f16>; 495defm V_MFMA_I32_4X4X4I8 : MAIInst<"v_mfma_i32_4x4x4i8", "I32_I32_X4", int_amdgcn_mfma_i32_4x4x4i8>; 496defm V_MFMA_F32_16X16X1F32 : MAIInst<"v_mfma_f32_16x16x1f32", "F32_F32_X16", int_amdgcn_mfma_f32_16x16x1f32>; 497defm V_MFMA_F32_16X16X4F32 : MAIInst<"v_mfma_f32_16x16x4f32", "F32_F32_X4", int_amdgcn_mfma_f32_16x16x4f32>; 498defm V_MFMA_F32_16X16X4F16 : MAIInst<"v_mfma_f32_16x16x4f16", "F32_V4F16_X16", int_amdgcn_mfma_f32_16x16x4f16>; 499defm V_MFMA_F32_16X16X16F16 : MAIInst<"v_mfma_f32_16x16x16f16", "F32_V4F16_X4", int_amdgcn_mfma_f32_16x16x16f16>; 500defm V_MFMA_I32_16X16X4I8 : MAIInst<"v_mfma_i32_16x16x4i8", "I32_I32_X16", int_amdgcn_mfma_i32_16x16x4i8>; 501defm V_MFMA_F32_32X32X1F32 : MAIInst<"v_mfma_f32_32x32x1f32", "F32_F32_X32", int_amdgcn_mfma_f32_32x32x1f32>; 502defm V_MFMA_F32_32X32X2F32 : MAIInst<"v_mfma_f32_32x32x2f32", "F32_F32_X16", int_amdgcn_mfma_f32_32x32x2f32>; 503defm V_MFMA_F32_32X32X4F16 : MAIInst<"v_mfma_f32_32x32x4f16", "F32_V4F16_X32", int_amdgcn_mfma_f32_32x32x4f16>; 504defm V_MFMA_F32_32X32X8F16 : MAIInst<"v_mfma_f32_32x32x8f16", "F32_V4F16_X16", int_amdgcn_mfma_f32_32x32x8f16>; 505defm V_MFMA_I32_32X32X4I8 : MAIInst<"v_mfma_i32_32x32x4i8", "I32_I32_X32", int_amdgcn_mfma_i32_32x32x4i8>; 506defm V_MFMA_I32_16X16X16I8 : MAIInst<"v_mfma_i32_16x16x16i8", "I32_I32_X4", int_amdgcn_mfma_i32_16x16x16i8>; 507defm V_MFMA_I32_32X32X8I8 : MAIInst<"v_mfma_i32_32x32x8i8", "I32_I32_X16", int_amdgcn_mfma_i32_32x32x8i8>; 508defm V_MFMA_F32_4X4X2BF16 : MAIInst<"v_mfma_f32_4x4x2bf16", "F32_V2I16_X4", int_amdgcn_mfma_f32_4x4x2bf16>; 509defm V_MFMA_F32_16X16X2BF16 : MAIInst<"v_mfma_f32_16x16x2bf16", "F32_V2I16_X16", int_amdgcn_mfma_f32_16x16x2bf16>; 510defm V_MFMA_F32_16X16X8BF16 : MAIInst<"v_mfma_f32_16x16x8bf16", "F32_V2I16_X4", int_amdgcn_mfma_f32_16x16x8bf16>; 511defm V_MFMA_F32_32X32X2BF16 : MAIInst<"v_mfma_f32_32x32x2bf16", "F32_V2I16_X32", int_amdgcn_mfma_f32_32x32x2bf16>; 512defm V_MFMA_F32_32X32X4BF16 : MAIInst<"v_mfma_f32_32x32x4bf16", "F32_V2I16_X16", int_amdgcn_mfma_f32_32x32x4bf16>; 513 514} // End SubtargetPredicate = HasMAIInsts 515 516let Predicates = [isGFX90APlus] in { 517 defm V_MFMA_F32_32X32X4BF16_1K : MAIInst<"v_mfma_f32_32x32x4bf16_1k", "F32_V4I16_X32", int_amdgcn_mfma_f32_32x32x4bf16_1k>; 518 defm V_MFMA_F32_16X16X4BF16_1K : MAIInst<"v_mfma_f32_16x16x4bf16_1k", "F32_V4I16_X16", int_amdgcn_mfma_f32_16x16x4bf16_1k>; 519 defm V_MFMA_F32_4X4X4BF16_1K : MAIInst<"v_mfma_f32_4x4x4bf16_1k", "F32_V4I16_X4", int_amdgcn_mfma_f32_4x4x4bf16_1k>; 520 defm V_MFMA_F32_32X32X8BF16_1K : MAIInst<"v_mfma_f32_32x32x8bf16_1k", "F32_V4I16_X16", int_amdgcn_mfma_f32_32x32x8bf16_1k>; 521 defm V_MFMA_F32_16X16X16BF16_1K : MAIInst<"v_mfma_f32_16x16x16bf16_1k", "F32_V4I16_X4", int_amdgcn_mfma_f32_16x16x16bf16_1k>; 522 523 defm V_MFMA_F64_16X16X4F64 : MAIInst<"v_mfma_f64_16x16x4f64", "F64_16X16X4F64", int_amdgcn_mfma_f64_16x16x4f64>; 524 defm V_MFMA_F64_4X4X4F64 : MAIInst<"v_mfma_f64_4x4x4f64", "F64_4X4X4F64", int_amdgcn_mfma_f64_4x4x4f64>; 525} // End Predicates = [isGFX90APlus] 526 527let SubtargetPredicate = HasPackedFP32Ops, isCommutable = 1 in { 528 defm V_PK_FMA_F32 : VOP3PInst<"v_pk_fma_f32", VOP3_Profile<VOP_V2F32_V2F32_V2F32_V2F32, VOP3_PACKED>, any_fma>; 529 defm V_PK_MUL_F32 : VOP3PInst<"v_pk_mul_f32", VOP3_Profile<VOP_V2F32_V2F32_V2F32, VOP3_PACKED>, any_fmul>; 530 defm V_PK_ADD_F32 : VOP3PInst<"v_pk_add_f32", VOP3_Profile<VOP_V2F32_V2F32_V2F32, VOP3_PACKED>, any_fadd>; 531 defm V_PK_MOV_B32 : VOP3PInst<"v_pk_mov_b32", VOP3_Profile<VOP_V2I32_V2I32_V2I32, VOP3_PACKED>>; 532} // End SubtargetPredicate = HasPackedFP32Ops, isCommutable = 1 533 534def : MnemonicAlias<"v_accvgpr_read", "v_accvgpr_read_b32">; 535def : MnemonicAlias<"v_accvgpr_write", "v_accvgpr_write_b32">; 536 537//===----------------------------------------------------------------------===// 538// Begin Real Encodings 539//===----------------------------------------------------------------------===// 540 541//===----------------------------------------------------------------------===// 542// GFX8 (VI) 543//===----------------------------------------------------------------------===// 544 545multiclass VOP3P_Real_vi<bits<7> op> { 546 def _vi : VOP3P_Real<!cast<VOP3_Pseudo>(NAME), SIEncodingFamily.VI>, 547 VOP3Pe <op, !cast<VOP3_Pseudo>(NAME).Pfl> { 548 let AssemblerPredicate = HasVOP3PInsts; 549 let DecoderNamespace = "GFX8"; 550 let VOP3P = 1; 551 } 552} 553 554multiclass VOP3P_Real_MAI<bits<7> op> { 555 def _vi : VOP3P_Real<!cast<VOP3_Pseudo>(NAME#"_e64"), SIEncodingFamily.VI>, 556 VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME#"_e64").Pfl, ?> { 557 let AssemblerPredicate = HasMAIInsts; 558 let DecoderNamespace = "GFX8"; 559 let Inst{14} = ?; // op_sel_hi(2) 560 let Inst{59} = ?; // op_sel_hi(0) 561 let Inst{60} = ?; // op_sel_hi(1) 562 } 563} 564 565let Constraints = "" in { 566multiclass VOP3P_Real_MFMA_gfx90a<bits<7> op> { 567 let SubtargetPredicate = isGFX90AOnly, 568 AssemblerPredicate = isGFX90AOnly, DecoderNamespace = "GFX90A" in { 569 def _gfx90a_acd : VOP3P_Real<!cast<VOP3_Pseudo>(NAME#"_e64"), SIEncodingFamily.GFX90A>, 570 VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME#"_e64").Pfl, 1>; 571 572 def _gfx90a_vcd : VOP3P_Real<!cast<VOP3_Pseudo>(NAME # "_vgprcd" # "_e64"), SIEncodingFamily.GFX90A>, 573 VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME # "_vgprcd" # "_e64").Pfl, 0>; 574 } // End AssemblerPredicate = isGFX90AOnly, DecoderNamespace = "GFX90A" 575} 576} 577 578multiclass VOP3P_Real_MFMA_gfx940_aliases<string NameFrom, string NameTo, string Op, 579 VOP3_Pseudo PS_ACD = !cast<VOP3_Pseudo>(Op # "_e64"), 580 VOP3_Pseudo PS_VCD = !cast<VOP3_Pseudo>(Op # "_vgprcd" # "_e64"), 581 VOPProfile Pfl_ACD = PS_ACD.Pfl, 582 VOPProfile Pfl_VCD = PS_VCD.Pfl> { 583 let Predicates = [isGFX940Plus] in { 584 foreach _ = BoolToList<!ne(NameFrom, NameTo)>.ret in { 585 def : InstAlias <NameTo # " " # PS_ACD.AsmOperands, 586 (!cast<VOP3P_Real>(Op # "_gfx940_acd") Pfl_ACD.DstRC:$vdst, 587 Pfl_ACD.Src0RC64:$src0, Pfl_ACD.Src1RC64:$src1, Pfl_ACD.Src2RC64:$src2, 588 cbsz:$cbsz, abid:$abid, blgp:$blgp)>, PredicateControl; 589 def : InstAlias <NameTo # " " # PS_VCD.AsmOperands, 590 (!cast<VOP3P_Real>(Op # "_gfx940_vcd") Pfl_VCD.DstRC:$vdst, 591 Pfl_VCD.Src0RC64:$src0, Pfl_VCD.Src1RC64:$src1, Pfl_VCD.Src2RC64:$src2, 592 cbsz:$cbsz, abid:$abid, blgp:$blgp)>, PredicateControl; 593 } 594 } // End Predicates = [isGFX940Plus] 595} 596 597multiclass VOP3P_Real_MFMA_gfx940<bits<7> op, string Name = !cast<VOP3_Pseudo>(NAME#"_e64").Mnemonic, 598 VOP3_Pseudo PS_ACD = !cast<VOP3_Pseudo>(NAME # "_e64"), 599 VOP3_Pseudo PS_VCD = !cast<VOP3_Pseudo>(NAME # "_vgprcd" # "_e64")> { 600 let SubtargetPredicate = isGFX940Plus, 601 AssemblerPredicate = isGFX940Plus, DecoderNamespace = "GFX9", 602 AsmString = Name # PS_ACD.AsmOperands, Constraints = "" in { 603 def _gfx940_acd : VOP3P_Real<PS_ACD, SIEncodingFamily.GFX940>, 604 VOP3Pe_MAI <op, PS_ACD.Pfl, 1>; 605 606 def _gfx940_vcd : VOP3P_Real<PS_VCD, SIEncodingFamily.GFX940>, 607 VOP3Pe_MAI <op, PS_VCD.Pfl, 0>; 608 } // End AssemblerPredicate = isGFX940Plus, DecoderNamespace = "GFX9" 609 610 defm : VOP3P_Real_MFMA_gfx940_aliases<Name, PS_ACD.Mnemonic, NAME>; 611 612 foreach _ = BoolToList<!ne(!subst("_1k", "", PS_ACD.Mnemonic), PS_ACD.Mnemonic)>.ret in 613 defm : VOP3P_Real_MFMA_gfx940_aliases<Name, !subst("_1k", "", PS_ACD.Mnemonic), NAME>; 614} 615 616multiclass VOP3P_Real_MFMA<bits<7> op, string GFX940Name = !cast<VOP3_Pseudo>(NAME#"_e64").Mnemonic> : 617 VOP3P_Real_MFMA_gfx90a <op>, 618 VOP3P_Real_MFMA_gfx940 <op, GFX940Name> { 619 def _vi : VOP3P_Real<!cast<VOP3_Pseudo>(NAME#"_e64"), SIEncodingFamily.VI>, 620 VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME#"_e64").Pfl, ?> { 621 let AssemblerPredicate = HasMAIInsts; 622 let DecoderNamespace = "GFX8"; 623 let Constraints = ""; 624 } 625} 626 627defm V_PK_MAD_I16 : VOP3P_Real_vi <0x00>; 628defm V_PK_MUL_LO_U16 : VOP3P_Real_vi <0x01>; 629defm V_PK_ADD_I16 : VOP3P_Real_vi <0x02>; 630defm V_PK_SUB_I16 : VOP3P_Real_vi <0x03>; 631defm V_PK_LSHLREV_B16 : VOP3P_Real_vi <0x04>; 632defm V_PK_LSHRREV_B16 : VOP3P_Real_vi <0x05>; 633defm V_PK_ASHRREV_I16 : VOP3P_Real_vi <0x06>; 634defm V_PK_MAX_I16 : VOP3P_Real_vi <0x07>; 635defm V_PK_MIN_I16 : VOP3P_Real_vi <0x08>; 636defm V_PK_MAD_U16 : VOP3P_Real_vi <0x09>; 637 638defm V_PK_ADD_U16 : VOP3P_Real_vi <0x0a>; 639defm V_PK_SUB_U16 : VOP3P_Real_vi <0x0b>; 640defm V_PK_MAX_U16 : VOP3P_Real_vi <0x0c>; 641defm V_PK_MIN_U16 : VOP3P_Real_vi <0x0d>; 642defm V_PK_FMA_F16 : VOP3P_Real_vi <0x0e>; 643defm V_PK_ADD_F16 : VOP3P_Real_vi <0x0f>; 644defm V_PK_MUL_F16 : VOP3P_Real_vi <0x10>; 645defm V_PK_MIN_F16 : VOP3P_Real_vi <0x11>; 646defm V_PK_MAX_F16 : VOP3P_Real_vi <0x12>; 647 648 649let SubtargetPredicate = HasMadMixInsts in { 650defm V_MAD_MIX_F32 : VOP3P_Real_vi <0x20>; 651defm V_MAD_MIXLO_F16 : VOP3P_Real_vi <0x21>; 652defm V_MAD_MIXHI_F16 : VOP3P_Real_vi <0x22>; 653} 654 655let SubtargetPredicate = HasFmaMixInsts in { 656let DecoderNamespace = "GFX9_DL" in { 657// The mad_mix instructions were renamed and their behaviors changed, 658// but the opcode stayed the same so we need to put these in a 659// different DecoderNamespace to avoid the ambiguity. 660defm V_FMA_MIX_F32 : VOP3P_Real_vi <0x20>; 661defm V_FMA_MIXLO_F16 : VOP3P_Real_vi <0x21>; 662defm V_FMA_MIXHI_F16 : VOP3P_Real_vi <0x22>; 663} 664} 665 666 667let SubtargetPredicate = HasDot2Insts in { 668 669defm V_DOT2_I32_I16 : VOP3P_Real_vi <0x26>; 670defm V_DOT2_U32_U16 : VOP3P_Real_vi <0x27>; 671 672} // End SubtargetPredicate = HasDot2Insts 673 674let SubtargetPredicate = HasDot7Insts in { 675 676defm V_DOT2_F32_F16 : VOP3P_Real_vi <0x23>; 677defm V_DOT4_U32_U8 : VOP3P_Real_vi <0x29>; 678defm V_DOT8_U32_U4 : VOP3P_Real_vi <0x2b>; 679 680} // End SubtargetPredicate = HasDot7Insts 681 682let SubtargetPredicate = HasDot1Insts in { 683 684defm V_DOT4_I32_I8 : VOP3P_Real_vi <0x28>; 685defm V_DOT8_I32_I4 : VOP3P_Real_vi <0x2a>; 686 687} // End SubtargetPredicate = HasDot1Insts 688 689let SubtargetPredicate = HasMAIInsts in { 690 691defm V_ACCVGPR_READ_B32 : VOP3P_Real_MAI <0x58>; 692defm V_ACCVGPR_WRITE_B32 : VOP3P_Real_MAI <0x59>; 693defm V_MFMA_F32_32X32X1F32 : VOP3P_Real_MFMA <0x40, "v_mfma_f32_32x32x1_2b_f32">; 694defm V_MFMA_F32_16X16X1F32 : VOP3P_Real_MFMA <0x41, "v_mfma_f32_16x16x1_4b_f32">; 695defm V_MFMA_F32_4X4X1F32 : VOP3P_Real_MFMA <0x42, "v_mfma_f32_4x4x1_16b_f32">; 696defm V_MFMA_F32_32X32X2F32 : VOP3P_Real_MFMA <0x44, "v_mfma_f32_32x32x2_f32">; 697defm V_MFMA_F32_16X16X4F32 : VOP3P_Real_MFMA <0x45, "v_mfma_f32_16x16x4_f32">; 698defm V_MFMA_F32_32X32X4F16 : VOP3P_Real_MFMA <0x48, "v_mfma_f32_32x32x4_2b_f16">; 699defm V_MFMA_F32_16X16X4F16 : VOP3P_Real_MFMA <0x49, "v_mfma_f32_16x16x4_4b_f16">; 700defm V_MFMA_F32_4X4X4F16 : VOP3P_Real_MFMA <0x4a, "v_mfma_f32_4x4x4_16b_f16">; 701defm V_MFMA_F32_32X32X8F16 : VOP3P_Real_MFMA <0x4c, "v_mfma_f32_32x32x8_f16">; 702defm V_MFMA_F32_16X16X16F16 : VOP3P_Real_MFMA <0x4d, "v_mfma_f32_16x16x16_f16">; 703defm V_MFMA_I32_32X32X4I8 : VOP3P_Real_MFMA <0x50, "v_mfma_i32_32x32x4_2b_i8">; 704defm V_MFMA_I32_16X16X4I8 : VOP3P_Real_MFMA <0x51, "v_mfma_i32_16x16x4_4b_i8">; 705defm V_MFMA_I32_4X4X4I8 : VOP3P_Real_MFMA <0x52, "v_mfma_i32_4x4x4_16b_i8">; 706 707defm V_MFMA_I32_16X16X16I8 : VOP3P_Real_MFMA <0x55>; 708defm V_MFMA_I32_32X32X8I8 : VOP3P_Real_MFMA <0x54>; 709defm V_MFMA_F32_32X32X2BF16 : VOP3P_Real_MFMA <0x68>; 710defm V_MFMA_F32_16X16X2BF16 : VOP3P_Real_MFMA <0x69>; 711defm V_MFMA_F32_4X4X2BF16 : VOP3P_Real_MFMA <0x6b>; 712defm V_MFMA_F32_32X32X4BF16 : VOP3P_Real_MFMA <0x6c>; 713defm V_MFMA_F32_16X16X8BF16 : VOP3P_Real_MFMA <0x6d>; 714 715} // End SubtargetPredicate = HasMAIInsts 716 717defm V_MFMA_F32_32X32X4BF16_1K : VOP3P_Real_MFMA_gfx90a <0x63>; 718defm V_MFMA_F32_16X16X4BF16_1K : VOP3P_Real_MFMA_gfx90a <0x64>; 719defm V_MFMA_F32_4X4X4BF16_1K : VOP3P_Real_MFMA_gfx90a <0x65>; 720defm V_MFMA_F32_32X32X8BF16_1K : VOP3P_Real_MFMA_gfx90a <0x66>; 721defm V_MFMA_F32_16X16X16BF16_1K : VOP3P_Real_MFMA_gfx90a <0x67>; 722defm V_MFMA_F64_16X16X4F64 : VOP3P_Real_MFMA_gfx90a <0x6e>; 723defm V_MFMA_F64_4X4X4F64 : VOP3P_Real_MFMA_gfx90a <0x6f>; 724 725defm V_MFMA_F32_32X32X4BF16_1K : VOP3P_Real_MFMA_gfx940 <0x5d, "v_mfma_f32_32x32x4_2b_bf16">; 726defm V_MFMA_F32_16X16X4BF16_1K : VOP3P_Real_MFMA_gfx940 <0x5e, "v_mfma_f32_16x16x4_4b_bf16">; 727defm V_MFMA_F32_4X4X4BF16_1K : VOP3P_Real_MFMA_gfx940 <0x5f, "v_mfma_f32_4x4x4_16b_bf16">; 728defm V_MFMA_F32_32X32X8BF16_1K : VOP3P_Real_MFMA_gfx940 <0x60, "v_mfma_f32_32x32x8_bf16">; 729defm V_MFMA_F32_16X16X16BF16_1K : VOP3P_Real_MFMA_gfx940 <0x61, "v_mfma_f32_16x16x16_bf16">; 730 731defm V_MFMA_F64_16X16X4F64 : VOP3P_Real_MFMA_gfx940 <0x6e, "v_mfma_f64_16x16x4_f64">; 732defm V_MFMA_F64_4X4X4F64 : VOP3P_Real_MFMA_gfx940 <0x6f, "v_mfma_f64_4x4x4_4b_f64">; 733 734let SubtargetPredicate = HasPackedFP32Ops in { 735 defm V_PK_FMA_F32 : VOP3P_Real_vi <0x30>; 736 defm V_PK_MUL_F32 : VOP3P_Real_vi <0x31>; 737 defm V_PK_ADD_F32 : VOP3P_Real_vi <0x32>; 738 defm V_PK_MOV_B32 : VOP3P_Real_vi <0x33>; 739} // End SubtargetPredicate = HasPackedFP32Ops 740 741//===----------------------------------------------------------------------===// 742// GFX10. 743//===----------------------------------------------------------------------===// 744 745let AssemblerPredicate = isGFX10Plus, DecoderNamespace = "GFX10", VOP3P = 1 in { 746 multiclass VOP3P_Real_gfx10<bits<7> op> { 747 def _gfx10 : VOP3P_Real<!cast<VOP3P_Pseudo>(NAME), SIEncodingFamily.GFX10>, 748 VOP3Pe_gfx10 <op, !cast<VOP3P_Pseudo>(NAME).Pfl>; 749 } 750} // End AssemblerPredicate = isGFX10Plus, DecoderNamespace = "GFX10", VOP3P = 1 751 752defm V_PK_MAD_I16 : VOP3P_Real_gfx10<0x00>; 753defm V_PK_MUL_LO_U16 : VOP3P_Real_gfx10<0x01>; 754defm V_PK_ADD_I16 : VOP3P_Real_gfx10<0x02>; 755defm V_PK_SUB_I16 : VOP3P_Real_gfx10<0x03>; 756defm V_PK_LSHLREV_B16 : VOP3P_Real_gfx10<0x04>; 757defm V_PK_LSHRREV_B16 : VOP3P_Real_gfx10<0x05>; 758defm V_PK_ASHRREV_I16 : VOP3P_Real_gfx10<0x06>; 759defm V_PK_MAX_I16 : VOP3P_Real_gfx10<0x07>; 760defm V_PK_MIN_I16 : VOP3P_Real_gfx10<0x08>; 761defm V_PK_MAD_U16 : VOP3P_Real_gfx10<0x09>; 762defm V_PK_ADD_U16 : VOP3P_Real_gfx10<0x0a>; 763defm V_PK_SUB_U16 : VOP3P_Real_gfx10<0x0b>; 764defm V_PK_MAX_U16 : VOP3P_Real_gfx10<0x0c>; 765defm V_PK_MIN_U16 : VOP3P_Real_gfx10<0x0d>; 766defm V_PK_FMA_F16 : VOP3P_Real_gfx10<0x0e>; 767defm V_PK_ADD_F16 : VOP3P_Real_gfx10<0x0f>; 768defm V_PK_MUL_F16 : VOP3P_Real_gfx10<0x10>; 769defm V_PK_MIN_F16 : VOP3P_Real_gfx10<0x11>; 770defm V_PK_MAX_F16 : VOP3P_Real_gfx10<0x12>; 771defm V_FMA_MIX_F32 : VOP3P_Real_gfx10<0x20>; 772defm V_FMA_MIXLO_F16 : VOP3P_Real_gfx10<0x21>; 773defm V_FMA_MIXHI_F16 : VOP3P_Real_gfx10<0x22>; 774 775let SubtargetPredicate = HasDot2Insts in { 776 777defm V_DOT2_I32_I16 : VOP3P_Real_gfx10 <0x14>; 778defm V_DOT2_U32_U16 : VOP3P_Real_gfx10 <0x15>; 779 780} // End SubtargetPredicate = HasDot2Insts 781 782let SubtargetPredicate = HasDot7Insts in { 783 784defm V_DOT2_F32_F16 : VOP3P_Real_gfx10 <0x13>; 785defm V_DOT4_U32_U8 : VOP3P_Real_gfx10 <0x17>; 786defm V_DOT8_U32_U4 : VOP3P_Real_gfx10 <0x19>; 787 788} // End SubtargetPredicate = HasDot7Insts 789 790let SubtargetPredicate = HasDot1Insts in { 791 792defm V_DOT4_I32_I8 : VOP3P_Real_gfx10 <0x16>; 793defm V_DOT8_I32_I4 : VOP3P_Real_gfx10 <0x18>; 794 795} // End SubtargetPredicate = HasDot1Insts 796