brintos / llvm-project-archived public Read only

1//===-- SIInstrInfo.td -----------------------------------------------------===//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 9class AMDGPUMnemonicAlias<string From, string To, string VariantName = "">10    : MnemonicAlias<From, To, VariantName>, PredicateControl;11 12// Except for the NONE field, this must be kept in sync with the13// SIEncodingFamily enum in SIInstrInfo.cpp and the columns of the14// getMCOpcodeGen table.15def SIEncodingFamily {16  int NONE = -1;17  int SI = 0;18  int VI = 1;19  int SDWA = 2;20  int SDWA9 = 3;21  int GFX80 = 4;22  int GFX9 = 5;23  int GFX10 = 6;24  int SDWA10 = 7;25  int GFX90A = 8;26  int GFX940 = 9;27  int GFX11 = 10;28  int GFX12 = 11;29  int GFX1250 = 12;30}31 32//===----------------------------------------------------------------------===//33// Subtarget info34//===----------------------------------------------------------------------===//35 36class GFXGen<Predicate pred, string dn, string suffix, int sub> {37  Predicate AssemblerPredicate = pred;38  string DecoderNamespace = dn;39  string Suffix = suffix;40  int Subtarget = sub;41}42 43def GFX1250Gen       : GFXGen<isGFX125xOnly, "GFX1250", "_gfx1250", SIEncodingFamily.GFX1250>;44def GFX12Not12_50Gen : GFXGen<isGFX12Not12_50, "GFX12", "_gfx12", SIEncodingFamily.GFX12>;45def GFX12Gen : GFXGen<isGFX12Only, "GFX12", "_gfx12", SIEncodingFamily.GFX12>;46def GFX11Gen : GFXGen<isGFX11Only, "GFX11", "_gfx11", SIEncodingFamily.GFX11>;47def GFX10Gen : GFXGen<isGFX10Only, "GFX10", "_gfx10", SIEncodingFamily.GFX10>;48 49//===----------------------------------------------------------------------===//50// SI DAG Nodes51//===----------------------------------------------------------------------===//52 53// CLAMP value between 0.0 and 1.0. NaN clamped to 0, following clamp output54// modifier behavior with dx10_enable.55def AMDGPUclamp : SDNode<"AMDGPUISD::CLAMP", SDTFPUnaryOp>;56 57def SDTSBufferLoad : SDTypeProfile<1, 3,58    [                    // vdata59     SDTCisVT<1, v4i32>, // rsrc60     SDTCisVT<2, i32>,   // offset(imm)61     SDTCisVT<3, i32>]>; // cachepolicy62 63def SIsbuffer_load : SDNode<"AMDGPUISD::SBUFFER_LOAD", SDTSBufferLoad,64                            [SDNPMayLoad, SDNPMemOperand]>;65 66def SIsbuffer_load_byte : SDNode<"AMDGPUISD::SBUFFER_LOAD_BYTE", SDTSBufferLoad,67                                 [SDNPMayLoad, SDNPMemOperand]>;68 69def SIsbuffer_load_ubyte70    : SDNode<"AMDGPUISD::SBUFFER_LOAD_UBYTE", SDTSBufferLoad,71             [SDNPMayLoad, SDNPMemOperand]>;72 73def SIsbuffer_load_short74    : SDNode<"AMDGPUISD::SBUFFER_LOAD_SHORT", SDTSBufferLoad,75             [SDNPMayLoad, SDNPMemOperand]>;76 77def SIsbuffer_load_ushort78    : SDNode<"AMDGPUISD::SBUFFER_LOAD_USHORT", SDTSBufferLoad,79             [SDNPMayLoad, SDNPMemOperand]>;80 81def SIds_ordered_count : SDNode<"AMDGPUISD::DS_ORDERED_COUNT",82  SDTypeProfile<1, 2, [SDTCisVT<0, i32>, SDTCisVT<1, i32>, SDTCisVT<2, i16>]>,83  [SDNPMayLoad, SDNPMayStore, SDNPMemOperand, SDNPHasChain, SDNPInGlue]84>;85 86def SDTAtomic2_f32 : SDTypeProfile<1, 2, [87  SDTCisSameAs<0,2>, SDTCisFP<0>, SDTCisPtrTy<1>88]>;89 90// load_d16_{lo|hi} ptr, tied_input91def SIload_d16 : SDTypeProfile<1, 2, [92  SDTCisPtrTy<1>,93  SDTCisSameAs<0, 2>94]>;95 96 97def SDTtbuffer_load : SDTypeProfile<1, 8,98  [                     // vdata99   SDTCisVT<1, v4i32>,  // rsrc100   SDTCisVT<2, i32>,    // vindex(VGPR)101   SDTCisVT<3, i32>,    // voffset(VGPR)102   SDTCisVT<4, i32>,    // soffset(SGPR)103   SDTCisVT<5, i32>,    // offset(imm)104   SDTCisVT<6, i32>,    // format(imm)105   SDTCisVT<7, i32>,    // cachepolicy, swizzled buffer(imm)106   SDTCisVT<8, i1>      // idxen(imm)107  ]>;108 109def SItbuffer_load :   SDNode<"AMDGPUISD::TBUFFER_LOAD_FORMAT", SDTtbuffer_load,110                              [SDNPMayLoad, SDNPMemOperand, SDNPHasChain]>;111def SItbuffer_load_d16 : SDNode<"AMDGPUISD::TBUFFER_LOAD_FORMAT_D16",112                                SDTtbuffer_load,113                                [SDNPMayLoad, SDNPMemOperand, SDNPHasChain]>;114 115def SDTtbuffer_store : SDTypeProfile<0, 9,116    [                     // vdata117     SDTCisVT<1, v4i32>,  // rsrc118     SDTCisVT<2, i32>,    // vindex(VGPR)119     SDTCisVT<3, i32>,    // voffset(VGPR)120     SDTCisVT<4, i32>,    // soffset(SGPR)121     SDTCisVT<5, i32>,    // offset(imm)122     SDTCisVT<6, i32>,    // format(imm)123     SDTCisVT<7, i32>,    // cachepolicy, swizzled buffer(imm)124     SDTCisVT<8, i1>      // idxen(imm)125    ]>;126 127def SItbuffer_store : SDNode<"AMDGPUISD::TBUFFER_STORE_FORMAT", SDTtbuffer_store,128                             [SDNPMayStore, SDNPMemOperand, SDNPHasChain]>;129def SItbuffer_store_d16 : SDNode<"AMDGPUISD::TBUFFER_STORE_FORMAT_D16",130                                SDTtbuffer_store,131                                [SDNPMayStore, SDNPMemOperand, SDNPHasChain]>;132 133def SDTBufferLoad : SDTypeProfile<1, 7,134    [                    // vdata135     SDTCisVT<1, v4i32>, // rsrc136     SDTCisVT<2, i32>,   // vindex(VGPR)137     SDTCisVT<3, i32>,   // voffset(VGPR)138     SDTCisVT<4, i32>,   // soffset(SGPR)139     SDTCisVT<5, i32>,   // offset(imm)140     SDTCisVT<6, i32>,   // cachepolicy, swizzled buffer(imm)141     SDTCisVT<7, i1>]>;  // idxen(imm)142 143def SIbuffer_load : SDNode <"AMDGPUISD::BUFFER_LOAD", SDTBufferLoad,144                            [SDNPMemOperand, SDNPHasChain, SDNPMayLoad]>;145def SIbuffer_load_ubyte : SDNode <"AMDGPUISD::BUFFER_LOAD_UBYTE", SDTBufferLoad,146                            [SDNPMemOperand, SDNPHasChain, SDNPMayLoad]>;147def SIbuffer_load_ushort : SDNode <"AMDGPUISD::BUFFER_LOAD_USHORT", SDTBufferLoad,148                            [SDNPMemOperand, SDNPHasChain, SDNPMayLoad]>;149def SIbuffer_load_byte : SDNode <"AMDGPUISD::BUFFER_LOAD_BYTE", SDTBufferLoad,150                            [SDNPMemOperand, SDNPHasChain, SDNPMayLoad]>;151def SIbuffer_load_short: SDNode <"AMDGPUISD::BUFFER_LOAD_SHORT", SDTBufferLoad,152                            [SDNPMemOperand, SDNPHasChain, SDNPMayLoad]>;153def SIbuffer_load_tfe : SDNode <"AMDGPUISD::BUFFER_LOAD_TFE", SDTBufferLoad,154                            [SDNPMemOperand, SDNPHasChain, SDNPMayLoad]>;155def SIbuffer_load_ubyte_tfe : SDNode <"AMDGPUISD::BUFFER_LOAD_UBYTE_TFE", SDTBufferLoad,156                            [SDNPMemOperand, SDNPHasChain, SDNPMayLoad]>;157def SIbuffer_load_ushort_tfe : SDNode <"AMDGPUISD::BUFFER_LOAD_USHORT_TFE", SDTBufferLoad,158                            [SDNPMemOperand, SDNPHasChain, SDNPMayLoad]>;159def SIbuffer_load_byte_tfe : SDNode <"AMDGPUISD::BUFFER_LOAD_BYTE_TFE", SDTBufferLoad,160                            [SDNPMemOperand, SDNPHasChain, SDNPMayLoad]>;161def SIbuffer_load_short_tfe: SDNode <"AMDGPUISD::BUFFER_LOAD_SHORT_TFE", SDTBufferLoad,162                            [SDNPMemOperand, SDNPHasChain, SDNPMayLoad]>;163def SIbuffer_load_format : SDNode <"AMDGPUISD::BUFFER_LOAD_FORMAT", SDTBufferLoad,164                            [SDNPMemOperand, SDNPHasChain, SDNPMayLoad]>;165def SIbuffer_load_format_tfe : SDNode <"AMDGPUISD::BUFFER_LOAD_FORMAT_TFE", SDTBufferLoad,166                               [SDNPMemOperand, SDNPHasChain, SDNPMayLoad]>;167def SIbuffer_load_format_d16 : SDNode <"AMDGPUISD::BUFFER_LOAD_FORMAT_D16",168                                SDTBufferLoad,169                                [SDNPMemOperand, SDNPHasChain, SDNPMayLoad]>;170 171def SDTBufferStore : SDTypeProfile<0, 8,172    [                    // vdata173     SDTCisVT<1, v4i32>, // rsrc174     SDTCisVT<2, i32>,   // vindex(VGPR)175     SDTCisVT<3, i32>,   // voffset(VGPR)176     SDTCisVT<4, i32>,   // soffset(SGPR)177     SDTCisVT<5, i32>,   // offset(imm)178     SDTCisVT<6, i32>,   // cachepolicy, swizzled buffer(imm)179     SDTCisVT<7, i1>]>;  // idxen(imm)180 181def SIbuffer_store : SDNode <"AMDGPUISD::BUFFER_STORE", SDTBufferStore,182                             [SDNPMayStore, SDNPMemOperand, SDNPHasChain]>;183def SIbuffer_store_byte: SDNode <"AMDGPUISD::BUFFER_STORE_BYTE",184                         SDTBufferStore,185                         [SDNPMayStore, SDNPMemOperand, SDNPHasChain]>;186def SIbuffer_store_short : SDNode <"AMDGPUISD::BUFFER_STORE_SHORT",187                           SDTBufferStore,188                           [SDNPMayStore, SDNPMemOperand, SDNPHasChain]>;189def SIbuffer_store_format : SDNode <"AMDGPUISD::BUFFER_STORE_FORMAT",190                            SDTBufferStore,191                            [SDNPMayStore, SDNPMemOperand, SDNPHasChain]>;192def SIbuffer_store_format_d16 : SDNode <"AMDGPUISD::BUFFER_STORE_FORMAT_D16",193                            SDTBufferStore,194                            [SDNPMayStore, SDNPMemOperand, SDNPHasChain]>;195 196multiclass SDBufferAtomic<string opcode> {197  def "" : SDNode <opcode,198    SDTypeProfile<1, 8,199         [SDTCisVT<2, v4i32>, // rsrc200         SDTCisVT<3, i32>,   // vindex(VGPR)201         SDTCisVT<4, i32>,   // voffset(VGPR)202         SDTCisVT<5, i32>,   // soffset(SGPR)203         SDTCisVT<6, i32>,   // offset(imm)204         SDTCisVT<7, i32>,   // cachepolicy(imm)205         SDTCisVT<8, i1>]>,  // idxen(imm)206    [SDNPMemOperand, SDNPHasChain, SDNPMayLoad, SDNPMayStore]207  >;208  def "_noret" : PatFrag<209    (ops node:$vdata_in, node:$rsrc, node:$vindex, node:$voffset, node:$soffset,210      node:$offset, node:$cachepolicy, node:$idxen),211    (!cast<SDNode>(NAME) node:$vdata_in, node:$rsrc, node:$vindex,212      node:$voffset, node:$soffset, node:$offset, node:$cachepolicy,213      node:$idxen)> {214    let HasNoUse = true;215  }216}217 218defm SIbuffer_atomic_swap : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_SWAP">;219defm SIbuffer_atomic_add : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_ADD">;220defm SIbuffer_atomic_sub : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_SUB">;221defm SIbuffer_atomic_smin : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_SMIN">;222defm SIbuffer_atomic_umin : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_UMIN">;223defm SIbuffer_atomic_smax : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_SMAX">;224defm SIbuffer_atomic_umax : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_UMAX">;225defm SIbuffer_atomic_and : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_AND">;226defm SIbuffer_atomic_or : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_OR">;227defm SIbuffer_atomic_xor : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_XOR">;228defm SIbuffer_atomic_inc : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_INC">;229defm SIbuffer_atomic_dec : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_DEC">;230defm SIbuffer_atomic_csub : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_CSUB">;231defm SIbuffer_atomic_fadd : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_FADD">;232defm SIbuffer_atomic_fmin : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_FMIN">;233defm SIbuffer_atomic_fmax : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_FMAX">;234defm SIbuffer_atomic_cond_sub_u32 : SDBufferAtomic <"AMDGPUISD::BUFFER_ATOMIC_COND_SUB_U32">;235 236def SIbuffer_atomic_cmpswap : SDNode <"AMDGPUISD::BUFFER_ATOMIC_CMPSWAP",237  SDTypeProfile<1, 9,238    [SDTCisVT<3, v4i32>, // rsrc239     SDTCisVT<4, i32>,   // vindex(VGPR)240     SDTCisVT<5, i32>,   // voffset(VGPR)241     SDTCisVT<6, i32>,   // soffset(SGPR)242     SDTCisVT<7, i32>,   // offset(imm)243     SDTCisVT<8, i32>,   // cachepolicy(imm)244     SDTCisVT<9, i1>]>,  // idxen(imm)245  [SDNPMemOperand, SDNPHasChain, SDNPMayLoad, SDNPMayStore]246>;247 248def SIbuffer_atomic_cmpswap_noret : PatFrag<249  (ops node:$src, node:$cmp, node:$rsrc, node:$vindex, node:$voffset,250    node:$soffset, node:$offset, node:$cachepolicy, node:$idxen),251  (SIbuffer_atomic_cmpswap node:$src, node:$cmp, node:$rsrc, node:$vindex,252    node:$voffset, node:$soffset, node:$offset, node:$cachepolicy,253    node:$idxen)> {254  let HasNoUse = true;255}256 257class SDGlobalAtomicNoRtn<string opcode, ValueType ty> : SDNode <opcode,258  SDTypeProfile<0, 2,259      [SDTCisPtrTy<0>,     // vaddr260       SDTCisVT<1, ty>]>,  // vdata261  [SDNPMemOperand, SDNPHasChain, SDNPMayLoad, SDNPMayStore]262>;263 264def SIpc_add_rel_offset : SDNode<"AMDGPUISD::PC_ADD_REL_OFFSET",265  SDTypeProfile<1, 2, [SDTCisVT<0, iPTR>, SDTCisSameAs<0,1>, SDTCisSameAs<0,2>]>266>;267 268def SIpc_add_rel_offset64 : SDNode<"AMDGPUISD::PC_ADD_REL_OFFSET64",269  SDTypeProfile<1, 1, [SDTCisVT<0, iPTR>, SDTCisSameAs<0,1>]>270>;271 272def SIlds : SDNode<"AMDGPUISD::LDS",273  SDTypeProfile<1, 1, [SDTCisVT<0, iPTR>, SDTCisSameAs<0,1>]>274>;275 276def SIload_d16_lo : SDNode<"AMDGPUISD::LOAD_D16_LO",277  SIload_d16,278  [SDNPMayLoad, SDNPMemOperand, SDNPHasChain]279>;280 281def SIload_d16_lo_u8 : SDNode<"AMDGPUISD::LOAD_D16_LO_U8",282  SIload_d16,283  [SDNPMayLoad, SDNPMemOperand, SDNPHasChain]284>;285 286def SIload_d16_lo_i8 : SDNode<"AMDGPUISD::LOAD_D16_LO_I8",287  SIload_d16,288  [SDNPMayLoad, SDNPMemOperand, SDNPHasChain]289>;290 291def SIload_d16_hi : SDNode<"AMDGPUISD::LOAD_D16_HI",292  SIload_d16,293  [SDNPMayLoad, SDNPMemOperand, SDNPHasChain]294>;295 296def SIload_d16_hi_u8 : SDNode<"AMDGPUISD::LOAD_D16_HI_U8",297  SIload_d16,298  [SDNPMayLoad, SDNPMemOperand, SDNPHasChain]299>;300 301def SIload_d16_hi_i8 : SDNode<"AMDGPUISD::LOAD_D16_HI_I8",302  SIload_d16,303  [SDNPMayLoad, SDNPMemOperand, SDNPHasChain]304>;305 306def SIdenorm_mode : SDNode<"AMDGPUISD::DENORM_MODE",307  SDTypeProfile<0 ,1, [SDTCisInt<0>]>,308  [SDNPHasChain, SDNPOptInGlue, SDNPOutGlue]309>;310 311 312// Optimize v_mfma_scale* instructions to avoid the scale if the313// scales are known 0.314class UnscaledMFMAOptimizationPat<SDPatternOperator intrin> : PatFrag<315  (ops node:$srca, node:$srcb, node:$srcc,316       node:$cbsz, node:$blgp),317  (intrin $srca, $srcb, $srcc, $cbsz, $blgp,318          srcvalue, 0, srcvalue, 0)319>;320 321def mfma_f32_16x16x128_f8f6f4 : UnscaledMFMAOptimizationPat<int_amdgcn_mfma_scale_f32_16x16x128_f8f6f4>;322def mfma_f32_32x32x64_f8f6f4 : UnscaledMFMAOptimizationPat<int_amdgcn_mfma_scale_f32_32x32x64_f8f6f4>;323 324//===----------------------------------------------------------------------===//325// ValueType helpers326//===----------------------------------------------------------------------===//327 328class isIntType<ValueType SrcVT> {329  bit ret = !and(SrcVT.isInteger, !ne(SrcVT, i1));330}331 332def SDTSBufferPrefetch : SDTypeProfile<0, 3,333    [SDTCisVT<0, v4i32>, // rsrc334     SDTCisVT<1, i32>,   // offset(imm)335     SDTCisVT<2, i32>]>; // length336 337def SIsbuffer_prefetch : SDNode<"AMDGPUISD::SBUFFER_PREFETCH_DATA", SDTSBufferPrefetch,338                                [SDNPMayLoad, SDNPMayStore, SDNPMemOperand, SDNPHasChain]>;339 340//===----------------------------------------------------------------------===//341// SDNodes PatFrags for loads/stores with a glue input.342// This is for SDNodes and PatFrag for local loads and stores to343// enable s_mov_b32 m0, -1 to be glued to the memory instructions.344//345// These mirror the regular load/store PatFrags and rely on special346// processing during Select() to add the glued copy.347//348//===----------------------------------------------------------------------===//349 350def AMDGPUld_glue : SDNode <"ISD::LOAD", SDTLoad,351  [SDNPHasChain, SDNPMayLoad, SDNPMemOperand, SDNPInGlue]352>;353 354def AMDGPUatomic_ld_glue : SDNode <"ISD::ATOMIC_LOAD", SDTAtomicLoad,355  [SDNPHasChain, SDNPMayLoad, SDNPMemOperand, SDNPInGlue]356>;357 358def unindexedload_glue : PatFrag <(ops node:$ptr), (AMDGPUld_glue node:$ptr)> {359  let IsLoad = 1;360  let IsUnindexed = 1;361}362 363def load_glue : PatFrag <(ops node:$ptr), (unindexedload_glue node:$ptr)> {364  let IsLoad = 1;365  let IsNonExtLoad = 1;366}367 368def atomic_load_nonext_glue :369  PatFrag<(ops node:$ptr), (AMDGPUatomic_ld_glue node:$ptr)> {370  let IsAtomic = true; // FIXME: Should be IsLoad and/or IsAtomic?371  let IsNonExtLoad = true;372}373 374def atomic_load_zext_glue :375  PatFrag<(ops node:$ptr), (AMDGPUatomic_ld_glue node:$ptr)> {376  let IsAtomic = true; // FIXME: Should be IsLoad and/or IsAtomic?377  let IsZeroExtLoad = true;378}379 380def atomic_load_sext_glue :381  PatFrag<(ops node:$ptr), (AMDGPUatomic_ld_glue node:$ptr)> {382  let IsAtomic = true; // FIXME: Should be IsLoad and/or IsAtomic?383  let IsSignExtLoad = true;384}385 386def atomic_load_aext_glue :387  PatFrag<(ops node:$ptr), (AMDGPUatomic_ld_glue node:$ptr)> {388  let IsAtomic = true; // FIXME: Should be IsLoad and/or IsAtomic?389  let IsAnyExtLoad = true;390}391 392def atomic_load_nonext_16_glue : PatFrag<(ops node:$ptr),393  (atomic_load_nonext_glue node:$ptr)> {394  let IsAtomic = 1;395  let MemoryVT = i16;396}397 398def atomic_load_nonext_32_glue : PatFrag<(ops node:$ptr),399  (atomic_load_nonext_glue node:$ptr)> {400  let IsAtomic = 1;401  let MemoryVT = i32;402}403 404def atomic_load_nonext_64_glue : PatFrag<(ops node:$ptr),405  (atomic_load_nonext_glue node:$ptr)> {406  let IsAtomic = 1;407  let MemoryVT = i64;408}409 410def atomic_load_zext_8_glue : PatFrag<(ops node:$ptr),411  (atomic_load_zext_glue node:$ptr)> {412  let IsAtomic = 1;413  let MemoryVT = i8;414}415 416def atomic_load_sext_8_glue : PatFrag<(ops node:$ptr),417  (atomic_load_sext_glue node:$ptr)> {418  let IsAtomic = 1;419  let MemoryVT = i8;420}421 422def atomic_load_aext_8_glue : PatFrag<(ops node:$ptr),423  (atomic_load_aext_glue node:$ptr)> {424  let IsAtomic = 1;425  let MemoryVT = i8;426}427 428def atomic_load_zext_16_glue : PatFrag<(ops node:$ptr),429  (atomic_load_zext_glue node:$ptr)> {430  let IsAtomic = 1;431  let MemoryVT = i16;432}433 434def atomic_load_sext_16_glue : PatFrag<(ops node:$ptr),435  (atomic_load_sext_glue node:$ptr)> {436  let IsAtomic = 1;437  let MemoryVT = i16;438}439 440def atomic_load_aext_16_glue : PatFrag<(ops node:$ptr),441  (atomic_load_aext_glue node:$ptr)> {442  let IsAtomic = 1;443  let MemoryVT = i16;444}445 446def extload_glue : PatFrag<(ops node:$ptr), (unindexedload_glue node:$ptr)> {447  let IsLoad = 1;448  let IsAnyExtLoad = 1;449}450 451def sextload_glue : PatFrag<(ops node:$ptr), (unindexedload_glue node:$ptr)> {452  let IsLoad = 1;453  let IsSignExtLoad = 1;454}455 456def zextload_glue : PatFrag<(ops node:$ptr), (unindexedload_glue node:$ptr)> {457  let IsLoad = 1;458  let IsZeroExtLoad = 1;459}460 461def extloadi8_glue : PatFrag<(ops node:$ptr), (extload_glue node:$ptr)> {462  let IsLoad = 1;463  let MemoryVT = i8;464}465 466def zextloadi8_glue : PatFrag<(ops node:$ptr), (zextload_glue node:$ptr)> {467  let IsLoad = 1;468  let MemoryVT = i8;469}470 471def extloadi16_glue : PatFrag<(ops node:$ptr), (extload_glue node:$ptr)> {472  let IsLoad = 1;473  let MemoryVT = i16;474}475 476def zextloadi16_glue : PatFrag<(ops node:$ptr), (zextload_glue node:$ptr)> {477  let IsLoad = 1;478  let MemoryVT = i16;479}480 481def sextloadi8_glue : PatFrag<(ops node:$ptr), (sextload_glue node:$ptr)> {482  let IsLoad = 1;483  let MemoryVT = i8;484}485 486def sextloadi16_glue : PatFrag<(ops node:$ptr), (sextload_glue node:$ptr)> {487  let IsLoad = 1;488  let MemoryVT = i16;489}490 491 492let IsLoad = 1, AddressSpaces = LoadAddress_local.AddrSpaces in {493def load_local_m0 : PatFrag<(ops node:$ptr), (load_glue node:$ptr)> {494  let IsNonExtLoad = 1;495}496 497def extloadi8_local_m0 : PatFrag<(ops node:$ptr), (extloadi8_glue node:$ptr)>;498def sextloadi8_local_m0 : PatFrag<(ops node:$ptr), (sextloadi8_glue node:$ptr)>;499def zextloadi8_local_m0 : PatFrag<(ops node:$ptr), (zextloadi8_glue node:$ptr)>;500 501def extloadi16_local_m0 : PatFrag<(ops node:$ptr), (extloadi16_glue node:$ptr)>;502def sextloadi16_local_m0 : PatFrag<(ops node:$ptr), (sextloadi16_glue node:$ptr)>;503def zextloadi16_local_m0 : PatFrag<(ops node:$ptr), (zextloadi16_glue node:$ptr)>;504} // End IsLoad = 1, , AddressSpaces = LoadAddress_local.AddrSpaces505 506def load_align8_local_m0 : PatFrag<(ops node:$ptr),507                                   (load_local_m0 node:$ptr)> {508  let IsLoad = 1;509  int MinAlignment = 8;510}511 512def load_align16_local_m0 : PatFrag<(ops node:$ptr),513                                   (load_local_m0 node:$ptr)> {514  let IsLoad = 1;515  int MinAlignment = 16;516}517 518let IsAtomic = 1, AddressSpaces = LoadAddress_local.AddrSpaces in {519def atomic_load_nonext_16_local_m0 : PatFrag<(ops node:$ptr),520                                      (atomic_load_nonext_16_glue node:$ptr)>;521def atomic_load_nonext_32_local_m0 : PatFrag<(ops node:$ptr),522                                      (atomic_load_nonext_32_glue node:$ptr)>;523def atomic_load_nonext_64_local_m0 : PatFrag<(ops node:$ptr),524                                       (atomic_load_nonext_64_glue node:$ptr)>;525 526def atomic_load_zext_8_local_m0 : PatFrag<(ops node:$ptr),527                                      (atomic_load_zext_8_glue node:$ptr)>;528def atomic_load_sext_8_local_m0 : PatFrag<(ops node:$ptr),529                                      (atomic_load_sext_8_glue node:$ptr)>;530def atomic_load_aext_8_local_m0 : PatFrag<(ops node:$ptr),531                                      (atomic_load_aext_8_glue node:$ptr)>;532def atomic_load_zext_16_local_m0 : PatFrag<(ops node:$ptr),533                                      (atomic_load_zext_16_glue node:$ptr)>;534def atomic_load_sext_16_local_m0 : PatFrag<(ops node:$ptr),535                                      (atomic_load_sext_16_glue node:$ptr)>;536def atomic_load_aext_16_local_m0 : PatFrag<(ops node:$ptr),537                                      (atomic_load_aext_16_glue node:$ptr)>;538} // End let AddressSpaces = LoadAddress_local.AddrSpaces539 540 541def AMDGPUst_glue : SDNode <"ISD::STORE", SDTStore,542  [SDNPHasChain, SDNPMayStore, SDNPMemOperand, SDNPInGlue]543>;544 545def AMDGPUatomic_st_glue : SDNode <"ISD::ATOMIC_STORE", SDTAtomicStore,546  [SDNPHasChain, SDNPMayStore, SDNPMemOperand, SDNPInGlue]547>;548 549def unindexedstore_glue : PatFrag<(ops node:$val, node:$ptr),550                                   (AMDGPUst_glue node:$val, node:$ptr)> {551  let IsStore = 1;552  let IsUnindexed = 1;553}554 555def store_glue : PatFrag<(ops node:$val, node:$ptr),556                         (unindexedstore_glue node:$val, node:$ptr)> {557  let IsStore = 1;558  let IsTruncStore = 0;559}560 561def truncstore_glue : PatFrag<(ops node:$val, node:$ptr),562  (unindexedstore_glue node:$val, node:$ptr)> {563  let IsStore = 1;564  let IsTruncStore = 1;565}566 567def truncstorei8_glue : PatFrag<(ops node:$val, node:$ptr),568                           (truncstore_glue node:$val, node:$ptr)> {569  let IsStore = 1;570  let MemoryVT = i8;571}572 573def truncstorei16_glue : PatFrag<(ops node:$val, node:$ptr),574                           (truncstore_glue node:$val, node:$ptr)> {575  let IsStore = 1;576  let MemoryVT = i16;577}578 579let IsStore = 1, AddressSpaces = StoreAddress_local.AddrSpaces in {580def store_local_m0 : PatFrag<(ops node:$val, node:$ptr),581                             (store_glue node:$val, node:$ptr)>;582def truncstorei8_local_m0 : PatFrag<(ops node:$val, node:$ptr),583                                    (truncstorei8_glue node:$val, node:$ptr)>;584def truncstorei16_local_m0 : PatFrag<(ops node:$val, node:$ptr),585                                    (truncstorei16_glue node:$val, node:$ptr)>;586}587 588def store_align8_local_m0 : PatFrag <(ops node:$value, node:$ptr),589                                     (store_local_m0 node:$value, node:$ptr)>,590                            Aligned<8> {591  let IsStore = 1;592}593 594def store_align16_local_m0 : PatFrag <(ops node:$value, node:$ptr),595                                     (store_local_m0 node:$value, node:$ptr)>,596                            Aligned<16> {597  let IsStore = 1;598}599 600let PredicateCode = [{return cast<MemSDNode>(N)->getAlign() < 4;}],601    GISelPredicateCode = [{return (*MI.memoperands_begin())->getAlign() < 4;}],602    AddressSpaces = [ AddrSpaces.Local ] in {603def load_align_less_than_4_local : PatFrag<(ops node:$ptr),604                                           (load_local node:$ptr)> {605  let IsLoad = 1;606  let IsNonExtLoad = 1;607}608 609def load_align_less_than_4_local_m0 : PatFrag<(ops node:$ptr),610                                              (load_local_m0 node:$ptr)> {611  let IsLoad = 1;612  let IsNonExtLoad = 1;613}614 615def store_align_less_than_4_local : PatFrag <(ops node:$value, node:$ptr),616                                             (store_local node:$value, node:$ptr)> {617  let IsStore = 1;618  let IsTruncStore = 0;619}620 621def store_align_less_than_4_local_m0 : PatFrag <(ops node:$value, node:$ptr),622                                                (store_local_m0 node:$value, node:$ptr)> {623  let IsStore = 1;624  let IsTruncStore = 0;625}626}627 628def atomic_store_8_glue : PatFrag <629  (ops node:$ptr, node:$value),630  (AMDGPUatomic_st_glue node:$ptr, node:$value)> {631  let IsAtomic = 1;632  let MemoryVT = i8;633}634 635def atomic_store_16_glue : PatFrag <636  (ops node:$ptr, node:$value),637  (AMDGPUatomic_st_glue node:$ptr, node:$value)> {638  let IsAtomic = 1;639  let MemoryVT = i16;640}641 642def atomic_store_32_glue : PatFrag <643  (ops node:$ptr, node:$value),644  (AMDGPUatomic_st_glue node:$ptr, node:$value)> {645  let IsAtomic = 1;646  let MemoryVT = i32;647}648 649def atomic_store_64_glue : PatFrag <650  (ops node:$ptr, node:$value),651  (AMDGPUatomic_st_glue node:$ptr, node:$value)> {652  let IsAtomic = 1;653  let MemoryVT = i64;654}655 656let IsAtomic = 1, AddressSpaces = StoreAddress_local.AddrSpaces in {657def atomic_store_8_local_m0 : PatFrag<(ops node:$val, node:$ptr),658                                       (atomic_store_8_glue node:$val, node:$ptr)>;659def atomic_store_16_local_m0 : PatFrag<(ops node:$val, node:$ptr),660                                       (atomic_store_16_glue node:$val, node:$ptr)>;661def atomic_store_32_local_m0 : PatFrag<(ops node:$val, node:$ptr),662                                       (atomic_store_32_glue node:$val, node:$ptr)>;663def atomic_store_64_local_m0 : PatFrag<(ops node:$val, node:$ptr),664                                       (atomic_store_64_glue node:$val, node:$ptr)>;665} // End let IsAtomic = 1, AddressSpaces = StoreAddress_local.AddrSpaces666 667 668//===----------------------------------------------------------------------===//669// SDNodes PatFrags for a16 loads and stores with 3 components.670// v3f16/v3i16 is widened to v4f16/v4i16, so we need to match on the memory671// load/store size.672//===----------------------------------------------------------------------===//673 674class mubuf_intrinsic_load<SDPatternOperator name, ValueType vt> : PatFrag <675  (ops node:$rsrc, node:$vindex, node:$voffset, node:$soffset, node:$offset,676            node:$auxiliary, node:$idxen),677  (name node:$rsrc, node:$vindex, node:$voffset, node:$soffset, node:$offset,678            node:$auxiliary, node:$idxen)> {679  let IsLoad = 1;680  let MemoryVT = vt;681}682 683class mubuf_intrinsic_store<SDPatternOperator name, ValueType vt> : PatFrag <684  (ops node:$vdata, node:$rsrc, node:$vindex, node:$voffset, node:$soffset, node:$offset,685            node:$auxiliary, node:$idxen),686  (name node:$vdata, node:$rsrc, node:$vindex, node:$voffset, node:$soffset, node:$offset,687            node:$auxiliary, node:$idxen)> {688  let IsStore = 1;689  let MemoryVT = vt;690}691 692class mtbuf_intrinsic_load<SDPatternOperator name, ValueType vt> : PatFrag <693  (ops node:$rsrc, node:$vindex, node:$voffset, node:$soffset, node:$offset,694            node:$format, node:$auxiliary, node:$idxen),695  (name node:$rsrc, node:$vindex, node:$voffset, node:$soffset, node:$offset,696            node:$format, node:$auxiliary, node:$idxen)> {697  let IsLoad = 1;698  let MemoryVT = vt;699}700 701class mtbuf_intrinsic_store<SDPatternOperator name, ValueType vt> : PatFrag <702  (ops node:$vdata, node:$rsrc, node:$vindex, node:$voffset, node:$soffset, node:$offset,703            node:$format, node:$auxiliary, node:$idxen),704  (name node:$vdata, node:$rsrc, node:$vindex, node:$voffset, node:$soffset, node:$offset,705            node:$format, node:$auxiliary, node:$idxen)> {706  let IsStore = 1;707  let MemoryVT = vt;708}709 710//===----------------------------------------------------------------------===//711// SDNodes PatFrags for d16 loads712//===----------------------------------------------------------------------===//713 714class LoadD16Frag <SDPatternOperator op> : PatFrag<715  (ops node:$ptr, node:$tied_in),716  (op node:$ptr, node:$tied_in)> {717  let IsLoad = 1;718}719 720foreach as = [ "global", "flat", "constant", "local", "private", "region" ] in {721let AddressSpaces = !cast<AddressSpaceList>("LoadAddress_"#as).AddrSpaces in {722 723def load_d16_hi_#as : LoadD16Frag <SIload_d16_hi>;724 725def az_extloadi8_d16_hi_#as : LoadD16Frag <SIload_d16_hi_u8> {726  let MemoryVT = i8;727}728 729def sextloadi8_d16_hi_#as : LoadD16Frag <SIload_d16_hi_i8> {730  let MemoryVT = i8;731}732 733def load_d16_lo_#as : LoadD16Frag <SIload_d16_lo>;734 735def az_extloadi8_d16_lo_#as : LoadD16Frag <SIload_d16_lo_u8> {736  let MemoryVT = i8;737}738 739def sextloadi8_d16_lo_#as : LoadD16Frag <SIload_d16_lo_i8> {740  let MemoryVT = i8;741}742 743} // End let AddressSpaces = ...744} // End foreach AddrSpace745 746def lshr_rev : PatFrag <747  (ops node:$src1, node:$src0),748  (srl $src0, $src1)749>;750 751def ashr_rev : PatFrag <752  (ops node:$src1, node:$src0),753  (sra $src0, $src1)754>;755 756def lshl_rev : PatFrag <757  (ops node:$src1, node:$src0),758  (shl $src0, $src1)759>;760 761def add_ctpop : PatFrag <762  (ops node:$src0, node:$src1),763  (add (ctpop $src0), $src1)764>;765 766def xnor : PatFrag <767  (ops node:$src0, node:$src1),768  (not (xor $src0, $src1))769>;770 771foreach I = 1-4 in {772def shl#I#_add : PatFrag <773  (ops node:$src0, node:$src1),774  (add (shl_oneuse $src0, (i32 I)), $src1)>;775}776 777multiclass SIAtomicM0Glue2 <string op_name, bit is_amdgpu = 0,778                            SDTypeProfile tc = SDTAtomic2,779                            bit IsInt = 1> {780 781  def _glue : SDNode <782    !if(is_amdgpu, "AMDGPUISD", "ISD")#"::ATOMIC_"#op_name, tc,783    [SDNPHasChain, SDNPMayStore, SDNPMayLoad, SDNPMemOperand, SDNPInGlue]784  >;785 786  let AddressSpaces = StoreAddress_local.AddrSpaces in {787 788    if IsInt then {789      defm _local_m0 : binary_atomic_op <!cast<SDNode>(NAME#"_glue")>;790      defm _local_m0 : noret_binary_atomic_op <!cast<SDNode>(NAME#"_glue")>;791    } else {792      defm _local_m0 : binary_atomic_op_fp <!cast<SDNode>(NAME#"_glue")>;793      defm _local_m0 : noret_binary_atomic_op_fp <!cast<SDNode>(NAME#"_glue")>;794     }795  }796 797  let AddressSpaces = StoreAddress_region.AddrSpaces in {798    if IsInt then {799      defm _region_m0 : binary_atomic_op <!cast<SDNode>(NAME#"_glue")>;800      defm _region_m0 : noret_binary_atomic_op <!cast<SDNode>(NAME#"_glue")>;801    } else {802      defm _region_m0 : binary_atomic_op_fp <!cast<SDNode>(NAME#"_glue")>;803      defm _region_m0 : noret_binary_atomic_op_fp <!cast<SDNode>(NAME#"_glue")>;804    }805  }806}807 808defm atomic_load_add : SIAtomicM0Glue2 <"LOAD_ADD">;809defm atomic_load_sub : SIAtomicM0Glue2 <"LOAD_SUB">;810defm atomic_load_uinc_wrap : SIAtomicM0Glue2 <"LOAD_UINC_WRAP">;811defm atomic_load_udec_wrap : SIAtomicM0Glue2 <"LOAD_UDEC_WRAP">;812defm atomic_load_and : SIAtomicM0Glue2 <"LOAD_AND">;813defm atomic_load_min : SIAtomicM0Glue2 <"LOAD_MIN">;814defm atomic_load_max : SIAtomicM0Glue2 <"LOAD_MAX">;815defm atomic_load_or : SIAtomicM0Glue2 <"LOAD_OR">;816defm atomic_load_xor : SIAtomicM0Glue2 <"LOAD_XOR">;817defm atomic_load_umin : SIAtomicM0Glue2 <"LOAD_UMIN">;818defm atomic_load_umax : SIAtomicM0Glue2 <"LOAD_UMAX">;819defm atomic_swap : SIAtomicM0Glue2 <"SWAP">;820defm atomic_load_fadd : SIAtomicM0Glue2 <"LOAD_FADD", 0, SDTAtomic2_f32, 0>;821defm atomic_load_fmin : SIAtomicM0Glue2 <"LOAD_FMIN", 0, SDTAtomic2_f32, 0>;822defm atomic_load_fmax : SIAtomicM0Glue2 <"LOAD_FMAX", 0, SDTAtomic2_f32, 0>;823 824def as_i1timm : SDNodeXForm<timm, [{825  return CurDAG->getTargetConstant(N->getZExtValue(), SDLoc(N), MVT::i1);826}]>;827 828def as_i1timm_zext : SDNodeXForm<timm, [{829  return CurDAG->getTargetConstant(N->getZExtValue(), SDLoc(N), MVT::i32);830}]>;831 832def as_i8imm : SDNodeXForm<imm, [{833  return CurDAG->getTargetConstant(N->getZExtValue(), SDLoc(N), MVT::i8);834}]>;835 836def as_i8timm : SDNodeXForm<timm, [{837  return CurDAG->getSignedTargetConstant(N->getSExtValue(), SDLoc(N), MVT::i16);838}]>;839 840def as_i16imm : SDNodeXForm<imm, [{841  return CurDAG->getSignedTargetConstant(N->getSExtValue(), SDLoc(N), MVT::i16);842}]>;843 844def as_i16timm : SDNodeXForm<timm, [{845  // Explicit cast, as this is used with both signed and unsigned immediates.846  return CurDAG->getSignedTargetConstant(int16_t(N->getSExtValue()), SDLoc(N),847                                         MVT::i16);848}]>;849 850def as_i32imm: SDNodeXForm<imm, [{851  return CurDAG->getSignedTargetConstant(N->getSExtValue(), SDLoc(N), MVT::i32);852}]>;853 854def as_i32timm: SDNodeXForm<timm, [{855  return CurDAG->getSignedTargetConstant(N->getSExtValue(), SDLoc(N), MVT::i32);856}]>;857 858def as_i64imm: SDNodeXForm<imm, [{859  return CurDAG->getSignedTargetConstant(N->getSExtValue(), SDLoc(N), MVT::i64);860}]>;861 862def cond_as_i32imm: SDNodeXForm<cond, [{863  return CurDAG->getTargetConstant(N->get(), SDLoc(N), MVT::i32);864}]>;865 866// Copied from the AArch64 backend:867def bitcast_fpimm_to_i16 : SDNodeXForm<fpimm, [{868return CurDAG->getTargetConstant(869  N->getValueAPF().bitcastToAPInt().getZExtValue(), SDLoc(N), MVT::i16);870}]>;871 872def bitcast_fpimm_to_i32 : SDNodeXForm<fpimm, [{873return CurDAG->getTargetConstant(874  N->getValueAPF().bitcastToAPInt().getZExtValue(), SDLoc(N), MVT::i32);875}]>;876 877def frameindex_to_targetframeindex : SDNodeXForm<frameindex, [{878  auto FI = cast<FrameIndexSDNode>(N);879  return CurDAG->getTargetFrameIndex(FI->getIndex(), MVT::i32);880}]>;881 882// Copied from the AArch64 backend:883def bitcast_fpimm_to_i64 : SDNodeXForm<fpimm, [{884return CurDAG->getTargetConstant(885  N->getValueAPF().bitcastToAPInt().getZExtValue(), SDLoc(N), MVT::i64);886}]>;887 888def as_hw_round_mode : SDNodeXForm<timm, [{889  // "round.towardzero" -> TowardZero 0        -> FP_ROUND_ROUND_TO_ZERO 3890  // "round.tonearest"  -> NearestTiesToEven 1 -> FP_ROUND_ROUND_TO_NEAREST 0891  // "round.upward"     -> TowardPositive 2    -> FP_ROUND_ROUND_TO_INF 1892  // "round.downward    -> TowardNegative 3    -> FP_ROUND_ROUND_TO_NEGINF 2893  return CurDAG->getTargetConstant((N->getSExtValue() + 3) % 4, SDLoc(N),894                                    MVT::i32);895}]>;896 897def SupportedRoundMode : TImmLeaf<i32, [{898  return Imm == (int)RoundingMode::TowardZero ||899         Imm == (int)RoundingMode::NearestTiesToEven ||900         Imm == (int)RoundingMode::TowardPositive ||901         Imm == (int)RoundingMode::TowardNegative;902}]>;903 904def VOP3PModsNeg : SDNodeXForm<timm, [{905  unsigned Mods = SISrcMods::OP_SEL_1;906  if (N->getZExtValue())907    Mods ^= SISrcMods::NEG;908  return CurDAG->getTargetConstant(Mods, SDLoc(N), MVT::i32);909}]>;910 911def VOP3PModsNegs : SDNodeXForm<timm, [{912  unsigned Mods = SISrcMods::OP_SEL_1;913  if (N->getZExtValue())914    Mods ^= (SISrcMods::NEG | SISrcMods::NEG_HI);915  return CurDAG->getTargetConstant(Mods, SDLoc(N), MVT::i32);916}]>;917 918def VOP3PModsNegAbs : SDNodeXForm<timm, [{919  unsigned Val = N->getZExtValue();920  unsigned Mods = SISrcMods::OP_SEL_1; // default: none921  if (Val == 1) // neg922    Mods ^= SISrcMods::NEG;923  if (Val == 2) // abs924    Mods ^= SISrcMods::ABS;925  if (Val == 3) // neg and abs926    Mods ^= (SISrcMods::NEG | SISrcMods::ABS);927  return CurDAG->getTargetConstant(Mods, SDLoc(N), MVT::i32);928}]>;929 930class bitextract_imm<int bitnum> : SDNodeXForm<imm, [{931  uint64_t Imm = N->getZExtValue();932  unsigned Bit = (Imm >> }] # bitnum # [{ ) & 1;933  return CurDAG->getTargetConstant(Bit, SDLoc(N), MVT::i1);934}]>;935 936def SIMM16bit : TImmLeaf <i32,937  [{return isInt<16>(Imm) || isUInt<16>(Imm);}],938  as_i16timm939>;940 941def i64imm_32bit : ImmLeaf<i64, [{942  return (Imm & 0xffffffffULL) == static_cast<uint64_t>(Imm);943}]>;944 945def InlineImm64 : IntImmLeaf<i64, [{946  return isInlineImmediate(Imm);947}]>;948 949def InlineImmFP32 : FPImmLeaf<f32, [{950  return isInlineImmediate(Imm);951}]>;952 953def InlineImmFP64 : FPImmLeaf<f64, [{954  return isInlineImmediate(Imm);955}]>;956 957 958class VGPRImm <dag frag> : PatLeaf<frag, [{959  return isVGPRImm(N);960}]> {961  let GISelPredicateCode = [{return true;}];962}963 964def NegateImm : SDNodeXForm<imm, [{965  return CurDAG->getSignedConstant(-N->getSExtValue(), SDLoc(N), MVT::i32);966}]>;967 968// TODO: When FP inline imm values work?969def NegSubInlineConst32 : ImmLeaf<i32, [{970  return Imm < -16 && Imm >= -64;971}], NegateImm>;972 973def NegSubInlineIntConst16 : ImmLeaf<i16, [{974  return Imm < -16 && Imm >= -64;975}], NegateImm>;976 977def ShiftAmt32Imm : ImmLeaf <i32, [{978  return Imm < 32;979}]>;980 981def fp16_zeros_high_16bits : PatLeaf<(f16 VGPR_32:$src), [{982  return fp16SrcZerosHighBits(N->getOpcode());983}]>;984 985def MFMALdScaleXForm : SDNodeXForm<timm, [{986  unsigned Val = N->getZExtValue();987  unsigned New = 0;988  if (Val & 0x1)989    New |= SISrcMods::OP_SEL_0;990  if (Val & 0x2)991    New |= SISrcMods::OP_SEL_1;992  return CurDAG->getTargetConstant(New, SDLoc(N), MVT::i32);993}]>;994 995def is_canonicalized : PatLeaf<(fAny srcvalue:$src), [{996  const SITargetLowering &Lowering =997      *static_cast<const SITargetLowering *>(getTargetLowering());998  return Lowering.isCanonicalized(*CurDAG, Op);999}]> {1000  let GISelPredicateCode = [{1001    const SITargetLowering *TLI = static_cast<const SITargetLowering *>(1002        MF.getSubtarget().getTargetLowering());1003    const MachineOperand &Dst = MI.getOperand(0);1004    assert(Dst.isDef());1005    return TLI->isCanonicalized(Dst.getReg(), MF);1006   }];1007}1008 1009//===----------------------------------------------------------------------===//1010// MUBUF/SMEM Patterns1011//===----------------------------------------------------------------------===//1012 1013def extract_cpol : SDNodeXForm<timm, [{1014  return CurDAG->getTargetConstant(1015      N->getZExtValue() & (Subtarget->getGeneration() >= AMDGPUSubtarget::GFX121016                               ? AMDGPU::CPol::ALL1017                               : AMDGPU::CPol::ALL_pregfx12),1018      SDLoc(N), MVT::i8);1019}]>;1020 1021def extract_swz : SDNodeXForm<timm, [{1022  const bool Swizzle =1023      N->getZExtValue() & (Subtarget->getGeneration() >= AMDGPUSubtarget::GFX121024                               ? AMDGPU::CPol::SWZ1025                               : AMDGPU::CPol::SWZ_pregfx12);1026  return CurDAG->getTargetConstant(Swizzle, SDLoc(N), MVT::i8);1027}]>;1028 1029def extract_cpol_set_glc : SDNodeXForm<timm, [{1030  const uint32_t cpol = N->getZExtValue() & (Subtarget->getGeneration() >= AMDGPUSubtarget::GFX121031                               ? AMDGPU::CPol::ALL1032                               : AMDGPU::CPol::ALL_pregfx12);1033  return CurDAG->getTargetConstant(cpol | AMDGPU::CPol::GLC, SDLoc(N), MVT::i8);1034}]>;1035 1036//===----------------------------------------------------------------------===//1037// Custom Operands1038//===----------------------------------------------------------------------===//1039 1040def SOPPBrTarget : CustomOperand<OtherVT> {1041  let PrintMethod = "printOperand";1042  let EncoderMethod = "getSOPPBrEncoding";1043  let DecoderMethod = "decodeSOPPBrTarget";1044  let OperandType = "OPERAND_PCREL";1045}1046 1047def si_ga : Operand<iPTR>;1048 1049def InterpSlot : CustomOperand<i32>;1050 1051// It appears to be necessary to create a separate operand for this to1052// be able to parse attr<num> with no space.1053def InterpAttr : CustomOperand<i32>;1054 1055def InterpAttrChan : ImmOperand<i32>;1056 1057def SplitBarrier : ImmOperand<i32> {1058  let OperandNamespace = "AMDGPU";1059  let OperandType = "OPERAND_INLINE_SPLIT_BARRIER_INT32";1060  let DecoderMethod = "decodeSplitBarrier";1061  let PrintMethod = "printOperand";1062}1063 1064// Pseudo-operand type. This is a pair of 32-bit inline constants1065// packed into a single 64-bit value.1066def AV_64_PSEUDO_IMM : Operand<i64> {1067  let OperandNamespace = "AMDGPU";1068  let OperandType = "OPERAND_INLINE_C_AV64_PSEUDO";1069}1070 1071def VReg32OrOffClass : AsmOperandClass {1072  let Name = "VReg32OrOff";1073  let ParserMethod = "parseVReg32OrOff";1074}1075 1076def SendMsg : CustomOperand<i32>;1077 1078def Swizzle : CustomOperand<i16, 1>;1079 1080def Endpgm : CustomOperand<i16, 1>;1081 1082def SWaitCnt : CustomOperand<i32>;1083 1084def DepCtr : CustomOperand<i32>;1085 1086def SDelayALU : CustomOperand<i32>;1087 1088include "SIInstrFormats.td"1089include "VIInstrFormats.td"1090 1091def BoolReg : AsmOperandClass {1092  let Name = "BoolReg";1093  let ParserMethod = "parseBoolReg";1094  let RenderMethod = "addRegOperands";1095}1096 1097class BoolRC : RegisterOperand<SReg_1> {1098  let ParserMatchClass = BoolReg;1099  let DecoderMethod = "decodeBoolReg";1100}1101 1102def SSrc_i1 : RegisterOperand<SReg_1_XEXEC> {1103  let ParserMatchClass = BoolReg;1104  let DecoderMethod = "decodeBoolReg";1105}1106 1107def VOPDstS64orS32 : BoolRC {1108  let PrintMethod = "printVOPDst";1109}1110 1111// SCSrc_i1 is the operand for pseudo instructions only.1112// Boolean immediates shall not be exposed to codegen instructions.1113def SCSrc_i1 : RegisterOperand<SReg_1_XEXEC> {1114  let OperandNamespace = "AMDGPU";1115  let OperandType = "OPERAND_REG_IMM_INT32";1116  let ParserMatchClass = BoolReg;1117  let DecoderMethod = "decodeBoolReg";1118}1119 1120// ===----------------------------------------------------------------------===//1121// ExpSrc* Special cases for exp src operands which are printed as1122// "off" depending on en operand.1123// ===----------------------------------------------------------------------===//1124 1125def ExpSrc0 : RegisterOperand<VGPR_32> {1126  let PrintMethod = "printExpSrc0";1127  let ParserMatchClass = VReg32OrOffClass;1128}1129 1130def ExpSrc1 : RegisterOperand<VGPR_32> {1131  let PrintMethod = "printExpSrc1";1132  let ParserMatchClass = VReg32OrOffClass;1133}1134 1135def ExpSrc2 : RegisterOperand<VGPR_32> {1136  let PrintMethod = "printExpSrc2";1137  let ParserMatchClass = VReg32OrOffClass;1138}1139 1140def ExpSrc3 : RegisterOperand<VGPR_32> {1141  let PrintMethod = "printExpSrc3";1142  let ParserMatchClass = VReg32OrOffClass;1143}1144 1145// FIXME: Should change class based on hasSDWAScalar to exclude SGPRs1146class SDWASrc<ValueType vt> : RegisterOperand<VS_32> {1147  let OperandNamespace = "AMDGPU";1148  string Type = !if(vt.isFP, "FP", "INT");1149  let OperandType = "OPERAND_REG_INLINE_C_"#Type#vt.Size;1150  let DecoderMethod = "decodeSDWASrc"#vt.Size;1151  let EncoderMethod = "getSDWASrcEncoding";1152}1153 1154def SDWASrc_i32 : SDWASrc<i32>;1155def SDWASrc_i16 : SDWASrc<i16>;1156def SDWASrc_f32 : SDWASrc<f32>;1157def SDWASrc_f16 : SDWASrc<f16>;1158 1159def SDWAVopcDst : BoolRC {1160  let OperandNamespace = "AMDGPU";1161  let OperandType = "OPERAND_SDWA_VOPC_DST";1162  let EncoderMethod = "getSDWAVopcDstEncoding";1163  let DecoderMethod = "decodeSDWAVopcDst";1164  let PrintMethod = "printVOPDst";1165}1166 1167class NamedIntOperand<string prefix, bit Optional = 1, string name = NAME>1168    : CustomOperand<i32, Optional, name> {1169  string Prefix = prefix;1170 1171  let PredicateMethod =1172    "getPredicate([](const AMDGPUOperand &Op) -> bool { "#1173    "return Op.isImmTy(AMDGPUOperand::"#ImmTy#"); })";1174 1175  string Validator = "[](int64_t V) { return true; }";1176  string ConvertMethod = "[](int64_t &V) { return "#Validator#"(V); }";1177  let ParserMethod =1178    "[this](OperandVector &Operands) -> ParseStatus { "#1179    "return parseIntWithPrefix(\""#Prefix#"\", Operands, "#1180    "AMDGPUOperand::"#ImmTy#", "#ConvertMethod#"); }";1181 1182  bit PrintInHex = 0;1183  bit AlwaysPrint = 0;1184  let PrintMethod = "[this](const MCInst *MI, unsigned OpNo, "1185                    "const MCSubtargetInfo &STI, raw_ostream &O) { "1186                    "printNamedInt(MI, OpNo, STI, O, \""#Prefix#"\", "#1187                    !if(PrintInHex, "true", "false")#", "#1188                    !if(AlwaysPrint, "true", "false")#"); }";1189}1190 1191class NamedBitOperand<string Id, string Name = NAME>1192    : CustomOperand<i1, 1, Name> {1193  let PredicateMethod = "isImmTy<AMDGPUOperand::"#ImmTy#">";1194  let ParserMethod =1195    "[this](OperandVector &Operands) -> ParseStatus { "#1196    "return parseNamedBit(\""#Id#"\", Operands, AMDGPUOperand::"#ImmTy#"); }";1197  let PrintMethod = "[this](const MCInst *MI, unsigned OpNo, "#1198    "const MCSubtargetInfo &STI, raw_ostream &O) { "#1199    "printNamedBit(MI, OpNo, O, \""#Id#"\"); }";1200}1201 1202class DefaultOperand<CustomOperand Op, int Value>1203  : OperandWithDefaultOps<Op.Type, (ops (Op.Type Value))>,1204    CustomOperandProps<1> {1205  let ParserMatchClass = Op.ParserMatchClass;1206  let PrintMethod = Op.PrintMethod;1207}1208 1209class SDWAOperand<string Id, string Name = NAME>1210    : CustomOperand<i32, 1, Name> {1211  let ParserMethod =1212    "[this](OperandVector &Operands) -> ParseStatus { "#1213    "return parseSDWASel(Operands, \""#Id#"\", AMDGPUOperand::"#ImmTy#"); }";1214}1215 1216class ArrayOperand0<string Id, string Name = NAME>1217  : OperandWithDefaultOps<i32, (ops (i32 0))>,1218    CustomOperandProps<1, Name> {1219  let ParserMethod =1220    "[this](OperandVector &Operands) -> ParseStatus { "#1221    "return parseOperandArrayWithPrefix(\""#Id#"\", Operands, "#1222    "AMDGPUOperand::"#ImmTy#"); }";1223}1224 1225let ImmTy = "ImmTyOffset" in1226def flat_offset : CustomOperand<i32, 1, "FlatOffset">;1227let PrintMethod = "printOffset" in1228def Offset : NamedIntOperand<"offset">;1229let Validator = "isUInt<8>" in {1230def Offset0 : NamedIntOperand<"offset0">;1231def Offset1 : NamedIntOperand<"offset1">;1232}1233 1234def gds : NamedBitOperand<"gds", "GDS">;1235 1236def omod : CustomOperand<i32, 1, "OModSI">;1237def omod0 : DefaultOperand<omod, 0>;1238 1239// We need to make the cases with a default of 0 distinct from no1240// default to help deal with some cases where the operand appears1241// before a mandatory operand.1242def Clamp : NamedBitOperand<"clamp">;1243def Clamp0 : DefaultOperand<Clamp, 0>;1244def highmod : NamedBitOperand<"high", "High">;1245 1246def CPol : CustomOperand<i32, 1>;1247def CPol_0 : DefaultOperand<CPol, 0>;1248def CPol_GLC1 : DefaultOperand<CPol, 1>;1249def CPol_GLC : ValuePredicatedOperand<CPol, "Op.getImm() & CPol::GLC">;1250def CPol_NonGLC : ValuePredicatedOperand<CPol, "!(Op.getImm() & CPol::GLC)", 1>;1251def CPol_GLC_WithDefault : DefaultOperand<CPol_GLC, !shl(1, CPolBit.GLC)>;1252def CPol_NonGLC_WithDefault : DefaultOperand<CPol_NonGLC, 0>;1253 1254def TFE : NamedBitOperand<"tfe">;1255def UNorm : NamedBitOperand<"unorm">;1256def DA : NamedBitOperand<"da">;1257def R128A16 : CustomOperand<i1, 1>;1258def A16 : NamedBitOperand<"a16">;1259def D16 : NamedBitOperand<"d16">;1260def LWE : NamedBitOperand<"lwe">;1261def exp_compr : NamedBitOperand<"compr", "ExpCompr">;1262def exp_vm : NamedBitOperand<"vm", "ExpVM">;1263 1264def FORMAT : CustomOperand<i8>;1265 1266let PrintInHex = 1 in1267def DMask : NamedIntOperand<"dmask">;1268 1269def Dim : CustomOperand<i8, /*optional=*/1>;1270 1271def dst_sel : SDWAOperand<"dst_sel", "SDWADstSel">;1272def src0_sel : SDWAOperand<"src0_sel", "SDWASrc0Sel">;1273def src1_sel : SDWAOperand<"src1_sel", "SDWASrc1Sel">;1274def dst_unused : CustomOperand<i32, 1, "SDWADstUnused">;1275 1276def op_sel0 : ArrayOperand0<"op_sel", "OpSel">;1277def op_sel_hi0 : ArrayOperand0<"op_sel_hi", "OpSelHi">;1278def neg_lo0 : ArrayOperand0<"neg_lo", "NegLo">;1279def neg_hi0 : ArrayOperand0<"neg_hi", "NegHi">;1280 1281def IndexKey32bit : CustomOperand<i32, 1>;1282def IndexKey16bit : CustomOperand<i32, 1>;1283def IndexKey8bit : CustomOperand<i32, 1>;1284 1285def dpp8 : CustomOperand<i32, 0, "DPP8">;1286def dpp_ctrl : CustomOperand<i32, 0, "DPPCtrl">;1287 1288let DefaultValue = "0xf", PrintInHex = 1, AlwaysPrint = 1 in {1289def DppRowMask : NamedIntOperand<"row_mask">;1290def DppBankMask : NamedIntOperand<"bank_mask">;1291}1292def DppBoundCtrl : NamedIntOperand<"bound_ctrl"> {1293  let ConvertMethod = "[this] (int64_t &BC) -> bool { return convertDppBoundCtrl(BC); }";1294  let PrintMethod = "printDppBoundCtrl";1295}1296 1297let DecoderMethod = "decodeDpp8FI", PrintMethod = "printDppFI" in1298def Dpp8FI : NamedIntOperand<"fi", 1, "DppFI">;1299let PrintMethod = "printDppFI" in1300def Dpp16FI : NamedIntOperand<"fi", 1, "DppFI">;1301 1302def blgp : CustomOperand<i32, 1, "BLGP">;1303def CBSZ : NamedIntOperand<"cbsz"> {1304  let Validator = "isUInt<3>";1305}1306def ABID : NamedIntOperand<"abid"> {1307  let Validator = "isUInt<4>";1308}1309def hwreg : CustomOperand<i32, 0, "Hwreg">;1310 1311def exp_tgt : CustomOperand<i32, 0, "ExpTgt">;1312 1313let AlwaysPrint = 1 in {1314def WaitVDST : NamedIntOperand<"wait_vdst"> {1315  let Validator = "isUInt<4>";1316}1317def WaitEXP : NamedIntOperand<"wait_exp"> {1318  let Validator = "isUInt<3>";1319}1320def WaitVAVDst : NamedIntOperand<"wait_va_vdst"> {1321  let Validator = "isUInt<4>";1322}1323def WaitVMVSrc : NamedIntOperand<"wait_vm_vsrc"> {1324  let Validator = "isUInt<1>";1325}1326} // End AlwaysPrint = 11327 1328def ByteSel : NamedIntOperand<"byte_sel"> {1329  let Validator = "isUInt<2>";1330}1331def ByteSel0 : DefaultOperand<ByteSel, 0>;1332 1333let PrintMethod = "printBitOp3" in1334def BitOp3 : NamedIntOperand<"bitop3">;1335def bitop3_0 : DefaultOperand<BitOp3, 0>;1336 1337def MatrixAFMT : CustomOperand<i32, 1, "MatrixAFMT">;1338def MatrixBFMT : CustomOperand<i32, 1, "MatrixBFMT">;1339 1340def MatrixAScale : CustomOperand<i32, 1, "MatrixAScale">;1341def MatrixBScale : CustomOperand<i32, 1, "MatrixBScale">;1342 1343def MatrixAScaleFmt : CustomOperand<i32, 1, "MatrixAScaleFmt">;1344def MatrixBScaleFmt : CustomOperand<i32, 1, "MatrixBScaleFmt">;1345 1346def MatrixAReuse : NamedBitOperand<"matrix_a_reuse">;1347def MatrixBReuse : NamedBitOperand<"matrix_b_reuse">;1348 1349def ScaleSel : NamedIntOperand<"scale_sel"> {1350  let Validator = "isUInt<4>";1351}1352 1353class KImmFPOperand<ValueType vt> : ImmOperand<vt> {1354  let OperandNamespace = "AMDGPU";1355  let OperandType = "OPERAND_KIMM"#vt.Size;1356  let PrintMethod = "printU"#vt.Size#"ImmOperand";1357  let DecoderMethod = "decodeOperand_KImmFP";1358}1359 1360// 32-bit VALU immediate operand that uses the constant bus.1361def KImmFP32 : KImmFPOperand<i32>;1362 1363// 32-bit VALU immediate operand with a 16-bit value that uses the1364// constant bus.1365def KImmFP16 : KImmFPOperand<i16>;1366 1367// 64-bit VALU immediate operand that uses the constant bus.1368def KImmFP64 : KImmFPOperand<i64> {1369  let DecoderMethod = "decodeOperand_KImmFP64";1370  let PrintMethod = "printFP64ImmOperand";1371}1372 1373class FPInputModsMatchClass <int opSize> : AsmOperandClass {1374  let Name = "RegOrImmWithFP"#opSize#"InputMods";1375  let ParserMethod = "parseRegOrImmWithFPInputMods";1376  let PredicateMethod = "isRegOrImmWithFP"#opSize#"InputMods";1377}1378 1379class FPVCSrcInputModsMatchClass <int opSize> : FPInputModsMatchClass <opSize> {1380  let Name = "RegOrInlineImmWithFP"#opSize#"InputMods";1381  let PredicateMethod = "isRegOrInlineImmWithFP"#opSize#"InputMods";1382}1383 1384class FPVRegSrcInputModsMatchClass <int opSize> : FPInputModsMatchClass <opSize> {1385  let Name = "VRegWithFP"#opSize#"InputMods";1386  let PredicateMethod = "isVRegWithFP"#opSize#"InputMods";1387}1388 1389def FP16InputModsMatchClass : FPInputModsMatchClass<16>;1390class FPT16InputModsMatchClass<bit IsFake16> : FPInputModsMatchClass<16> {1391  let Name = !if(IsFake16, "RegOrImmWithFPFake16InputMods",1392                 "RegOrImmWithFPT16InputMods");1393  let PredicateMethod = "isRegOrImmWithFPT16InputMods<" #1394                        !if(IsFake16, "true", "false") # ">";1395}1396def FP32InputModsMatchClass : FPInputModsMatchClass<32>;1397def FP64InputModsMatchClass : FPInputModsMatchClass<64>;1398 1399class FP16VCSrcInputModsMatchClass<bit IsFake16>1400    : FPVCSrcInputModsMatchClass<16> {1401  let Name = !if(IsFake16, "RegOrInlineImmWithFPFake16InputMods",1402                 "RegOrInlineImmWithFPT16InputMods");1403  let PredicateMethod = "isRegOrInlineImmWithFP16InputMods<" #1404                        !if(IsFake16, "true", "false") # ">";1405}1406def FP32VCSrcInputModsMatchClass : FPVCSrcInputModsMatchClass<32>;1407def FP64VCSrcInputModsMatchClass : FPVCSrcInputModsMatchClass<64>;1408 1409def FP32VRegSrcInputModsMatchClass : FPVRegSrcInputModsMatchClass<32>;1410def FP64VRegSrcInputModsMatchClass : FPVRegSrcInputModsMatchClass<64>;1411 1412class InputMods <AsmOperandClass matchClass> : Operand <i32> {1413  let OperandNamespace = "AMDGPU";1414  let OperandType = "OPERAND_INPUT_MODS";1415  let ParserMatchClass = matchClass;1416}1417 1418class FPInputMods <FPInputModsMatchClass matchClass> : InputMods <matchClass> {1419  let PrintMethod = "printOperandAndFPInputMods";1420}1421 1422def FP16InputMods : FPInputMods<FP16InputModsMatchClass>;1423class FPT16InputMods<bit IsFake16> : FPInputMods<FPT16InputModsMatchClass<IsFake16>> {1424  let EncoderMethod = "getMachineOpValueT16";1425}1426def FP32InputMods : FPInputMods<FP32InputModsMatchClass>;1427def FP32T16DstInputMods : FPInputMods<FP32InputModsMatchClass> {1428  let EncoderMethod = "getMachineOpValueT16";1429}1430def FP64InputMods : FPInputMods<FP64InputModsMatchClass>;1431 1432class FPT16VCSrcInputMods<bit IsFake16 = 1>1433  : FPInputMods<FP16VCSrcInputModsMatchClass<IsFake16>> {1434  let EncoderMethod = "getMachineOpValueT16";1435}1436def FP32VCSrcInputMods : FPInputMods<FP32VCSrcInputModsMatchClass>;1437def FP64VCSrcInputMods : FPInputMods<FP64VCSrcInputModsMatchClass>;1438 1439def FP32VRegSrcInputMods : FPInputMods<FP32VRegSrcInputModsMatchClass>;1440def FP64VRegSrcInputMods : FPInputMods<FP64VRegSrcInputModsMatchClass>;1441 1442class IntInputModsMatchClass <int opSize> : AsmOperandClass {1443  let Name = "RegOrImmWithInt"#opSize#"InputMods";1444  let ParserMethod = "parseRegOrImmWithIntInputMods";1445  let PredicateMethod = "isRegOrImmWithInt"#opSize#"InputMods";1446}1447class IntVCSrcInputModsMatchClass <int opSize> : IntInputModsMatchClass <opSize> {1448  let Name = "RegOrInlineImmWithInt"#opSize#"InputMods";1449  let PredicateMethod = "isRegOrInlineImmWithInt"#opSize#"InputMods";1450}1451class IntT16InputModsMatchClass<bit IsFake16> : IntInputModsMatchClass<16> {1452  let Name = !if(IsFake16, "RegOrImmWithIntFake16InputMods",1453                 "RegOrImmWithIntT16InputMods");1454  let PredicateMethod = "isRegOrImmWithIntT16InputMods<" #1455                        !if(IsFake16, "true", "false") # ">";1456}1457def Int32InputModsMatchClass : IntInputModsMatchClass<32>;1458def Int64InputModsMatchClass : IntInputModsMatchClass<64>;1459def Int32VCSrcInputModsMatchClass : IntVCSrcInputModsMatchClass<32>;1460class IntT16VCSrcInputModsMatchClass<bit IsFake16> : IntInputModsMatchClass<16> {1461  let Name = !if(IsFake16, "RegOrInlineImmWithIntFake16InputMods",1462                 "RegOrInlineImmWithIntT16InputMods");1463  let PredicateMethod = "isRegOrInlineImmWithIntT16InputMods<" #1464                        !if(IsFake16, "true", "false") # ">";1465}1466 1467class IntInputMods <IntInputModsMatchClass matchClass> : InputMods <matchClass> {1468  let PrintMethod = "printOperandAndIntInputMods";1469}1470class IntT16InputMods<bit IsFake16> : IntInputMods<IntT16InputModsMatchClass<IsFake16>> {1471  let EncoderMethod = "getMachineOpValueT16";1472}1473def Int32InputMods : IntInputMods<Int32InputModsMatchClass>;1474def Int32T16DstInputMods : IntInputMods<Int32InputModsMatchClass> {1475  let EncoderMethod = "getMachineOpValueT16";1476}1477def Int64InputMods : IntInputMods<Int64InputModsMatchClass>;1478def Int32VCSrcInputMods : IntInputMods<Int32VCSrcInputModsMatchClass>;1479class IntT16VCSrcInputMods<bit IsFake16 = 1>1480    : IntInputMods<IntT16VCSrcInputModsMatchClass<IsFake16>> {1481  let EncoderMethod = "getMachineOpValueT16";1482}1483 1484class OpSelModsMatchClass : AsmOperandClass {1485  let Name = "OpSelMods";1486  let ParserMethod = "parseRegOrImm";1487  let PredicateMethod = "isRegOrImm";1488}1489 1490def IntOpSelModsMatchClass : OpSelModsMatchClass;1491def IntOpSelMods : InputMods<IntOpSelModsMatchClass>;1492 1493class FPSDWAInputModsMatchClass <int opSize> : AsmOperandClass {1494  let Name = "SDWAWithFP"#opSize#"InputMods";1495  let ParserMethod = "parseRegOrImmWithFPInputMods";1496  let PredicateMethod = "isSDWAFP"#opSize#"Operand";1497}1498 1499def FP16SDWAInputModsMatchClass : FPSDWAInputModsMatchClass<16>;1500def FP32SDWAInputModsMatchClass : FPSDWAInputModsMatchClass<32>;1501 1502class FPSDWAInputMods <FPSDWAInputModsMatchClass matchClass> :1503  InputMods <matchClass> {1504  let PrintMethod = "printOperandAndFPInputMods";1505}1506 1507def FP16SDWAInputMods : FPSDWAInputMods<FP16SDWAInputModsMatchClass>;1508def FP32SDWAInputMods : FPSDWAInputMods<FP32SDWAInputModsMatchClass>;1509 1510def FPVRegInputModsMatchClass : AsmOperandClass {1511  let Name = "VRegWithFPInputMods";1512  let ParserMethod = "parseRegWithFPInputMods";1513  let PredicateMethod = "isVRegWithInputMods";1514}1515 1516def FPVRegInputMods : InputMods <FPVRegInputModsMatchClass> {1517  let PrintMethod = "printOperandAndFPInputMods";1518}1519 1520def FPVRegT16DstInputMods : InputMods <FPVRegInputModsMatchClass> {1521  let PrintMethod = "printOperandAndFPInputMods";1522  let EncoderMethod = "getMachineOpValueT16";1523}1524 1525class FPT16_Lo128VRegInputModsMatchClass<bit IsFake16> : AsmOperandClass {1526  let Name = !if(IsFake16, "Fake16_Lo128VRegWithFPInputMods",1527                 "T16_Lo128VRegWithFPInputMods");1528  let ParserMethod = "parseRegWithFPInputMods";1529  let PredicateMethod = "isT16_Lo128VRegWithInputMods<" #1530                        !if(IsFake16, "true", "false") # ">";1531}1532 1533class FPT16VRegInputModsMatchClass<bit IsFake16> : AsmOperandClass {1534  let Name = !if(IsFake16, "Fake16VRegWithFPInputMods",1535                 "T16VRegWithFPInputMods");1536  let ParserMethod = "parseRegWithFPInputMods";1537  let PredicateMethod = "isT16VRegWithInputMods<" #1538                        !if(IsFake16, "true", "false") # ">";1539}1540 1541class FPT16_Lo128VRegInputMods<bit IsFake16 = 1>1542    : InputMods <FPT16_Lo128VRegInputModsMatchClass<IsFake16>> {1543  let PrintMethod = "printOperandAndFPInputMods";1544  let EncoderMethod = "getMachineOpValueT16Lo128";1545}1546 1547class FPT16VRegInputMods<bit IsFake16 = 1>1548    : InputMods <FPT16VRegInputModsMatchClass<IsFake16>> {1549  let PrintMethod = "printOperandAndFPInputMods";1550  let EncoderMethod = "getMachineOpValueT16";1551}1552 1553class IntSDWAInputModsMatchClass <int opSize> : AsmOperandClass {1554  let Name = "SDWAWithInt"#opSize#"InputMods";1555  let ParserMethod = "parseRegOrImmWithIntInputMods";1556  let PredicateMethod = "isSDWAInt"#opSize#"Operand";1557}1558 1559def Int16SDWAInputModsMatchClass : IntSDWAInputModsMatchClass<16>;1560def Int32SDWAInputModsMatchClass : IntSDWAInputModsMatchClass<32>;1561def Bin32SDWAInputModsMatchClass : IntSDWAInputModsMatchClass<32> {1562  let Name = "SDWAWithBin32InputMods";1563  let ParserMethod = "parseRegOrImm";1564}1565 1566class IntSDWAInputMods <IntSDWAInputModsMatchClass matchClass> :1567  InputMods <matchClass> {1568  let PrintMethod = "printOperandAndIntInputMods";1569}1570 1571def Int16SDWAInputMods : IntSDWAInputMods<Int16SDWAInputModsMatchClass>;1572def Int32SDWAInputMods : IntSDWAInputMods<Int32SDWAInputModsMatchClass>;1573def Bin32SDWAInputMods : IntSDWAInputMods<Bin32SDWAInputModsMatchClass>;1574 1575def IntVRegInputModsMatchClass : AsmOperandClass {1576  let Name = "VRegWithIntInputMods";1577  let ParserMethod = "parseRegWithIntInputMods";1578  let PredicateMethod = "isVRegWithInputMods";1579}1580 1581class IntT16_Lo128VRegInputModsMatchClass<bit IsFake16 = 1> : AsmOperandClass {1582  let Name = !if(IsFake16, "Fake16_Lo128VRegWithIntInputMods",1583                 "T16_Lo128VRegWithIntInputMods");1584  let ParserMethod = "parseRegWithIntInputMods";1585  let PredicateMethod = "isT16_Lo128VRegWithInputMods<" #1586                        !if(IsFake16, "true", "false") # ">";1587}1588 1589class IntT16VRegInputModsMatchClass<bit IsFake16 = 1> : AsmOperandClass {1590  let Name = !if(IsFake16, "Fake16VRegWithIntInputMods",1591                 "T16VRegWithIntInputMods");1592  let ParserMethod = "parseRegWithIntInputMods";1593  let PredicateMethod = "isT16VRegWithInputMods<" #1594                        !if(IsFake16, "true", "false") # ">";1595}1596 1597class IntT16_Lo128VRegInputMods<bit IsFake16 = 1>1598    : InputMods <IntT16_Lo128VRegInputModsMatchClass<IsFake16>> {1599  let PrintMethod = "printOperandAndIntInputMods";1600  let EncoderMethod = "getMachineOpValueT16Lo128";1601}1602 1603class IntT16VRegInputMods<bit IsFake16 = 1>1604    : InputMods <IntT16VRegInputModsMatchClass<IsFake16>> {1605  let PrintMethod = "printOperandAndIntInputMods";1606  let EncoderMethod = "getMachineOpValueT16";1607}1608 1609def IntVRegInputMods : InputMods <IntVRegInputModsMatchClass> {1610  let PrintMethod = "printOperandAndIntInputMods";1611}1612 1613def IntVRegT16DstInputMods : InputMods <IntVRegInputModsMatchClass> {1614  let PrintMethod = "printOperandAndIntInputMods";1615  let EncoderMethod = "getMachineOpValueT16";1616}1617 1618class PackedFPInputModsMatchClass <int opSize> : AsmOperandClass {1619  let Name = "PackedFP"#opSize#"InputMods";1620  let ParserMethod = "parseRegOrImmWithFPInputMods";1621  let PredicateMethod = "isPackedFP"#opSize#"InputMods";1622}1623 1624class PackedVGPRFPInputModsMatchClass <int opSize> : PackedFPInputModsMatchClass<opSize> {1625  let PredicateMethod = "isPackedVGPRFP"#opSize#"InputMods";1626}1627 1628class PackedIntInputModsMatchClass <int opSize> : AsmOperandClass {1629  let Name = "PackedInt"#opSize#"InputMods";1630  let ParserMethod = "parseRegOrImm";1631  let PredicateMethod = "isRegOrImm";1632//  let PredicateMethod = "isPackedInt"#opSize#"InputMods";1633}1634 1635def PackedF16InputModsMatchClass : PackedFPInputModsMatchClass<16>;1636def PackedI16InputModsMatchClass : PackedIntInputModsMatchClass<16>;1637def PackedVGPRF32InputModsMatchClass : PackedVGPRFPInputModsMatchClass<32>;1638 1639class PackedFPInputMods <PackedFPInputModsMatchClass matchClass> : InputMods <matchClass> {1640  let PrintMethod = "printOperandAndFPInputMods";1641}1642 1643class PackedIntInputMods <PackedIntInputModsMatchClass matchClass> : InputMods <matchClass> {1644  //let PrintMethod = "printPackedIntInputMods";1645}1646 1647def PackedF16InputMods : PackedFPInputMods<PackedF16InputModsMatchClass>;1648def PackedI16InputMods : PackedIntInputMods<PackedI16InputModsMatchClass>;1649def PackedVGPRF32InputMods : PackedFPInputMods<PackedVGPRF32InputModsMatchClass>;1650 1651def MFMALdScaleModifierOp : TImmLeaf<i32, [{1652  return isUInt<2>(Imm);1653}], MFMALdScaleXForm>;1654 1655//===----------------------------------------------------------------------===//1656// Complex patterns1657//===----------------------------------------------------------------------===//1658 1659def DS1Addr1Offset : ComplexPattern<iPTR, 2, "SelectDS1Addr1Offset">;1660def DS64Bit4ByteAligned : ComplexPattern<iPTR, 3, "SelectDS64Bit4ByteAligned">;1661def DS128Bit8ByteAligned : ComplexPattern<iPTR, 3, "SelectDS128Bit8ByteAligned">;1662 1663def MOVRELOffset : ComplexPattern<iPTR, 2, "SelectMOVRELOffset">;1664 1665def VOP3Mods0 : ComplexPattern<untyped, 4, "SelectVOP3Mods0">;1666 1667// Modifiers for floating point instructions.1668def VOP3Mods  : ComplexPattern<untyped, 2, "SelectVOP3Mods">;1669 1670// VOP3 modifiers used for instructions that do not read canonicalized1671// floating point values (i.e. integer operations with FP source1672// modifiers)1673def VOP3ModsNonCanonicalizing : ComplexPattern<untyped, 2,1674  "SelectVOP3ModsNonCanonicalizing">;1675 1676def VOP3NoMods : ComplexPattern<untyped, 1, "SelectVOP3NoMods">;1677 1678def VOP3OMods : ComplexPattern<untyped, 3, "SelectVOP3OMods">;1679 1680def VOP3PMods  : ComplexPattern<untyped, 2, "SelectVOP3PMods">;1681 1682def VOP3PModsDOT  : ComplexPattern<untyped, 2, "SelectVOP3PModsDOT">;1683def WMMAOpSelVOP3PMods  : ComplexPattern<untyped, 1, "SelectWMMAOpSelVOP3PMods">;1684 1685def WMMAModsF32NegAbs  : ComplexPattern<untyped, 2, "SelectWMMAModsF32NegAbs">;1686def WMMAModsF16Neg  : ComplexPattern<untyped, 2, "SelectWMMAModsF16Neg">;1687def WMMAModsF16NegAbs  : ComplexPattern<untyped, 2, "SelectWMMAModsF16NegAbs">;1688def WMMAVISrc  : ComplexPattern<untyped, 1, "SelectWMMAVISrc">;1689def SWMMACIndex8  : ComplexPattern<untyped, 2, "SelectSWMMACIndex8">;1690def SWMMACIndex16  : ComplexPattern<untyped, 2, "SelectSWMMACIndex16">;1691def SWMMACIndex32  : ComplexPattern<untyped, 2, "SelectSWMMACIndex32">;1692 1693def VOP3OpSel  : ComplexPattern<untyped, 2, "SelectVOP3OpSel">;1694 1695def VOP3OpSelMods  : ComplexPattern<untyped, 2, "SelectVOP3OpSelMods">;1696 1697def VOP3PMadMixModsExt : ComplexPattern<untyped, 2, "SelectVOP3PMadMixModsExt">;1698def VOP3PMadMixMods : ComplexPattern<untyped, 2, "SelectVOP3PMadMixMods">;1699def VOP3PMadMixBF16ModsExt : ComplexPattern<untyped, 2, "SelectVOP3PMadMixBF16ModsExt">;1700def VOP3PMadMixBF16Mods : ComplexPattern<untyped, 2, "SelectVOP3PMadMixBF16Mods">;1701 1702def VINTERPMods  : ComplexPattern<untyped, 2, "SelectVINTERPMods">;1703def VINTERPModsHi  : ComplexPattern<untyped, 2, "SelectVINTERPModsHi">;1704 1705//===----------------------------------------------------------------------===//1706// SI assembler operands1707//===----------------------------------------------------------------------===//1708 1709def SIOperand {1710  int ZERO = 0x80;1711  int VCC = 0x6A;1712  int FLAT_SCR = 0x68;1713}1714 1715// This should be kept in sync with SISrcMods enum1716def SRCMODS {1717  int NONE = 0;1718  int NEG = 1;1719  int ABS = 2;1720  int NEG_ABS = 3;1721 1722  int NEG_HI = ABS;1723  int OP_SEL_0 = 4;1724  int OP_SEL_1 = 8;1725  int DST_OP_SEL = 8;1726}1727 1728def DSTCLAMP {1729  int NONE = 0;1730  int ENABLE = 1;1731}1732 1733def DSTOMOD {1734  int NONE = 0;1735}1736 1737def HWREG {1738  int MODE = 1;1739  int STATUS = 2;1740  int TRAPSTS = 3;1741  int HW_ID = 4;1742  int GPR_ALLOC = 5;1743  int LDS_ALLOC = 6;1744  int IB_STS = 7;1745  int MEM_BASES = 15;1746  int TBA_LO = 16;1747  int TBA_HI = 17;1748  int TMA_LO = 18;1749  int TMA_HI = 19;1750  int FLAT_SCR_LO = 20;1751  int FLAT_SCR_HI = 21;1752  int XNACK_MASK = 22;1753  int POPS_PACKER = 25;1754  int SHADER_CYCLES = 29;1755}1756 1757class getHwRegImm<int Reg, int Offset = 0, int Size = 32> {1758  int ret = !and(!or(Reg,1759                     !shl(Offset, 6),1760                     !shl(!add(Size, -1), 11)), 65535);1761}1762 1763//===----------------------------------------------------------------------===//1764//1765// SI Instruction multiclass helpers.1766//1767// Instructions with _32 take 32-bit operands.1768// Instructions with _64 take 64-bit operands.1769//1770// VOP_* instructions can use either a 32-bit or 64-bit encoding.  The 32-bit1771// encoding is the standard encoding, but instruction that make use of1772// any of the instruction modifiers must use the 64-bit encoding.1773//1774// Instructions with _e32 use the 32-bit encoding.1775// Instructions with _e64 use the 64-bit encoding.1776//1777//===----------------------------------------------------------------------===//1778 1779class SIMCInstr <string pseudo, int subtarget> {1780  string PseudoInstr = pseudo;1781  int Subtarget = subtarget;1782}1783 1784//===----------------------------------------------------------------------===//1785// Vector ALU classes1786//===----------------------------------------------------------------------===//1787 1788class getNumSrcArgs<ValueType Src0, ValueType Src1, ValueType Src2> {1789  int ret =1790    !if (!eq(Src0, untyped),      0,1791      !if (!eq(Src1, untyped),    1,   // VOP11792         !if (!eq(Src2, untyped), 2,   // VOP21793                                  3))); // VOP31794}1795 1796// Returns the register class to use for the destination of VOP[123C]1797// instructions for the given VT.1798class getVALUDstForVT<ValueType VT, bit IsTrue16 = 0, bit IsVOP3Encoding = 0> {1799  defvar op16 = !if(IsTrue16, !if (IsVOP3Encoding, VOPDstOperand_t16,1800                                   VOPDstOperand_t16Lo128),1801                    VOPDstOperand<VGPR_32>);1802  RegisterOperand ret = !cond(!eq(VT.Size, 1024) : VOPDstOperand<VReg_1024_AlignTarget>,1803                              !eq(VT.Size, 512)  : VOPDstOperand<VReg_512_AlignTarget>,1804                              !eq(VT.Size, 256)  : VOPDstOperand<VReg_256_AlignTarget>,1805                              !eq(VT.Size, 192)  : VOPDstOperand<VReg_192_AlignTarget>,1806                              !eq(VT.Size, 128)  : VOPDstOperand<VReg_128_AlignTarget>,1807                              !eq(VT.Size, 96)   : VOPDstOperand<VReg_96_AlignTarget>,1808                              !eq(VT.Size, 64)   : VOPDstOperand<VReg_64_AlignTarget>,1809                              !eq(VT.Size, 32)   : VOPDstOperand<VGPR_32>,1810                              !eq(VT.Size, 16)   : op16,1811                              1                  : VOPDstS64orS32); // else VT == i11812}1813 1814class getVALUDstForVT_fake16<ValueType VT> {1815  RegisterOperand ret = !if(!eq(VT.Size, 32), VOPDstOperand<VGPR_32>,1816                          !if(!eq(VT.Size, 128), VOPDstOperand<VReg_128_AlignTarget>,1817                            !if(!eq(VT.Size, 64), VOPDstOperand<VReg_64_AlignTarget>,1818                              !if(!eq(VT.Size, 16), VOPDstOperand<VGPR_32_Lo128>,1819                              VOPDstS64orS32)))); // else VT == i11820}1821 1822// Returns the register class to use for the destination of VOP[12C]1823// instructions with SDWA extension1824class getSDWADstForVT<ValueType VT> {1825  RegisterOperand ret = !if(!eq(VT.Size, 1),1826                            SDWAVopcDst, // VOPC1827                            VOPDstOperand<VGPR_32>); // VOP1/2 32-bit dst1828}1829 1830// Returns the register class to use for source 0 of VOP[12C]1831// instructions for the given VT.1832class getVOPSrc0ForVT<ValueType VT, bit IsTrue16, bit IsFake16 = 1> {1833  RegisterOperand ret =1834  !cond(!eq(VT, i64)    : VSrc_b64,1835        !eq(VT, f64)    : VSrc_f64,1836        !eq(VT, i32)    : VSrc_b32,1837        !eq(VT, f32)    : VSrc_f32,1838        !eq(VT, i16)    : !if(IsTrue16,1839                              !if(IsFake16, VSrcFake16_b16_Lo128, VSrcT_b16_Lo128),1840                              VSrc_b16),1841        !eq(VT, f16)    : !if(IsTrue16,1842                              !if(IsFake16, VSrcFake16_f16_Lo128, VSrcT_f16_Lo128),1843                              VSrc_f16),1844        !eq(VT, bf16)   : !if(IsTrue16,1845                              !if(IsFake16, VSrcFake16_bf16_Lo128, VSrcT_bf16_Lo128),1846                              VSrc_bf16),1847        !eq(VT, v2i16)  : VSrc_v2b16,1848        !eq(VT, v2f16)  : VSrc_v2f16,1849        !eq(VT, v2bf16) : VSrc_v2bf16,1850        !eq(VT, v4f16)  : AVSrc_64,1851        !eq(VT, v4bf16) : AVSrc_64,1852        1               : VSrc_b32);1853}1854 1855// Returns the register class to use for source VGPR, SGPR or inline constant1856// for the given VT.1857class getVCSrcForVT<ValueType VT> {1858  RegisterOperand ret =1859    !if(VT.isFP,1860      !if(!eq(VT.Size, 64),1861         VCSrc_f64,1862         !cond(!eq(VT, f16)    : VCSrc_f16,1863               !eq(VT, bf16)   : VCSrc_bf16,1864               !eq(VT, v2f16)  : VCSrc_v2f16,1865               !eq(VT, v2bf16) : VCSrc_v2bf16,1866               1 : VCSrc_f32)1867       ),1868       !if(!eq(VT.Size, 64),1869          VCSrc_b64,1870          !if(!eq(VT, i16),1871             VCSrc_b16,1872             !if(!eq(VT, v2i16),1873                VCSrc_v2b16,1874                VCSrc_b321875             )1876          )1877       )1878    );1879}1880 1881class getSOPSrcForVT<ValueType VT> {1882  RegisterOperand ret = !if(!eq(VT.Size, 64), SSrc_b64, SSrc_b32);1883}1884 1885// Returns the vreg register operand to use for source operand given VT.1886// This should only be used for a target instruction's ins list.1887class getVregSrcForVT<ValueType VT, bit IsTrue16 = 0, bit IsFake16 = 1> {1888  RegisterOperand ret =1889  !cond(!eq(VT.Size, 512) : RegisterOperand<VReg_512_AlignTarget>,1890        !eq(VT.Size, 192) : RegisterOperand<VReg_192_AlignTarget>,1891        !eq(VT.Size, 128) : RegisterOperand<VReg_128_AlignTarget>,1892        !eq(VT.Size, 96)  : RegisterOperand<VReg_96_AlignTarget>,1893        !eq(VT.Size, 64)  : RegisterOperand<VReg_64_AlignTarget>,1894        !eq(VT.Size, 48)  : RegisterOperand<VReg_64_AlignTarget>,1895        !eq(VT.Size, 16)  : !if(IsTrue16,1896                                !if(IsFake16, VGPROp_32_Lo128, VGPROp_16_Lo128),1897                                RegisterOperand<VGPR_32>),1898        1                 : RegisterOperand<VGPR_32>);1899}1900 1901// Returns a concrete vgpr register class to use for a value type VT,1902// which exists separately from a real instruction use.1903class getVregClassForVT<ValueType VT, bit IsTrue16 = 0, bit IsFake16 = 1> {1904  RegisterClass ret =1905  !cond(!eq(VT.Size, 512) : VReg_512,1906        !eq(VT.Size, 192) : VReg_192,1907        !eq(VT.Size, 128) : VReg_128,1908        !eq(VT.Size, 96)  : VReg_96,1909        !eq(VT.Size, 64)  : VReg_64,1910        !eq(VT.Size, 48)  : VReg_64,1911        !eq(VT.Size, 16)  : !if(IsTrue16,1912                                !if(IsFake16, VGPR_32_Lo128, VGPR_16_Lo128),1913                                VGPR_32),1914        1                 : VGPR_32);1915}1916 1917class getSDWASrcForVT <ValueType VT> {1918  RegisterOperand retFlt = !if(!eq(VT.Size, 16), SDWASrc_f16, SDWASrc_f32);1919  RegisterOperand retInt = !if(!eq(VT.Size, 16), SDWASrc_i16, SDWASrc_i32);1920  RegisterOperand ret = !if(VT.isFP, retFlt, retInt);1921}1922 1923// Returns the register class to use for sources of VOP3 instructions for the1924// given VT.1925class getVOP3SrcForVT<ValueType VT, bit IsTrue16 = 0> {1926  RegisterOperand ret =1927  !cond(!eq(VT, f64)       : VSrc_f64,1928        !eq(VT, f32)       : VSrc_f32,1929        !eq(VT, f16)       : !if(IsTrue16, VSrcT_f16, VSrc_f16),1930        !eq(VT, bf16)      : !if(IsTrue16, VSrcT_bf16, VSrc_bf16),1931        !eq(VT, i16)       : !if(IsTrue16, VSrcT_b16, VSrc_b16),1932        !eq(VT, i1)        : SSrc_i1,1933        !eq(VT, v2f32)     : VSrc_v2f32,1934        !eq(VT, v2i32)     : VSrc_v2b32,1935        !eq(VT, v2f16)     : VSrc_v2f16,1936        !eq(VT, v2bf16)    : VSrc_v2bf16,1937        !eq(VT, v2i16)     : VSrc_v2b16,1938        !eq(VT, v4f16)     : AVSrc_64,1939        !eq(VT, v4bf16)    : AVSrc_64,1940        !eq(VT.Size, 1024) : VRegSrc_1024,1941        !eq(VT.Size, 512)  : VRegSrc_512,1942        !eq(VT.Size, 384)  : VRegSrc_384,1943        !eq(VT.Size, 256)  : VRegSrc_256,1944        !eq(VT.Size, 192)  : VRegSrc_192,1945        !eq(VT.Size, 128)  : VRegSrc_128,1946        !eq(VT.Size, 96)   : VRegSrc_96,1947        !eq(VT.Size, 64)   : VSrc_b64,1948        1                  : VSrc_b32);1949}1950 1951// VGPR only VOP3 src with 9 bit encoding1952class getVOP3VRegSrcForVT<ValueType VT> {1953  RegisterOperand ret = !cond(!eq(VT.Size, 1024) : VRegSrc_1024,1954                              !eq(VT.Size, 512)  : VRegSrc_512,1955                              !eq(VT.Size, 384)  : VRegSrc_384,1956                              !eq(VT.Size, 256)  : VRegSrc_256,1957                              !eq(VT.Size, 192)  : VRegSrc_192,1958                              !eq(VT.Size, 128)  : VRegSrc_128,1959                              !eq(VT.Size, 96)   : VRegSrc_96,1960                              !eq(VT.Size, 64)   : VRegSrc_64,1961                              1 : VRegSrc_32);1962}1963 1964// VGPR only VOP3 src with 8 bit encoding e.g. VOP3DPP src0.1965class getVGPRSrcForVT<ValueType VT, bit IsTrue16 = 0, bit IsFake16 = 0> {1966  RegisterOperand ret =1967   !cond(!eq(VT.Size, 128) : VGPROp_128,1968         !eq(VT.Size, 96)  : VGPROp_96,1969         !eq(VT.Size, 64)  : VGPROp_64,1970         !eq(VT.Size, 48)  : VGPROp_64,1971         !eq(VT.Size, 16)  : !if(IsTrue16,1972                                 !if(IsFake16, VGPROp_32,1973                                               VGPROp_16),1974                                 VGPROp_32),1975         1                 : VGPROp_32);1976}1977 1978// Src2 of VOP3 DPP instructions cannot be a literal1979class getVOP3DPPSrcForVT<ValueType VT, bit IsFake16 = 1> {1980  RegisterOperand ret =1981  !cond(!eq(VT, i1)     : SSrc_i1,1982        !eq(VT, i16)    : !if (IsFake16, VCSrc_b16, VCSrcT_b16),1983        !eq(VT, i64)    : VCSrc_b64,1984        !eq(VT, f16)    : !if (IsFake16, VCSrc_f16, VCSrcT_f16),1985        !eq(VT, bf16)   : !if (IsFake16, VCSrc_bf16, VCSrcT_bf16),1986        !eq(VT, v2i16)  : VCSrc_v2b16,1987        !eq(VT, v2f16)  : VCSrc_v2f16,1988        !eq(VT, v2bf16) : VCSrc_v2bf16,1989        !eq(VT, f32)    : VCSrc_f32,1990        !eq(VT, f64)    : VCSrc_f64,1991        !eq(VT, v2i32)  : VCSrc_v2b32,1992        1               : VCSrc_b32);1993}1994 1995// Float or packed int1996class isModifierType<ValueType SrcVT> {1997  bit ret = !or(!eq(SrcVT, f16),1998                !eq(SrcVT, bf16),1999                !eq(SrcVT, f32),2000                !eq(SrcVT, f64),2001                !eq(SrcVT, v2f16),2002                !eq(SrcVT, v2i16),2003                !eq(SrcVT, v2bf16),2004                !eq(SrcVT, v2f32),2005                !eq(SrcVT, v2i32),2006                !eq(SrcVT, v4f16),2007                !eq(SrcVT, v4i16),2008                !eq(SrcVT, v4bf16),2009                !eq(SrcVT, v4f32),2010                !eq(SrcVT, v4i32),2011                !eq(SrcVT, v8f16),2012                !eq(SrcVT, v8i16),2013                !eq(SrcVT, v8bf16),2014                !eq(SrcVT, v8f32),2015                !eq(SrcVT, v8i32),2016                !eq(SrcVT, v16f16),2017                !eq(SrcVT, v16i16),2018                !eq(SrcVT, v16bf16));2019}2020 2021// Return type of input modifiers operand for specified input operand.2022// True16: If the destination is a 16-bit value, the src0 modifier must hold2023// dst's opsel bit. Use a dummy value for DstVT if getting the mod for a src operand besides 0.2024// 64-bit src types are not implemented for True16 dst.2025class getSrc0Mod <ValueType VT, ValueType DstVT, bit IsTrue16 = 0, bit IsFake16 = 1> {2026  defvar T16Dst =  !if(!eq(VT.Size, 64),2027                     !if(VT.isFP, FP64InputMods, Int64InputMods),2028                     !if(!eq(VT.Size, 16),2029                         !if(VT.isFP, !if(IsTrue16, FPT16InputMods<IsFake16>, FP16InputMods),2030                                      !if(IsTrue16, IntT16InputMods<IsFake16>, IntOpSelMods)),2031                         !if(VT.isFP, FP32T16DstInputMods, Int32T16DstInputMods)));2032  defvar Normal =  !if(!eq(VT.Size, 64),2033                     !if(VT.isFP, FP64InputMods, Int64InputMods),2034                     !if(!eq(VT.Size, 16),2035                         !if(VT.isFP, !if(IsTrue16, FPT16InputMods<IsFake16>, FP16InputMods),2036                                      !if(IsTrue16, IntT16InputMods<IsFake16>, IntOpSelMods)),2037                         !if(VT.isFP, FP32InputMods, Int32InputMods)));2038  Operand ret = !if(!and(IsTrue16, !eq(DstVT.Size, 16)), T16Dst, Normal);2039}2040 2041class getSrcMod<ValueType VT, bit IsTrue16 = 0, bit IsFake16 = 1> : getSrc0Mod<VT, f128/*Dummy Arg*/, IsTrue16, IsFake16>;2042 2043// Return type of input modifiers operand specified input operand for DPP2044class getSrcModDPP <ValueType VT> {2045  Operand ret = !if(VT.isFP, FPVRegInputMods, IntVRegInputMods);2046}2047 2048class getSrcModDPP_t16 <ValueType VT, bit IsFake16 = 1> {2049  Operand ret =2050      !if (VT.isFP,2051           !if (!or(!eq(VT, f16), !eq(VT, bf16)),2052                FPT16_Lo128VRegInputMods<IsFake16>, FPVRegInputMods),2053           !if (!eq(VT, i16),2054                IntT16_Lo128VRegInputMods<IsFake16>, IntVRegInputMods));2055}2056 2057// Return type of input modifiers operand for specified input operand for DPP2058// or VOPD3.2059class getSrcModVOP3VC <ValueType VT, bit IsFake16 = 1> {2060  Operand ret =2061      !if (VT.isFP,2062           !if (!or(!eq(VT, f16), !eq(VT, bf16)),2063                FPT16VCSrcInputMods<IsFake16>,2064                !if (!eq(VT, f64), FP64VCSrcInputMods,2065                     FP32VCSrcInputMods)),2066           !if (!eq(VT, i16),2067                IntT16VCSrcInputMods<IsFake16>,2068                Int32VCSrcInputMods));2069}2070 2071// Return type of input modifiers operand for specified input operand for DPP2072// True16: If the destination is a 16-bit value, the src0 modifier must hold2073// dst's opsel bit. Use a dummy value for DstVT if getting the mod for a src operand besides 0.2074// 64-bit src types are not implemented for True16 dst.2075class getSrc0ModVOP3DPP <ValueType VT, ValueType DstVT, bit IsFake16 = 1> {2076  defvar T16Dst =2077      !if (VT.isFP,2078           !if (!or(!eq(VT, f16), !eq(VT, bf16)),2079                FPT16VRegInputMods<IsFake16>, FPVRegT16DstInputMods),2080           !if (!eq(VT, i16), IntT16VRegInputMods<IsFake16>,2081                IntVRegT16DstInputMods));2082  defvar Normal =2083      !if (VT.isFP,2084           !if (!or(!eq(VT, f16), !eq(VT, bf16)),2085                FPT16VRegInputMods<IsFake16>, FPVRegInputMods),2086           !if (!eq(VT, i16),2087                IntT16VRegInputMods<IsFake16>,2088                IntVRegInputMods));2089  Operand ret = !if(!and(!not(IsFake16), !eq(DstVT.Size, 16)), T16Dst, Normal);2090}2091 2092// Return type of input modifiers operand for specified input operand for VGPR2093// only operands (VOPD3 vsrc1 and vsrc2).2094class getSrcModVOP3V <ValueType VT> {2095  Operand ret =2096      !if (!eq(VT, f64), FP64VRegSrcInputMods,2097           FP32VRegSrcInputMods);2098}2099 2100// Return type of input modifiers operand specified input operand for SDWA2101class getSrcModSDWA <ValueType VT> {2102  Operand ret = !if(!eq(VT, f16), FP16SDWAInputMods,2103                !if(!eq(VT, f32), FP32SDWAInputMods,2104                !if(!eq(VT, i16), Int16SDWAInputMods,2105                !if(!eq(VT, bf16), FP16SDWAInputMods,2106                Int32SDWAInputMods))));2107}2108 2109// Returns the input arguments for VOP[12C] instructions for the given SrcVT.2110class getIns32 <RegisterOperand Src0RC, RegisterOperand Src1RC, int NumSrcArgs> {2111  dag ret = !if(!eq(NumSrcArgs, 1), (ins Src0RC:$src0),               // VOP12112            !if(!eq(NumSrcArgs, 2), (ins Src0RC:$src0, Src1RC:$src1), // VOP22113                                    (ins)));2114}2115 2116// Returns the input arguments for VOP3 instructions for the given SrcVT.2117class getIns64 <RegisterOperand Src0RC, RegisterOperand Src1RC,2118                RegisterOperand Src2RC, int NumSrcArgs,2119                bit HasClamp, bit HasModifiers, bit HasSrc2Mods, bit HasOMod,2120                Operand Src0Mod, Operand Src1Mod, Operand Src2Mod,2121                bit HasFP8ByteSel = 0, bit HasFP8DstByteSel = 0,2122                bit HasBitOp3 = 0> {2123  dag src0 = !if(!ge(NumSrcArgs, 1),2124                 !if (HasModifiers,2125                      (ins Src0Mod:$src0_modifiers, Src0RC:$src0),2126                      (ins Src0RC:$src0)),2127                 (ins));2128  dag src1 = !if(!ge(NumSrcArgs, 2),2129                 !if (HasModifiers,2130                      (ins Src1Mod:$src1_modifiers, Src1RC:$src1),2131                      (ins Src1RC:$src1)),2132                 (ins));2133  dag src2 = !if(!ge(NumSrcArgs, 3),2134                 !if (HasSrc2Mods,2135                      (ins Src2Mod:$src2_modifiers, Src2RC:$src2),2136                      (ins Src2RC:$src2)),2137                 (ins));2138  // If there is vdst_in after clamp with HasFP8DstByteSel we cannot use2139  // Clamp0 with default value, all default operands must be at the end.2140  dag clamp = !if(HasClamp, !if(HasFP8DstByteSel, (ins Clamp:$clamp),2141                                                  (ins Clamp0:$clamp)),2142                            (ins));2143  dag omod = !if(HasOMod, (ins omod0:$omod), (ins));2144  dag bytesel = !if(HasFP8ByteSel,2145                    !con(!if(HasFP8DstByteSel, (ins VGPR_32:$vdst_in), (ins)),2146                         (ins ByteSel0:$byte_sel)),2147                    (ins));2148  dag bitop3 = !if(HasBitOp3, (ins bitop3_0:$bitop3), (ins));2149 2150  dag ret = !con(src0, src1, src2, clamp, omod, bytesel, bitop3);2151}2152 2153class getInsVOP3Base<RegisterOperand Src0RC, RegisterOperand Src1RC,2154                RegisterOperand Src2RC, int NumSrcArgs,2155                bit HasClamp, bit HasModifiers, bit HasSrc2Mods, bit HasOMod,2156                Operand Src0Mod, Operand Src1Mod, Operand Src2Mod, bit HasOpSel,2157                bit HasFP8ByteSel = 0, bit HasFP8DstByteSel = 0, bit HasBitOp3 = 0> {2158  // getInst64 handles clamp and omod. implicit mutex between vop3p and omod2159  dag base = getIns64 <Src0RC, Src1RC, Src2RC, NumSrcArgs,2160                HasClamp, HasModifiers, HasSrc2Mods, HasOMod,2161                Src0Mod, Src1Mod, Src2Mod, HasFP8ByteSel, HasFP8DstByteSel>.ret;2162  dag opsel = (ins op_sel0:$op_sel);2163  dag bitop3 = (ins bitop3_0:$bitop3);2164  dag ret = !con(base, !if(HasBitOp3, bitop3, (ins)), !if(HasOpSel, opsel, (ins)));2165}2166 2167class getInsVOP3P <RegisterOperand Src0RC, RegisterOperand Src1RC,2168                   RegisterOperand Src2RC, int NumSrcArgs, bit HasClamp, bit HasOpSel,2169                   bit HasNeg,2170                   Operand Src0Mod, Operand Src1Mod, Operand Src2Mod> {2171  dag base = getInsVOP3Base<Src0RC, Src1RC, Src2RC, NumSrcArgs,2172                    HasClamp, 1/*HasModifiers*/, 1/*HasSrc2Mods*/,2173                    0/*HasOMod*/, Src0Mod, Src1Mod, Src2Mod, HasOpSel>.ret;2174 2175  dag vop3pOpsel = (ins op_sel_hi0:$op_sel_hi);2176  dag vop3p_neg = !if(HasNeg, (ins neg_lo0:$neg_lo, neg_hi0:$neg_hi), (ins));2177 2178  dag vop3pFields = !con(!if(HasOpSel, vop3pOpsel, (ins)), vop3p_neg);2179  dag ret = !con(base, vop3pFields);2180}2181 2182class getInsVOP3OpSel <RegisterOperand Src0RC, RegisterOperand Src1RC,2183                       RegisterOperand Src2RC, int NumSrcArgs,2184                       bit HasClamp, bit HasOMod,2185                       Operand Src0Mod, Operand Src1Mod, Operand Src2Mod,2186                       bit HasFP8ByteSel = 0, bit HasFP8DstByteSel = 0, bit HasBitOp3 = 0> {2187  dag ret = getInsVOP3Base<Src0RC, Src1RC,2188                    Src2RC, NumSrcArgs,2189                    HasClamp, 1/*HasModifiers*/, 1/*HasSrc2Mods*/, HasOMod,2190                    Src0Mod, Src1Mod, Src2Mod, /*HasOpSel=*/1,2191                    HasFP8ByteSel, HasFP8DstByteSel, HasBitOp3>.ret;2192}2193 2194class getInsDPPBase <RegisterOperand OldRC, RegisterOperand Src0RC, RegisterOperand Src1RC,2195                     RegisterOperand Src2RC, int NumSrcArgs, bit HasModifiers,2196                     Operand Src0Mod, Operand Src1Mod, Operand Src2Mod, bit HasOld> {2197  dag ret = !if(!eq(NumSrcArgs, 0),2198                // VOP1 without input operands (V_NOP)2199                (ins ),2200                !con(2201                  !if(HasOld ,(ins OldRC:$old), (ins)),2202                  !if (!eq(NumSrcArgs, 1),2203                    !if (HasModifiers,2204                      // VOP1_DPP with modifiers2205                      (ins Src0Mod:$src0_modifiers, Src0RC:$src0)2206                    /* else */,2207                      // VOP1_DPP without modifiers2208                      (ins Src0RC:$src0)2209                    /* endif */),2210                  !if (!eq(NumSrcArgs, 2),2211                    !if (HasModifiers,2212                      // VOP2_DPP with modifiers2213                      (ins Src0Mod:$src0_modifiers, Src0RC:$src0,2214                       Src1Mod:$src1_modifiers, Src1RC:$src1)2215                    /* else */,2216                      // VOP2_DPP without modifiers2217                      (ins Src0RC:$src0, Src1RC:$src1)2218                    )2219                    /* NumSrcArgs == 3, VOP3 */,2220                    !if (HasModifiers,2221                      // VOP3_DPP with modifiers2222                      (ins Src0Mod:$src0_modifiers, Src0RC:$src0,2223                       Src1Mod:$src1_modifiers, Src1RC:$src1,2224                       Src2Mod:$src2_modifiers, Src2RC:$src2)2225                    /* else */,2226                      // VOP3_DPP without modifiers2227                      (ins Src0RC:$src0, Src1RC:$src1,2228                       Src2RC:$src2)2229                      )2230                    )2231                  )2232                )2233            );2234}2235 2236class getInsDPP <RegisterOperand OldRC, RegisterOperand Src0RC, RegisterOperand Src1RC,2237                 RegisterOperand Src2RC, int NumSrcArgs, bit HasModifiers,2238                 Operand Src0Mod, Operand Src1Mod, Operand Src2Mod, bit HasOld = 1> {2239  dag ret = !con(getInsDPPBase<OldRC, Src0RC, Src1RC, Src2RC, NumSrcArgs,2240                           HasModifiers, Src0Mod, Src1Mod, Src2Mod, HasOld>.ret,2241                 (ins dpp_ctrl:$dpp_ctrl, DppRowMask:$row_mask,2242                      DppBankMask:$bank_mask, DppBoundCtrl:$bound_ctrl));2243}2244 2245class getInsDPP16 <RegisterOperand OldRC, RegisterOperand Src0RC, RegisterOperand Src1RC,2246                   RegisterOperand Src2RC, int NumSrcArgs, bit HasModifiers,2247                   Operand Src0Mod, Operand Src1Mod, Operand Src2Mod, bit HasOld = 1> {2248  dag ret = !con(getInsDPP<OldRC, Src0RC, Src1RC, Src2RC, NumSrcArgs,2249                           HasModifiers, Src0Mod, Src1Mod, Src2Mod, HasOld>.ret,2250                 (ins Dpp16FI:$fi));2251}2252 2253class getInsDPP8 <RegisterOperand OldRC, RegisterOperand Src0RC, RegisterOperand Src1RC,2254                  RegisterOperand Src2RC, int NumSrcArgs, bit HasModifiers,2255                  Operand Src0Mod, Operand Src1Mod, Operand Src2Mod, bit HasOld = 1> {2256  dag ret = !con(getInsDPPBase<OldRC, Src0RC, Src1RC, Src2RC, NumSrcArgs,2257                           HasModifiers, Src0Mod, Src1Mod, Src2Mod, HasOld>.ret,2258                 (ins dpp8:$dpp8, Dpp8FI:$fi));2259}2260 2261class getInsVOP3DPPBase<dag VOP3Base, RegisterOperand OldRC, int NumSrcArgs, bit HasOld> {2262  dag old = ( ins OldRC:$old );2263  dag base = VOP3Base;2264  dag ret =  !con(2265                !if(!and(HasOld,!ne(NumSrcArgs, 0)), old, (ins)),2266                base2267              );2268}2269 2270class getInsVOP3DPP<dag VOP3Base, RegisterOperand OldRC, int NumSrcArgs, bit HasOld = 1> {2271  dag ret = !con(getInsVOP3DPPBase<VOP3Base,OldRC,NumSrcArgs,HasOld>.ret,2272                 (ins dpp_ctrl:$dpp_ctrl, DppRowMask:$row_mask,2273                      DppBankMask:$bank_mask, DppBoundCtrl:$bound_ctrl));2274}2275 2276class getInsVOP3DPP16<dag VOP3Base, RegisterOperand OldRC, int NumSrcArgs, bit HasOld = 1> {2277  dag ret = !con(getInsVOP3DPP<VOP3Base,OldRC,NumSrcArgs,HasOld>.ret,2278                 (ins Dpp16FI:$fi));2279}2280 2281class getInsVOP3DPP8<dag VOP3Base, RegisterOperand OldRC, int NumSrcArgs, bit HasOld = 1> {2282  dag ret = !con(getInsVOP3DPPBase<VOP3Base,OldRC,NumSrcArgs,HasOld>.ret,2283                 (ins dpp8:$dpp8, Dpp8FI:$fi));2284}2285 2286// Ins for SDWA2287class getInsSDWA <RegisterOperand Src0RC, RegisterOperand Src1RC, int NumSrcArgs,2288                  bit HasSDWAOMod, Operand Src0Mod, Operand Src1Mod,2289                  ValueType DstVT> {2290 2291  dag ret = !if(!eq(NumSrcArgs, 0),2292               // VOP1 without input operands (V_NOP)2293               (ins),2294            !if(!eq(NumSrcArgs, 1),2295               // VOP12296               !if(!not(HasSDWAOMod),2297                  // VOP1_SDWA without omod2298                  (ins Src0Mod:$src0_modifiers, Src0RC:$src0,2299                       Clamp:$clamp,2300                       dst_sel:$dst_sel, dst_unused:$dst_unused,2301                       src0_sel:$src0_sel),2302                  // VOP1_SDWA with omod2303                  (ins Src0Mod:$src0_modifiers, Src0RC:$src0,2304                       Clamp:$clamp, omod:$omod,2305                       dst_sel:$dst_sel, dst_unused:$dst_unused,2306                       src0_sel:$src0_sel)),2307            !if(!eq(NumSrcArgs, 2),2308               !if(!eq(DstVT.Size, 1),2309                  // VOPC_SDWA2310                  (ins Src0Mod:$src0_modifiers, Src0RC:$src0,2311                       Src1Mod:$src1_modifiers, Src1RC:$src1,2312                       Clamp:$clamp, src0_sel:$src0_sel, src1_sel:$src1_sel),2313                  // VOP2_SDWA2314                  !if(!not(HasSDWAOMod),2315                     // VOP2_SDWA without omod2316                     (ins Src0Mod:$src0_modifiers, Src0RC:$src0,2317                          Src1Mod:$src1_modifiers, Src1RC:$src1,2318                          Clamp:$clamp,2319                          dst_sel:$dst_sel, dst_unused:$dst_unused,2320                          src0_sel:$src0_sel, src1_sel:$src1_sel),2321                     // VOP2_SDWA with omod2322                     (ins Src0Mod:$src0_modifiers, Src0RC:$src0,2323                          Src1Mod:$src1_modifiers, Src1RC:$src1,2324                          Clamp:$clamp, omod:$omod,2325                          dst_sel:$dst_sel, dst_unused:$dst_unused,2326                          src0_sel:$src0_sel, src1_sel:$src1_sel))),2327            (ins)/* endif */)));2328}2329 2330// Ins for VOPD32331class getInsVOPD3<RegisterOperand Src0VOPD3, RegisterOperand Src1VOPD3, RegisterOperand Src2VOPD3,2332                  Operand Src0Mod, Operand Src1Mod, Operand Src2Mod,2333                  bit HasSrc1, bit HasSrc2, bit HasModifiers, bit IsCompY> {2334  dag Src0 = !if(HasModifiers,2335                 !if(IsCompY, (ins Src0Mod:$src0Y_modifiers, Src0VOPD3:$src0Y),2336                              (ins Src0Mod:$src0X_modifiers, Src0VOPD3:$src0X)),2337                 !if(IsCompY, (ins Src0VOPD3:$src0Y), (ins Src0VOPD3:$src0X)));2338  dag Src1 = !if(HasModifiers,2339                 !if(IsCompY, (ins Src1Mod:$vsrc1Y_modifiers, Src1VOPD3:$vsrc1Y),2340                              (ins Src1Mod:$vsrc1X_modifiers, Src1VOPD3:$vsrc1X)),2341                 !if(IsCompY, (ins Src1VOPD3:$vsrc1Y), (ins Src1VOPD3:$vsrc1X)));2342  dag Src2 = !if(HasModifiers,2343                 !if(IsCompY, (ins Src2Mod:$vsrc2Y_modifiers, Src2VOPD3:$vsrc2Y),2344                              (ins Src2Mod:$vsrc2X_modifiers, Src2VOPD3:$vsrc2X)),2345                 !if(IsCompY, (ins Src2VOPD3:$vsrc2Y), (ins Src2VOPD3:$vsrc2X)));2346  dag ret = !con(Src0,2347                 !if(HasSrc1, Src1, (ins)),2348                 !if(HasSrc2, Src2, (ins)));2349}2350 2351// Outs for DPP2352class getOutsDPP <bit HasDst, ValueType DstVT, RegisterOperand DstRCDPP> {2353  dag ret = !if(HasDst,2354                !if(!eq(DstVT.Size, 1),2355                    (outs), // no dst for VOPC, we use "vcc"-token as dst in SDWA VOPC instructions2356                    (outs DstRCDPP:$vdst)),2357                (outs)); // V_NOP2358}2359 2360// Outs for SDWA2361class getOutsSDWA <bit HasDst, ValueType DstVT, RegisterOperand DstRCSDWA> {2362  dag ret = !if(HasDst,2363                !if(!eq(DstVT.Size, 1),2364                    (outs DstRCSDWA:$sdst),2365                    (outs DstRCSDWA:$vdst)),2366                (outs)); // V_NOP2367}2368 2369// Returns the assembly string for the inputs and outputs of a VOP[12C]2370// instruction.2371class getAsm32 <bit HasDst, int NumSrcArgs, ValueType DstVT = i32> {2372  string dst = !if(!eq(DstVT.Size, 1), "$sdst", "$vdst"); // use $sdst for VOPC2373  string src0 = ", $src0";2374  string src1 = ", $src1";2375  string src2 = ", $src2";2376  string ret = !if(HasDst, dst, "") #2377               !if(!eq(NumSrcArgs, 1), src0, "") #2378               !if(!eq(NumSrcArgs, 2), src0#src1, "") #2379               !if(!eq(NumSrcArgs, 3), src0#src1#src2, "");2380}2381 2382class getAsmVOPDPart <int NumSrcArgs, string XorY, bit HasVOPD3Src2 = 0, bit HasModifiers = 0> {2383  string mods = !if(HasModifiers, "_modifiers", "");2384  string dst = "$vdst" # XorY;2385  string src0 = ", $src0" # XorY # mods;2386  string src1 = ", $vsrc1" # XorY # mods;2387  string src2 = ", $vsrc2" # XorY # mods;2388  string ret = dst #2389               !if(!ge(NumSrcArgs, 1), src0, "") #2390               !if(!ge(NumSrcArgs, 2), src1, "") #2391               !if(HasVOPD3Src2, src2, "");2392}2393 2394// Returns the assembly string for the inputs and outputs of a VOP3P2395// instruction.2396class getAsmVOP3P <bit HasDst, int NumSrcArgs, bit HasNeg,2397                   bit HasClamp, bit HasOpSel> {2398  string dst = !if(HasDst, "$vdst"# !if(!gt(NumSrcArgs, 0), ",", ""), "");2399  string src0 = !if(!eq(NumSrcArgs, 1), " $src0", " $src0,");2400  string src1 = !if(!eq(NumSrcArgs, 1), "",2401                   !if(!eq(NumSrcArgs, 2), " $src1",2402                                           " $src1,"));2403  string src2 = !if(!eq(NumSrcArgs, 3), " $src2", "");2404 2405  string mods = !if(HasNeg, "$neg_lo$neg_hi", "");2406  string clamp = !if(HasClamp, "$clamp", "");2407  string opsel = !if(HasOpSel, "$op_sel$op_sel_hi", "");2408 2409  // Each modifier is printed as an array of bits for each operand, so2410  // all operands are printed as part of src0_modifiers.2411  string ret = dst#src0#src1#src2#opsel#mods#clamp;2412}2413 2414class getAsmDPP <bit HasDst, int NumSrcArgs, bit HasModifiers, ValueType DstVT = i32> {2415  string dst = !if(HasDst,2416                   !if(!eq(DstVT.Size, 1),2417                       "$sdst",2418                       "$vdst"),2419                    ""); // use $sdst for VOPC2420  string src0 = !if(!eq(NumSrcArgs, 1), "$src0_modifiers", "$src0_modifiers,");2421  string src1 = !if(!eq(NumSrcArgs, 1), "",2422                   !if(!eq(NumSrcArgs, 2), " $src1_modifiers",2423                                           " $src1_modifiers,"));2424  string args = !if(!not(HasModifiers),2425                     getAsm32<0, NumSrcArgs, DstVT>.ret,2426                     ", "#src0#src1);2427  string ret = dst#args#" $dpp_ctrl$row_mask$bank_mask$bound_ctrl";2428}2429 2430class getAsmDPP16 <bit HasDst, int NumSrcArgs, bit HasModifiers, ValueType DstVT = i32> {2431  string ret = getAsmDPP<HasDst, NumSrcArgs, HasModifiers, DstVT>.ret#"$fi";2432}2433 2434class getAsmDPP8 <bit HasDst, int NumSrcArgs, bit HasModifiers, ValueType DstVT = i32>2435  : getAsmDPP<HasDst, NumSrcArgs, HasModifiers, DstVT>{2436  let ret = dst#args#" $dpp8$fi";2437}2438 2439class getAsmVOP3Base <int NumSrcArgs, bit HasDst, bit HasClamp,2440                       bit HasOpSel, bit HasOMod, bit IsVOP3P,2441                       bit HasNeg, bit Src0HasMods,2442                       bit Src1HasMods, bit Src2HasMods, ValueType DstVT = i32,2443                       bit HasByteSel = 0, bit HasBitOp3 = 0> {2444  string dst = !if(HasDst,2445                   !if(!eq(DstVT.Size, 1),2446                       "$sdst",2447                       "$vdst"),2448                    ""); // use $sdst for VOPC2449  string src0nomods = !if(!eq(NumSrcArgs, 1), "$src0", "$src0,");2450  string src1nomods = !if(!eq(NumSrcArgs, 1), "",2451                    !if(!eq(NumSrcArgs, 2), " $src1",2452                                            " $src1,"));2453  string src2nomods = !if(!eq(NumSrcArgs, 3), " $src2", "");2454 2455  string src0mods = !if(!eq(NumSrcArgs, 1), "$src0_modifiers", "$src0_modifiers,");2456  string src1mods = !if(!eq(NumSrcArgs, 1), "",2457                    !if(!eq(NumSrcArgs, 2), " $src1_modifiers",2458                                            " $src1_modifiers,"));2459  string src2mods = !if(!eq(NumSrcArgs, 3), " $src2_modifiers", "");2460 2461  string src0 = !if(Src0HasMods, src0mods, src0nomods);2462  string src1 = !if(Src1HasMods, src1mods, src1nomods);2463  string src2 = !if(Src2HasMods, src2mods, src2nomods);2464  string opsel = !if(HasOpSel, "$op_sel", "");2465  string bytesel = !if(HasByteSel, "$byte_sel", "");2466  string bitop3 = !if(HasBitOp3, "$bitop3", "");2467  string 3PMods = !if(IsVOP3P,2468                      !if(HasOpSel, "$op_sel_hi", "")2469                        #!if(HasNeg, "$neg_lo$neg_hi", ""),2470                      "");2471  string clamp = !if(HasClamp, "$clamp", "");2472  string omod = !if(HasOMod, "$omod", "");2473 2474  string ret = dst#!if(!eq(NumSrcArgs,0),2475                       "",2476                       !if(HasDst,", ", "")#src0#src1#src2#bitop3#opsel#bytesel#3PMods#clamp#omod);2477}2478 2479class getAsmVOP3DPP<string base> {2480  string ret = base # " $dpp_ctrl$row_mask$bank_mask$bound_ctrl";2481}2482 2483class getAsmVOP3DPP16<string base> {2484  string ret = getAsmVOP3DPP<base>.ret # "$fi";2485}2486 2487class getAsmVOP3DPP8<string base> {2488  string ret = base # " $dpp8$fi";2489}2490 2491 2492class getAsmSDWA <bit HasDst, int NumSrcArgs, ValueType DstVT = i32> {2493  string dst = !if(HasDst,2494                   !if(!eq(DstVT.Size, 1),2495                       " vcc", // use vcc token as dst for VOPC instructions2496                       "$vdst"),2497                    "");2498  string src0 = "$src0_modifiers";2499  string src1 = "$src1_modifiers";2500  string args = !if(!eq(NumSrcArgs, 0),2501                    "",2502                    !if(!eq(NumSrcArgs, 1),2503                        ", "#src0#"$clamp",2504                        ", "#src0#", "#src1#"$clamp"2505                     )2506                );2507  string sdwa = !if(!eq(NumSrcArgs, 0),2508                    "",2509                    !if(!eq(NumSrcArgs, 1),2510                        " $dst_sel $dst_unused $src0_sel",2511                        !if(!eq(DstVT.Size, 1),2512                            " $src0_sel $src1_sel", // No dst_sel and dst_unused for VOPC2513                            " $dst_sel $dst_unused $src0_sel $src1_sel"2514                        )2515                    )2516                );2517  string ret = dst#args#sdwa;2518}2519 2520class getAsmSDWA9 <bit HasDst, bit HasOMod, int NumSrcArgs,2521                   ValueType DstVT = i32> {2522  string dst = !if(HasDst,2523                   !if(!eq(DstVT.Size, 1),2524                       "$sdst", // VOPC2525                       "$vdst"), // VOP1/22526                    "");2527  string src0 = "$src0_modifiers";2528  string src1 = "$src1_modifiers";2529  string out_mods = !if(!not(HasOMod), "$clamp", "$clamp$omod");2530  string args = !if(!eq(NumSrcArgs, 0), "",2531                    !if(!eq(NumSrcArgs, 1),2532                        ", "#src0,2533                        ", "#src0#", "#src12534                     )2535                );2536  string sdwa = !if(!eq(NumSrcArgs, 0), "",2537                    !if(!eq(NumSrcArgs, 1),2538                        out_mods#" $dst_sel $dst_unused $src0_sel",2539                        !if(!eq(DstVT.Size, 1),2540                            " $src0_sel $src1_sel", // No dst_sel, dst_unused and output modifiers for VOPC2541                            out_mods#" $dst_sel $dst_unused $src0_sel $src1_sel"2542                        )2543                    )2544                );2545  string ret = dst#args#sdwa;2546}2547 2548class getHas64BitOps <int NumSrcArgs, ValueType DstVT, ValueType Src0VT,2549                      ValueType Src1VT> {2550  bit ret = !if(!eq(NumSrcArgs, 3),2551                0,2552                !if(!eq(DstVT.Size, 64),2553                    1,2554                    !if(!eq(Src0VT.Size, 64),2555                        1,2556                        !if(!eq(Src1VT.Size, 64),2557                            1,2558                            02559                        )2560                    )2561                )2562            );2563}2564 2565class getHasSDWA <int NumSrcArgs, ValueType DstVT = i32, ValueType Src0VT = i32,2566                  ValueType Src1VT = i32> {2567  bit ret = !if(!eq(NumSrcArgs, 3),2568                0, // NumSrcArgs == 3 - No SDWA for VOP32569                !if(!eq(DstVT.Size, 64),2570                    0, // 64-bit dst - No SDWA for 64-bit operands2571                    !if(!eq(Src0VT.Size, 64),2572                        0, // 64-bit src02573                        !if(!eq(Src1VT.Size, 64),2574                            0, // 64-bit src22575                            12576                        )2577                    )2578                )2579            );2580}2581 2582class getHasDPP <int NumSrcArgs> {2583  bit ret = !if(!eq(NumSrcArgs, 3),2584                0, // NumSrcArgs == 3 - No DPP for VOP32585                1);2586}2587 2588class getHasExt32BitDPP <int NumSrcArgs, ValueType DstVT = i32, ValueType Src0VT = i32,2589                 ValueType Src1VT = i32> {2590  bit ret = !and(getHasDPP<NumSrcArgs>.ret,2591                 !not(getHas64BitOps<NumSrcArgs, DstVT, Src0VT, Src1VT>.ret));2592}2593 2594class getHasExt64BitDPP <int NumSrcArgs, ValueType DstVT = i32, ValueType Src0VT = i32,2595                 ValueType Src1VT = i32> {2596  bit ret = !and(getHasDPP<NumSrcArgs>.ret,2597                 getHas64BitOps<NumSrcArgs, DstVT, Src0VT, Src1VT>.ret);2598}2599 2600// Function that checks if instruction supports DPP and SDWA2601class getHasExt <int NumSrcArgs, ValueType DstVT = i32, ValueType Src0VT = i32,2602                 ValueType Src1VT = i32> {2603  bit ret = !or(getHasDPP<NumSrcArgs>.ret,2604                getHasSDWA<NumSrcArgs, DstVT, Src0VT, Src1VT>.ret);2605}2606 2607class getAlign2RegOp<RegisterOperand RC> {2608  RegisterOperand ret =2609    !cond(!eq(RC, VGPROp_16) : VGPROp_16,2610          !eq(RC, VGPROp_32) : VGPROp_32,2611 2612          !eq(RC, VGPROp_64) : VGPROp_64_Align2,2613          !eq(RC, VGPROp_64_Align1) : VGPROp_64_Align2,2614          !eq(RC, VGPROp_64_Align2) : VGPROp_64_Align2,2615 2616          !eq(RC, VGPROp_96) : VGPROp_96_Align2,2617          !eq(RC, VGPROp_96_Align1) : VGPROp_96_Align2,2618          !eq(RC, VGPROp_96_Align2) : VGPROp_96_Align2,2619 2620          !eq(RC, VGPROp_128) : VGPROp_128_Align2,2621          !eq(RC, VGPROp_128_Align1) : VGPROp_128_Align2,2622          !eq(RC, VGPROp_128_Align2) : VGPROp_128_Align2,2623 2624          !eq(RC, VGPROp_160) : VGPROp_160_Align2,2625          !eq(RC, VGPROp_160_Align1) : VGPROp_160_Align2,2626          !eq(RC, VGPROp_160_Align2) : VGPROp_160_Align2,2627 2628          !eq(RC, VGPROp_1024) : VGPROp_1024_Align2,2629          !eq(RC, VGPROp_1024_Align1) : VGPROp_1024_Align2,2630          !eq(RC, VGPROp_1024_Align2) : VGPROp_1024_Align2,2631 2632          !eq(RC, AVLdSt_32) : AVLdSt_32,2633          !eq(RC, AVLdSt_64_Align1) : AVLdSt_64_Align2,2634          !eq(RC, AVLdSt_64) : AVLdSt_64_Align2,2635 2636          !eq(RC, AVLdSt_96) : AVLdSt_96_Align2,2637          !eq(RC, AVLdSt_96_Align1) : AVLdSt_96_Align1,2638          !eq(RC, AVLdSt_96_Align1) : AVLdSt_96_Align2,2639 2640          !eq(RC, AVLdSt_128) : AVLdSt_128_Align2,2641          !eq(RC, AVLdSt_128_Align1) : AVLdSt_128_Align2,2642          !eq(RC, AVLdSt_128_Align2) : AVLdSt_128_Align2,2643 2644          !eq(RC, AVLdSt_160) : AVLdSt_160_Align2,2645          !eq(RC, AVLdSt_160_Align1) : AVLdSt_160_Align2,2646          !eq(RC, AVLdSt_160_Align2) : AVLdSt_160_Align2);2647}2648 2649class getEquivalentAGPROperand<RegisterOperand RC> {2650  defvar Size = !cast<SIRegisterClassLike>(RC.RegClass).Size;2651  RegisterOperand ret =2652    !cond(!eq(Size, 32)   : RegisterOperand<AGPR_32>,2653          !eq(Size, 64)   : RegisterOperand<AReg_64>,2654          !eq(Size, 96)   : RegisterOperand<AReg_96>,2655          !eq(Size, 128)  : RegisterOperand<AReg_128>,2656          !eq(Size, 160)  : RegisterOperand<AReg_160>,2657          !eq(Size, 1024) : RegisterOperand<AReg_1024>);2658}2659 2660class getEquivalentVGPROperand<RegisterOperand RC> {2661  defvar Size = !cast<SIRegisterClassLike>(RC.RegClass).Size;2662  RegisterOperand ret =2663    !cond(2664      !eq(RC, VGPROp_32) : VGPROp_32,2665      !eq(RC, VGPROp_64) : VGPROp_64,2666 2667      !eq(RC, AVLdSt_32) : VGPROp_32,2668      !eq(RC, AVLdSt_64) : VGPROp_64,2669      !eq(RC, AVLdSt_96) : VGPROp_96,2670      !eq(RC, AVLdSt_128) : VGPROp_128,2671      !eq(RC, AVLdSt_160) : VGPROp_160,2672      !eq(RC, AVLdSt_1024) : VGPROp_1024,2673 2674      !eq(RC, AVLdSt_64_Align2) : VGPROp_64_Align2,2675      !eq(RC, AVLdSt_96_Align2) : VGPROp_96_Align2,2676      !eq(RC, AVLdSt_128_Align2) : VGPROp_128_Align2,2677      !eq(RC, AVLdSt_160_Align2) : VGPROp_160_Align2,2678      !eq(RC, AVLdSt_1024_Align2) : VGPROp_1024_Align2,2679 2680      !eq(RC, AVLdSt_64_Align1) : VGPROp_64_Align1,2681      !eq(RC, AVLdSt_96_Align1) : VGPROp_96_Align1,2682      !eq(RC, AVLdSt_128_Align1) : VGPROp_128_Align1,2683      !eq(RC, AVLdSt_160_Align1) : VGPROp_160_Align1,2684      !eq(RC, AVLdSt_1024_Align1) : VGPROp_1024_Align1);2685}2686 2687 2688class getHasVOP3DPP <ValueType DstVT = i32, ValueType Src0VT = i32,2689                 ValueType Src1VT = i32, ValueType Src2VT = i32> {2690  bit ret =    !if(!eq(DstVT.Size, 64),2691                    0, // 64-bit dst No DPP for 64-bit operands2692                    !if(!eq(Src0VT.Size, 64),2693                        0, // 64-bit src02694                        !if(!eq(Src1VT.Size, 64),2695                            0, // 64-bit src12696                            !if(!eq(Src2VT.Size, 64),2697                                0, // 64-bit src22698                                12699                            )2700                        )2701                    )2702                );2703}2704 2705 2706def PatGenMode {2707  int NoPattern = 0;2708  int Pattern   = 1;2709}2710 2711class VOPProfile <list<ValueType> _ArgVT, bit _EnableClamp = 0> {2712 2713  field list<ValueType> ArgVT = _ArgVT;2714  field bit EnableClamp = _EnableClamp;2715  field bit IsTrue16 = 0;2716  field bit IsRealTrue16 = 0;2717 2718  field ValueType DstVT = ArgVT[0];2719  field ValueType Src0VT = ArgVT[1];2720  field ValueType Src1VT = ArgVT[2];2721  field ValueType Src2VT = ArgVT[3];2722  field RegisterOperand DstRC = getVALUDstForVT<DstVT>.ret;2723  field RegisterOperand DstRCDPP = DstRC;2724  field RegisterOperand DstRC64 = DstRC;2725  field RegisterOperand DstRCVOP3DPP = DstRC64;2726  field RegisterOperand DstRCSDWA = getSDWADstForVT<DstVT>.ret;2727  field RegisterOperand Src0RC32 = getVOPSrc0ForVT<Src0VT, IsTrue16>.ret;2728  field RegisterOperand Src1RC32 = getVregSrcForVT<Src1VT>.ret;2729  field RegisterOperand Src0RC64 = getVOP3SrcForVT<Src0VT>.ret;2730  field RegisterOperand Src1RC64 = getVOP3SrcForVT<Src1VT>.ret;2731  field RegisterOperand Src2RC64 = getVOP3SrcForVT<Src2VT>.ret;2732  field RegisterOperand Src0DPP = getVregSrcForVT<Src0VT>.ret;2733  field RegisterOperand Src1DPP = getVregSrcForVT<Src1VT>.ret;2734  field RegisterOperand Src2DPP = getVregSrcForVT<Src2VT>.ret;2735  field RegisterOperand Src0VOP3DPP = getVGPRSrcForVT<Src0VT>.ret;2736  field RegisterOperand Src1VOP3DPP = getVOP3DPPSrcForVT<Src1VT>.ret;2737  field RegisterOperand Src2VOP3DPP = getVOP3DPPSrcForVT<Src2VT>.ret;2738  field RegisterOperand Src0SDWA = getSDWASrcForVT<Src0VT>.ret;2739  field RegisterOperand Src1SDWA = getSDWASrcForVT<Src0VT>.ret;2740  field Operand Src0Mod = getSrc0Mod<Src0VT, DstVT>.ret;2741  field Operand Src1Mod = getSrcMod<Src1VT>.ret;2742  field Operand Src2Mod = getSrcMod<Src2VT>.ret;2743  field Operand Src0ModDPP = getSrcModDPP<Src0VT>.ret;2744  field Operand Src1ModDPP = getSrcModDPP<Src1VT>.ret;2745  field Operand Src2ModDPP = getSrcModDPP<Src2VT>.ret;2746  field Operand Src0ModVOP3DPP = getSrc0ModVOP3DPP<Src0VT, DstVT>.ret;2747  field Operand Src1ModVOP3DPP = getSrcModVOP3VC<Src1VT>.ret;2748  field Operand Src2ModVOP3DPP = getSrcModVOP3VC<Src2VT>.ret;2749  field Operand Src0ModSDWA = getSrcModSDWA<Src0VT>.ret;2750  field Operand Src1ModSDWA = getSrcModSDWA<Src1VT>.ret;2751  field RegisterOperand Src0VOPD3 = getVCSrcForVT<Src0VT>.ret;2752  field RegisterOperand Src1VOPD3 = getVregSrcForVT<Src1VT>.ret;2753  field RegisterOperand Src2VOPD3 = getVregSrcForVT<Src2VT>.ret;2754  field Operand Src0ModVOPD3 = getSrcModVOP3VC<Src0VT>.ret;2755  field Operand Src1ModVOPD3 = getSrcModVOP3V<Src1VT>.ret;2756  field Operand Src2ModVOPD3 = getSrcModVOP3V<Src2VT>.ret;2757 2758 2759  field bit IsMAI = 0;2760  field bit IsVOP3P = 0;2761  field bit IsDOT = 0;2762  field bit IsSingle = 0;2763  field bit IsWMMA = 0;2764  field bit IsSWMMAC = 0;2765 2766  field bit HasFP8SrcByteSel = 0;2767  field bit HasFP8DstByteSel = 0;2768  field bit HasFP4DstByteSel = 0;2769  field bit HasFP8ByteSel = !or(HasFP8SrcByteSel, HasFP8DstByteSel);2770  field bit HasBitOp3 = 0;2771 2772  field bit HasDst = !ne(DstVT, untyped);2773  field bit HasDst32 = HasDst;2774  field bit EmitDst = HasDst; // force dst encoding, see v_movreld_b32 special case2775  field bit EmitDstSel = EmitDst;2776  field int NumSrcArgs = getNumSrcArgs<Src0VT, Src1VT, Src2VT>.ret;2777  field bit HasSrc0 = !ne(Src0VT, untyped);2778  field bit HasSrc1 = !ne(Src1VT, untyped);2779  field bit HasSrc2 = !ne(Src2VT, untyped);2780 2781  field bit HasSrc0FloatMods = Src0VT.isFP;2782  field bit HasSrc1FloatMods = Src1VT.isFP;2783  field bit HasSrc2FloatMods = Src2VT.isFP;2784 2785  field bit HasSrc0IntMods = isIntType<Src0VT>.ret;2786  field bit HasSrc1IntMods = isIntType<Src1VT>.ret;2787  field bit HasSrc2IntMods = isIntType<Src2VT>.ret;2788 2789  field bit HasClamp = !or(isModifierType<Src0VT>.ret, EnableClamp);2790  field bit HasSDWAClamp = EmitDst;2791  field bit HasFPClamp = !and(DstVT.isFP, HasClamp);2792  field bit HasIntClamp = !if(DstVT.isFP, 0, HasClamp);2793  field bit HasClampLo = HasClamp;2794  field bit HasClampHi = !and(DstVT.isVector, HasClamp);2795  field bit HasHigh = 0;2796 2797  field bit IsPacked = Src0VT.isVector;2798  field bit HasOpSel = IsPacked;2799  field bit HasOMod = !if(IsVOP3P, 0, DstVT.isFP);2800  field bit HasSDWAOMod = DstVT.isFP;2801 2802  field bit HasModifiers = !or(isModifierType<Src0VT>.ret,2803                               isModifierType<Src1VT>.ret,2804                               isModifierType<Src2VT>.ret,2805                               HasOMod);2806  field bit HasNeg = HasModifiers;2807  field bit HasMatrixFMT = 0;2808  field bit HasMatrixScale = 0;2809  field bit HasMatrixReuse = 0;2810 2811  field bit HasSrc0Mods = HasModifiers;2812  field bit HasSrc1Mods = !if(HasModifiers, !or(HasSrc1FloatMods, HasSrc1IntMods), 0);2813  field bit HasSrc2Mods = !if(HasModifiers, !or(HasSrc2FloatMods, HasSrc2IntMods), 0);2814 2815  field bit HasExt = getHasExt<NumSrcArgs, DstVT, Src0VT, Src1VT>.ret;2816  field bit HasExtVOP3DPP = getHasVOP3DPP<DstVT, Src0VT, Src1VT, Src2VT>.ret;2817  field bit HasExtDPP = !or(getHasDPP<NumSrcArgs>.ret, HasExtVOP3DPP);2818  field bit HasExt32BitDPP = getHasExt32BitDPP<NumSrcArgs, DstVT, Src0VT, Src1VT>.ret;2819  field bit HasExt64BitDPP = getHasExt64BitDPP<NumSrcArgs, DstVT, Src0VT, Src1VT>.ret;2820  field bit HasExtSDWA = getHasSDWA<NumSrcArgs, DstVT, Src0VT, Src1VT>.ret;2821  field bit HasExtSDWA9 = HasExtSDWA;2822  field int NeedPatGen = PatGenMode.NoPattern;2823 2824  field Operand Src0PackedMod = !if(HasSrc0FloatMods, PackedF16InputMods, PackedI16InputMods);2825  field Operand Src1PackedMod = !if(HasSrc1FloatMods, PackedF16InputMods, PackedI16InputMods);2826  field Operand Src2PackedMod = !if(HasSrc2FloatMods, PackedF16InputMods, PackedI16InputMods);2827 2828  field dag Outs = !if(HasDst,(outs DstRC:$vdst),(outs));2829 2830  // VOP3b instructions are a special case with a second explicit2831  // output. This is manually overridden for them.2832  field dag Outs32 = Outs;2833  field dag Outs64 = !if(HasDst,(outs DstRC64:$vdst),(outs));2834  field dag OutsDPP = getOutsDPP<HasDst, DstVT, DstRCDPP>.ret;2835  field dag OutsDPP8 = OutsDPP;2836  field dag OutsVOP3DPP = getOutsDPP<HasDst, DstVT, DstRCVOP3DPP>.ret;2837  field dag OutsVOP3DPP8 = OutsVOP3DPP;2838  field dag OutsSDWA = getOutsSDWA<HasDst, DstVT, DstRCSDWA>.ret;2839 2840  field dag Ins32 = getIns32<Src0RC32, Src1RC32, NumSrcArgs>.ret;2841  field dag Ins64 = getIns64<Src0RC64, Src1RC64, Src2RC64, NumSrcArgs,2842                             HasClamp, HasModifiers, HasSrc2Mods,2843                             HasOMod, Src0Mod, Src1Mod, Src2Mod,2844                             HasFP8ByteSel, HasFP8DstByteSel, HasBitOp3>.ret;2845  field dag InsVOP3P = getInsVOP3P<Src0RC64, Src1RC64, Src2RC64,2846                                   NumSrcArgs, HasClamp, HasOpSel, HasNeg,2847                                   Src0PackedMod, Src1PackedMod, Src2PackedMod>.ret;2848  field dag InsVOP3OpSel = getInsVOP3OpSel<Src0RC64, Src1RC64, Src2RC64,2849                                NumSrcArgs, HasClamp, HasOMod,2850                                Src0Mod, Src1Mod, Src2Mod,2851                                HasFP8ByteSel, HasFP8DstByteSel, HasBitOp3>.ret;2852  field dag InsDPP = !if(HasExtDPP,2853                         getInsDPP<DstRCDPP, Src0DPP, Src1DPP, Src2DPP, NumSrcArgs,2854                                   HasModifiers, Src0ModDPP, Src1ModDPP, Src2ModDPP>.ret,2855                         (ins));2856  field dag InsDPP16 = getInsDPP16<DstRCDPP, Src0DPP, Src1DPP, Src2DPP, NumSrcArgs,2857                                   HasModifiers, Src0ModDPP, Src1ModDPP, Src2ModDPP>.ret;2858  field dag InsDPP8 = getInsDPP8<DstRCDPP, Src0DPP, Src1DPP, Src2DPP,2859                                 NumSrcArgs, HasModifiers,2860                                 Src0ModDPP, Src1ModDPP, Src2ModDPP>.ret;2861  defvar InsVOP3DPPBase = getInsVOP3Base<Src0VOP3DPP, Src1VOP3DPP,2862                  Src2VOP3DPP, NumSrcArgs, HasClamp, HasModifiers, HasSrc2Mods, HasOMod,2863                  Src0ModVOP3DPP, Src1ModVOP3DPP, Src2ModVOP3DPP, HasOpSel,2864                  HasFP8ByteSel, HasFP8DstByteSel, HasBitOp3>.ret;2865  defvar InsVOP3PDPPBase = getInsVOP3P<Src0VOP3DPP, Src1VOP3DPP,2866                  Src2VOP3DPP, NumSrcArgs, HasClamp, HasOpSel, HasNeg,2867                  Src0ModVOP3DPP, Src1ModVOP3DPP, Src2ModVOP3DPP>.ret;2868 2869  field dag InsVOP3Base = !if(IsVOP3P, InsVOP3PDPPBase, InsVOP3DPPBase);2870 2871  field dag InsVOP3DPP = getInsVOP3DPP<InsVOP3Base, DstRCVOP3DPP, NumSrcArgs>.ret;2872  field dag InsVOP3DPP16 = getInsVOP3DPP16<InsVOP3Base, DstRCVOP3DPP, NumSrcArgs>.ret;2873  field dag InsVOP3DPP8 = getInsVOP3DPP8<InsVOP3Base, DstRCVOP3DPP, NumSrcArgs>.ret;2874  field dag InsSDWA = getInsSDWA<Src0SDWA, Src1SDWA, NumSrcArgs,2875                                 HasSDWAOMod, Src0ModSDWA, Src1ModSDWA,2876                                 DstVT>.ret;2877  field dag InsVOPDX = (ins Src0RC32:$src0X, Src1RC32:$vsrc1X);2878  // It is a slight misnomer to use the f32 operand type for non-float2879  // operands, but this operand type will only be used if the other dual2880  // component is FMAAK or FMAMK2881  field dag InsVOPDX_immX = (ins !if(!eq(Src0VT.Size, 32), VSrc_f32, VSrc_f16):$src0X, VGPR_32:$vsrc1X);2882  field dag InsVOPDY = (ins Src0RC32:$src0Y, Src1RC32:$vsrc1Y);2883  field bit HasVOPD3Src2 = HasSrc2;2884  field dag InsVOPD3X = getInsVOPD3<Src0VOPD3, Src1VOPD3, Src2VOPD3,2885                                    Src0ModVOPD3, Src1ModVOPD3, Src2ModVOPD3,2886                                    HasSrc1, HasVOPD3Src2, HasModifiers, 0>.ret;2887  field dag InsVOPD3Y = getInsVOPD3<Src0VOPD3, Src1VOPD3, Src2VOPD3,2888                                    Src0ModVOPD3, Src1ModVOPD3, Src2ModVOPD3,2889                                    HasSrc1, HasVOPD3Src2, HasModifiers, 1>.ret;2890 2891  field string Asm32 = getAsm32<HasDst, NumSrcArgs, DstVT>.ret;2892  field string AsmDPP = !if(HasExtDPP,2893                            getAsmDPP<HasDst, NumSrcArgs, HasModifiers, DstVT>.ret, "");2894  field string AsmDPP16 = getAsmDPP16<HasDst, NumSrcArgs, HasModifiers, DstVT>.ret;2895  // DPP8 encoding has no fields for modifiers, and it is enforced by setting2896  // the asm operand name via this HasModifiers flag2897  field string AsmDPP8 = getAsmDPP8<HasDst, NumSrcArgs, 0 /*HasModifiers*/, DstVT>.ret;2898  field string AsmVOP3Base = getAsmVOP3Base<NumSrcArgs, HasDst, HasClamp,2899   HasOpSel, HasOMod, IsVOP3P, HasNeg, HasSrc0FloatMods, HasSrc1FloatMods,2900   HasSrc2FloatMods, DstVT, HasFP8ByteSel, HasBitOp3>.ret;2901  field string Asm64 = AsmVOP3Base;2902  field string AsmVOP3P = getAsmVOP3P<HasDst, NumSrcArgs, HasNeg, HasClamp, HasOpSel>.ret;2903  field string AsmVOP3DPP = getAsmVOP3DPP<AsmVOP3Base>.ret;2904  field string AsmVOP3DPP16 = getAsmVOP3DPP16<AsmVOP3Base>.ret;2905  field string AsmVOP3DPP8 = getAsmVOP3DPP8<AsmVOP3Base>.ret;2906  field string AsmSDWA = getAsmSDWA<HasDst, NumSrcArgs, DstVT>.ret;2907  field string AsmSDWA9 = getAsmSDWA9<HasDst, HasSDWAOMod, NumSrcArgs, DstVT>.ret;2908  field string AsmVOPDX = getAsmVOPDPart<NumSrcArgs, "X">.ret;2909  field string AsmVOPDY = getAsmVOPDPart<NumSrcArgs, "Y">.ret;2910  field string AsmVOPD3X = getAsmVOPDPart<NumSrcArgs, "X", HasVOPD3Src2, HasModifiers>.ret;2911  field string AsmVOPD3Y = getAsmVOPDPart<NumSrcArgs, "Y", HasVOPD3Src2, HasModifiers>.ret;2912  field string TieRegDPP = "$old";2913  field bit IsSMFMAC = false;2914  field bit HasAbid = !and(IsMAI, HasSrc1);2915}2916 2917  class VOP_NO_EXT <VOPProfile p> : VOPProfile <p.ArgVT> {2918  let HasExt = 0;2919  let HasExtDPP = 0;2920  let HasExtVOP3DPP = 0;2921  let HasExt32BitDPP = 0;2922  let HasExt64BitDPP = 0;2923  let HasExtSDWA = 0;2924  let HasExtSDWA9 = 0;2925}2926 2927class VOP_PAT_GEN <VOPProfile p, int mode=PatGenMode.NoPattern> : VOPProfile <p.ArgVT> {2928  let NeedPatGen = mode;2929}2930 2931// VOPC_Profile_t16, VOPC_NoSdst_Profile_t16, VOPC_Class_Profile_t16,2932// VOPC_Class_NoSdst_Profile_t16, and  VOP_MAC_F16_t16 do not inherit from this2933// class, so copy changes to this class in those profiles2934class VOPProfile_True16<VOPProfile P> : VOPProfile<P.ArgVT> {2935  let IsTrue16 = 1;2936  let IsRealTrue16 = 1;2937 2938  let HasOpSel = 1;2939  let HasModifiers = 1; // All instructions at least have OpSel.2940 2941  // Most DstVT are 16-bit, but not all.2942  let DstRC = getVALUDstForVT<DstVT, 1 /*IsTrue16*/, 0 /*IsVOP3Encoding*/>.ret;2943  let Src0RC32 = getVOPSrc0ForVT<Src0VT, 1 /*IsTrue16*/, 0 /*IsFake16*/>.ret;2944  let Src1RC32 = getVregSrcForVT<Src1VT, 1 /*IsTrue16*/, 0 /*IsFake16*/>.ret;2945  let Src0DPP = getVregSrcForVT<Src0VT, 1 /*IsTrue16*/, 0 /*IsFake16*/>.ret;2946  let Src1DPP = getVregSrcForVT<Src1VT, 1 /*IsTrue16*/, 0 /*IsFake16*/>.ret;2947  let Src2DPP = getVregSrcForVT<Src2VT, 1 /*IsTrue16*/, 0 /*IsFake16*/>.ret;2948  let Src0ModDPP = getSrcModDPP_t16<Src0VT, 0 /*IsFake16*/>.ret;2949  let Src1ModDPP = getSrcModDPP_t16<Src1VT, 0 /*IsFake16*/>.ret;2950  let Src2ModDPP = getSrcModDPP_t16<Src2VT, 0 /*IsFake16*/>.ret;2951  let Src0VOP3DPP = !if (!eq(Src0VT.Size, 16), VGPROp_16, VGPROp_32);2952  let Src1VOP3DPP = getVOP3DPPSrcForVT<Src1VT, 0 /*IsFake16*/>.ret;2953  let Src2VOP3DPP = getVOP3DPPSrcForVT<Src2VT, 0 /*IsFake16*/>.ret;2954  let Src0ModVOP3DPP = getSrc0ModVOP3DPP<Src0VT, DstVT, 0/*IsFake16*/>.ret;2955  let Src1ModVOP3DPP = getSrcModVOP3VC<Src1VT, 0 /*IsFake16*/>.ret;2956  let Src2ModVOP3DPP = getSrcModVOP3VC<Src2VT, 0 /*IsFake16*/>.ret;2957 2958  let DstRC64 = getVALUDstForVT<DstVT, 1 /*IsTrue16*/, 1 /*IsVOP3Encoding*/>.ret;2959  let Src0RC64 = getVOP3SrcForVT<Src0VT, 1 /*IsTrue16*/>.ret;2960  let Src1RC64 = getVOP3SrcForVT<Src1VT, 1 /*IsTrue16*/>.ret;2961  let Src2RC64 = getVOP3SrcForVT<Src2VT, 1 /*IsTrue16*/>.ret;2962  let Src0Mod = getSrc0Mod<Src0VT, DstVT, 1/*IsTrue16*/, 0/*IsFake16*/>.ret;2963  let Src1Mod = getSrcMod<Src1VT, 1 /*IsTrue16*/, 0/*IsFake16*/>.ret;2964  let Src2Mod = getSrcMod<Src2VT, 1 /*IsTrue16*/, 0/*IsFake16*/>.ret;2965}2966 2967class VOPProfile_Fake16<VOPProfile P> : VOPProfile<P.ArgVT> {2968  let IsTrue16 = 1;2969  // Most DstVT are 16-bit, but not all2970  let DstRC = getVALUDstForVT_fake16<DstVT>.ret;2971  let DstRC64 = getVALUDstForVT<DstVT>.ret;2972  let Src1RC32 = getVregSrcForVT<Src1VT, 1/*IsTrue16*/, 1/*IsFake16*/>.ret;2973  let Src0DPP = getVregSrcForVT<Src0VT, 1/*IsTrue16*/, 1/*IsFake16*/>.ret;2974  let Src1DPP = getVregSrcForVT<Src1VT, 1/*IsTrue16*/, 1/*IsFake16*/>.ret;2975  let Src2DPP = getVregSrcForVT<Src2VT, 1/*IsTrue16*/, 1/*IsFake16*/>.ret;2976  let Src0ModDPP = getSrcModDPP_t16<Src0VT, 1/*IsFake16*/>.ret;2977  let Src1ModDPP = getSrcModDPP_t16<Src1VT, 1/*IsFake16*/>.ret;2978  let Src2ModDPP = getSrcModDPP_t16<Src2VT, 1/*IsFake16*/>.ret;2979  let Src0Mod = getSrc0Mod<Src0VT, DstVT, 1/*IsTrue16*/, 1/*IsFake16*/>.ret;2980  let Src1Mod = getSrcMod<Src1VT, 1 /*IsTrue16*/, 1/*IsFake16*/>.ret;2981  let Src2Mod = getSrcMod<Src2VT, 1 /*IsTrue16*/, 1/*IsFake16*/>.ret;2982  let Src1VOP3DPP = getVOP3DPPSrcForVT<Src1VT, 1 /*IsFake16*/>.ret;2983  let Src2VOP3DPP = getVOP3DPPSrcForVT<Src2VT, 1 /*IsFake16*/>.ret;2984  let Src0ModVOP3DPP = getSrc0ModVOP3DPP<Src0VT, DstVT, 1/*IsFake16*/>.ret;2985  let Src1ModVOP3DPP = getSrcModVOP3VC<Src1VT, 1 /*IsFake16*/>.ret;2986  let Src2ModVOP3DPP = getSrcModVOP3VC<Src2VT, 1 /*IsFake16*/>.ret;2987}2988 2989def VOP_F16_F16 : VOPProfile<[f16, f16, untyped, untyped]>;2990def VOP_F16_I16 : VOPProfile <[f16, i16, untyped, untyped]>;2991def VOP_I16_F16 : VOPProfile <[i16, f16, untyped, untyped]>;2992def VOP_I16_I16 : VOPProfile <[i16, i16, untyped, untyped]>;2993def VOP_BF16_BF16 : VOPProfile<[bf16, bf16, untyped, untyped]>;2994def VOP1_I16_I32 :  VOPProfile<[i16, i32, untyped, untyped]>;2995def VOP_I16_V2F16 : VOPProfile<[i16, v2f16, untyped, untyped]>;2996 2997def VOP_F16_F16_F16 : VOPProfile <[f16, f16, f16, untyped]>;2998def VOP_F16_F16_I16 : VOPProfile <[f16, f16, i16, untyped]>;2999def VOP_F16_F16_I32 : VOPProfile <[f16, f16, i32, untyped]>;3000def VOP_I16_I16_I16 : VOPProfile <[i16, i16, i16, untyped]>;3001def VOP_I16_I16_I16_ARITH : VOPProfile <[i16, i16, i16, untyped], /*EnableClamp=*/1>;3002def VOP_BF16_BF16_BF16 : VOPProfile <[bf16, bf16, bf16, untyped]>;3003 3004def VOP_I16_I16_I16_I16 : VOPProfile <[i16, i16, i16, i16, untyped]>;3005def VOP_F16_F16_F16_F16 : VOPProfile <[f16, f16, f16, f16, untyped]>;3006def VOP_BF16_BF16_BF16_BF16 : VOPProfile <[bf16, bf16, bf16, bf16, untyped]>;3007 3008def VOP_I32_I16_I16_I32 : VOPProfile <[i32, i16, i16, i32, untyped]>;3009def VOP_I32_I16 : VOPProfile <[i32, i16, untyped, untyped]>;3010def VOP_I16_I32 : VOPProfile <[i16, i32, untyped, untyped]>;3011 3012def VOP_V2F16_V2F16_V2F16 : VOPProfile <[v2f16, v2f16, v2f16, untyped]>;3013def VOP_V2I16_V2I16_V2I16 : VOPProfile <[v2i16, v2i16, v2i16, untyped]>;3014def VOP_V2BF16_V2BF16_V2BF16 : VOPProfile <[v2bf16, v2bf16, v2bf16, untyped]>;3015def VOP_B32_F16_F16 : VOPProfile <[i32, f16, f16, untyped]>;3016 3017def VOP_V2F16_V2F16_V2F16_V2F16 : VOPProfile <[v2f16, v2f16, v2f16, v2f16]>;3018def VOP_V2I16_V2I16_V2I16_V2I16 : VOPProfile <[v2i16, v2i16, v2i16, v2i16]>;3019def VOP_V2BF16_V2BF16_V2BF16_V2BF16 : VOPProfile <[v2bf16, v2bf16, v2bf16, v2bf16]>;3020def VOP_V2I16_F32_F32 : VOPProfile <[v2i16, f32, f32, untyped]>;3021def VOP_V2I16_I32_I32 : VOPProfile <[v2i16, i32, i32, untyped]>;3022 3023def VOP_F16_V2F16_V2F16_F16 : VOPProfile <[f16, v2f16, v2f16, f16]>;3024def VOP_BF16_V2BF16_V2BF16_BF16: VOPProfile <[bf16, v2bf16, v2bf16, bf16]>;3025def VOP_F32_V2BF16_V2BF16_F32 : VOPProfile <[f32, v2bf16, v2bf16, f32]>;3026 3027def VOP_F32_V2F16_V2F16_V2F16 : VOPProfile <[f32, v2f16, v2f16, v2f16]>;3028 3029def VOP_NONE : VOPProfile <[untyped, untyped, untyped, untyped]>;3030 3031def VOP_F32_F32 : VOPProfile <[f32, f32, untyped, untyped]>;3032def VOP_F32_F64 : VOPProfile <[f32, f64, untyped, untyped]>;3033def VOP_F32_I32 : VOPProfile <[f32, i32, untyped, untyped]>;3034def VOP_F64_F32 : VOPProfile <[f64, f32, untyped, untyped]>;3035def VOP_F64_F64 : VOPProfile <[f64, f64, untyped, untyped]>;3036def VOP_F64_I32 : VOPProfile <[f64, i32, untyped, untyped]>;3037def VOP_I32_F32 : VOPProfile <[i32, f32, untyped, untyped]>;3038def VOP_I32_F64 : VOPProfile <[i32, f64, untyped, untyped]>;3039def VOP_I32_I32 : VOPProfile <[i32, i32, untyped, untyped]>;3040def VOP_F16_F32 : VOPProfile <[f16, f32, untyped, untyped]>;3041def VOP_F32_F16 : VOPProfile <[f32, f16, untyped, untyped]>;3042def VOP_I64_I64 : VOPProfile <[i64, i64, untyped, untyped]>;3043def VOP_F32_BF16 : VOPProfile <[f32, bf16, untyped, untyped]>;3044 3045def VOP_F32_F32_F16 : VOPProfile <[f32, f32, f16, untyped]>;3046def VOP_F32_F32_F32 : VOPProfile <[f32, f32, f32, untyped]>;3047def VOP_F32_F32_I32 : VOPProfile <[f32, f32, i32, untyped]>;3048def VOP_F64_F64_F64 : VOPProfile <[f64, f64, f64, untyped]>;3049def VOP_F64_F64_I32 : VOPProfile <[f64, f64, i32, untyped]>;3050def VOP_I32_F32_F32 : VOPProfile <[i32, f32, f32, untyped]>;3051def VOP_I32_F32_I32 : VOPProfile <[i32, f32, i32, untyped]>;3052def VOP_I32_I32_I32 : VOPProfile <[i32, i32, i32, untyped]>;3053def VOP_I16_F32_F32 : VOPProfile <[i16, f32, f32, untyped]>;3054def VOP_I32_I32_I32_ARITH : VOPProfile <[i32, i32, i32, untyped], /*EnableClamp=*/1>;3055def VOP_I64_I64_I64_ARITH : VOPProfile <[i64, i64, i64, untyped], /*EnableClamp=*/1>;3056def VOP_V2F16_F32_F32 : VOPProfile <[v2f16, f32, f32, untyped]>;3057def VOP_F32_F16_F16_F16 : VOPProfile <[f32, f16, f16, f16]>;3058def VOP_F32_BF16_BF16_BF16 : VOPProfile <[f32, bf16, bf16, bf16]>;3059def VOP_V2BF16_F32_F32 : VOPProfile <[v2bf16, f32, f32, untyped]>;3060def VOP_V32F32_V6I32_F32 : VOPProfile <[v32f32, v6i32, f32, untyped]>;3061def VOP_V32F16_V6I32_F32 : VOPProfile <[v32f16, v6i32, f32, untyped]>;3062def VOP_V32BF16_V6I32_F32 : VOPProfile <[v32bf16, v6i32, f32, untyped]>;3063def VOP_V2BF16_F32_F32_I32 : VOPProfile <[v2bf16, f32, f32, i32]>;3064def VOP_V2F16_F32_F32_I32 : VOPProfile <[v2f16, f32, f32, i32]>;3065def VOP_V6I32_V32F16_F32 : VOPProfile<[v6i32, v32f16, f32, untyped]>;3066def VOP_V6I32_V32BF16_F32 : VOPProfile<[v6i32, v32bf16, f32, untyped]>;3067def VOP_V3I32_V16F16_F32 : VOPProfile<[v3i32, v16f16, f32, untyped]>;3068def VOP_V3I32_V16BF16_F32 : VOPProfile<[v3i32, v16bf16, f32, untyped]>;3069def VOP_V3I32_V16F32_F32 : VOPProfile<[v3i32, v16f32, f32, untyped]>;3070def VOP_V6I32_V16F32_V16F32_F32 : VOPProfile<[v6i32, v16f32, v16f32, f32]>;3071def VOP_V2F16_I32_F32 : VOPProfile<[v2f16, i32, f32, untyped]>;3072def VOP_V2I16_F32_F32_F32 : VOPProfile<[v2i16, f32, f32, f32]>;3073def VOP_V2I16_V2F16_F32 : VOPProfile<[v2i16, v2f16, f32, untyped]>;3074def VOP_V2I16_V2BF16_F32 : VOPProfile<[v2i16, v2bf16, f32, untyped]>;3075def VOP_I32_F32_F32_F32 : VOPProfile<[i32, f32, f32, f32]>;3076def VOP_I32_V2F32_I32_F32 : VOPProfile<[i32, v2f32, i32, f32]>;3077def VOP_I32_V2F16_F32_F32 : VOPProfile<[i32, v2f16, f32, f32]>;3078def VOP_I32_V2BF16_F32_F32: VOPProfile<[i32, v2bf16, f32, f32]>;3079def VOP_BF16_F32_I32 : VOPProfile<[bf16, f32, i32, untyped]>;3080def VOP_F16_F32_I32 : VOPProfile<[f16, f32, i32, untyped]>;3081def VOP_I32_BF16_I32_F32 : VOPProfile<[i32, bf16, i32, f32]>;3082def VOP_I32_F16_I32_F32 : VOPProfile<[i32, f16, i32, f32]>;3083def VOP_V16F16_V3I32_I32  : VOPProfile<[v16f16, v3i32, i32, untyped]>;3084def VOP_V16BF16_V3I32_I32 : VOPProfile<[v16bf16, v3i32, i32, untyped]>;3085def VOP_V8F16_V2I32_I32  : VOPProfile<[v8f16, v2i32, i32, untyped]>;3086def VOP_V8BF16_V2I32_I32 : VOPProfile<[v8bf16, v2i32, i32, untyped]>;3087def VOP_V8F16_I32_I32  : VOPProfile<[v8f16, i32, i32, untyped]>;3088def VOP_V8BF16_I32_I32 : VOPProfile<[v8bf16, i32, i32, untyped]>;3089def VOP_V16F32_V3I32_I32 : VOPProfile<[v16f32, v3i32, i32, untyped]>;3090def VOP_V8F32_V2I32_I32 : VOPProfile<[v8f32, v2i32, i32, untyped]>;3091def VOP_V8F32_I32_I32 : VOPProfile<[v8f32, i32, i32, untyped]>;3092def VOP_V2I32_V8BF16_F32 : VOPProfile<[v2i32, v8bf16, f32, untyped]>;3093def VOP_V2I32_V8F16_F32 : VOPProfile<[v2i32, v8f16, f32, untyped]>;3094def VOP_V2I32_V8F32_F32 : VOPProfile<[v2i32, v8f32, f32, untyped]>;3095def VOP_I32_V8F32_F32 : VOPProfile<[i32, v8f32, f32, untyped]>;3096def VOP_I32_V8F16_F32 : VOPProfile<[i32, v8f16, f32, untyped]>;3097def VOP_I32_V8BF16_F32 : VOPProfile<[i32, v8bf16, f32, untyped]>;3098def VOP_I32_F32_I32_F32 : VOPProfile<[i32, f32, i32, f32]>;3099 3100def VOP_V6I32_V32BF16_I32_F32 : VOPProfile<[v6i32, v32bf16, i32, f32]>;3101def VOP_V6I32_V32F16_I32_F32 : VOPProfile<[v6i32, v32f16, i32, f32]>;3102def VOP_V6I32_V32F32_I32_F32 : VOPProfile<[v6i32, v32f32, i32, f32]>;3103def VOP_V3I32_V16F16_I32_F32 : VOPProfile<[v3i32, v16f16, i32, f32]>;3104def VOP_V3I32_V16BF16_I32_F32 : VOPProfile<[v3i32, v16bf16, i32, f32]>;3105def VOP_V3I32_V16F32_I32_F32 : VOPProfile<[v3i32, v16f32, i32, f32]>;3106def VOP_V2I32_V8BF16_I32_F32 : VOPProfile<[v2i32, v8bf16, i32, f32]>;3107def VOP_V2I32_V8F16_I32_F32 : VOPProfile<[v2i32, v8f16, i32, f32]>;3108def VOP_V2I32_V8F32_I32_F32 : VOPProfile<[v2i32, v8f32, i32, f32]>;3109def VOP_I32_V8F32_I32_F32 : VOPProfile<[i32, v8f32, i32, f32]>;3110def VOP_I32_V8F16_I32_F32 : VOPProfile<[i32, v8f16, i32, f32]>;3111def VOP_I32_V8BF16_I32_F32 : VOPProfile<[i32, v8bf16, i32, f32]>;3112 3113def VOP_I64_I64_I32 : VOPProfile <[i64, i64, i32, untyped]>;3114def VOP_I64_I32_I64 : VOPProfile <[i64, i32, i64, untyped]>;3115def VOP_I64_I64_I64 : VOPProfile <[i64, i64, i64, untyped]>;3116 3117def VOP_F16_F32_F16_F32 : VOPProfile <[f16, f32, f16, f32]>;3118def VOP_F32_F32_F16_F16 : VOPProfile <[f32, f32, f16, f16]>;3119def VOP_F32_F32_F32_F32 : VOPProfile <[f32, f32, f32, f32]>;3120def VOP_F64_F64_F64_F64 : VOPProfile <[f64, f64, f64, f64]>;3121def VOP_I32_I32_I32_I32 : VOPProfile <[i32, i32, i32, i32]>;3122def VOP_I32_I32_I32_I16 : VOPProfile <[i32, i32, i32, i16]>;3123def VOP_I64_I32_I32_I64 : VOPProfile <[i64, i32, i32, i64]>;3124def VOP_I32_F32_I32_I32 : VOPProfile <[i32, f32, i32, i32]>;3125def VOP_I64_I64_I32_I64 : VOPProfile <[i64, i64, i32, i64]>;3126def VOP_V4I32_I64_I32_V4I32 : VOPProfile <[v4i32, i64, i32, v4i32]>;3127def VOP_I16_I32_I32_I32 : VOPProfile <[i16, i32, i32, i32]>;3128 3129def VOP_F32_V2F16_V2F16_F32 : VOPProfile <[f32, v2f16, v2f16, f32]>;3130def VOP_I32_V2I16_V2I16_I32 : VOPProfile <[i32, v2i16, v2i16, i32]>;3131 3132def VOP_V4F32_F32_F32_V4F32       : VOPProfile <[v4f32,  f32,   f32,   v4f32]>;3133def VOP_V16F32_F32_F32_V16F32     : VOPProfile <[v16f32, f32,   f32,   v16f32]>;3134def VOP_V32F32_F32_F32_V32F32     : VOPProfile <[v32f32, f32,   f32,   v32f32]>;3135def VOP_V4F32_V4F16_V4F16_V4F32   : VOPProfile <[v4f32,  v4f16, v4f16, v4f32]>;3136def VOP_V16F32_V4F16_V4F16_V16F32 : VOPProfile <[v16f32, v4f16, v4f16, v16f32]>;3137def VOP_V32F32_V4F16_V4F16_V32F32 : VOPProfile <[v32f32, v4f16, v4f16, v32f32]>;3138def VOP_V4F32_V2I16_V2I16_V4F32   : VOPProfile <[v4f32,  v2i16, v2i16, v4f32]>;3139def VOP_V16F32_V2I16_V2I16_V16F32 : VOPProfile <[v16f32, v2i16, v2i16, v16f32]>;3140def VOP_V32F32_V2I16_V2I16_V32F32 : VOPProfile <[v32f32, v2i16, v2i16, v32f32]>;3141def VOP_V4I32_I32_I32_V4I32       : VOPProfile <[v4i32,  i32,   i32,   v4i32]>;3142def VOP_V16I32_I32_I32_V16I32     : VOPProfile <[v16i32, i32,   i32,   v16i32]>;3143def VOP_V32I32_I32_I32_V32I32     : VOPProfile <[v32i32, i32,   i32,   v32i32]>;3144 3145def VOP_V4F64_F64_F64_V4F64 : VOPProfile <[v4f64, f64, f64, v4f64]>;3146def VOP_V1F64_F64_F64_V1F64 : VOPProfile <[v1f64, f64, f64, v1f64]>;3147 3148def VOP_V2F32_V2F32_V2F32_V2F32   : VOPProfile <[v2f32,  v2f32, v2f32, v2f32]>;3149def VOP_V2F32_V2F32_V2F32         : VOPProfile <[v2f32,  v2f32, v2f32, untyped]>;3150def VOP_V2I32_V2I32_V2I32         : VOPProfile <[v2i32,  v2i32, v2i32, untyped]>;3151def VOP_V4F32_V4I16_V4I16_V4F32   : VOPProfile <[v4f32,  v4i16, v4i16, v4f32]>;3152def VOP_V16F32_V4I16_V4I16_V16F32 : VOPProfile <[v16f32, v4i16, v4i16, v16f32]>;3153def VOP_V32F32_V4I16_V4I16_V32F32 : VOPProfile <[v32f32, v4i16, v4i16, v32f32]>;3154 3155def VOP_V4I32_I64_I64_V4I32       : VOPProfile <[v4i32,  i64,   i64,   v4i32]>;3156def VOP_V16I32_I64_I64_V16I32     : VOPProfile <[v16i32, i64,   i64,   v16i32]>;3157def VOP_V4F32_V2F32_V2F32_V4F32   : VOPProfile <[v4f32,  v2f32, v2f32, v4f32]>;3158def VOP_V16F32_V2F32_V2F32_V16F32 : VOPProfile <[v16f32, v2f32, v2f32, v16f32]>;3159def VOP_V4F32_I64_I64_V4F32       : VOPProfile <[v4f32,  i64,   i64,   v4f32]>;3160def VOP_V16F32_I64_I64_V16F32     : VOPProfile <[v16f32, i64,   i64,   v16f32]>;3161 3162def VOP_V4F32_V4F16_V8F16_I32     : VOPProfile <[v4f32,  v4f16, v8f16, i32]>;3163def VOP_V4F32_V8F16_V16F16_I32    : VOPProfile <[v4f32,  v8f16, v16f16, i32]>;3164def VOP_V4F32_V8BF16_V16BF16_I32  : VOPProfile <[v4f32,  v8bf16, v16bf16, i32]>;3165def VOP_V16F32_V4F16_V8F16_I32    : VOPProfile <[v16f32, v4f16, v8f16, i32]>;3166def VOP_V16F32_V8F16_V16F16_I32   : VOPProfile <[v16f32, v8f16, v16f16, i32]>;3167def VOP_V16F32_V8BF16_V16BF16_I32 : VOPProfile <[v16f32, v8bf16, v16bf16, i32]>;3168def VOP_V4F32_V4I16_V8I16_I32     : VOPProfile <[v4f32,  v4i16, v8i16, i32]>;3169def VOP_V16F32_V4I16_V8I16_I32    : VOPProfile <[v16f32, v4i16, v8i16, i32]>;3170def VOP_V4I32_V2I32_V4I32_I32     : VOPProfile <[v4i32,  v2i32, v4i32, i32]>;3171def VOP_V16I32_V2I32_V4I32_I32    : VOPProfile <[v16i32, v2i32, v4i32, i32]>;3172def VOP_V4F32_V2I32_V4I32_I32     : VOPProfile <[v4f32,  v2i32, v4i32, i32]>;3173def VOP_V16F32_V2I32_V4I32_I32    : VOPProfile <[v16f32, v2i32, v4i32, i32]>;3174def VOP_V4I32_V4I32_V8I32_I32     : VOPProfile <[v4i32,  v4i32, v8i32, i32]>;3175def VOP_V16I32_V4I32_V8I32_I32    : VOPProfile <[v16i32, v4i32, v8i32, i32]>;3176def VOP_V4F32_V4I32_V8I32_I32     : VOPProfile <[v4f32, v4i32, v8i32, i32]>;3177def VOP_V16F32_V4I32_V8I32_I32    : VOPProfile <[v16f32, v4i32, v8i32, i32]>;3178 3179def VOP_V4F32_V8F16_V8F16_V4F32   : VOPProfile <[v4f32,  v8f16, v8f16, v4f32]>;3180def VOP_V16F32_V8F16_V8F16_V16F32 : VOPProfile <[v16f32, v8f16, v8f16, v16f32]>;3181def VOP_V16F32_V8BF16_V8BF16_V16F32 : VOPProfile <[v16f32, v8bf16, v8bf16, v16f32]>;3182def VOP_V4F32_V8BF16_V8BF16_V4F32 : VOPProfile <[v4f32, v8bf16, v8bf16, v4f32]>;3183def VOP_V4F32_V8I32_V8I32_V4F32   : VOPProfile <[v4f32,  v8i32, v8i32, v4f32]>;3184 3185def VOP_V4F32_V8I32_V6I32_V4F32   : VOPProfile <[v4f32,  v8i32, v6i32, v4f32]>;3186def VOP_V4F32_V6I32_V8I32_V4F32   : VOPProfile <[v4f32,  v6i32, v8i32, v4f32]>;3187def VOP_V4F32_V6I32_V6I32_V4F32   : VOPProfile <[v4f32,  v6i32, v6i32, v4f32]>;3188 3189def VOP_V4F32_V8I32_V4I32_V4F32   : VOPProfile <[v4f32,  v8i32, v4i32, v4f32]>;3190def VOP_V4F32_V4I32_V8I32_V4F32   : VOPProfile <[v4f32,  v4i32, v8i32, v4f32]>;3191def VOP_V4F32_V6I32_V4I32_V4F32   : VOPProfile <[v4f32,  v6i32, v4i32, v4f32]>;3192def VOP_V4F32_V4I32_V6I32_V4F32   : VOPProfile <[v4f32,  v4i32, v6i32, v4f32]>;3193def VOP_V4F32_V4I32_V4I32_V4F32   : VOPProfile <[v4f32,  v4i32, v4i32, v4f32]>;3194 3195def VOP_V16F32_V8I32_V8I32_V16F32   : VOPProfile <[v16f32, v8i32, v8i32, v16f32]>;3196def VOP_V16F32_V8I32_V6I32_V16F32   : VOPProfile <[v16f32, v8i32, v6i32, v16f32]>;3197def VOP_V16F32_V6I32_V8I32_V16F32   : VOPProfile <[v16f32, v6i32, v8i32, v16f32]>;3198def VOP_V16F32_V6I32_V6I32_V16F32   : VOPProfile <[v16f32, v6i32, v6i32, v16f32]>;3199 3200def VOP_V16F32_V8I32_V4I32_V16F32   : VOPProfile <[v16f32, v8i32, v4i32, v16f32]>;3201def VOP_V16F32_V4I32_V8I32_V16F32   : VOPProfile <[v16f32, v4i32, v8i32, v16f32]>;3202def VOP_V16F32_V6I32_V4I32_V16F32   : VOPProfile <[v16f32, v6i32, v4i32, v16f32]>;3203def VOP_V16F32_V4I32_V6I32_V16F32   : VOPProfile <[v16f32, v4i32, v6i32, v16f32]>;3204def VOP_V16F32_V4I32_V4I32_V16F32   : VOPProfile <[v16f32, v4i32, v4i32, v16f32]>;3205 3206def VOP_V4I32_V4I32_V4I32_V4I32     : VOPProfile <[v4i32,  v4i32, v4i32, v4i32]>;3207def VOP_V16I32_V4I32_V4I32_V16I32   : VOPProfile <[v16i32,  v4i32, v4i32, v16i32]>;3208 3209 3210class Commutable_REV <string revOp, bit isOrig> {3211  string RevOp = revOp;3212  bit IsOrig = isOrig;3213}3214 3215//===----------------------------------------------------------------------===//3216// Interpolation opcodes3217//===----------------------------------------------------------------------===//3218 3219class VINTRPDstOperand <RegisterClassLike rc> : RegisterOperand <rc, "printVINTRPDst">;3220 3221class VINTRP_Pseudo <string opName, dag outs, dag ins, list<dag> pattern> :3222  VINTRPCommon <outs, ins, "", pattern>,3223  SIMCInstr<opName, SIEncodingFamily.NONE> {3224  let isPseudo = 1;3225  let isCodeGenOnly = 1;3226}3227 3228// FIXME-GFX10: WIP.3229class VINTRP_Real_si <bits <2> op, string opName, dag outs, dag ins,3230                      string asm, int encodingFamily> :3231  VINTRPCommon <outs, ins, asm, []>,3232  VINTRPe <op>,3233  SIMCInstr<opName, encodingFamily> {3234}3235 3236class VINTRP_Real_vi <bits <2> op, string opName, dag outs, dag ins,3237                      string asm> :3238  VINTRPCommon <outs, ins, asm, []>,3239  VINTRPe_vi <op>,3240  SIMCInstr<opName, SIEncodingFamily.VI> {3241  let AssemblerPredicate = isGFX8GFX9;3242  let DecoderNamespace = "GFX8";3243}3244 3245// FIXME-GFX10: WIP.3246multiclass VINTRP_m <bits <2> op, dag outs, dag ins, string asm,3247                     list<dag> pattern = []> {3248  def "" : VINTRP_Pseudo <NAME, outs, ins, pattern>;3249 3250  let AssemblerPredicate = isGFX6GFX7, DecoderNamespace = "GFX6GFX7" in {3251    def _si : VINTRP_Real_si <op, NAME, outs, ins, asm, SIEncodingFamily.SI>;3252  } // End AssemblerPredicate = isGFX6GFX7, DecoderNamespace = "GFX6GFX7"3253 3254  def _vi : VINTRP_Real_vi <op, NAME, outs, ins, asm>;3255 3256  let AssemblerPredicate = isGFX10Only, DecoderNamespace = "GFX10" in {3257    def _gfx10 : VINTRP_Real_si<op, NAME, outs, ins, asm, SIEncodingFamily.GFX10>;3258  } // End AssemblerPredicate = isGFX10Only, DecoderNamespace = "GFX10"3259}3260 3261//===----------------------------------------------------------------------===//3262// Vector instruction mappings3263//===----------------------------------------------------------------------===//3264 3265// Maps an opcode in e32 form to its e64 equivalent3266def getVOPe64 : InstrMapping {3267  let FilterClass = "VOP";3268  let RowFields = ["OpName"];3269  let ColFields = ["Size", "VOP3"];3270  let KeyCol = ["4", "0"];3271  let ValueCols = [["8", "1"]];3272}3273 3274// Maps an opcode in e64 form to its e32 equivalent3275def getVOPe32 : InstrMapping {3276  let FilterClass = "VOP";3277  let RowFields = ["OpName"];3278  let ColFields = ["Size", "VOP3"];3279  let KeyCol = ["8", "1"];3280  let ValueCols = [["4", "0"]];3281}3282 3283// Maps ordinary instructions to their SDWA counterparts3284def getSDWAOp : InstrMapping {3285  let FilterClass = "VOP";3286  let RowFields = ["OpName"];3287  let ColFields = ["AsmVariantName"];3288  let KeyCol = ["Default"];3289  let ValueCols = [["SDWA"]];3290}3291 3292// Maps SDWA instructions to their ordinary counterparts3293def getBasicFromSDWAOp : InstrMapping {3294  let FilterClass = "VOP";3295  let RowFields = ["OpName"];3296  let ColFields = ["AsmVariantName"];3297  let KeyCol = ["SDWA"];3298  let ValueCols = [["Default"]];3299}3300 3301// Maps ordinary instructions to their DPP counterparts3302def getDPPOp32 : InstrMapping {3303  let FilterClass = "VOP";3304  let RowFields = ["OpName"];3305  let ColFields = ["AsmVariantName"];3306  let KeyCol = ["Default"];3307  let ValueCols = [["DPP"]];3308}3309 3310def getDPPOp64 : InstrMapping {3311  let FilterClass = "VOP";3312  let RowFields = ["OpName"];3313  let ColFields = ["AsmVariantName"];3314  let KeyCol = ["VOP3"];3315  let ValueCols = [["VOP3_DPP"]];3316}3317 3318// Maps an commuted opcode to its original version3319def getCommuteOrig : InstrMapping {3320  let FilterClass = "Commutable_REV";3321  let RowFields = ["RevOp"];3322  let ColFields = ["IsOrig"];3323  let KeyCol = ["0"];3324  let ValueCols = [["1"]];3325}3326 3327// Maps an original opcode to its commuted version3328def getCommuteRev : InstrMapping {3329  let FilterClass = "Commutable_REV";3330  let RowFields = ["RevOp"];3331  let ColFields = ["IsOrig"];3332  let KeyCol = ["1"];3333  let ValueCols = [["0"]];3334}3335 3336def getMCOpcodeGen : InstrMapping {3337  let FilterClass = "SIMCInstr";3338  let RowFields = ["PseudoInstr"];3339  let ColFields = ["Subtarget"];3340  let KeyCol = [!cast<string>(SIEncodingFamily.NONE)];3341  // These columns must be kept in sync with the SIEncodingFamily enumeration.3342  let ValueCols = [[!cast<string>(SIEncodingFamily.SI)],3343                   [!cast<string>(SIEncodingFamily.VI)],3344                   [!cast<string>(SIEncodingFamily.SDWA)],3345                   [!cast<string>(SIEncodingFamily.SDWA9)],3346                   // GFX80 encoding is added to work around a multiple matching3347                   // issue for buffer instructions with unpacked d16 data. This3348                   // does not actually change the encoding, and thus may be3349                   // removed later.3350                   [!cast<string>(SIEncodingFamily.GFX80)],3351                   [!cast<string>(SIEncodingFamily.GFX9)],3352                   [!cast<string>(SIEncodingFamily.GFX10)],3353                   [!cast<string>(SIEncodingFamily.SDWA10)],3354                   [!cast<string>(SIEncodingFamily.GFX90A)],3355                   [!cast<string>(SIEncodingFamily.GFX940)],3356                   [!cast<string>(SIEncodingFamily.GFX11)],3357                   [!cast<string>(SIEncodingFamily.GFX12)],3358                   [!cast<string>(SIEncodingFamily.GFX1250)]];3359}3360 3361// Get equivalent SOPK instruction.3362def getSOPKOp : InstrMapping {3363  let FilterClass = "SOPKInstTable";3364  let RowFields = ["BaseCmpOp"];3365  let ColFields = ["IsSOPK"];3366  let KeyCol = ["0"];3367  let ValueCols = [["1"]];3368}3369 3370def getAddr64Inst : InstrMapping {3371  let FilterClass = "MUBUFAddr64Table";3372  let RowFields = ["OpName"];3373  let ColFields = ["IsAddr64"];3374  let KeyCol = ["0"];3375  let ValueCols = [["1"]];3376}3377 3378def getIfAddr64Inst : InstrMapping {3379  let FilterClass = "MUBUFAddr64Table";3380  let RowFields = ["OpName"];3381  let ColFields = ["IsAddr64"];3382  let KeyCol = ["1"];3383  let ValueCols = [["1"]];3384}3385 3386// Maps a GLOBAL to its SADDR form.3387def getGlobalSaddrOp : InstrMapping {3388  let FilterClass = "GlobalSaddrTable";3389  let RowFields = ["SaddrOp"];3390  let ColFields = ["IsSaddr"];3391  let KeyCol = ["0"];3392  let ValueCols = [["1"]];3393}3394 3395// Maps a GLOBAL SADDR to its VADDR form.3396def getGlobalVaddrOp : InstrMapping {3397  let FilterClass = "GlobalSaddrTable";3398  let RowFields = ["SaddrOp"];3399  let ColFields = ["IsSaddr"];3400  let KeyCol = ["1"];3401  let ValueCols = [["0"]];3402}3403 3404// Maps a v_cmpx opcode with sdst to opcode without sdst.3405def getVCMPXNoSDstOp : InstrMapping {3406  let FilterClass = "VCMPXNoSDstTable";3407  let RowFields = ["NoSDstOp"];3408  let ColFields = ["HasSDst"];3409  let KeyCol = ["1"];3410  let ValueCols = [["0"]];3411}3412 3413// Maps a SOPP to a SOPP with S_NOP3414def getSOPPWithRelaxation : InstrMapping {3415  let FilterClass = "SOPPRelaxTable";3416  let RowFields = ["KeyName"];3417  let ColFields = ["IsRelaxed"];3418  let KeyCol = ["0"];3419  let ValueCols = [["1"]];3420}3421 3422// Maps flat scratch opcodes by addressing modes3423def getFlatScratchInstSTfromSS : InstrMapping {3424  let FilterClass = "FlatScratchInst";3425  let RowFields = ["SVOp"];3426  let ColFields = ["Mode"];3427  let KeyCol = ["SS"];3428  let ValueCols = [["ST"]];3429}3430 3431def getFlatScratchInstSSfromSV : InstrMapping {3432  let FilterClass = "FlatScratchInst";3433  let RowFields = ["SVOp"];3434  let ColFields = ["Mode"];3435  let KeyCol = ["SV"];3436  let ValueCols = [["SS"]];3437}3438 3439def getFlatScratchInstSVfromSVS : InstrMapping {3440  let FilterClass = "FlatScratchInst";3441  let RowFields = ["SVOp"];3442  let ColFields = ["Mode"];3443  let KeyCol = ["SVS"];3444  let ValueCols = [["SV"]];3445}3446 3447def getFlatScratchInstSVfromSS : InstrMapping {3448  let FilterClass = "FlatScratchInst";3449  let RowFields = ["SVOp"];3450  let ColFields = ["Mode"];3451  let KeyCol = ["SS"];3452  let ValueCols = [["SV"]];3453}3454 3455def getMFMAEarlyClobberOp : InstrMapping {3456  let FilterClass = "MFMATable";3457  let RowFields = ["FMAOp"];3458  let ColFields = ["IsMac"];3459  let KeyCol = ["1"];3460  let ValueCols = [["0"]];3461}3462 3463// Map from an mfma using VGPRs to one using AGPRs.3464def getMFMASrcCVDstAGPROp : InstrMapping {3465  let FilterClass = "MFMATable";3466  let RowFields = ["AGPROp"];3467  let ColFields = ["MFMAKind"];3468  let KeyCol = ["VGPR"];3469  let ValueCols = [["AGPR"]];3470}3471 3472// Maps an v_cmp instruction to its v_cmpx equivalent.3473def getVCMPXOpFromVCMP : InstrMapping {3474  let FilterClass = "VCMPVCMPXTable";3475  let RowFields = ["VCMPOp"];3476  let ColFields = ["IsVCMPX"];3477  let KeyCol = ["0"];3478  let ValueCols = [["1"]];3479}3480 3481// Map encoded mfma(_scale)?_f8f6f4 instructions depending on the3482// number of registers required for the used format.3483def getMFMA_F8F6F4_WithSize : GenericTable {3484  let FilterClass = "MFMA_F8F6F4_WithSizeTable";3485  let CppTypeName = "MFMA_F8F6F4_Info";3486  let Fields = [ "Opcode", "F8F8Opcode", "NumRegsSrcA", "NumRegsSrcB" ];3487  let PrimaryKey = [ "NumRegsSrcA", "NumRegsSrcB", "F8F8Opcode" ];3488  let PrimaryKeyName = "getMFMA_F8F6F4_InstWithNumRegs" ;3489}3490 3491def isMFMA_F8F6F4Table : GenericTable {3492  let FilterClass = "MFMA_F8F6F4_WithSizeTable";3493  let CppTypeName = "MFMA_F8F6F4_Info";3494//  let Fields = [ "Opcode" ];3495  let Fields = [ "Opcode", "F8F8Opcode", "NumRegsSrcA", "NumRegsSrcB" ];3496  let PrimaryKey = [ "Opcode" ];3497  let PrimaryKeyName = "isMFMA_F8F6F4" ;3498}3499 3500def FP4FP8DstByteSelTable : GenericTable {3501  let FilterClass = "VOP3_Pseudo";3502  let CppTypeName = "FP4FP8DstByteSelInfo";3503  let Fields = ["Opcode", "HasFP8DstByteSel", "HasFP4DstByteSel"];3504 3505  let PrimaryKey = ["Opcode"];3506  let PrimaryKeyName = "getFP4FP8DstByteSelHelper";3507}3508 3509def VOPDComponentTable : GenericTable {3510  let FilterClass = "VOPD_Component";3511  let CppTypeName = "VOPDComponentInfo";3512  let Fields = ["BaseVOP", "VOPDOp", "CanBeVOPDX", "CanBeVOPD3X"];3513  let PrimaryKey = ["BaseVOP"];3514  let PrimaryKeyName = "getVOPDComponentHelper";3515}3516 3517def getVOPDBaseFromComponent : SearchIndex {3518  let Table = VOPDComponentTable;3519  let Key = ["VOPDOp"];3520}3521 3522def VOPDPairs : GenericTable {3523  let FilterClass = "VOPD_Base";3524  let CppTypeName = "VOPDInfo";3525  let Fields = ["Opcode", "OpX", "OpY", "SubTgt", "VOPD3"];3526  let PrimaryKey = ["Opcode"];3527  let PrimaryKeyName = "getVOPDOpcodeHelper";3528}3529 3530def getVOPDInfoFromComponentOpcodes : SearchIndex {3531  let Table = VOPDPairs;3532  let Key = ["OpX", "OpY", "SubTgt", "VOPD3"];3533}3534 3535include "SIInstructions.td"3536 3537include "DSInstructions.td"3538include "MIMGInstructions.td"3539