brintos / llvm-project-archived public Read only

1//===-- AMDGPU.td - AMDGPU Tablegen files --------*- tablegen -*-===//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 9include "llvm/TableGen/SearchableTable.td"10include "llvm/Target/Target.td"11include "AMDGPUFeatures.td"12include "AMDGPUPredicateControl.td"13 14def p0 : PtrValueType<i64, 0>;15def p1 : PtrValueType<i64, 1>;16def p2 : PtrValueType<i32, 2>;17def p3 : PtrValueType<i32, 3>;18def p4 : PtrValueType<i64, 4>;19def p5 : PtrValueType<i32, 5>;20def p6 : PtrValueType<i32, 6>;21 22//===------------------------------------------------------------===//23// Subtarget Features (device properties)24//===------------------------------------------------------------===//25 26def FeatureFastFMAF32 : SubtargetFeature<"fast-fmaf",27  "FastFMAF32",28  "true",29  "Assuming f32 fma is at least as fast as mul + add"30>;31 32def FeatureFastDenormalF32 : SubtargetFeature<"fast-denormal-f32",33  "FastDenormalF32",34  "true",35  "Enabling denormals does not cause f32 instructions to run at f64 rates"36>;37 38def FeatureMIMG_R128 : SubtargetFeature<"mimg-r128",39  "MIMG_R128",40  "true",41  "Support 128-bit texture resources"42>;43 44def HalfRate64Ops : SubtargetFeature<"half-rate-64-ops",45  "HalfRate64Ops",46  "true",47  "Most fp64 instructions are half rate instead of quarter"48>;49 50def FullRate64Ops : SubtargetFeature<"full-rate-64-ops",51  "FullRate64Ops",52  "true",53  "Most fp64 instructions are full rate"54>;55 56def FeatureFlatAddressSpace : SubtargetFeature<"flat-address-space",57  "FlatAddressSpace",58  "true",59  "Support flat address space"60>;61 62def FeatureFlatInstOffsets : SubtargetFeature<"flat-inst-offsets",63  "FlatInstOffsets",64  "true",65  "Flat instructions have immediate offset addressing mode"66>;67 68def FeatureFlatGlobalInsts : SubtargetFeature<"flat-global-insts",69  "FlatGlobalInsts",70  "true",71  "Have global_* flat memory instructions",72  [FeatureFlatAddressSpace]73>;74 75def FeatureFlatScratchInsts : SubtargetFeature<"flat-scratch-insts",76  "FlatScratchInsts",77  "true",78  "Have scratch_* flat memory instructions",79  [FeatureFlatAddressSpace]80>;81 82def FeatureScalarFlatScratchInsts : SubtargetFeature<"scalar-flat-scratch-insts",83  "ScalarFlatScratchInsts",84  "true",85  "Have s_scratch_* flat memory instructions"86>;87 88def FeatureEnableFlatScratch : SubtargetFeature<"enable-flat-scratch",89  "EnableFlatScratch",90  "true",91  "Use scratch_* flat memory instructions to access scratch"92>;93 94def FeatureFlatGVSMode : SubtargetFeature<"flat-gvs-mode",95  "FlatGVSMode",96  "true",97  "Have GVS addressing mode with flat_* instructions",98  [FeatureFlatAddressSpace]99>;100 101def FeatureAddNoCarryInsts : SubtargetFeature<"add-no-carry-insts",102  "AddNoCarryInsts",103  "true",104  "Have VALU add/sub instructions without carry out"105>;106 107def FeatureUnalignedBufferAccess : SubtargetFeature<"unaligned-buffer-access",108  "UnalignedBufferAccess",109  "true",110  "Hardware supports unaligned global loads and stores"111>;112 113def FeatureTrapHandler: SubtargetFeature<"trap-handler",114  "TrapHandler",115  "true",116  "Trap handler support"117>;118 119def FeatureUnalignedScratchAccess : SubtargetFeature<"unaligned-scratch-access",120  "UnalignedScratchAccess",121  "true",122  "Support unaligned scratch loads and stores"123>;124 125def FeatureUnalignedDSAccess : SubtargetFeature<"unaligned-ds-access",126  "UnalignedDSAccess",127  "true",128  "Hardware supports unaligned local and region loads and stores"129>;130 131def FeatureRelaxedBufferOOBMode : SubtargetFeature<"relaxed-buffer-oob-mode",132  "RelaxedBufferOOBMode",133  "true",134  "Disable strict out-of-bounds buffer guarantees. An OOB access may potentially cause an adjacent access to be treated as if it were also OOB"135>;136 137def FeatureApertureRegs : SubtargetFeature<"aperture-regs",138  "HasApertureRegs",139  "true",140  "Has Memory Aperture Base and Size Registers"141>;142 143def FeatureMadMixInsts : SubtargetFeature<"mad-mix-insts",144  "HasMadMixInsts",145  "true",146  "Has v_mad_mix_f32, v_mad_mixlo_f16, v_mad_mixhi_f16 instructions"147>;148 149def FeatureFmaMixInsts : SubtargetFeature<"fma-mix-insts",150  "HasFmaMixInsts",151  "true",152  "Has v_fma_mix_f32, v_fma_mixlo_f16, v_fma_mixhi_f16 instructions"153>;154 155def FeatureFmaMixBF16Insts : SubtargetFeature<"fma-mix-bf16-insts",156  "HasFmaMixBF16Insts",157  "true",158  "Has v_fma_mix_f32_bf16, v_fma_mixlo_bf16, v_fma_mixhi_bf16 instructions"159>;160 161def FeatureIEEEMinimumMaximumInsts : SubtargetFeature<"ieee-minimum-maximum-insts",162  "HasIEEEMinimumMaximumInsts",163  "true",164  "Has v_minimum/maximum_f16/f32/f64, v_minimummaximum/maximumminimum_f16/f32 and v_pk_minimum/maximum_f16 instructions"165>;166 167def FeatureMinimum3Maximum3F32 : SubtargetFeature<"minimum3-maximum3-f32",168  "HasMinimum3Maximum3F32",169  "true",170  "Has v_minimum3_f32 and v_maximum3_f32 instructions"171>;172 173def FeatureMinimum3Maximum3F16 : SubtargetFeature<"minimum3-maximum3-f16",174  "HasMinimum3Maximum3F16",175  "true",176  "Has v_minimum3_f16 and v_maximum3_f16 instructions"177>;178 179def FeatureMin3Max3PKF16 : SubtargetFeature<"min3-max3-pkf16",180  "HasMin3Max3PKF16",181  "true",182  "Has v_pk_min3_num_f16 and v_pk_max3_num_f16 instructions"183>;184 185def FeatureMinimum3Maximum3PKF16 : SubtargetFeature<"minimum3-maximum3-pkf16",186  "HasMinimum3Maximum3PKF16",187  "true",188  "Has v_pk_minimum3_f16 and v_pk_maximum3_f16 instructions"189>;190 191def FeatureSupportsXNACK : SubtargetFeature<"xnack-support",192  "SupportsXNACK",193  "true",194  "Hardware supports XNACK"195>;196 197// XNACK is disabled if SH_MEM_CONFIG.ADDRESS_MODE = GPUVM on chips that support198// XNACK. The current default kernel driver setting is:199// - graphics ring: XNACK disabled200// - compute ring: XNACK enabled201//202// If XNACK is enabled, the VMEM latency can be worse.203// If XNACK is disabled, the 2 SGPRs can be used for general purposes.204def FeatureXNACK : SubtargetFeature<"xnack",205  "EnableXNACK",206  "true",207  "Enable XNACK support"208>;209 210def FeatureTgSplit : SubtargetFeature<"tgsplit",211  "EnableTgSplit",212  "true",213  "Enable threadgroup split execution"214>;215 216def FeatureCuMode : SubtargetFeature<"cumode",217  "EnableCuMode",218  "true",219  "Enable CU wavefront execution mode"220>;221 222def FeaturePreciseMemory223    : SubtargetFeature<"precise-memory", "EnablePreciseMemory",224                       "true", "Enable precise memory mode">;225 226def FeatureSGPRInitBug : SubtargetFeature<"sgpr-init-bug",227  "SGPRInitBug",228  "true",229  "VI SGPR initialization bug requiring a fixed SGPR allocation size"230>;231 232def FeatureUserSGPRInit16Bug : SubtargetFeature<"user-sgpr-init16-bug",233  "UserSGPRInit16Bug",234  "true",235  "Bug requiring at least 16 user+system SGPRs to be enabled"236>;237 238def FeatureLdsMisalignedBug : SubtargetFeature<"lds-misaligned-bug",239  "LDSMisalignedBug",240  "true",241  "Some GFX10 bug with multi-dword LDS and flat access that is not naturally aligned in WGP mode"242>;243 244def FeatureMFMAInlineLiteralBug : SubtargetFeature<"mfma-inline-literal-bug",245  "HasMFMAInlineLiteralBug",246  "true",247  "MFMA cannot use inline literal as SrcC"248>;249 250def FeatureVcmpxPermlaneHazard : SubtargetFeature<"vcmpx-permlane-hazard",251  "HasVcmpxPermlaneHazard",252  "true",253  "TODO: describe me"254>;255 256def FeatureVMEMtoScalarWriteHazard : SubtargetFeature<"vmem-to-scalar-write-hazard",257  "HasVMEMtoScalarWriteHazard",258  "true",259  "VMEM instruction followed by scalar writing to EXEC mask, M0 or SGPR leads to incorrect execution."260>;261 262def FeatureSMEMtoVectorWriteHazard : SubtargetFeature<"smem-to-vector-write-hazard",263  "HasSMEMtoVectorWriteHazard",264  "true",265  "s_load_dword followed by v_cmp page faults"266>;267 268def FeatureInstFwdPrefetchBug : SubtargetFeature<"inst-fwd-prefetch-bug",269  "HasInstFwdPrefetchBug",270  "true",271  "S_INST_PREFETCH instruction causes shader to hang"272>;273 274def FeatureVmemPrefInsts : SubtargetFeature<"vmem-pref-insts",275  "HasVmemPrefInsts",276  "true",277  "Has flat_prefect_b8 and global_prefetch_b8 instructions"278>;279 280def FeatureSafeSmemPrefetch : SubtargetFeature<"safe-smem-prefetch",281  "HasSafeSmemPrefetch",282  "true",283  "SMEM prefetches do not fail on illegal address"284>;285 286def FeatureSafeCUPrefetch : SubtargetFeature<"safe-cu-prefetch",287  "HasSafeCUPrefetch",288  "true",289  "VMEM CU scope prefetches do not fail on illegal address"290>;291 292def FeatureVcmpxExecWARHazard : SubtargetFeature<"vcmpx-exec-war-hazard",293  "HasVcmpxExecWARHazard",294  "true",295  "V_CMPX WAR hazard on EXEC (V_CMPX issue ONLY)"296>;297 298def FeatureLdsBranchVmemWARHazard : SubtargetFeature<"lds-branch-vmem-war-hazard",299  "HasLdsBranchVmemWARHazard",300  "true",301  "Switching between LDS and VMEM-tex not waiting VM_VSRC=0"302>;303 304class FeatureMaxHardClauseLength<int size> : SubtargetFeature<305  "max-hard-clause-length-"#size,306  "MaxHardClauseLength",307  !cast<string>(size),308  "Maximum number of instructions in an explicit S_CLAUSE is "#size309>;310 311/// Work around a hardware bug on some chips that can be triggered312/// under certain circumstances when clauses are longer than 32 operations.313def FeatureMaxHardClauseLength32 : FeatureMaxHardClauseLength<32>;314/// While the S_CLAUSE instruction permits encoding clause lengths up to 64,315/// hardware documentation for gfx10+ indicates that 63 is the maximum316/// permitted clause length.317def FeatureMaxHardClauseLength63 : FeatureMaxHardClauseLength<63>;318 319def FeatureNSAtoVMEMBug : SubtargetFeature<"nsa-to-vmem-bug",320  "HasNSAtoVMEMBug",321  "true",322  "MIMG-NSA followed by VMEM fail if EXEC_LO or EXEC_HI equals zero"323>;324 325def FeatureNSAClauseBug : SubtargetFeature<"nsa-clause-bug",326  "HasNSAClauseBug",327  "true",328  "MIMG-NSA in a hard clause has unpredictable results on GFX10.1"329>;330 331def FeatureFlatSegmentOffsetBug : SubtargetFeature<"flat-segment-offset-bug",332  "HasFlatSegmentOffsetBug",333  "true",334  "GFX10 bug where inst_offset is ignored when flat instructions access global memory"335>;336 337def FeatureNegativeScratchOffsetBug : SubtargetFeature<"negative-scratch-offset-bug",338  "NegativeScratchOffsetBug",339  "true",340  "Negative immediate offsets in scratch instructions with an SGPR offset page fault on GFX9"341>;342 343def FeatureNegativeUnalignedScratchOffsetBug : SubtargetFeature<"negative-unaligned-scratch-offset-bug",344  "NegativeUnalignedScratchOffsetBug",345  "true",346  "Scratch instructions with a VGPR offset and a negative immediate offset that is not a multiple of 4 read wrong memory on GFX10"347>;348 349def FeatureOffset3fBug : SubtargetFeature<"offset-3f-bug",350  "HasOffset3fBug",351  "true",352  "Branch offset of 3f hardware bug"353>;354 355def FeatureImageStoreD16Bug : SubtargetFeature<"image-store-d16-bug",356  "HasImageStoreD16Bug",357  "true",358  "Image Store D16 hardware bug"359>;360 361def FeatureImageGather4D16Bug : SubtargetFeature<"image-gather4-d16-bug",362  "HasImageGather4D16Bug",363  "true",364  "Image Gather4 D16 hardware bug"365>;366 367def FeatureMADIntraFwdBug : SubtargetFeature<"mad-intra-fwd-bug",368  "HasMADIntraFwdBug",369  "true",370  "MAD_U64/I64 intra instruction forwarding bug"371>;372 373def FeatureMSAALoadDstSelBug : SubtargetFeature<"msaa-load-dst-sel-bug",374  "HasMSAALoadDstSelBug",375  "true",376  "MSAA loads not honoring dst_sel bug"377>;378 379def FeaturePrivEnabledTrap2NopBug : SubtargetFeature<"priv-enabled-trap2-nop-bug",380  "HasPrivEnabledTrap2NopBug",381  "true",382  "Hardware that runs with PRIV=1 interpreting 's_trap 2' as a nop bug"383>;384 385class SubtargetFeatureLDSBankCount <int Value> : SubtargetFeature <386  "ldsbankcount"#Value,387  "LDSBankCount",388  !cast<string>(Value),389  "The number of LDS banks per compute unit."390>;391 392def FeatureLDSBankCount16 : SubtargetFeatureLDSBankCount<16>;393def FeatureLDSBankCount32 : SubtargetFeatureLDSBankCount<32>;394 395def FeatureGCN3Encoding : SubtargetFeature<"gcn3-encoding",396  "GCN3Encoding",397  "true",398  "Encoding format for VI"399>;400 401def FeatureCIInsts : SubtargetFeature<"ci-insts",402  "CIInsts",403  "true",404  "Additional instructions for CI+"405>;406 407def FeatureGFX8Insts : SubtargetFeature<"gfx8-insts",408  "GFX8Insts",409  "true",410  "Additional instructions for GFX8+"411>;412 413def FeatureGFX9Insts : SubtargetFeature<"gfx9-insts",414  "GFX9Insts",415  "true",416  "Additional instructions for GFX9+"417>;418 419def FeatureRequiresAlignedVGPRs : SubtargetFeature<"vgpr-align2",420  "RequiresAlignVGPR",421  "true",422  "VGPR and AGPR tuple operands require even alignment"423>;424 425def FeatureGFX90AInsts : SubtargetFeature<"gfx90a-insts",426  "GFX90AInsts",427  "true",428  "Additional instructions for GFX90A+"429  // [HasAtomicFMinFMaxF64GlobalInsts, HasAtomicFMinFMaxF64FlatInsts] // TODO430>;431 432def FeatureGFX940Insts : SubtargetFeature<"gfx940-insts",433  "GFX940Insts",434  "true",435  "Additional instructions for GFX940+"436>;437 438def FeaturePermlane16Swap : SubtargetFeature<"permlane16-swap",439  "HasPermlane16Swap",440  "true",441  "Has v_permlane16_swap_b32 instructions"442>;443 444def FeaturePermlane32Swap : SubtargetFeature<"permlane32-swap",445  "HasPermlane32Swap",446  "true",447  "Has v_permlane32_swap_b32 instructions"448>;449 450def FeatureFP8ConversionScaleInsts : SubtargetFeature<"fp8-cvt-scale-insts",451  "HasFP8ConversionScaleInsts",452  "true",453  "Has fp8 conversion scale instructions"454>;455 456def FeatureBF8ConversionScaleInsts : SubtargetFeature<"bf8-cvt-scale-insts",457  "HasBF8ConversionScaleInsts",458  "true",459  "Has bf8 conversion scale instructions"460>;461 462def FeatureFP4ConversionScaleInsts : SubtargetFeature<"fp4-cvt-scale-insts",463  "HasFP4ConversionScaleInsts",464  "true",465  "Has fp4 conversion scale instructions"466>;467 468def FeatureFP6BF6ConversionScaleInsts : SubtargetFeature<"fp6bf6-cvt-scale-insts",469  "HasFP6BF6ConversionScaleInsts",470  "true",471  "Has fp6 and bf6 conversion scale instructions"472>;473 474def FeatureF16BF16ToFP6BF6ConversionScaleInsts : SubtargetFeature<"f16bf16-to-fp6bf6-cvt-scale-insts",475  "HasF16BF16ToFP6BF6ConversionScaleInsts",476  "true",477  "Has f16bf16 to fp6bf6 conversion scale instructions"478>;479 480def FeatureF32ToF16BF16ConversionSRInsts : SubtargetFeature<"f32-to-f16bf16-cvt-sr-insts",481  "HasF32ToF16BF16ConversionSRInsts",482  "true",483  "Has f32 to f16bf16 conversion scale instructions"484>;485 486def FeatureAshrPkInsts : SubtargetFeature<"ashr-pk-insts",487  "HasAshrPkInsts",488  "true",489  "Has Arithmetic Shift Pack instructions"490>;491 492def FeatureCvtPkF16F32Inst : SubtargetFeature<"cvt-pk-f16-f32-inst",493  "HasCvtPkF16F32Inst",494  "true",495  "Has cvt_pk_f16_f32 instruction"496>;497 498def FeatureGFX950Insts : SubtargetFeature<"gfx950-insts",499  "GFX950Insts",500  "true",501  "Additional instructions for GFX950+",502  [FeaturePermlane16Swap,503   FeaturePermlane32Swap,504   FeatureAshrPkInsts,505   FeatureFP8ConversionScaleInsts,506   FeatureBF8ConversionScaleInsts,507   FeatureFP4ConversionScaleInsts,508   FeatureFP6BF6ConversionScaleInsts,509   FeatureF16BF16ToFP6BF6ConversionScaleInsts,510   FeatureF32ToF16BF16ConversionSRInsts,511   FeatureCvtPkF16F32Inst,512   FeatureMinimum3Maximum3F32,513   FeatureMinimum3Maximum3PKF16,514   ]515>;516 517def FeatureGFX10Insts : SubtargetFeature<"gfx10-insts",518  "GFX10Insts",519  "true",520  "Additional instructions for GFX10+"521>;522 523def FeatureGFX11Insts : SubtargetFeature<"gfx11-insts",524  "GFX11Insts",525  "true",526  "Additional instructions for GFX11+"527>;528 529def FeatureGFX12Insts : SubtargetFeature<"gfx12-insts",530  "GFX12Insts",531  "true",532  "Additional instructions for GFX12+"533>;534 535def FeatureGFX1250Insts : SubtargetFeature<"gfx1250-insts",536  "GFX1250Insts",537  "true",538  "Additional instructions for GFX1250+"539>;540 541def FeatureGFX10_3Insts : SubtargetFeature<"gfx10-3-insts",542  "GFX10_3Insts",543  "true",544  "Additional instructions for GFX10.3"545>;546 547def FeatureGFX7GFX8GFX9Insts : SubtargetFeature<"gfx7-gfx8-gfx9-insts",548  "GFX7GFX8GFX9Insts",549  "true",550  "Instructions shared in GFX7, GFX8, GFX9"551>;552 553def FeatureSMemRealTime : SubtargetFeature<"s-memrealtime",554  "HasSMemRealTime",555  "true",556  "Has s_memrealtime instruction"557>;558 559def FeatureInv2PiInlineImm : SubtargetFeature<"inv-2pi-inline-imm",560  "HasInv2PiInlineImm",561  "true",562  "Has 1 / (2 * pi) as inline immediate"563>;564 565def Feature16BitInsts : SubtargetFeature<"16-bit-insts",566  "Has16BitInsts",567  "true",568  "Has i16/f16 instructions"569>;570 571def FeatureTrue16BitInsts : SubtargetFeature<"true16",572  "HasTrue16BitInsts",573  "true",574  "True 16-bit operand instructions"575>;576 577def FeatureRealTrue16Insts : SubtargetFeature<"real-true16",578  "EnableRealTrue16Insts",579  "true",580  "Use true 16-bit registers"581>;582 583def FeatureD16Writes32BitVgpr : SubtargetFeature<"d16-write-vgpr32",584  "EnableD16Writes32BitVgpr",585  "true",586  "D16 instructions potentially have 32-bit data dependencies"587>;588 589def FeatureBF16TransInsts : SubtargetFeature<"bf16-trans-insts",590  "HasBF16TransInsts",591  "true",592  "Has bf16 transcendental instructions"593>;594 595def FeatureBF16ConversionInsts : SubtargetFeature<"bf16-cvt-insts",596  "HasBF16ConversionInsts",597  "true",598  "Has bf16 conversion instructions"599>;600 601def FeatureBF16PackedInsts : SubtargetFeature<"bf16-pk-insts",602  "HasBF16PackedInsts",603  "true",604  "Has bf16 packed instructions (fma, add, mul, max, min)"605>;606 607def FeatureVOP3P : SubtargetFeature<"vop3p",608  "HasVOP3PInsts",609  "true",610  "Has VOP3P packed instructions"611>;612 613def FeatureMovrel : SubtargetFeature<"movrel",614  "HasMovrel",615  "true",616  "Has v_movrel*_b32 instructions"617>;618 619def FeatureVGPRIndexMode : SubtargetFeature<"vgpr-index-mode",620  "HasVGPRIndexMode",621  "true",622  "Has VGPR mode register indexing"623>;624 625def FeatureScalarDwordx3Loads : SubtargetFeature<"scalar-dwordx3-loads",626  "HasScalarDwordx3Loads",627  "true",628  "Has 96-bit scalar load instructions"629>;630 631def FeatureScalarStores : SubtargetFeature<"scalar-stores",632  "HasScalarStores",633  "true",634  "Has store scalar memory instructions"635>;636 637def FeatureScalarAtomics : SubtargetFeature<"scalar-atomics",638  "HasScalarAtomics",639  "true",640  "Has atomic scalar memory instructions"641>;642 643def FeatureSDWA : SubtargetFeature<"sdwa",644  "HasSDWA",645  "true",646  "Support SDWA (Sub-DWORD Addressing) extension"647>;648 649def FeatureSDWAOmod : SubtargetFeature<"sdwa-omod",650  "HasSDWAOmod",651  "true",652  "Support OMod with SDWA (Sub-DWORD Addressing) extension"653>;654 655def FeatureSDWAScalar : SubtargetFeature<"sdwa-scalar",656  "HasSDWAScalar",657  "true",658  "Support scalar register with SDWA (Sub-DWORD Addressing) extension"659>;660 661def FeatureSDWASdst : SubtargetFeature<"sdwa-sdst",662  "HasSDWASdst",663  "true",664  "Support scalar dst for VOPC with SDWA (Sub-DWORD Addressing) extension"665>;666 667def FeatureSDWAMac : SubtargetFeature<"sdwa-mav",668  "HasSDWAMac",669  "true",670  "Support v_mac_f32/f16 with SDWA (Sub-DWORD Addressing) extension"671>;672 673def FeatureSDWAOutModsVOPC : SubtargetFeature<"sdwa-out-mods-vopc",674  "HasSDWAOutModsVOPC",675  "true",676  "Support clamp for VOPC with SDWA (Sub-DWORD Addressing) extension"677>;678 679def FeatureDPP : SubtargetFeature<"dpp",680  "HasDPP",681  "true",682  "Support DPP (Data Parallel Primitives) extension"683>;684 685// DPP8 allows arbitrary cross-lane swizzling within groups of 8 lanes.686def FeatureDPP8 : SubtargetFeature<"dpp8",687  "HasDPP8",688  "true",689  "Support DPP8 (Data Parallel Primitives) extension"690>;691 692def FeatureDPALU_DPP : SubtargetFeature<"dpp-64bit",693  "HasDPALU_DPP",694  "true",695  "Support DPP (Data Parallel Primitives) extension in DP ALU"696>;697 698def FeatureDPPSrc1SGPR : SubtargetFeature<"dpp-src1-sgpr",699  "HasDPPSrc1SGPR",700  "true",701  "Support SGPR for Src1 of DPP instructions"702>;703 704def FeaturePackedFP32Ops : SubtargetFeature<"packed-fp32-ops",705  "HasPackedFP32Ops",706  "true",707  "Support packed fp32 instructions"708>;709 710def FeatureR128A16 : SubtargetFeature<"r128-a16",711  "HasR128A16",712  "true",713  "Support gfx9-style A16 for 16-bit coordinates/gradients/lod/clamp/mip image operands, where a16 is aliased with r128"714>;715 716def FeatureA16 : SubtargetFeature<"a16",717  "HasA16",718  "true",719  "Support A16 for 16-bit coordinates/gradients/lod/clamp/mip image operands"720>;721 722def FeatureG16 : SubtargetFeature<"g16",723  "HasG16",724  "true",725  "Support G16 for 16-bit gradient image operands"726>;727 728def FeatureNSAEncoding : SubtargetFeature<"nsa-encoding",729  "HasNSAEncoding",730  "true",731  "Support NSA encoding for image instructions"732>;733 734def FeaturePartialNSAEncoding : SubtargetFeature<"partial-nsa-encoding",735  "HasPartialNSAEncoding",736  "true",737  "Support partial NSA encoding for image instructions"738>;739 740def FeatureImageInsts : SubtargetFeature<"image-insts",741  "HasImageInsts",742  "true",743  "Support image instructions"744>;745 746def FeatureExtendedImageInsts : SubtargetFeature<"extended-image-insts",747  "HasExtendedImageInsts",748  "true",749  "Support mips != 0, lod != 0, gather4, and get_lod"750>;751 752def FeatureGFX10_AEncoding : SubtargetFeature<"gfx10_a-encoding",753  "GFX10_AEncoding",754  "true",755  "Has BVH ray tracing instructions"756>;757 758def FeatureGFX10_BEncoding : SubtargetFeature<"gfx10_b-encoding",759  "GFX10_BEncoding",760  "true",761  "Encoding format GFX10_B"762>;763 764def FeatureIntClamp : SubtargetFeature<"int-clamp-insts",765  "HasIntClamp",766  "true",767  "Support clamp for integer destination"768>;769 770def FeatureUnpackedD16VMem : SubtargetFeature<"unpacked-d16-vmem",771  "HasUnpackedD16VMem",772  "true",773  "Has unpacked d16 vmem instructions"774>;775 776def FeatureDLInsts : SubtargetFeature<"dl-insts",777  "HasDLInsts",778  "true",779  "Has v_fmac_f32 and v_xnor_b32 instructions"780>;781 782def FeatureFmacF64Inst : SubtargetFeature<"fmacf64-inst",783  "HasFmacF64Inst",784  "true",785  "Has v_fmac_f64 instruction"786>;787 788def FeatureDot1Insts : SubtargetFeature<"dot1-insts",789  "HasDot1Insts",790  "true",791  "Has v_dot4_i32_i8 and v_dot8_i32_i4 instructions"792>;793 794def FeatureDot2Insts : SubtargetFeature<"dot2-insts",795  "HasDot2Insts",796  "true",797  "Has v_dot2_i32_i16, v_dot2_u32_u16 instructions"798>;799 800def FeatureDot3Insts : SubtargetFeature<"dot3-insts",801  "HasDot3Insts",802  "true",803  "Has v_dot8c_i32_i4 instruction"804>;805 806def FeatureDot4Insts : SubtargetFeature<"dot4-insts",807  "HasDot4Insts",808  "true",809  "Has v_dot2c_i32_i16 instruction"810>;811 812def FeatureDot5Insts : SubtargetFeature<"dot5-insts",813  "HasDot5Insts",814  "true",815  "Has v_dot2c_f32_f16 instruction"816>;817 818def FeatureDot6Insts : SubtargetFeature<"dot6-insts",819  "HasDot6Insts",820  "true",821  "Has v_dot4c_i32_i8 instruction"822>;823 824def FeatureDot7Insts : SubtargetFeature<"dot7-insts",825  "HasDot7Insts",826  "true",827  "Has v_dot4_u32_u8, v_dot8_u32_u4 instructions"828>;829 830def FeatureDot8Insts : SubtargetFeature<"dot8-insts",831  "HasDot8Insts",832  "true",833  "Has v_dot4_i32_iu8, v_dot8_i32_iu4 instructions"834>;835 836def FeatureDot9Insts : SubtargetFeature<"dot9-insts",837  "HasDot9Insts",838  "true",839  "Has v_dot2_f16_f16, v_dot2_bf16_bf16 instructions"840>;841 842def FeatureDot10Insts : SubtargetFeature<"dot10-insts",843  "HasDot10Insts",844  "true",845  "Has v_dot2_f32_f16 instruction"846>;847 848def FeatureDot11Insts : SubtargetFeature<"dot11-insts",849  "HasDot11Insts",850  "true",851  "Has v_dot4_f32_fp8_fp8, v_dot4_f32_fp8_bf8, v_dot4_f32_bf8_fp8, v_dot4_f32_bf8_bf8 instructions"852>;853 854def FeatureDot12Insts : SubtargetFeature<"dot12-insts",855  "HasDot12Insts",856  "true",857  "Has v_dot2_f32_bf16 instructions"858>;859 860def FeatureDot13Insts : SubtargetFeature<"dot13-insts",861  "HasDot13Insts",862  "true",863  "Has v_dot2c_f32_bf16 instructions"864>;865 866 867def FeatureMAIInsts : SubtargetFeature<"mai-insts",868  "HasMAIInsts",869  "true",870  "Has mAI instructions"871>;872 873def FeatureFP8Insts : SubtargetFeature<"fp8-insts",874  "HasFP8Insts",875  "true",876  "Has fp8 and bf8 instructions"877>;878 879def FeatureFP8ConversionInsts : SubtargetFeature<"fp8-conversion-insts",880  "HasFP8ConversionInsts",881  "true",882  "Has fp8 and bf8 conversion instructions"883>;884 885def FeatureFP8E5M3Insts : SubtargetFeature<"fp8e5m3-insts",886  "HasFP8E5M3Insts",887  "true",888  "Has fp8 e5m3 format support"889>;890 891def FeatureCvtFP8VOP1Bug : SubtargetFeature<"cvt-fp8-vop1-bug",892  "HasCvtFP8Vop1Bug",893  "true",894  "FP8/BF8 VOP1 form of conversion to F32 is unreliable",895  [FeatureFP8ConversionInsts]896>;897 898def FeaturePkFmacF16Inst : SubtargetFeature<"pk-fmac-f16-inst",899  "HasPkFmacF16Inst",900  "true",901  "Has v_pk_fmac_f16 instruction"902>;903 904def FeatureCubeInsts : SubtargetFeature<"cube-insts",905  "HasCubeInsts",906  "true",907  "Has v_cube* instructions"908>; 909 910def FeatureLerpInst : SubtargetFeature<"lerp-inst",911  "HasLerpInst",912  "true",913  "Has v_lerp_u8 instruction"914>; 915 916def FeatureSadInsts : SubtargetFeature<"sad-insts",917  "HasSadInsts",918  "true",919  "Has v_sad* instructions"920>; 921 922def FeatureQsadInsts : SubtargetFeature<"qsad-insts",923  "HasQsadInsts",924  "true",925  "Has v_qsad* instructions"926>; 927 928def FeatureCvtNormInsts : SubtargetFeature<"cvt-norm-insts",929  "HasCvtNormInsts",930  "true",931  "Has v_cvt_norm* instructions"932>;933 934def FeatureCvtPkNormVOP2Insts : SubtargetFeature<"cvt-pknorm-vop2-insts",935  "HasCvtPkNormVOP2Insts",936  "true",937  "Has v_cvt_pk_norm_*f32 instructions/Has v_cvt_pk_norm_*_f16 instructions"938>;939 940def FeatureCvtPkNormVOP3Insts : SubtargetFeature<"cvt-pknorm-vop3-insts",941  "HasCvtPkNormVOP3Insts",942  "true",943  "Has v_cvt_pk_norm_*f32 instructions/Has v_cvt_pk_norm_*_f16 instructions"944>;945 946def FeatureAtomicDsPkAdd16Insts : SubtargetFeature<"atomic-ds-pk-add-16-insts",947  "HasAtomicDsPkAdd16Insts",948  "true",949  "Has ds_pk_add_bf16, ds_pk_add_f16, ds_pk_add_rtn_bf16, "950  "ds_pk_add_rtn_f16 instructions"951>;952 953def FeatureAtomicFlatPkAdd16Insts : SubtargetFeature<"atomic-flat-pk-add-16-insts",954  "HasAtomicFlatPkAdd16Insts",955  "true",956  "Has flat_atomic_pk_add_f16 and flat_atomic_pk_add_bf16 instructions"957>;958 959def FeatureAtomicFaddRtnInsts : SubtargetFeature<"atomic-fadd-rtn-insts",960  "HasAtomicFaddRtnInsts",961  "true",962  "Has buffer_atomic_add_f32 and global_atomic_add_f32 instructions that "963  "return original value",964  [FeatureFlatGlobalInsts]965>;966 967def FeatureAtomicFMinFMaxF32GlobalInsts : SubtargetFeature<"atomic-fmin-fmax-global-f32",968  "HasAtomicFMinFMaxF32GlobalInsts",969  "true",970  "Has global/buffer instructions for atomicrmw fmin/fmax for float"971>;972 973def FeatureAtomicFMinFMaxF64GlobalInsts : SubtargetFeature<"atomic-fmin-fmax-global-f64",974  "HasAtomicFMinFMaxF64GlobalInsts",975  "true",976  "Has global/buffer instructions for atomicrmw fmin/fmax for float"977>;978 979def FeatureAtomicFMinFMaxF32FlatInsts : SubtargetFeature<"atomic-fmin-fmax-flat-f32",980  "HasAtomicFMinFMaxF32FlatInsts",981  "true",982  "Has flat memory instructions for atomicrmw fmin/fmax for float",983  [FeatureFlatAddressSpace]984>;985 986def FeatureAtomicFMinFMaxF64FlatInsts : SubtargetFeature<"atomic-fmin-fmax-flat-f64",987  "HasAtomicFMinFMaxF64FlatInsts",988  "true",989  "Has flat memory instructions for atomicrmw fmin/fmax for double",990  [FeatureFlatAddressSpace]991>;992 993def FeatureAtomicFaddNoRtnInsts : SubtargetFeature<"atomic-fadd-no-rtn-insts",994  "HasAtomicFaddNoRtnInsts",995  "true",996  "Has buffer_atomic_add_f32 and global_atomic_add_f32 instructions that "997  "don't return original value",998  [FeatureFlatGlobalInsts]999>;1000 1001def FeatureAtomicBufferGlobalPkAddF16NoRtnInsts1002  : SubtargetFeature<"atomic-buffer-global-pk-add-f16-no-rtn-insts",1003  "HasAtomicBufferGlobalPkAddF16NoRtnInsts",1004  "true",1005  "Has buffer_atomic_pk_add_f16 and global_atomic_pk_add_f16 instructions that "1006  "don't return original value",1007  [FeatureFlatGlobalInsts]1008>;1009 1010def FeatureAtomicBufferGlobalPkAddF16Insts : SubtargetFeature<"atomic-buffer-global-pk-add-f16-insts",1011 "HasAtomicBufferGlobalPkAddF16Insts",1012 "true",1013 "Has buffer_atomic_pk_add_f16 and global_atomic_pk_add_f16 instructions that "1014 "can return original value",1015 [FeatureFlatGlobalInsts]1016>;1017 1018def FeatureAtomicGlobalPkAddBF16Inst : SubtargetFeature<"atomic-global-pk-add-bf16-inst",1019 "HasAtomicGlobalPkAddBF16Inst",1020 "true",1021 "Has global_atomic_pk_add_bf16 instruction",1022 [FeatureFlatGlobalInsts]1023>;1024 1025def FeatureAtomicBufferPkAddBF16Inst : SubtargetFeature<"atomic-buffer-pk-add-bf16-inst",1026 "HasAtomicBufferPkAddBF16Inst",1027 "true",1028 "Has buffer_atomic_pk_add_bf16 instruction"1029>;1030 1031def FeatureAtomicCSubNoRtnInsts : SubtargetFeature<"atomic-csub-no-rtn-insts",1032  "HasAtomicCSubNoRtnInsts",1033  "true",1034  "Has buffer_atomic_csub and global_atomic_csub instructions that don't "1035  "return original value"1036>;1037 1038def FeatureFlatAtomicFaddF32Inst1039  : SubtargetFeature<"flat-atomic-fadd-f32-inst",1040  "HasFlatAtomicFaddF32Inst",1041  "true",1042  "Has flat_atomic_add_f32 instruction",1043  [FeatureFlatAddressSpace]1044>;1045 1046def FeatureFlatBufferGlobalAtomicFaddF64Inst1047  : SubtargetFeature<"flat-buffer-global-fadd-f64-inst",1048  "HasFlatBufferGlobalAtomicFaddF64Inst",1049  "true",1050  "Has flat, buffer, and global instructions for f64 atomic fadd"1051>;1052 1053def FeatureMemoryAtomicFAddF32DenormalSupport1054  : SubtargetFeature<"memory-atomic-fadd-f32-denormal-support",1055  "HasMemoryAtomicFaddF32DenormalSupport",1056  "true",1057  "global/flat/buffer atomic fadd for float supports denormal handling"1058>;1059 1060def FeatureAgentScopeFineGrainedRemoteMemoryAtomics1061  : SubtargetFeature<"agent-scope-fine-grained-remote-memory-atomics",1062  "HasAgentScopeFineGrainedRemoteMemoryAtomics",1063  "true",1064  "Agent (device) scoped atomic operations, excluding those directly "1065  "supported by PCIe (i.e. integer atomic add, exchange, and "1066  "compare-and-swap), are functional for allocations in host or peer "1067  "device memory."1068>;1069 1070def FeatureEmulatedSystemScopeAtomics1071  : SubtargetFeature<"emulated-system-scope-atomics",1072  "HasEmulatedSystemScopeAtomics",1073  "true",1074  "System scope atomics unsupported by the PCI-e are emulated in HW via CAS "1075  "loop and functional."1076>;1077 1078def FeatureDefaultComponentZero : SubtargetFeature<"default-component-zero",1079  "HasDefaultComponentZero",1080  "true",1081  "BUFFER/IMAGE store instructions set unspecified components to zero (before GFX12)"1082>;1083 1084def FeatureDefaultComponentBroadcast : SubtargetFeature<"default-component-broadcast",1085  "HasDefaultComponentBroadcast",1086  "true",1087  "BUFFER/IMAGE store instructions set unspecified components to x component (GFX12)"1088>;1089 1090def FeatureSupportsSRAMECC : SubtargetFeature<"sramecc-support",1091  "SupportsSRAMECC",1092  "true",1093  "Hardware supports SRAMECC"1094>;1095 1096def FeatureSRAMECC : SubtargetFeature<"sramecc",1097  "EnableSRAMECC",1098  "true",1099  "Enable SRAMECC"1100>;1101 1102def FeatureNoSdstCMPX : SubtargetFeature<"no-sdst-cmpx",1103  "HasNoSdstCMPX",1104  "true",1105  "V_CMPX does not write VCC/SGPR in addition to EXEC"1106>;1107 1108def FeatureVscnt : SubtargetFeature<"vscnt",1109  "HasVscnt",1110  "true",1111  "Has separate store vscnt counter"1112>;1113 1114def FeatureGetWaveIdInst : SubtargetFeature<"get-wave-id-inst",1115  "HasGetWaveIdInst",1116  "true",1117  "Has s_get_waveid_in_workgroup instruction"1118>;1119 1120def FeatureSMemTimeInst : SubtargetFeature<"s-memtime-inst",1121  "HasSMemTimeInst",1122  "true",1123  "Has s_memtime instruction"1124>;1125 1126def FeatureShaderCyclesRegister : SubtargetFeature<"shader-cycles-register",1127  "HasShaderCyclesRegister",1128  "true",1129  "Has SHADER_CYCLES hardware register"1130>;1131 1132def FeatureShaderCyclesHiLoRegisters : SubtargetFeature<"shader-cycles-hi-lo-registers",1133  "HasShaderCyclesHiLoRegisters",1134  "true",1135  "Has SHADER_CYCLES_HI/LO hardware registers"1136>;1137 1138def FeatureMadMacF32Insts : SubtargetFeature<"mad-mac-f32-insts",1139  "HasMadMacF32Insts",1140  "true",1141  "Has v_mad_f32/v_mac_f32/v_madak_f32/v_madmk_f32 instructions"1142>;1143 1144def FeatureDsSrc2Insts : SubtargetFeature<"ds-src2-insts",1145  "HasDsSrc2Insts",1146  "true",1147  "Has ds_*_src2 instructions"1148>;1149 1150def FeatureVOP3Literal : SubtargetFeature<"vop3-literal",1151  "HasVOP3Literal",1152  "true",1153  "Can use one literal in VOP3"1154>;1155 1156def FeatureNoDataDepHazard : SubtargetFeature<"no-data-dep-hazard",1157  "HasNoDataDepHazard",1158  "true",1159  "Does not need SW waitstates"1160>;1161 1162// Allocate 1536 VGPRs for wave32 and 768 VGPRs for wave641163// with allocation granularity 24 for wave32 and 12 for wave641164def Feature1_5xVGPRs : SubtargetFeature<"allocate1_5xvgprs",1165  "Has1_5xVGPRs",1166  "true",1167  "Has 50% more physical VGPRs and 50% larger allocation granule"1168>;1169 1170 1171def FeatureVOPD : SubtargetFeature<"vopd",1172  "HasVOPDInsts",1173  "true",1174  "Has VOPD dual issue wave32 instructions"1175>;1176 1177def FeatureVALUTransUseHazard : SubtargetFeature<"valu-trans-use-hazard",1178  "HasVALUTransUseHazard",1179  "true",1180  "Hazard when TRANS instructions are closely followed by a use of the result"1181>;1182 1183def FeatureSALUFloatInsts : SubtargetFeature<"salu-float",1184  "HasSALUFloatInsts",1185  "true",1186  "Has SALU floating point instructions"1187>;1188 1189def FeaturePseudoScalarTrans : SubtargetFeature<"pseudo-scalar-trans",1190  "HasPseudoScalarTrans",1191  "true",1192  "Has Pseudo Scalar Transcendental instructions"1193>;1194 1195def FeatureHasRestrictedSOffset : SubtargetFeature<"restricted-soffset",1196  "HasRestrictedSOffset",1197  "true",1198  "Has restricted SOffset (immediate not supported)."1199>;1200 1201def FeatureRequiredExportPriority : SubtargetFeature<"required-export-priority",1202  "HasRequiredExportPriority",1203  "true",1204  "Export priority must be explicitly manipulated on GFX11.5"1205>;1206 1207def FeatureVmemWriteVgprInOrder : SubtargetFeature<"vmem-write-vgpr-in-order",1208  "HasVmemWriteVgprInOrder",1209  "true",1210  "VMEM instructions of the same type write VGPR results in order"1211>;1212 1213def FeatureBitOp3Insts : SubtargetFeature<"bitop3-insts",1214  "HasBitOp3Insts",1215  "true",1216  "Has v_bitop3_b32/v_bitop3_b16 instructions"1217>;1218 1219def FeatureTanhInsts : SubtargetFeature<"tanh-insts",1220  "HasTanhInsts",1221  "true",1222  "Has v_tanh_f32/f16 instructions"1223>;1224 1225def FeatureTensorCvtLutInsts : SubtargetFeature<"tensor-cvt-lut-insts",1226  "HasTensorCvtLutInsts",1227  "true",1228  "Has v_perm_pk16* instructions"1229>;1230 1231def FeatureTransposeLoadF4F6Insts : SubtargetFeature<"transpose-load-f4f6-insts",1232  "HasTransposeLoadF4F6Insts",1233  "true",1234  "Has ds_load_tr4/tr6 and global_load_tr4/tr6 instructions"1235>;1236 1237def FeaturePrngInst : SubtargetFeature<"prng-inst",1238  "HasPrngInst",1239  "true",1240  "Has v_prng_b32 instruction"1241>;1242 1243def FeatureBVHDualAndBVH8Insts : SubtargetFeature<"bvh-dual-bvh-8-insts",1244  "HasBVHDualAndBVH8Insts",1245  "true",1246  "Has image_bvh_dual_intersect_ray and image_bvh8_intersect_ray instructions"1247>;1248 1249def FeaturePointSampleAccel : SubtargetFeature<"point-sample-accel",1250  "HasPointSampleAccel",1251  "true",1252  "Has point sample acceleration feature"1253>;1254 1255def Feature64BitLiterals : SubtargetFeature<"64-bit-literals",1256  "Has64BitLiterals",1257  "true",1258  "Can use 64-bit literals with single DWORD instructions"1259>;1260 1261def Feature1024AddressableVGPRs : SubtargetFeature<"1024-addressable-vgprs",1262  "Has1024AddressableVGPRs",1263  "true",1264  "Has 1024 addressable VGPRs"1265>;1266 1267def FeatureWaitXcnt : SubtargetFeature<"wait-xcnt",1268  "HasWaitXcnt",1269  "true",1270  "Has s_wait_xcnt instruction"1271>;1272 1273def FeatureSetPrioIncWgInst : SubtargetFeature<"setprio-inc-wg-inst",1274  "HasSetPrioIncWgInst",1275  "true",1276  "Has s_setprio_inc_wg instruction."1277>;1278 1279//===------------------------------------------------------------===//1280// Subtarget Features (options and debugging)1281//===------------------------------------------------------------===//1282 1283// Ugly hack to accomodate assembling modules with mixed1284// wavesizes. Ideally we would have a mapping symbol in assembly which1285// would keep track of which sections of code should be treated as1286// wave32 and wave64. Instead what users do is assemble with both1287// wavesizes enabled. We translate this into this special mode so this1288// only influences assembler behavior and nothing else.1289def FeatureAssemblerPermissiveWavesize : SubtargetFeature<1290  "assembler-permissive-wavesize",1291  "AssemblerPermissiveWavesize",1292  "true",1293  "allow parsing wave32 and wave64 variants of instructions"1294>;1295 1296class FeatureMaxPrivateElementSize<int size> : SubtargetFeature<1297  "max-private-element-size-"#size,1298  "MaxPrivateElementSize",1299  !cast<string>(size),1300  "Maximum private access size may be "#size1301>;1302 1303def FeatureMaxPrivateElementSize4 : FeatureMaxPrivateElementSize<4>;1304def FeatureMaxPrivateElementSize8 : FeatureMaxPrivateElementSize<8>;1305def FeatureMaxPrivateElementSize16 : FeatureMaxPrivateElementSize<16>;1306 1307def FeatureDumpCode : SubtargetFeature <"DumpCode",1308  "DumpCode",1309  "true",1310  "Dump MachineInstrs in the CodeEmitter"1311>;1312 1313def FeatureDumpCodeLower : SubtargetFeature <"dumpcode",1314  "DumpCode",1315  "true",1316  "Dump MachineInstrs in the CodeEmitter"1317>;1318 1319// XXX - This should probably be removed once enabled by default1320def FeatureEnableLoadStoreOpt : SubtargetFeature <"load-store-opt",1321  "EnableLoadStoreOpt",1322  "true",1323  "Enable SI load/store optimizer pass"1324>;1325 1326// Performance debugging feature. Allow using DS instruction immediate1327// offsets even if the base pointer can't be proven to be base. On SI,1328// base pointer values that won't give the same result as a 16-bit add1329// are not safe to fold, but this will override the conservative test1330// for the base pointer.1331def FeatureEnableUnsafeDSOffsetFolding : SubtargetFeature <1332  "unsafe-ds-offset-folding",1333  "EnableUnsafeDSOffsetFolding",1334  "true",1335  "Force using DS instruction immediate offsets on SI"1336>;1337 1338def FeatureEnableSIScheduler : SubtargetFeature<"si-scheduler",1339  "EnableSIScheduler",1340  "true",1341  "Enable SI Machine Scheduler"1342>;1343 1344def FeatureEnableDS128 : SubtargetFeature<"enable-ds128",1345  "EnableDS128",1346  "true",1347  "Use ds_{read|write}_b128"1348>;1349 1350// Sparse texture support requires that all result registers are zeroed when1351// PRTStrictNull is set to true. This feature is turned on for all architectures1352// but is enabled as a feature in case there are situations where PRTStrictNull1353// is disabled by the driver.1354def FeatureEnablePRTStrictNull : SubtargetFeature<"enable-prt-strict-null",1355  "EnablePRTStrictNull",1356  "true",1357  "Enable zeroing of result registers for sparse texture fetches"1358>;1359 1360// Unless +-flat-for-global is specified, turn on FlatForGlobal for1361// all OS-es on VI and newer hardware to avoid assertion failures due1362// to missing ADDR64 variants of MUBUF instructions.1363// FIXME: moveToVALU should be able to handle converting addr64 MUBUF1364// instructions.1365 1366def FeatureFlatForGlobal : SubtargetFeature<"flat-for-global",1367  "FlatForGlobal",1368  "true",1369  "Force to generate flat instruction for global"1370>;1371 1372def FeatureAutoWaitcntBeforeBarrier : SubtargetFeature <1373  "auto-waitcnt-before-barrier",1374  "AutoWaitcntBeforeBarrier",1375  "true",1376  "Hardware automatically inserts waitcnt before barrier"1377>;1378 1379def FeatureBackOffBarrier : SubtargetFeature <"back-off-barrier",1380  "BackOffBarrier",1381  "true",1382  "Hardware supports backing off s_barrier if an exception occurs"1383>;1384 1385def FeatureTrigReducedRange : SubtargetFeature<"trig-reduced-range",1386  "HasTrigReducedRange",1387  "true",1388  "Requires use of fract on arguments to trig instructions"1389>;1390 1391def FeatureKernargPreload : SubtargetFeature <"kernarg-preload",1392  "KernargPreload",1393  "true",1394  "Hardware supports preloading of kernel arguments in user SGPRs."1395>;1396 1397// Alignment enforcement is controlled by a configuration register:1398// SH_MEM_CONFIG.alignment_mode1399def FeatureUnalignedAccessMode : SubtargetFeature<"unaligned-access-mode",1400  "UnalignedAccessMode",1401  "true",1402  "Enable unaligned global, local and region loads and stores if the hardware"1403  " supports it"1404>;1405 1406def FeaturePackedTID : SubtargetFeature<"packed-tid",1407  "HasPackedTID",1408  "true",1409  "Workitem IDs are packed into v0 at kernel launch"1410>;1411 1412def FeatureArchitectedFlatScratch : SubtargetFeature<"architected-flat-scratch",1413  "HasArchitectedFlatScratch",1414  "true",1415  "Flat Scratch register is a readonly SPI initialized architected register"1416>;1417 1418def FeatureArchitectedSGPRs : SubtargetFeature<"architected-sgprs",1419  "HasArchitectedSGPRs",1420  "true",1421  "Enable the architected SGPRs"1422>;1423 1424def FeatureGDS : SubtargetFeature<"gds",1425  "HasGDS",1426  "true",1427  "Has Global Data Share"1428>;1429 1430def FeatureGWS : SubtargetFeature<"gws",1431  "HasGWS",1432  "true",1433  "Has Global Wave Sync"1434>;1435 1436def FeatureRequiresCOV6 : SubtargetFeature<"requires-cov6",1437  "RequiresCOV6",1438  "true",1439  "Target Requires Code Object V6"1440>;1441 1442def FeatureXF32Insts : SubtargetFeature<"xf32-insts",1443   "HasXF32Insts",1444   "true",1445   "Has instructions that support xf32 format, such as "1446   "v_mfma_f32_16x16x8_xf32 and v_mfma_f32_32x32x4_xf32"1447 >;1448 1449def FeatureGloballyAddressableScratch : SubtargetFeature<1450  "globally-addressable-scratch",1451  "HasGloballyAddressableScratch",1452  "true",1453  "FLAT instructions can access scratch memory for any thread in any wave"1454>;1455 1456// Enable the use of SCRATCH_STORE/LOAD_BLOCK instructions for saving and1457// restoring the callee-saved registers.1458def FeatureUseBlockVGPROpsForCSR : SubtargetFeature<"block-vgpr-csr",1459  "UseBlockVGPROpsForCSR",1460  "true",1461  "Use block load/store for VGPR callee saved registers"1462>;1463 1464def FeatureLshlAddU64Inst1465    : SubtargetFeature<"lshl-add-u64-inst", "HasLshlAddU64Inst", "true",1466                       "Has v_lshl_add_u64 instruction">;1467 1468def FeatureAddSubU64Insts1469    : SubtargetFeature<"add-sub-u64-insts", "HasAddSubU64Insts", "true",1470                       "Has v_add_u64 and v_sub_u64 instructions">;1471 1472def FeatureMadU32Inst : SubtargetFeature<"mad-u32-inst", "HasMadU32Inst",1473                                         "true", "Has v_mad_u32 instruction">;1474 1475def FeatureAddMinMaxInsts : SubtargetFeature<"add-min-max-insts",1476  "HasAddMinMaxInsts",1477  "true",1478  "Has v_add_{min|max}_{i|u}32 instructions"1479>;1480 1481def FeaturePkAddMinMaxInsts : SubtargetFeature<"pk-add-min-max-insts",1482  "HasPkAddMinMaxInsts",1483  "true",1484  "Has v_pk_add_{min|max}_{i|u}16 instructions"1485>;1486 1487def FeatureMemToLDSLoad : SubtargetFeature<"vmem-to-lds-load-insts",1488  "HasVMemToLDSLoad",1489  "true",1490  "The platform has memory to lds instructions (global_load w/lds bit set, buffer_load w/lds bit set or global_load_lds. This does not include scratch_load_lds."1491>;1492 1493def FeatureLdsBarrierArriveAtomic : SubtargetFeature< "lds-barrier-arrive-atomic",1494  "HasLdsBarrierArriveAtomic",1495  "true",1496  "Has LDS barrier-arrive atomic instructions"1497>;1498 1499def Feature45BitNumRecordsBufferResource : SubtargetFeature< "45-bit-num-records-buffer-resource",1500  "Has45BitNumRecordsBufferResource",1501  "true",1502  "The buffer resource (V#) supports 45-bit num_records"1503>;1504 1505def FeatureClusters : SubtargetFeature< "clusters",1506  "HasClusters",1507  "true",1508  "Has clusters of workgroups support"1509>;1510 1511def FeatureWaitsBeforeSystemScopeStores : SubtargetFeature<1512  "waits-before-system-scope-stores",1513  "RequiresWaitsBeforeSystemScopeStores",1514  "true",1515  "Target requires waits for loads and atomics before system scope stores"1516>;1517 1518// Dummy feature used to disable assembler instructions.1519def FeatureDisable : SubtargetFeature<"",1520  "FeatureDisable","true",1521  "Dummy feature to disable assembler instructions"1522>;1523 1524//===----------------------------------------------------------------------===//1525 1526class GCNSubtargetFeatureGeneration <string Value,1527                                     string FeatureName,1528                                     list<SubtargetFeature> Implies> :1529        SubtargetFeatureGeneration <Value, FeatureName, "GCNSubtarget", Implies>;1530 1531def FeatureSouthernIslands : GCNSubtargetFeatureGeneration<"SOUTHERN_ISLANDS",1532    "southern-islands",1533  [FeatureFP64, FeatureAddressableLocalMemorySize32768, FeatureMIMG_R128,1534  FeatureWavefrontSize64, FeatureSMemTimeInst, FeatureMadMacF32Insts,1535  FeatureDsSrc2Insts, FeatureLDSBankCount32, FeatureMovrel,1536  FeatureTrigReducedRange, FeatureExtendedImageInsts, FeatureImageInsts,1537  FeatureGDS, FeatureGWS, FeatureDefaultComponentZero,1538  FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF64GlobalInsts,1539  FeatureVmemWriteVgprInOrder, FeatureCubeInsts, FeatureLerpInst,1540  FeatureSadInsts, FeatureCvtPkNormVOP2Insts1541  ]1542>;1543 1544def FeatureSeaIslands : GCNSubtargetFeatureGeneration<"SEA_ISLANDS",1545    "sea-islands",1546  [FeatureFP64, FeatureAddressableLocalMemorySize65536, FeatureMIMG_R128,1547  FeatureWavefrontSize64, FeatureFlatAddressSpace,1548  FeatureCIInsts, FeatureMovrel, FeatureTrigReducedRange,1549  FeatureGFX7GFX8GFX9Insts, FeatureSMemTimeInst, FeatureMadMacF32Insts,1550  FeatureDsSrc2Insts, FeatureExtendedImageInsts, FeatureUnalignedBufferAccess,1551  FeatureImageInsts, FeatureGDS, FeatureGWS, FeatureDefaultComponentZero,1552  FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF64GlobalInsts,1553  FeatureAtomicFMinFMaxF32FlatInsts, FeatureAtomicFMinFMaxF64FlatInsts,1554  FeatureVmemWriteVgprInOrder, FeatureCubeInsts, FeatureLerpInst,1555  FeatureSadInsts, FeatureQsadInsts, FeatureCvtPkNormVOP2Insts1556  ]1557>;1558 1559def FeatureVolcanicIslands : GCNSubtargetFeatureGeneration<"VOLCANIC_ISLANDS",1560  "volcanic-islands",1561  [FeatureFP64, FeatureAddressableLocalMemorySize65536, FeatureMIMG_R128,1562   FeatureWavefrontSize64, FeatureFlatAddressSpace,1563   FeatureGCN3Encoding, FeatureCIInsts, Feature16BitInsts,1564   FeatureSMemRealTime, FeatureVGPRIndexMode, FeatureMovrel,1565   FeatureScalarStores, FeatureInv2PiInlineImm,1566   FeatureSDWA, FeatureSDWAOutModsVOPC, FeatureSDWAMac, FeatureDPP,1567   FeatureIntClamp, FeatureTrigReducedRange, FeatureGFX8Insts,1568   FeatureGFX7GFX8GFX9Insts, FeatureSMemTimeInst, FeatureMadMacF32Insts,1569   FeatureDsSrc2Insts, FeatureExtendedImageInsts, FeatureFastDenormalF32,1570   FeatureUnalignedBufferAccess, FeatureImageInsts, FeatureGDS, FeatureGWS,1571   FeatureDefaultComponentZero, FeatureVmemWriteVgprInOrder, FeatureCubeInsts, 1572   FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts, 1573   FeatureCvtPkNormVOP2Insts1574  ]1575>;1576 1577def FeatureGFX9 : GCNSubtargetFeatureGeneration<"GFX9",1578  "gfx9",1579  [FeatureFP64,1580   FeatureWavefrontSize64, FeatureFlatAddressSpace,1581   FeatureGCN3Encoding, FeatureCIInsts, Feature16BitInsts,1582   FeatureSMemRealTime, FeatureScalarStores, FeatureInv2PiInlineImm,1583   FeatureApertureRegs, FeatureGFX9Insts, FeatureVOP3P, FeatureVGPRIndexMode,1584   FeatureFastFMAF32, FeatureDPP, FeatureIntClamp,1585   FeatureSDWA, FeatureSDWAOmod, FeatureSDWAScalar, FeatureSDWASdst,1586   FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts,1587   FeatureAddNoCarryInsts, FeatureGFX8Insts, FeatureGFX7GFX8GFX9Insts,1588   FeatureScalarFlatScratchInsts, FeatureScalarAtomics, FeatureR128A16,1589   FeatureA16, FeatureSMemTimeInst, FeatureFastDenormalF32, FeatureSupportsXNACK,1590   FeatureUnalignedBufferAccess, FeatureUnalignedScratchAccess,1591   FeatureUnalignedDSAccess, FeatureNegativeScratchOffsetBug, FeatureGWS,1592   FeatureDefaultComponentZero,FeatureVmemWriteVgprInOrder, FeatureMemToLDSLoad,1593   FeatureCubeInsts, FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts, 1594   FeatureCvtNormInsts, FeatureCvtPkNormVOP2Insts, 1595   FeatureCvtPkNormVOP3Insts1596  ]1597>;1598 1599def FeatureGFX10 : GCNSubtargetFeatureGeneration<"GFX10",1600  "gfx10",1601  [FeatureFP64, FeatureAddressableLocalMemorySize65536, FeatureMIMG_R128,1602   FeatureFlatAddressSpace,1603   FeatureCIInsts, Feature16BitInsts,1604   FeatureSMemRealTime, FeatureInv2PiInlineImm,1605   FeatureApertureRegs, FeatureGFX9Insts, FeatureGFX10Insts, FeatureVOP3P,1606   FeatureMovrel, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp,1607   FeatureSDWA, FeatureSDWAOmod, FeatureSDWAScalar, FeatureSDWASdst,1608   FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts,1609   FeatureAddNoCarryInsts, FeatureFmaMixInsts, FeatureGFX8Insts,1610   FeatureNoSdstCMPX, FeatureVscnt,1611   FeatureVOP3Literal, FeatureDPP8, FeatureExtendedImageInsts,1612   FeatureNoDataDepHazard, FeaturePkFmacF16Inst,1613   FeatureA16, FeatureSMemTimeInst, FeatureFastDenormalF32, FeatureG16,1614   FeatureUnalignedBufferAccess, FeatureUnalignedScratchAccess,1615   FeatureUnalignedDSAccess, FeatureImageInsts, FeatureGDS, FeatureGWS,1616   FeatureDefaultComponentZero, FeatureMaxHardClauseLength63,1617   FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF64GlobalInsts,1618   FeatureAtomicFMinFMaxF32FlatInsts, FeatureAtomicFMinFMaxF64FlatInsts,1619   FeatureVmemWriteVgprInOrder, FeatureMemToLDSLoad, FeatureCubeInsts, 1620   FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts, 1621   FeatureCvtNormInsts, FeatureCvtPkNormVOP2Insts, 1622   FeatureCvtPkNormVOP3Insts1623  ]1624>;1625 1626def FeatureGFX11 : GCNSubtargetFeatureGeneration<"GFX11",1627  "gfx11",1628  [FeatureFP64, FeatureAddressableLocalMemorySize65536, FeatureMIMG_R128,1629   FeatureFlatAddressSpace, Feature16BitInsts,1630   FeatureInv2PiInlineImm, FeatureApertureRegs,1631   FeatureCIInsts, FeatureGFX8Insts, FeatureGFX9Insts, FeatureGFX10Insts,1632   FeatureGFX10_AEncoding, FeatureGFX10_BEncoding, FeatureGFX10_3Insts,1633   FeatureGFX11Insts, FeatureVOP3P, FeatureVOPD, FeatureTrue16BitInsts,1634   FeatureMovrel, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp,1635   FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts,1636   FeatureAddNoCarryInsts, FeatureFmaMixInsts,1637   FeatureNoSdstCMPX, FeatureVscnt,1638   FeatureVOP3Literal, FeatureDPP8, FeatureExtendedImageInsts,1639   FeatureNoDataDepHazard, FeaturePkFmacF16Inst,1640   FeatureA16, FeatureFastDenormalF32, FeatureG16,1641   FeatureUnalignedBufferAccess, FeatureUnalignedScratchAccess,1642   FeatureUnalignedDSAccess, FeatureGDS, FeatureGWS,1643   FeatureDefaultComponentZero, FeatureMaxHardClauseLength32,1644   FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF32FlatInsts,1645   FeatureVmemWriteVgprInOrder, FeatureCubeInsts, FeatureLerpInst,1646   FeatureSadInsts, FeatureQsadInsts, FeatureCvtNormInsts, 1647   FeatureCvtPkNormVOP2Insts, FeatureCvtPkNormVOP3Insts1648  ]1649>;1650 1651def FeatureGFX12 : GCNSubtargetFeatureGeneration<"GFX12",1652  "gfx12",1653  [FeatureFP64, FeatureMIMG_R128,1654   FeatureFlatAddressSpace, Feature16BitInsts,1655   FeatureInv2PiInlineImm, FeatureApertureRegs,1656   FeatureCIInsts, FeatureGFX8Insts, FeatureGFX9Insts, FeatureGFX10Insts,1657   FeatureGFX10_AEncoding, FeatureGFX10_BEncoding, FeatureGFX10_3Insts,1658   FeatureGFX11Insts, FeatureGFX12Insts, FeatureVOP3P, FeatureVOPD,1659   FeatureMovrel, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp,1660   FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts,1661   FeatureAddNoCarryInsts, FeatureFmaMixInsts,1662   FeatureNoSdstCMPX, FeatureVscnt,1663   FeatureVOP3Literal, FeatureDPP8,1664   FeatureNoDataDepHazard, FeaturePkFmacF16Inst,1665   FeatureA16, FeatureFastDenormalF32, FeatureG16,1666   FeatureUnalignedBufferAccess, FeatureUnalignedScratchAccess,1667   FeatureUnalignedDSAccess, FeatureTrue16BitInsts,1668   FeatureDefaultComponentBroadcast, FeatureMaxHardClauseLength32,1669   FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF32FlatInsts,1670   FeatureIEEEMinimumMaximumInsts, FeatureMinimum3Maximum3F32,1671   FeatureMinimum3Maximum3F16, FeatureAgentScopeFineGrainedRemoteMemoryAtomics1672  ]1673>;1674 1675//===----------------------------------------------------------------------===//1676 1677class FeatureSet<list<SubtargetFeature> Features_> {1678  list<SubtargetFeature> Features = Features_;1679}1680 1681def FeatureISAVersion6_0_0 : FeatureSet<[FeatureSouthernIslands,1682   FeatureFastFMAF32,1683   HalfRate64Ops,1684   FeatureLDSBankCount32]>;1685 1686def FeatureISAVersion6_0_1 : FeatureSet<1687  [FeatureSouthernIslands,1688   FeatureLDSBankCount32]>;1689 1690def FeatureISAVersion6_0_2 : FeatureSet<1691  [FeatureSouthernIslands,1692   FeatureLDSBankCount32]>;1693 1694def FeatureISAVersion7_0_0 : FeatureSet<1695  [FeatureSeaIslands,1696   FeatureLDSBankCount32]>;1697 1698def FeatureISAVersion7_0_1 : FeatureSet<1699  [FeatureSeaIslands,1700   HalfRate64Ops,1701   FeatureLDSBankCount32,1702   FeatureFastFMAF32]>;1703 1704def FeatureISAVersion7_0_2 : FeatureSet<1705  [FeatureSeaIslands,1706   FeatureLDSBankCount16,1707   FeatureFastFMAF32]>;1708 1709def FeatureISAVersion7_0_3 : FeatureSet<1710  [FeatureSeaIslands,1711   FeatureLDSBankCount16]>;1712 1713def FeatureISAVersion7_0_4 : FeatureSet<1714  [FeatureSeaIslands,1715   FeatureLDSBankCount32]>;1716 1717def FeatureISAVersion7_0_5 : FeatureSet<1718  [FeatureSeaIslands,1719   FeatureLDSBankCount16]>;1720 1721def FeatureISAVersion8_0_Common : FeatureSet<1722  [FeatureVolcanicIslands,1723   FeatureLDSBankCount32,1724   FeatureUnpackedD16VMem]>;1725 1726def FeatureISAVersion8_0_1 : FeatureSet<1727  !listconcat(FeatureISAVersion8_0_Common.Features,1728    [FeatureFastFMAF32,1729     HalfRate64Ops,1730     FeatureSupportsXNACK])>;1731 1732def FeatureISAVersion8_0_2 : FeatureSet<1733  !listconcat(FeatureISAVersion8_0_Common.Features,1734    [FeatureSGPRInitBug])>;1735 1736def FeatureISAVersion8_0_3 : FeatureSet<1737  !listconcat(FeatureISAVersion8_0_Common.Features,1738    [])>;1739 1740def FeatureISAVersion8_0_5 : FeatureSet<1741  !listconcat(FeatureISAVersion8_0_Common.Features,1742    [FeatureSGPRInitBug])>;1743 1744def FeatureISAVersion8_1_0 : FeatureSet<1745  [FeatureVolcanicIslands,1746   FeatureLDSBankCount16,1747   FeatureSupportsXNACK,1748   FeatureImageStoreD16Bug,1749   FeatureImageGather4D16Bug]>;1750 1751def FeatureISAVersion9_0_Common : FeatureSet<1752  [FeatureGFX9,1753   FeatureAddressableLocalMemorySize65536,1754   FeatureLDSBankCount32,1755   FeatureImageInsts,1756   FeatureMadMacF32Insts]>;1757 1758def FeatureISAVersion9_0_Consumer_Common : FeatureSet<1759  !listconcat(FeatureISAVersion9_0_Common.Features,1760    [FeatureImageGather4D16Bug,1761     FeatureDsSrc2Insts,1762     FeatureExtendedImageInsts,1763     FeatureGDS])>;1764 1765def FeatureISAVersion9_Generic : FeatureSet<1766  !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,1767    [FeatureRequiresCOV6])>;1768 1769def FeatureISAVersion9_0_MI_Common : FeatureSet<1770  !listconcat(FeatureISAVersion9_0_Common.Features,1771    [FeatureFmaMixInsts,1772     FeatureDLInsts,1773     FeatureDot1Insts,1774     FeatureDot2Insts,1775     FeatureDot3Insts,1776     FeatureDot4Insts,1777     FeatureDot5Insts,1778     FeatureDot6Insts,1779     FeatureDot7Insts,1780     FeatureDot10Insts,1781     FeatureMAIInsts,1782     FeaturePkFmacF16Inst,1783     FeatureAtomicFaddNoRtnInsts,1784     FeatureSupportsSRAMECC])>;1785 1786def FeatureISAVersion9_0_0 : FeatureSet<1787  !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,1788    [FeatureMadMixInsts])>;1789 1790def FeatureISAVersion9_0_2 : FeatureSet<1791  !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,1792    [FeatureMadMixInsts])>;1793 1794def FeatureISAVersion9_0_4 : FeatureSet<1795  !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,1796    [FeatureFmaMixInsts])>;1797 1798def FeatureISAVersion9_0_6 : FeatureSet<1799  !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,1800    [HalfRate64Ops,1801     FeatureFmaMixInsts,1802     FeatureDLInsts,1803     FeatureDot1Insts,1804     FeatureDot2Insts,1805     FeatureDot7Insts,1806     FeatureDot10Insts,1807     FeatureSupportsSRAMECC])>;1808 1809def FeatureISAVersion9_0_8 : FeatureSet<1810  !listconcat(FeatureISAVersion9_0_MI_Common.Features,1811    [FeatureGDS,1812     HalfRate64Ops,1813     FeatureDsSrc2Insts,1814     FeatureExtendedImageInsts,1815     FeatureAtomicBufferGlobalPkAddF16NoRtnInsts,1816     FeatureMFMAInlineLiteralBug,1817     FeatureImageGather4D16Bug])>;1818 1819def FeatureISAVersion9_0_9 : FeatureSet<1820  !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,1821    [FeatureMadMixInsts])>;1822 1823def FeatureISAVersion9_0_A : FeatureSet<1824  !listconcat(FeatureISAVersion9_0_MI_Common.Features,1825    [FeatureGFX90AInsts,1826     FeatureRequiresAlignedVGPRs,1827     FeatureFmacF64Inst,1828     FeatureDPALU_DPP,1829     FeaturePackedFP32Ops,1830     FeatureAtomicFaddRtnInsts,1831     FeatureAtomicBufferGlobalPkAddF16Insts,1832     FeaturePackedTID,1833     FullRate64Ops,1834     FeatureBackOffBarrier,1835     FeatureKernargPreload,1836     FeatureAtomicFMinFMaxF64GlobalInsts,1837     FeatureAtomicFMinFMaxF64FlatInsts,1838     FeatureFlatBufferGlobalAtomicFaddF64Inst1839     ])>;1840 1841def FeatureISAVersion9_0_C : FeatureSet<1842  !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,1843    [FeatureMadMixInsts])>;1844 1845def FeatureISAVersion9_4_Common : FeatureSet<1846  [FeatureGFX9,1847   FeatureGFX90AInsts,1848   FeatureGFX940Insts,1849   FeatureRequiresAlignedVGPRs,1850   FeatureFmaMixInsts,1851   FeatureLDSBankCount32,1852   FeatureDLInsts,1853   FeatureFmacF64Inst,1854   FeatureDot1Insts,1855   FeatureDot2Insts,1856   FeatureDot3Insts,1857   FeatureDot4Insts,1858   FeatureDot5Insts,1859   FeatureDot6Insts,1860   FeatureDot7Insts,1861   FeatureDot10Insts,1862   FeatureAtomicDsPkAdd16Insts,1863   FeatureAtomicFlatPkAdd16Insts,1864   FeatureDPALU_DPP,1865   FeaturePackedFP32Ops,1866   FeatureMAIInsts,1867   FeaturePkFmacF16Inst,1868   FeatureAtomicFaddRtnInsts,1869   FeatureAtomicFaddNoRtnInsts,1870   FeatureAtomicBufferGlobalPkAddF16Insts,1871   FeatureAtomicGlobalPkAddBF16Inst,1872   FeatureFlatAtomicFaddF32Inst,1873   FeatureSupportsSRAMECC,1874   FeaturePackedTID,1875   FeatureArchitectedFlatScratch,1876   FullRate64Ops,1877   FeatureBackOffBarrier,1878   FeatureKernargPreload,1879   FeatureAtomicFMinFMaxF64GlobalInsts,1880   FeatureAtomicFMinFMaxF64FlatInsts,1881   FeatureAgentScopeFineGrainedRemoteMemoryAtomics,1882   FeatureMemoryAtomicFAddF32DenormalSupport,1883   FeatureFlatBufferGlobalAtomicFaddF64Inst,1884   FeatureLshlAddU64Inst,1885   ]>;1886 1887def FeatureISAVersion9_5_Common : FeatureSet<1888  !listconcat(FeatureISAVersion9_4_Common.Features,1889  [FeatureAddressableLocalMemorySize163840,1890   FeatureFP8Insts,1891   FeatureFP8ConversionInsts,1892   FeatureGFX950Insts,1893   FeaturePrngInst,1894   FeatureBF16ConversionInsts,1895   FeatureBitOp3Insts,1896   FeatureFP8ConversionScaleInsts,1897   FeatureBF8ConversionScaleInsts,1898   FeatureFP4ConversionScaleInsts,1899   FeatureFP6BF6ConversionScaleInsts,1900   FeatureDot12Insts,1901   FeatureDot13Insts,1902   FeatureAtomicBufferPkAddBF16Inst1903   ])>;1904 1905def FeatureISAVersion9_4_2 : FeatureSet<1906  !listconcat(FeatureISAVersion9_4_Common.Features,1907    [1908      FeatureAddressableLocalMemorySize65536,1909      FeatureFP8Insts,1910      FeatureFP8ConversionInsts,1911      FeatureCvtFP8VOP1Bug,1912      FeatureXF32Insts1913    ])>;1914 1915def FeatureISAVersion9_4_Generic : FeatureSet<1916  !listconcat(FeatureISAVersion9_4_Common.Features,1917    [FeatureAddressableLocalMemorySize65536,1918     FeatureRequiresCOV6])>;1919 1920def FeatureISAVersion9_5_0 : FeatureSet<FeatureISAVersion9_5_Common.Features>;1921 1922def FeatureISAVersion10_Common : FeatureSet<1923  [FeatureGFX10,1924   FeatureLDSBankCount32,1925   FeatureDLInsts,1926   FeatureNSAEncoding,1927   FeatureBackOffBarrier]>;1928 1929def FeatureISAVersion10_1_Common : FeatureSet<1930  !listconcat(FeatureISAVersion10_Common.Features,1931    [FeatureScalarStores,1932     FeatureScalarAtomics,1933     FeatureScalarFlatScratchInsts,1934     FeatureGetWaveIdInst,1935     FeatureMadMacF32Insts,1936     FeatureDsSrc2Insts,1937     FeatureLdsMisalignedBug,1938     FeatureSupportsXNACK,1939     // gfx101x bugs1940     FeatureVcmpxPermlaneHazard,1941     FeatureVMEMtoScalarWriteHazard,1942     FeatureSMEMtoVectorWriteHazard,1943     FeatureInstFwdPrefetchBug,1944     FeatureVcmpxExecWARHazard,1945     FeatureLdsBranchVmemWARHazard,1946     FeatureNSAtoVMEMBug,1947     FeatureNSAClauseBug,1948     FeatureOffset3fBug,1949     FeatureFlatSegmentOffsetBug,1950     FeatureNegativeUnalignedScratchOffsetBug])>;1951 1952def FeatureISAVersion10_1_Generic : FeatureSet<1953  !listconcat(FeatureISAVersion10_1_Common.Features,1954    [FeatureRequiresCOV6])>;1955 1956def FeatureISAVersion10_1_0 : FeatureSet<1957  !listconcat(FeatureISAVersion10_1_Common.Features,1958    [])>;1959 1960def FeatureISAVersion10_1_1 : FeatureSet<1961  !listconcat(FeatureISAVersion10_1_Common.Features,1962    [FeatureDot1Insts,1963     FeatureDot2Insts,1964     FeatureDot5Insts,1965     FeatureDot6Insts,1966     FeatureDot7Insts,1967     FeatureDot10Insts])>;1968 1969def FeatureISAVersion10_1_2 : FeatureSet<1970  !listconcat(FeatureISAVersion10_1_Common.Features,1971    [FeatureDot1Insts,1972     FeatureDot2Insts,1973     FeatureDot5Insts,1974     FeatureDot6Insts,1975     FeatureDot7Insts,1976     FeatureDot10Insts])>;1977 1978def FeatureISAVersion10_1_3 : FeatureSet<1979  !listconcat(FeatureISAVersion10_1_Common.Features,1980    [FeatureGFX10_AEncoding])>;1981 1982def FeatureISAVersion10_3_0 : FeatureSet<1983  !listconcat(FeatureISAVersion10_Common.Features,1984    [FeatureGFX10_AEncoding,1985     FeatureGFX10_BEncoding,1986     FeatureGFX10_3Insts,1987     FeatureDot1Insts,1988     FeatureDot2Insts,1989     FeatureDot5Insts,1990     FeatureDot6Insts,1991     FeatureDot7Insts,1992     FeatureDot10Insts,1993     FeatureShaderCyclesRegister])>;1994 1995def FeatureISAVersion10_3_Generic: FeatureSet<1996  !listconcat(FeatureISAVersion10_3_0.Features,1997    [FeatureRequiresCOV6])>;1998 1999def FeatureISAVersion11_Common : FeatureSet<2000  [FeatureGFX11,2001   FeatureBackOffBarrier,2002   FeatureLDSBankCount32,2003   FeatureDLInsts,2004   FeatureDot5Insts,2005   FeatureDot7Insts,2006   FeatureDot8Insts,2007   FeatureDot9Insts,2008   FeatureDot10Insts,2009   FeatureDot12Insts,2010   FeatureNSAEncoding,2011   FeaturePartialNSAEncoding,2012   FeatureShaderCyclesRegister,2013   FeatureArchitectedFlatScratch,2014   FeatureAtomicFaddRtnInsts,2015   FeatureAtomicFaddNoRtnInsts,2016   FeatureFlatAtomicFaddF32Inst,2017   FeatureImageInsts,2018   FeaturePackedTID,2019   FeatureVcmpxPermlaneHazard,2020   FeatureMemoryAtomicFAddF32DenormalSupport,2021   FeatureRealTrue16Insts,2022   FeatureD16Writes32BitVgpr,2023]>;2024 2025// There are few workarounds that need to be2026// added to all targets. This pessimizes codegen2027// a bit on the generic GFX11 target.2028def FeatureISAVersion11_Generic: FeatureSet<2029  !listconcat(FeatureISAVersion11_Common.Features,2030    [FeatureMSAALoadDstSelBug,2031     FeatureVALUTransUseHazard,2032     FeatureUserSGPRInit16Bug,2033     FeatureMADIntraFwdBug,2034     FeaturePrivEnabledTrap2NopBug,2035     FeatureRequiresCOV6,2036     FeatureRequiredExportPriority])>;2037 2038def FeatureISAVersion11_0_Common : FeatureSet<2039  !listconcat(FeatureISAVersion11_Common.Features,2040    [FeatureMSAALoadDstSelBug,2041     FeatureVALUTransUseHazard,2042     FeatureMADIntraFwdBug,2043     FeaturePrivEnabledTrap2NopBug])>;2044 2045def FeatureISAVersion11_0_0 : FeatureSet<2046  !listconcat(FeatureISAVersion11_0_Common.Features,2047    [Feature1_5xVGPRs,2048     FeatureUserSGPRInit16Bug])>;2049 2050def FeatureISAVersion11_0_1 : FeatureSet<2051  !listconcat(FeatureISAVersion11_0_Common.Features,2052    [Feature1_5xVGPRs])>;2053 2054def FeatureISAVersion11_0_2 : FeatureSet<2055  !listconcat(FeatureISAVersion11_0_Common.Features,2056    [FeatureUserSGPRInit16Bug])>;2057 2058def FeatureISAVersion11_0_3 : FeatureSet<2059  !listconcat(FeatureISAVersion11_0_Common.Features,2060    [])>;2061 2062def FeatureISAVersion11_5_Common : FeatureSet<2063  !listconcat(FeatureISAVersion11_Common.Features,2064    [FeatureSALUFloatInsts,2065     FeatureDPPSrc1SGPR,2066     FeatureRequiredExportPriority])>;2067 2068def FeatureISAVersion11_5_0 : FeatureSet<2069  !listconcat(FeatureISAVersion11_5_Common.Features,2070    [FeaturePointSampleAccel])>;2071 2072def FeatureISAVersion11_5_1 : FeatureSet<2073  !listconcat(FeatureISAVersion11_5_Common.Features,2074    [Feature1_5xVGPRs,2075     FeaturePointSampleAccel])>;2076 2077def FeatureISAVersion11_5_2 : FeatureSet<2078  !listconcat(FeatureISAVersion11_5_Common.Features,2079    [FeaturePointSampleAccel])>;2080 2081def FeatureISAVersion11_5_3 : FeatureSet<2082  !listconcat(FeatureISAVersion11_5_Common.Features,2083    [])>;2084 2085def FeatureISAVersion12 : FeatureSet<2086  [FeatureGFX12,2087   FeatureBackOffBarrier,2088   FeatureAddressableLocalMemorySize65536,2089   FeatureLDSBankCount32,2090   FeatureDLInsts,2091   FeatureDot7Insts,2092   FeatureDot8Insts,2093   FeatureDot9Insts,2094   FeatureDot10Insts,2095   FeatureDot11Insts,2096   FeatureDot12Insts,2097   FeatureNSAEncoding,2098   FeaturePartialNSAEncoding,2099   FeatureShaderCyclesHiLoRegisters,2100   FeatureArchitectedFlatScratch,2101   FeatureArchitectedSGPRs,2102   FeatureAtomicFaddRtnInsts,2103   FeatureAtomicFaddNoRtnInsts,2104   FeatureAtomicDsPkAdd16Insts,2105   FeatureAtomicFlatPkAdd16Insts,2106   FeatureAtomicBufferGlobalPkAddF16Insts,2107   FeatureAtomicGlobalPkAddBF16Inst,2108   FeatureAtomicBufferPkAddBF16Inst,2109   FeatureFlatAtomicFaddF32Inst,2110   FeatureImageInsts,2111   FeatureExtendedImageInsts,2112   FeatureFP8ConversionInsts,2113   FeatureIEEEMinimumMaximumInsts,2114   FeaturePackedTID,2115   FeatureVcmpxPermlaneHazard,2116   FeatureSALUFloatInsts,2117   FeaturePseudoScalarTrans,2118   FeatureHasRestrictedSOffset,2119   FeatureScalarDwordx3Loads,2120   FeatureDPPSrc1SGPR,2121   FeatureMaxHardClauseLength32,2122   Feature1_5xVGPRs,2123   FeatureMemoryAtomicFAddF32DenormalSupport,2124   FeatureBVHDualAndBVH8Insts,2125   FeatureWaitsBeforeSystemScopeStores,2126   FeatureD16Writes32BitVgpr,2127   FeatureCubeInsts, 2128   FeatureLerpInst, 2129   FeatureSadInsts,2130   FeatureQsadInsts, 2131   FeatureCvtNormInsts, 2132   FeatureCvtPkNormVOP2Insts, 2133   FeatureCvtPkNormVOP3Insts2134   ]>;2135 2136def FeatureISAVersion12_50_Common : FeatureSet<2137  [FeatureGFX12,2138   FeatureGFX1250Insts,2139   FeatureRequiresAlignedVGPRs,2140   FeatureAddressableLocalMemorySize327680,2141   FeatureCuMode,2142   Feature1024AddressableVGPRs,2143   Feature64BitLiterals,2144   FeatureLDSBankCount32,2145   FeatureDLInsts,2146   FeatureFmacF64Inst,2147   FeaturePackedFP32Ops,2148   FeatureDot7Insts,2149   FeatureDot8Insts,2150   FeatureWavefrontSize32,2151   FeatureShaderCyclesHiLoRegisters,2152   FeatureArchitectedFlatScratch,2153   FeatureArchitectedSGPRs,2154   FeatureFlatGVSMode,2155   FeatureAtomicFaddRtnInsts,2156   FeatureAtomicFaddNoRtnInsts,2157   FeatureAtomicDsPkAdd16Insts,2158   FeatureAtomicFlatPkAdd16Insts,2159   FeatureAtomicBufferGlobalPkAddF16Insts,2160   FeatureAtomicGlobalPkAddBF16Inst,2161   FeatureAtomicBufferPkAddBF16Inst,2162   FeatureFlatAtomicFaddF32Inst,2163   FeatureFP8ConversionInsts,2164   FeatureFP8E5M3Insts,2165   FeaturePackedTID,2166   FeatureVcmpxPermlaneHazard,2167   FeatureSALUFloatInsts,2168   FeaturePseudoScalarTrans,2169   FeatureHasRestrictedSOffset,2170   FeatureScalarDwordx3Loads,2171   FeatureDPPSrc1SGPR,2172   FeatureBitOp3Insts,2173   FeatureTanhInsts,2174   FeatureTensorCvtLutInsts,2175   FeatureTransposeLoadF4F6Insts,2176   FeatureBF16TransInsts,2177   FeatureBF16ConversionInsts,2178   FeatureBF16PackedInsts,2179   FeatureCvtPkF16F32Inst,2180   FeatureFmaMixBF16Insts,2181   FeatureMin3Max3PKF16,2182   FeatureMinimum3Maximum3PKF16,2183   FeaturePrngInst,2184   FeaturePermlane16Swap,2185   FeatureAshrPkInsts,2186   FeatureSupportsSRAMECC,2187   FeatureMaxHardClauseLength63,2188   FeatureWaitXcnt,2189   FeatureAtomicFMinFMaxF64GlobalInsts,2190   FeatureAtomicFMinFMaxF64FlatInsts,2191   FeatureFlatBufferGlobalAtomicFaddF64Inst,2192   FeatureMemoryAtomicFAddF32DenormalSupport,2193   FeatureEmulatedSystemScopeAtomics,2194   FeatureGloballyAddressableScratch,2195   FeatureKernargPreload,2196   FeatureVmemPrefInsts,2197   FeatureLshlAddU64Inst,2198   FeatureAddSubU64Insts,2199   FeatureMadU32Inst,2200   FeatureAddMinMaxInsts,2201   FeaturePkAddMinMaxInsts,2202   FeatureLdsBarrierArriveAtomic,2203   FeatureSetPrioIncWgInst,2204   Feature45BitNumRecordsBufferResource,2205   FeatureSupportsXNACK,2206   FeatureXNACK,2207   FeatureClusters,2208   FeatureD16Writes32BitVgpr,2209]>;2210 2211def FeatureISAVersion12_50 : FeatureSet<2212  !listconcat(FeatureISAVersion12_50_Common.Features,2213  [FeatureCubeInsts, 2214   FeatureLerpInst, 2215   FeatureSadInsts, 2216   FeatureQsadInsts, 2217   FeatureCvtNormInsts, 2218   FeatureCvtPkNormVOP2Insts, 2219   FeatureCvtPkNormVOP3Insts])>;2220 2221def FeatureISAVersion12_51 : FeatureSet<2222  !listconcat(FeatureISAVersion12_50.Features,2223  [FeatureDPALU_DPP])>;2224 2225def FeatureISAVersion12_Generic: FeatureSet<2226  !listconcat(FeatureISAVersion12.Features,2227    [FeatureRequiresCOV6])>;2228 2229//===----------------------------------------------------------------------===//2230 2231def AMDGPUInstrInfo : InstrInfo {2232  let guessInstructionProperties = 1;2233}2234 2235def AMDGPUAsmParser : AsmParser {2236  // Some of the R600 registers have the same name, so this crashes.2237  // For example T0_XYZW and T0_XY both have the asm name T0.2238  let ShouldEmitMatchRegisterName = 0;2239 2240  // Call the custom operand parser for all operands.2241  let OperandParserMethod = "parseCustomOperand";2242  let CallCustomParserForAllOperands = true;2243}2244 2245def AMDGPUAsmWriter : AsmWriter {2246  int PassSubtarget = 1;2247}2248 2249def AMDGPUAsmVariants {2250  string Default = "Default";2251  int Default_ID = 0;2252  string VOP3 = "VOP3";2253  int VOP3_ID = 1;2254  string SDWA = "SDWA";2255  int SDWA_ID = 2;2256  string SDWA9 = "SDWA9";2257  int SDWA9_ID = 3;2258  string DPP = "DPP";2259  int DPP_ID = 4;2260  string VOP3_DPP = "VOP3_DPP";2261  int VOP3_DPP_ID = 5;2262  string Disable = "Disable";2263  int Disable_ID = 6;2264}2265 2266def DefaultAMDGPUAsmParserVariant : AsmParserVariant {2267  let Variant = AMDGPUAsmVariants.Default_ID;2268  let Name = AMDGPUAsmVariants.Default;2269}2270 2271def VOP3AsmParserVariant : AsmParserVariant {2272  let Variant = AMDGPUAsmVariants.VOP3_ID;2273  let Name = AMDGPUAsmVariants.VOP3;2274}2275 2276def SDWAAsmParserVariant : AsmParserVariant {2277  let Variant = AMDGPUAsmVariants.SDWA_ID;2278  let Name = AMDGPUAsmVariants.SDWA;2279}2280 2281def SDWA9AsmParserVariant : AsmParserVariant {2282  let Variant = AMDGPUAsmVariants.SDWA9_ID;2283  let Name = AMDGPUAsmVariants.SDWA9;2284}2285 2286def DPPAsmParserVariant : AsmParserVariant {2287  let Variant = AMDGPUAsmVariants.DPP_ID;2288  let Name = AMDGPUAsmVariants.DPP;2289}2290 2291def VOP3_DPPAsmParserVariant : AsmParserVariant {2292  let Variant = AMDGPUAsmVariants.VOP3_DPP_ID;2293  let Name = AMDGPUAsmVariants.VOP3_DPP;2294}2295 2296def AMDGPU : Target {2297  // Pull in Instruction Info:2298  let InstructionSet = AMDGPUInstrInfo;2299  let AssemblyParsers = [AMDGPUAsmParser];2300  let AssemblyParserVariants = [DefaultAMDGPUAsmParserVariant,2301                                VOP3AsmParserVariant,2302                                SDWAAsmParserVariant,2303                                SDWA9AsmParserVariant,2304                                DPPAsmParserVariant,2305                                VOP3_DPPAsmParserVariant];2306  let AssemblyWriters = [AMDGPUAsmWriter];2307  let AllowRegisterRenaming = 1;2308}2309 2310// Dummy Instruction itineraries for pseudo instructions2311def ALU_NULL : FuncUnit;2312def NullALU : InstrItinClass;2313 2314//===----------------------------------------------------------------------===//2315// Predicate helper class2316//===----------------------------------------------------------------------===//2317 2318def isGFX6 :2319  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::SOUTHERN_ISLANDS">,2320  AssemblerPredicate<(all_of FeatureSouthernIslands)>;2321 2322def isGFX6GFX7 :2323  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::SOUTHERN_ISLANDS ||"2324            "Subtarget->getGeneration() == AMDGPUSubtarget::SEA_ISLANDS">,2325  AssemblerPredicate<(all_of (not FeatureGCN3Encoding), (not FeatureGFX10Insts))>;2326 2327def isGFX6GFX7GFX10 :2328  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::SOUTHERN_ISLANDS ||"2329            "Subtarget->getGeneration() == AMDGPUSubtarget::SEA_ISLANDS ||"2330            "Subtarget->getGeneration() == AMDGPUSubtarget::GFX10">,2331  AssemblerPredicate<(all_of (not FeatureGCN3Encoding), (not FeatureGFX11Insts))>;2332 2333def isGFX6GFX7GFX10Plus :2334  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::SOUTHERN_ISLANDS ||"2335            "Subtarget->getGeneration() == AMDGPUSubtarget::SEA_ISLANDS ||"2336            "Subtarget->getGeneration() >= AMDGPUSubtarget::GFX10">,2337  AssemblerPredicate<(all_of (not FeatureGCN3Encoding))>;2338 2339def isGFX7Only :2340  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::SEA_ISLANDS">,2341  AssemblerPredicate<(all_of (not FeatureGCN3Encoding), FeatureCIInsts, (not FeatureGFX10Insts))>;2342 2343def isGFX7GFX10 :2344  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::SEA_ISLANDS ||"2345            "Subtarget->getGeneration() == AMDGPUSubtarget::GFX10">,2346  AssemblerPredicate<(all_of (not FeatureGCN3Encoding), FeatureCIInsts, (not FeatureGFX11Insts))>;2347 2348def isGFX7GFX10GFX11 :2349  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::SEA_ISLANDS ||"2350            "Subtarget->getGeneration() == AMDGPUSubtarget::GFX10 ||"2351            "Subtarget->getGeneration() == AMDGPUSubtarget::GFX11">,2352  AssemblerPredicate<(all_of (not FeatureGCN3Encoding), FeatureCIInsts)>;2353 2354def isGFX7GFX8GFX9 :2355  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::SEA_ISLANDS ||"2356            "Subtarget->getGeneration() == AMDGPUSubtarget::VOLCANIC_ISLANDS ||"2357            "Subtarget->getGeneration() == AMDGPUSubtarget::GFX9">,2358  AssemblerPredicate<(all_of FeatureGFX7GFX8GFX9Insts)>;2359 2360def isGFX6GFX7GFX8GFX9 :2361  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::SOUTHERN_ISLANDS ||"2362            "Subtarget->getGeneration() == AMDGPUSubtarget::SEA_ISLANDS ||"2363            "Subtarget->getGeneration() == AMDGPUSubtarget::VOLCANIC_ISLANDS ||"2364            "Subtarget->getGeneration() == AMDGPUSubtarget::GFX9">,2365  AssemblerPredicate<(all_of (not FeatureGFX10Insts))>;2366 2367def isGFX6GFX7GFX8GFX9NotGFX90A :2368  Predicate<"!Subtarget->hasGFX90AInsts() &&"2369            "(Subtarget->getGeneration() == AMDGPUSubtarget::SOUTHERN_ISLANDS ||"2370            " Subtarget->getGeneration() == AMDGPUSubtarget::SEA_ISLANDS ||"2371            " Subtarget->getGeneration() == AMDGPUSubtarget::VOLCANIC_ISLANDS ||"2372            " Subtarget->getGeneration() == AMDGPUSubtarget::GFX9)">,2373  AssemblerPredicate<(all_of (not FeatureGFX10Insts), (not FeatureGFX90AInsts))>;2374 2375def isGFX6GFX7GFX8GFX9GFX10 :2376  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::SOUTHERN_ISLANDS ||"2377            "Subtarget->getGeneration() == AMDGPUSubtarget::SEA_ISLANDS ||"2378            "Subtarget->getGeneration() == AMDGPUSubtarget::VOLCANIC_ISLANDS ||"2379            "Subtarget->getGeneration() == AMDGPUSubtarget::GFX9 ||"2380            "Subtarget->getGeneration() == AMDGPUSubtarget::GFX10">,2381  AssemblerPredicate<(all_of (not FeatureGFX11Insts))>;2382 2383def isNotGFX12Plus :2384  Predicate<"Subtarget->getGeneration() <= AMDGPUSubtarget::GFX11">,2385  AssemblerPredicate<(all_of (not FeatureGFX12Insts))>;2386 2387def isGFX7GFX8GFX9GFX10 :2388  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::SEA_ISLANDS ||"2389            "Subtarget->getGeneration() == AMDGPUSubtarget::VOLCANIC_ISLANDS ||"2390            "Subtarget->getGeneration() == AMDGPUSubtarget::GFX9 ||"2391            "Subtarget->getGeneration() == AMDGPUSubtarget::GFX10">,2392  AssemblerPredicate<(all_of FeatureCIInsts, (not FeatureGFX11Insts))>;2393 2394def isGFX8GFX9GFX10GFX11 :2395  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::VOLCANIC_ISLANDS ||"2396            "Subtarget->getGeneration() == AMDGPUSubtarget::GFX9 ||"2397            "Subtarget->getGeneration() == AMDGPUSubtarget::GFX10 ||"2398            "Subtarget->getGeneration() == AMDGPUSubtarget::GFX11">,2399  AssemblerPredicate<(all_of FeatureGFX8Insts, (not FeatureGFX12Insts))>;2400 2401def isGFX7Plus :2402  Predicate<"Subtarget->getGeneration() >= AMDGPUSubtarget::SEA_ISLANDS">,2403  AssemblerPredicate<(all_of FeatureCIInsts)>;2404 2405def isGFX8Plus :2406  Predicate<"Subtarget->getGeneration() >= AMDGPUSubtarget::VOLCANIC_ISLANDS">,2407  AssemblerPredicate<(all_of FeatureGFX8Insts)>;2408 2409def isGFX8Only : Predicate<"Subtarget->getGeneration() =="2410                           "AMDGPUSubtarget::VOLCANIC_ISLANDS">,2411  AssemblerPredicate <(all_of FeatureVolcanicIslands)>;2412 2413def isGFX9Plus :2414  Predicate<"Subtarget->getGeneration() >= AMDGPUSubtarget::GFX9">,2415  AssemblerPredicate<(all_of FeatureGFX9Insts)>;2416 2417def isNotGFX9Plus :2418  Predicate<"Subtarget->getGeneration() < AMDGPUSubtarget::GFX9">;2419 2420def isGFX9Only : Predicate <2421  "Subtarget->getGeneration() == AMDGPUSubtarget::GFX9">,2422  AssemblerPredicate<(all_of FeatureGCN3Encoding, FeatureGFX9Insts)>;2423 2424def isGCN3ExcludingGFX90A :2425  Predicate<"Subtarget->isGCN3Encoding() && !Subtarget->hasGFX90AInsts()">,2426  AssemblerPredicate<(all_of FeatureGCN3Encoding, (not FeatureGFX90AInsts))>;2427 2428def isGFX90APlus :2429  Predicate<"Subtarget->hasGFX90AInsts()">,2430  AssemblerPredicate<(all_of FeatureGFX90AInsts)>;2431 2432def isNotGFX90APlus :2433  Predicate<"!Subtarget->hasGFX90AInsts()">,2434  AssemblerPredicate<(all_of (not FeatureGFX90AInsts))>;2435 2436def isGFX8GFX9NotGFX90A :2437  Predicate<"!Subtarget->hasGFX90AInsts() &&"2438            "(Subtarget->getGeneration() == AMDGPUSubtarget::VOLCANIC_ISLANDS ||"2439            " Subtarget->getGeneration() == AMDGPUSubtarget::GFX9)">,2440  AssemblerPredicate<(all_of FeatureGFX8Insts, FeatureGCN3Encoding, (not FeatureGFX90AInsts))>;2441 2442def isGFX90AOnly :2443  Predicate<"Subtarget->hasGFX90AInsts() && !Subtarget->hasGFX940Insts()">,2444  AssemblerPredicate<(all_of FeatureGFX90AInsts, (not FeatureGFX940Insts))>;2445 2446def isGFX908orGFX90A :2447  Predicate<"Subtarget->hasMAIInsts() && !Subtarget->hasGFX940Insts()">,2448  AssemblerPredicate<(all_of FeatureMAIInsts, (not FeatureGFX940Insts))>;2449 2450def isGFX940Plus :2451  Predicate<"Subtarget->hasGFX940Insts()">,2452  AssemblerPredicate<(all_of FeatureGFX940Insts)>;2453 2454def isNotGFX940Plus :2455  Predicate<"!Subtarget->hasGFX940Insts()">,2456  AssemblerPredicate<(all_of (not FeatureGFX940Insts))>;2457 2458def HasGFX950Insts :2459  Predicate<"Subtarget->hasGFX950Insts()">,2460  AssemblerPredicate<(all_of FeatureGFX950Insts)>;2461 2462def HasPermlane16Swap :2463  Predicate<"Subtarget->hasPermlane16Swap()">,2464  AssemblerPredicate<(all_of FeaturePermlane16Swap)>;2465 2466def HasPermlane32Swap :2467  Predicate<"Subtarget->hasPermlane32Swap()">,2468  AssemblerPredicate<(all_of FeaturePermlane32Swap)>;2469 2470def isGFX8GFX9NotGFX940 :2471  Predicate<"!Subtarget->hasGFX940Insts() &&"2472            "(Subtarget->getGeneration() == AMDGPUSubtarget::VOLCANIC_ISLANDS ||"2473            " Subtarget->getGeneration() == AMDGPUSubtarget::GFX9)">,2474  AssemblerPredicate<(all_of FeatureGFX8Insts, FeatureGCN3Encoding, (not FeatureGFX940Insts))>;2475 2476def isGFX8GFX9 :2477  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::VOLCANIC_ISLANDS ||"2478            "Subtarget->getGeneration() == AMDGPUSubtarget::GFX9">,2479  AssemblerPredicate<(all_of FeatureGFX8Insts, FeatureGCN3Encoding)>;2480 2481def isGFX10Only :2482  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::GFX10">,2483  AssemblerPredicate<(all_of FeatureGFX10Insts, (not FeatureGFX11Insts))>;2484 2485def isGFX10Plus :2486  Predicate<"Subtarget->getGeneration() >= AMDGPUSubtarget::GFX10">,2487  AssemblerPredicate<(all_of FeatureGFX10Insts)>;2488 2489def isGFX10GFX11 :2490  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::GFX10 ||"2491            "Subtarget->getGeneration() == AMDGPUSubtarget::GFX11">,2492  AssemblerPredicate<(all_of FeatureGFX10Insts, (not FeatureGFX12Insts))>;2493 2494def isGFX10Before1030 :2495  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::GFX10 &&"2496            "!Subtarget->hasGFX10_3Insts()">,2497  AssemblerPredicate<(all_of FeatureGFX10Insts,(not FeatureGFX10_3Insts))>;2498 2499def isGFX9GFX10 :2500  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::GFX9 ||"2501            "Subtarget->getGeneration() == AMDGPUSubtarget::GFX10">,2502  AssemblerPredicate<(all_of FeatureGFX9Insts, (not FeatureGFX11Insts))>;2503 2504def isGFX9GFX10GFX11 :2505  Predicate<"Subtarget->getGeneration() >= AMDGPUSubtarget::GFX9 &&"2506            "Subtarget->getGeneration() < AMDGPUSubtarget::GFX12">,2507  AssemblerPredicate<(all_of FeatureGFX9Insts, (not FeatureGFX12Insts))>;2508 2509def isGFX8GFX9GFX10 :2510  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::VOLCANIC_ISLANDS ||"2511            "Subtarget->getGeneration() == AMDGPUSubtarget::GFX9 ||"2512            "Subtarget->getGeneration() == AMDGPUSubtarget::GFX10">,2513  AssemblerPredicate<(all_of FeatureGFX8Insts, (not FeatureGFX11Insts))>;2514 2515def isGFX11Only :2516  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::GFX11">,2517  AssemblerPredicate<(all_of FeatureGFX11Insts, (not FeatureGFX12Insts))>;2518 2519def isGFX11Plus :2520  Predicate<"Subtarget->getGeneration() >= AMDGPUSubtarget::GFX11">,2521  AssemblerPredicate<(all_of FeatureGFX11Insts)>;2522 2523def isGFX12Only :2524  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::GFX12">,2525  AssemblerPredicate<(all_of FeatureGFX12Insts)>;2526 2527def isGFX12Not12_50 :2528  Predicate<"Subtarget->getGeneration() == AMDGPUSubtarget::GFX12 && !Subtarget->hasGFX1250Insts()">,2529  AssemblerPredicate<(all_of FeatureGFX12Insts, (not FeatureGFX1250Insts))>;2530 2531def isGFX12Plus :2532  Predicate<"Subtarget->getGeneration() >= AMDGPUSubtarget::GFX12">,2533  AssemblerPredicate<(all_of FeatureGFX12Insts)>;2534 2535def isGFX12PlusNot12_50 :2536  Predicate<"Subtarget->getGeneration() >= AMDGPUSubtarget::GFX12 && !Subtarget->hasGFX1250Insts()">,2537  AssemblerPredicate<(all_of FeatureGFX12Insts, (not FeatureGFX1250Insts))>;2538 2539def isGFX125xOnly :2540  Predicate<"Subtarget->hasGFX1250Insts()">,2541  AssemblerPredicate<(all_of FeatureGFX1250Insts)>;2542 2543def isGFX1250Plus :2544  Predicate<"Subtarget->hasGFX1250Insts()">,2545  AssemblerPredicate<(all_of FeatureGFX1250Insts)>;2546 2547def isNotGFX1250Plus :2548  Predicate<"!Subtarget->hasGFX1250Insts()">,2549  AssemblerPredicate<(all_of (not FeatureGFX1250Insts))>;2550 2551def isGFX940orGFX1250 :2552  Predicate<"Subtarget->hasGFX940Insts() || Subtarget->hasGFX1250Insts()">,2553  AssemblerPredicate<(any_of FeatureGFX940Insts, FeatureGFX1250Insts)>;2554 2555def HasIEEEMinimumMaximumInsts :2556  Predicate<"Subtarget->hasIEEEMinimumMaximumInsts()">,2557  AssemblerPredicate<(all_of FeatureIEEEMinimumMaximumInsts)>;2558 2559def HasMinimum3Maximum3F32 :2560  Predicate<"Subtarget->hasMinimum3Maximum3F32()">,2561  AssemblerPredicate<(all_of FeatureMinimum3Maximum3F32)>;2562 2563def HasMinimum3Maximum3F16 :2564  Predicate<"Subtarget->hasMinimum3Maximum3F16()">,2565  AssemblerPredicate<(all_of FeatureMinimum3Maximum3F16)>;2566 2567def HasMin3Max3PKF16 :2568  Predicate<"Subtarget->hasMin3Max3PKF16()">,2569  AssemblerPredicate<(all_of FeatureMin3Max3PKF16)>;2570 2571def HasMinimum3Maximum3PKF16 :2572  Predicate<"Subtarget->hasMinimum3Maximum3PKF16()">,2573  AssemblerPredicate<(all_of FeatureMinimum3Maximum3PKF16)>;2574 2575 2576def HasFlatAddressSpace : Predicate<"Subtarget->hasFlatAddressSpace()">,2577  AssemblerPredicate<(all_of FeatureFlatAddressSpace)>;2578 2579def HasFlatBufferGlobalAtomicFaddF64Inst :2580  Predicate<"Subtarget->hasFlatBufferGlobalAtomicFaddF64Inst()">,2581  AssemblerPredicate<(any_of FeatureFlatBufferGlobalAtomicFaddF64Inst)>;2582 2583def HasAtomicFMinFMaxF32GlobalInsts :2584  Predicate<"Subtarget->hasAtomicFMinFMaxF32GlobalInsts()">,2585  AssemblerPredicate<(any_of FeatureAtomicFMinFMaxF32GlobalInsts)>;2586 2587def HasAtomicFMinFMaxF64GlobalInsts :2588  Predicate<"Subtarget->hasAtomicFMinFMaxF64GlobalInsts()">,2589  AssemblerPredicate<(any_of FeatureAtomicFMinFMaxF64GlobalInsts)>;2590 2591def HasAtomicFMinFMaxF32FlatInsts :2592  Predicate<"Subtarget->hasAtomicFMinFMaxF32FlatInsts()">,2593  AssemblerPredicate<(any_of FeatureAtomicFMinFMaxF32FlatInsts)>;2594 2595def HasAtomicFMinFMaxF64FlatInsts :2596  Predicate<"Subtarget->hasAtomicFMinFMaxF64FlatInsts()">,2597  AssemblerPredicate<(any_of FeatureAtomicFMinFMaxF64FlatInsts)>;2598 2599def HasLdsAtomicAddF64 :2600  Predicate<"Subtarget->hasLdsAtomicAddF64()">,2601  AssemblerPredicate<(any_of FeatureGFX90AInsts, FeatureGFX1250Insts)>;2602 2603def HasFlatGlobalInsts : Predicate<"Subtarget->hasFlatGlobalInsts()">,2604  AssemblerPredicate<(all_of FeatureFlatGlobalInsts)>;2605def HasFlatScratchInsts : Predicate<"Subtarget->hasFlatScratchInsts()">,2606  AssemblerPredicate<(all_of FeatureFlatScratchInsts)>;2607def HasScalarFlatScratchInsts : Predicate<"Subtarget->hasScalarFlatScratchInsts()">,2608  AssemblerPredicate<(all_of FeatureScalarFlatScratchInsts)>;2609def HasD16LoadStore : Predicate<"Subtarget->hasD16LoadStore()">,2610  AssemblerPredicate<(all_of FeatureGFX9Insts)>;2611 2612def HasFlatScratchSTMode : Predicate<"Subtarget->hasFlatScratchSTMode()">,2613  AssemblerPredicate<(any_of FeatureGFX10_3Insts, FeatureGFX940Insts)>;2614def HasFlatScratchSVSMode : Predicate<"Subtarget->hasFlatScratchSVSMode()">,2615  AssemblerPredicate<(any_of FeatureGFX940Insts, FeatureGFX11Insts)>;2616 2617def HasFlatGVSMode : Predicate<"Subtarget->hasFlatGVSMode()">,2618  AssemblerPredicate<(all_of FeatureFlatGVSMode)>;2619 2620def HasGFX10_AEncoding : Predicate<"Subtarget->hasGFX10_AEncoding()">,2621  AssemblerPredicate<(all_of FeatureGFX10_AEncoding)>;2622 2623def HasGFX10_BEncoding : Predicate<"Subtarget->hasGFX10_BEncoding()">,2624  AssemblerPredicate<(all_of FeatureGFX10_BEncoding)>;2625 2626def HasUnpackedD16VMem : Predicate<"Subtarget->hasUnpackedD16VMem()">,2627  AssemblerPredicate<(all_of FeatureUnpackedD16VMem)>;2628def HasPackedD16VMem : Predicate<"!Subtarget->hasUnpackedD16VMem()">,2629  AssemblerPredicate<(all_of (not FeatureUnpackedD16VMem))>;2630 2631def HasRestrictedSOffset : Predicate<"Subtarget->hasRestrictedSOffset()">,2632  AssemblerPredicate<(all_of FeatureHasRestrictedSOffset)>;2633def HasUnrestrictedSOffset : Predicate<"!Subtarget->hasRestrictedSOffset()">,2634  AssemblerPredicate<(all_of (not FeatureHasRestrictedSOffset))>;2635 2636def D16PreservesUnusedBits :2637  Predicate<"Subtarget->d16PreservesUnusedBits()">,2638  AssemblerPredicate<(all_of FeatureGFX9Insts, (not FeatureSRAMECC))>;2639 2640def LDSRequiresM0Init : Predicate<"Subtarget->ldsRequiresM0Init()">;2641def NotLDSRequiresM0Init : Predicate<"!Subtarget->ldsRequiresM0Init()">;2642 2643def HasMTBUFInsts : Predicate<"Subtarget->hasMTBUFInsts()">,2644  AssemblerPredicate<(all_of (not FeatureGFX1250Insts))>;2645 2646def HasFormattedMUBUFInsts : Predicate<"Subtarget->hasFormattedMUBUFInsts()">,2647  AssemblerPredicate<(all_of (not FeatureGFX1250Insts))>;2648 2649def HasExportInsts : Predicate<"Subtarget->hasExportInsts()">,2650  AssemblerPredicate<(all_of (not FeatureGFX90AInsts), (not FeatureGFX1250Insts))>;2651 2652def HasVINTERPEncoding : Predicate<"Subtarget->hasVINTERPEncoding()">,2653  AssemblerPredicate<(all_of FeatureGFX11Insts, (not FeatureGFX1250Insts))>;2654 2655def HasDSAddTid : Predicate<"Subtarget->getGeneration() >= AMDGPUSubtarget::GFX9">,2656  AssemblerPredicate<(all_of FeatureGFX9Insts)>;2657 2658def HasLDSFPAtomicAddF32 : Predicate<"Subtarget->hasLDSFPAtomicAddF32()">,2659  AssemblerPredicate<(all_of FeatureGFX8Insts)>;2660 2661def HasAddNoCarryInsts : Predicate<"Subtarget->hasAddNoCarry()">,2662  AssemblerPredicate<(all_of FeatureAddNoCarryInsts)>;2663 2664def NotHasAddNoCarryInsts : Predicate<"!Subtarget->hasAddNoCarry()">;2665 2666def HasXNACKEnabled : Predicate<"Subtarget->isXNACKEnabled()">;2667 2668def Has16BitInsts : Predicate<"Subtarget->has16BitInsts()">,2669  AssemblerPredicate<(all_of Feature16BitInsts)>;2670 2671def HasTrue16BitInsts : Predicate<"Subtarget->hasTrue16BitInsts()">,2672  AssemblerPredicate<(all_of FeatureTrue16BitInsts)>;2673def NotHasTrue16BitInsts : True16PredicateClass<"!Subtarget->hasTrue16BitInsts()">,2674  AssemblerPredicate<(all_of (not FeatureTrue16BitInsts))>;2675 2676// Control use of True16 instructions. The real True16 instructions are2677// True16 instructions as they are defined in the ISA. Fake True162678// instructions have the same encoding as real ones but syntactically2679// only allow 32-bit registers in operands and use low halves thereof.2680def UseRealTrue16Insts : True16PredicateClass<"Subtarget->useRealTrue16Insts()">,2681  AssemblerPredicate<(all_of FeatureTrue16BitInsts, FeatureRealTrue16Insts)>;2682def NotUseRealTrue16Insts : True16PredicateClass<"!Subtarget->useRealTrue16Insts()">,2683  AssemblerPredicate<(not (all_of FeatureTrue16BitInsts, FeatureRealTrue16Insts))>;2684def UseFakeTrue16Insts : True16PredicateClass<"Subtarget->hasTrue16BitInsts() && "2685                                              "!Subtarget->useRealTrue16Insts()">,2686  AssemblerPredicate<(all_of FeatureTrue16BitInsts)>;2687  // FIXME When we default to RealTrue16 instead of Fake, change the line as follows.2688  // AssemblerPredicate<(all_of FeatureTrue16BitInsts, (not FeatureRealTrue16Insts))>;2689 2690def UseTrue16WithSramECC : True16PredicateClass<"Subtarget->useRealTrue16Insts() && "2691                                                "!Subtarget->d16PreservesUnusedBits()">;2692 2693def HasD16Writes32BitVgpr: Predicate<"Subtarget->hasD16Writes32BitVgpr()">,2694  AssemblerPredicate<(all_of FeatureTrue16BitInsts, FeatureRealTrue16Insts, FeatureD16Writes32BitVgpr)>;2695def NotHasD16Writes32BitVgpr: Predicate<"!Subtarget->hasD16Writes32BitVgpr()">,2696  AssemblerPredicate<(all_of FeatureTrue16BitInsts, FeatureRealTrue16Insts, (not FeatureD16Writes32BitVgpr))>;2697 2698def HasBF16TransInsts : Predicate<"Subtarget->hasBF16TransInsts()">,2699  AssemblerPredicate<(all_of FeatureBF16TransInsts)>;2700 2701def HasBF16ConversionInsts : Predicate<"Subtarget->hasBF16ConversionInsts()">,2702  AssemblerPredicate<(all_of FeatureBF16ConversionInsts)>;2703 2704def HasBF16PackedInsts : Predicate<"Subtarget->hasBF16PackedInsts()">,2705  AssemblerPredicate<(all_of FeatureBF16PackedInsts)>;2706 2707def HasVOP3PInsts : Predicate<"Subtarget->hasVOP3PInsts()">,2708  AssemblerPredicate<(all_of FeatureVOP3P)>;2709 2710def NotHasMed3_16 : Predicate<"!Subtarget->hasMed3_16()">;2711def HasMed3_16 : Predicate<"Subtarget->hasMed3_16()">;2712 2713def HasMinMaxDenormModes : Predicate<"Subtarget->supportsMinMaxDenormModes()">;2714def NotHasMinMaxDenormModes : Predicate<"!Subtarget->supportsMinMaxDenormModes()">;2715 2716def HasFminFmaxLegacy : Predicate<"Subtarget->hasFminFmaxLegacy()">;2717 2718def HasSDWA : Predicate<"Subtarget->hasSDWA()">;2719 2720def HasSDWA8 : Predicate<"Subtarget->hasSDWA()">,2721  AssemblerPredicate<(all_of (not FeatureGFX9Insts), FeatureSDWA)>;2722 2723def HasSDWA9 :2724  Predicate<"Subtarget->hasSDWA()">,2725  AssemblerPredicate<(all_of FeatureGCN3Encoding, FeatureGFX9Insts,FeatureSDWA)>;2726 2727def HasSDWA10 :2728  Predicate<"Subtarget->hasSDWA()">,2729  AssemblerPredicate<(all_of (not FeatureGCN3Encoding), FeatureGFX10Insts, FeatureSDWA)>;2730 2731def HasDPP : Predicate<"Subtarget->hasDPP()">,2732  AssemblerPredicate<(all_of FeatureGCN3Encoding, FeatureDPP)>;2733 2734def HasDPP8 : Predicate<"Subtarget->hasDPP8()">,2735  AssemblerPredicate<(all_of (not FeatureGCN3Encoding), FeatureGFX10Insts, FeatureDPP8)>;2736 2737def HasDPALU_DPP : Predicate<"Subtarget->hasDPALU_DPP()">,2738  AssemblerPredicate<(all_of FeatureDPALU_DPP)>;2739 2740def HasPackedFP32Ops : Predicate<"Subtarget->hasPackedFP32Ops()">,2741  AssemblerPredicate<(all_of FeaturePackedFP32Ops)>;2742 2743def HasPkMovB32 : Predicate<"Subtarget->hasPkMovB32()">,2744  AssemblerPredicate<(all_of FeatureGFX90AInsts)>;2745 2746def HasFmaakFmamkF32Insts :2747  Predicate<"Subtarget->hasFmaakFmamkF32Insts()">,2748  AssemblerPredicate<(any_of FeatureGFX10Insts, FeatureGFX940Insts)>;2749 2750def HasFmaakFmamkF64Insts :2751  Predicate<"Subtarget->hasFmaakFmamkF64Insts()">,2752  AssemblerPredicate<(any_of FeatureGFX1250Insts)>;2753 2754def HasAddMinMaxInsts :2755  Predicate<"Subtarget->hasAddMinMaxInsts()">,2756  AssemblerPredicate<(any_of FeatureAddMinMaxInsts)>;2757 2758def HasPkAddMinMaxInsts :2759  Predicate<"Subtarget->hasPkAddMinMaxInsts()">,2760  AssemblerPredicate<(any_of FeaturePkAddMinMaxInsts)>;2761 2762def HasPkMinMax3Insts :2763  Predicate<"Subtarget->hasPkMinMax3Insts()">,2764  AssemblerPredicate<(any_of FeatureGFX1250Insts)>;2765 2766def HasSGetShaderCyclesInst :2767  Predicate<"Subtarget->hasSGetShaderCyclesInst()">,2768  AssemblerPredicate<(any_of FeatureGFX1250Insts)>;2769 2770def HasImageInsts : Predicate<"Subtarget->hasImageInsts()">,2771  AssemblerPredicate<(all_of FeatureImageInsts)>;2772 2773def HasExtendedImageInsts : Predicate<"Subtarget->hasExtendedImageInsts()">,2774  AssemblerPredicate<(all_of FeatureExtendedImageInsts)>;2775 2776def HasR128A16 : Predicate<"Subtarget->hasR128A16()">,2777  AssemblerPredicate<(all_of FeatureR128A16)>;2778 2779def HasA16 : Predicate<"Subtarget->hasA16()">,2780  AssemblerPredicate<(all_of FeatureA16)>;2781 2782def HasG16 : Predicate<"Subtarget->hasG16()">,2783  AssemblerPredicate<(all_of FeatureG16)>;2784 2785def HasDPP16 : Predicate<"Subtarget->hasDPP()">,2786  AssemblerPredicate<(all_of (not FeatureGCN3Encoding), FeatureGFX10Insts, FeatureDPP)>;2787 2788def HasIntClamp : Predicate<"Subtarget->hasIntClamp()">,2789  AssemblerPredicate<(all_of FeatureIntClamp)>;2790 2791def HasMadMixInsts : Predicate<"Subtarget->hasMadMixInsts()">,2792  AssemblerPredicate<(all_of FeatureMadMixInsts)>;2793 2794def HasScalarStores : Predicate<"Subtarget->hasScalarStores()">,2795  AssemblerPredicate<(all_of FeatureScalarStores)>;2796 2797def HasScalarAtomics : Predicate<"Subtarget->hasScalarAtomics()">,2798  AssemblerPredicate<(all_of FeatureScalarAtomics)>;2799 2800def HasNoSdstCMPX : Predicate<"Subtarget->hasNoSdstCMPX()">,2801  AssemblerPredicate<(all_of FeatureNoSdstCMPX)>;2802 2803def HasSdstCMPX : Predicate<"!Subtarget->hasNoSdstCMPX()">,2804  AssemblerPredicate<(all_of (not FeatureNoSdstCMPX))>;2805 2806def has16BankLDS : Predicate<"Subtarget->getLDSBankCount() == 16">;2807def has32BankLDS : Predicate<"Subtarget->getLDSBankCount() == 32">;2808def HasVGPRIndexMode : Predicate<"Subtarget->hasVGPRIndexMode()">,2809                      AssemblerPredicate<(all_of FeatureVGPRIndexMode)>;2810def HasMovrel : Predicate<"Subtarget->hasMovrel()">,2811                AssemblerPredicate<(all_of FeatureMovrel)>;2812 2813def HasFmaMixInsts : Predicate<"Subtarget->hasFmaMixInsts()">,2814  AssemblerPredicate<(all_of FeatureFmaMixInsts)>;2815 2816def HasFmaMixBF16Insts : Predicate<"Subtarget->hasFmaMixBF16Insts()">,2817  AssemblerPredicate<(all_of FeatureFmaMixBF16Insts)>;2818 2819def HasDLInsts : Predicate<"Subtarget->hasDLInsts()">,2820  AssemblerPredicate<(all_of FeatureDLInsts)>;2821 2822def HasFmacF64Inst : Predicate<"Subtarget->hasFmacF64Inst()">,2823  AssemblerPredicate<(all_of FeatureFmacF64Inst)>;2824 2825def HasDot1Insts : Predicate<"Subtarget->hasDot1Insts()">,2826  AssemblerPredicate<(all_of FeatureDot1Insts)>;2827 2828def HasDot2Insts : Predicate<"Subtarget->hasDot2Insts()">,2829  AssemblerPredicate<(all_of FeatureDot2Insts)>;2830 2831def HasDot3Insts : Predicate<"Subtarget->hasDot3Insts()">,2832  AssemblerPredicate<(all_of FeatureDot3Insts)>;2833 2834def HasDot4Insts : Predicate<"Subtarget->hasDot4Insts()">,2835  AssemblerPredicate<(all_of FeatureDot4Insts)>;2836 2837def HasDot5Insts : Predicate<"Subtarget->hasDot5Insts()">,2838  AssemblerPredicate<(all_of FeatureDot5Insts)>;2839 2840def HasDot6Insts : Predicate<"Subtarget->hasDot6Insts()">,2841  AssemblerPredicate<(all_of FeatureDot6Insts)>;2842 2843def HasDot7Insts : Predicate<"Subtarget->hasDot7Insts()">,2844  AssemblerPredicate<(all_of FeatureDot7Insts)>;2845 2846def HasDot8Insts : Predicate<"Subtarget->hasDot8Insts()">,2847  AssemblerPredicate<(all_of FeatureDot8Insts)>;2848 2849def HasDot9Insts : Predicate<"Subtarget->hasDot9Insts()">,2850  AssemblerPredicate<(all_of FeatureDot9Insts)>;2851 2852def HasDot10Insts : Predicate<"Subtarget->hasDot10Insts()">,2853  AssemblerPredicate<(all_of FeatureDot10Insts)>;2854 2855def HasDot11Insts : Predicate<"Subtarget->hasDot11Insts()">,2856  AssemblerPredicate<(all_of FeatureDot11Insts)>;2857 2858def HasDot12Insts : Predicate<"Subtarget->hasDot12Insts()">,2859  AssemblerPredicate<(all_of FeatureDot12Insts)>;2860 2861def HasDot13Insts : Predicate<"Subtarget->hasDot13Insts()">,2862  AssemblerPredicate<(all_of FeatureDot13Insts)>;2863 2864def HasGetWaveIdInst : Predicate<"Subtarget->hasGetWaveIdInst()">,2865  AssemblerPredicate<(all_of FeatureGetWaveIdInst)>;2866 2867def HasMAIInsts : Predicate<"Subtarget->hasMAIInsts()">,2868  AssemblerPredicate<(all_of FeatureMAIInsts)>;2869 2870def NotHasMAIInsts : Predicate<"!Subtarget->hasMAIInsts()">,2871  AssemblerPredicate<(all_of (not FeatureMAIInsts))>;2872 2873def HasSMemRealTime : Predicate<"Subtarget->hasSMemRealTime()">,2874  AssemblerPredicate<(all_of FeatureSMemRealTime)>;2875 2876def HasSMemTimeInst : Predicate<"Subtarget->hasSMemTimeInst()">,2877  AssemblerPredicate<(all_of FeatureSMemTimeInst)>;2878 2879def HasShaderCyclesRegister : Predicate<"Subtarget->hasShaderCyclesRegister()">,2880  AssemblerPredicate<(all_of FeatureShaderCyclesRegister)>;2881 2882def HasShaderCyclesHiLoRegisters : Predicate<"Subtarget->hasShaderCyclesHiLoRegisters()">;2883 2884def HasFP8Insts : Predicate<"Subtarget->hasFP8Insts()">,2885  AssemblerPredicate<(all_of FeatureFP8Insts)>;2886 2887def HasFP8ConversionInsts : Predicate<"Subtarget->hasFP8ConversionInsts()">,2888  AssemblerPredicate<(all_of FeatureFP8ConversionInsts)>;2889 2890def HasCubeInsts : Predicate<"Subtarget->hasCubeInsts()">,2891  AssemblerPredicate<(all_of FeatureCubeInsts)>;2892 2893def HasLerpInst : Predicate<"Subtarget->hasLerpInst()">,2894  AssemblerPredicate<(all_of FeatureLerpInst)>;2895 2896def HasSadInsts : Predicate<"Subtarget->hasSadInsts()">,2897  AssemblerPredicate<(all_of FeatureSadInsts)>;2898 2899def HasQsadInsts : Predicate<"Subtarget->hasQsadInsts()">,2900  AssemblerPredicate<(all_of FeatureQsadInsts)>;2901 2902def HasCvtNormInsts : Predicate<"Subtarget->hasCvtNormInsts()">,2903  AssemblerPredicate<(all_of FeatureCvtNormInsts)>;2904 2905def HasCvtPkNormVOP2Insts : Predicate<"Subtarget->hasCvtPkNormVOP2Insts()">,2906  AssemblerPredicate<(all_of FeatureCvtPkNormVOP2Insts)>;2907 2908def HasCvtPkNormVOP3Insts : Predicate<"Subtarget->hasCvtPkNormVOP3Insts()">,2909  AssemblerPredicate<(all_of FeatureCvtPkNormVOP3Insts)>;2910 2911def HasFP8E5M3Insts : Predicate<"Subtarget->hasFP8E5M3Insts()">,2912  AssemblerPredicate<(all_of FeatureFP8E5M3Insts)>;2913 2914def NotHasFP8E5M3Insts : Predicate<"!Subtarget->hasFP8E5M3Insts()">,2915  AssemblerPredicate<(all_of (not FeatureFP8E5M3Insts))>;2916 2917def HasPkFmacF16Inst : Predicate<"Subtarget->hasPkFmacF16Inst()">,2918  AssemblerPredicate<(all_of FeaturePkFmacF16Inst)>;2919 2920def HasMadMacF32Insts : Predicate<"Subtarget->hasMadMacF32Insts()">,2921  AssemblerPredicate<(all_of FeatureMadMacF32Insts)>;2922 2923def HasFmaLegacy32 : Predicate<"Subtarget->hasGFX10_3Insts()">,2924  AssemblerPredicate<(any_of FeatureGFX10_3Insts)>;2925 2926def HasAtomicDsPkAdd16Insts : Predicate<"Subtarget->hasAtomicDsPkAdd16Insts()">,2927  AssemblerPredicate<(any_of FeatureAtomicDsPkAdd16Insts)>;2928 2929def HasAtomicFlatPkAdd16Insts : Predicate<"Subtarget->hasAtomicFlatPkAdd16Insts()">,2930  AssemblerPredicate<(any_of FeatureAtomicFlatPkAdd16Insts)>;2931 2932def HasAtomicFaddRtnInsts : Predicate<"Subtarget->hasAtomicFaddRtnInsts()">,2933  AssemblerPredicate<(all_of FeatureAtomicFaddRtnInsts)>;2934def HasAtomicFaddNoRtnInsts : Predicate<"Subtarget->hasAtomicFaddNoRtnInsts()">,2935  AssemblerPredicate<(all_of FeatureAtomicFaddNoRtnInsts)>;2936def HasAtomicBufferGlobalPkAddF16NoRtnInsts2937  : Predicate<"Subtarget->hasAtomicBufferGlobalPkAddF16NoRtnInsts() || Subtarget->hasAtomicBufferGlobalPkAddF16Insts()">,2938  AssemblerPredicate<(any_of FeatureAtomicBufferGlobalPkAddF16NoRtnInsts, FeatureAtomicBufferGlobalPkAddF16Insts)>;2939def HasAtomicBufferGlobalPkAddF16Insts2940  : Predicate<"Subtarget->hasAtomicBufferGlobalPkAddF16Insts()">,2941  AssemblerPredicate<(all_of FeatureAtomicBufferGlobalPkAddF16Insts)>;2942def HasAtomicGlobalPkAddBF16Inst2943  : Predicate<"Subtarget->hasAtomicGlobalPkAddBF16Inst()">,2944    AssemblerPredicate<(all_of FeatureAtomicGlobalPkAddBF16Inst)>;2945def HasAtomicBufferPkAddBF16Inst2946  : Predicate<"Subtarget->hasAtomicBufferPkAddBF16Inst()">,2947    AssemblerPredicate<(all_of FeatureAtomicBufferPkAddBF16Inst)>;2948def HasFlatAtomicFaddF32Inst2949  : Predicate<"Subtarget->hasFlatAtomicFaddF32Inst()">,2950  AssemblerPredicate<(all_of FeatureFlatAtomicFaddF32Inst)>;2951 2952def HasDefaultComponentZero2953  : Predicate<"Subtarget->hasDefaultComponentZero()">,2954  AssemblerPredicate<(all_of FeatureDefaultComponentZero)>;2955def HasDefaultComponentBroadcast2956  : Predicate<"Subtarget->hasDefaultComponentBroadcast()">,2957  AssemblerPredicate<(all_of FeatureDefaultComponentBroadcast)>;2958 2959def HasDsSrc2Insts : Predicate<"!Subtarget->hasDsSrc2Insts()">,2960  AssemblerPredicate<(all_of FeatureDsSrc2Insts)>;2961 2962def HasAddPC64Inst : Predicate<"Subtarget->hasAddPC64Inst()">,2963  AssemblerPredicate<(any_of FeatureGFX1250Insts)>;2964 2965def EnableFlatScratch : Predicate<"Subtarget->enableFlatScratch()">;2966 2967def DisableFlatScratch : Predicate<"!Subtarget->enableFlatScratch()">;2968 2969def HasUnalignedAccessMode : Predicate<"Subtarget->hasUnalignedAccessMode()">,2970  AssemblerPredicate<(all_of FeatureUnalignedAccessMode)>;2971 2972def HasMADIntraFwdBug : Predicate<"Subtarget->hasMADIntraFwdBug()">;2973 2974def HasNotMADIntraFwdBug : Predicate<"!Subtarget->hasMADIntraFwdBug()">;2975 2976def HasSALUFloatInsts : Predicate<"Subtarget->hasSALUFloatInsts()">,2977  AssemblerPredicate<(all_of FeatureSALUFloatInsts)>;2978 2979def NotHasSALUFloatInsts : Predicate<"!Subtarget->hasSALUFloatInsts()">,2980  AssemblerPredicate<(all_of (not FeatureSALUFloatInsts))>;2981 2982def HasPseudoScalarTrans : Predicate<"Subtarget->hasPseudoScalarTrans()">,2983  AssemblerPredicate<(all_of FeaturePseudoScalarTrans)>;2984 2985def HasBitOp3Insts : Predicate<"Subtarget->hasBitOp3Insts()">,2986  AssemblerPredicate<(all_of FeatureBitOp3Insts)>;2987 2988def HasTanhInsts : Predicate<"Subtarget->hasTanhInsts()">,2989  AssemblerPredicate<(all_of FeatureTanhInsts)>;2990 2991def HasTensorCvtLutInsts : Predicate<"Subtarget->hasTensorCvtLutInsts()">,2992  AssemblerPredicate<(all_of FeatureTensorCvtLutInsts)>;2993 2994def HasTransposeLoadF4F6Insts : Predicate<"Subtarget->hasTransposeLoadF4F6Insts()">,2995  AssemblerPredicate<(all_of FeatureTransposeLoadF4F6Insts)>;2996 2997def HasPrngInst : Predicate<"Subtarget->hasPrngInst()">,2998  AssemblerPredicate<(all_of FeaturePrngInst)>;2999 3000def HasBVHDualAndBVH8Insts : Predicate<"Subtarget->hasBVHDualAndBVH8Insts()">,3001  AssemblerPredicate<(all_of FeatureBVHDualAndBVH8Insts)>;3002 3003def Has64BitLiterals : Predicate<"Subtarget->has64BitLiterals()">,3004  AssemblerPredicate<(all_of Feature64BitLiterals)>;3005 3006def Has1024AddressableVGPRs : Predicate<"Subtarget->has1024AddressableVGPRs()">,3007  AssemblerPredicate<(all_of Feature1024AddressableVGPRs)>;3008 3009def HasWaitXcnt : Predicate<"Subtarget->hasWaitXcnt()">,3010  AssemblerPredicate<(all_of FeatureWaitXcnt)>;3011 3012def HasFP8ConversionScaleInsts : Predicate<"Subtarget->hasFP8ConversionScaleInsts()">,3013  AssemblerPredicate<(all_of FeatureFP8ConversionScaleInsts)>;3014 3015def HasBF8ConversionScaleInsts : Predicate<"Subtarget->hasBF8ConversionScaleInsts()">,3016  AssemblerPredicate<(all_of FeatureBF8ConversionScaleInsts)>;3017 3018def HasFP4ConversionScaleInsts : Predicate<"Subtarget->hasFP4ConversionScaleInsts()">,3019  AssemblerPredicate<(all_of FeatureFP4ConversionScaleInsts)>;3020 3021def HasFP6BF6ConversionScaleInsts : Predicate<"Subtarget->hasFP6BF6ConversionScaleInsts()">,3022  AssemblerPredicate<(all_of FeatureFP6BF6ConversionScaleInsts)>;3023 3024def HasF16BF16ToFP6BF6ConversionScaleInsts : Predicate<"Subtarget->hasF16BF16ToFP6BF6ConversionScaleInsts()">,3025  AssemblerPredicate<(all_of FeatureF16BF16ToFP6BF6ConversionScaleInsts)>;3026 3027def HasCvtPkF16F32Inst : Predicate<"Subtarget->hasCvtPkF16F32Inst()">,3028  AssemblerPredicate<(all_of FeatureCvtPkF16F32Inst)>;3029 3030def HasF32ToF16BF16ConversionSRInsts : Predicate<"Subtarget->hasF32ToF16BF16ConversionSRInsts()">,3031  AssemblerPredicate<(all_of FeatureF32ToF16BF16ConversionSRInsts)>;3032 3033def HasGDS : Predicate<"Subtarget->hasGDS()">;3034 3035def HasGWS : Predicate<"Subtarget->hasGWS()">;3036 3037def HasCvtFP8VOP1Bug : Predicate<"Subtarget->hasCvtFP8VOP1Bug()">;3038def HasNoCvtFP8VOP1Bug : Predicate<"!Subtarget->hasCvtFP8VOP1Bug()">;3039 3040def HasAtomicCSubNoRtnInsts : Predicate<"Subtarget->hasAtomicCSubNoRtnInsts()">;3041 3042def HasScalarDwordx3Loads : Predicate<"Subtarget->hasScalarDwordx3Loads()">;3043 3044def HasXF32Insts : Predicate<"Subtarget->hasXF32Insts()">,3045   AssemblerPredicate<(all_of FeatureXF32Insts)>;3046 3047def HasVmemPrefInsts : Predicate<"Subtarget->hasVmemPrefInsts()">,3048  AssemblerPredicate<(all_of FeatureVmemPrefInsts)>;3049 3050def HasAshrPkInsts : Predicate<"Subtarget->hasAshrPkInsts()">,3051  AssemblerPredicate<(all_of FeatureAshrPkInsts)>;3052 3053def HasLshlAddU64Inst : Predicate<"Subtarget->hasLshlAddU64Inst()">,3054                        AssemblerPredicate<(all_of FeatureLshlAddU64Inst)>;3055 3056def HasAddSubU64Insts : Predicate<"Subtarget->hasAddSubU64Insts()">,3057                        AssemblerPredicate<(all_of FeatureAddSubU64Insts)>;3058 3059def HasMadU32Inst : Predicate<"Subtarget->hasMadU32Inst()">,3060                    AssemblerPredicate<(all_of FeatureMadU32Inst)>;3061 3062def HasLdsBarrierArriveAtomic : Predicate<"Subtarget->hasLdsBarrierArriveAtomic()">,3063  AssemblerPredicate<(all_of FeatureLdsBarrierArriveAtomic)>;3064 3065def HasSetPrioIncWgInst : Predicate<"Subtarget->hasSetPrioIncWgInst()">,3066 AssemblerPredicate<(all_of FeatureSetPrioIncWgInst)>;3067 3068def NeedsAlignedVGPRs : Predicate<"Subtarget->needsAlignedVGPRs()">,3069                      AssemblerPredicate<(all_of FeatureRequiresAlignedVGPRs)>;3070 3071def NotNeedsAlignedVGPRs : Predicate<"!Subtarget->needsAlignedVGPRs()">,3072                      AssemblerPredicate<(all_of (not FeatureRequiresAlignedVGPRs))>;3073 3074def isWave32 : Predicate<"Subtarget->isWave32()">,3075  AssemblerPredicate <(any_of FeatureWavefrontSize32,3076                              FeatureAssemblerPermissiveWavesize)>;3077def isWave64 : Predicate<"Subtarget->isWave64()">,3078  AssemblerPredicate <(any_of FeatureWavefrontSize64,3079                              FeatureAssemblerPermissiveWavesize)>;3080 3081def isWave32Strict : Predicate<"Subtarget->isWave32()">,3082  AssemblerPredicate <(all_of FeatureWavefrontSize32)>;3083def isWave64Strict : Predicate<"Subtarget->isWave64()">,3084  AssemblerPredicate <(all_of FeatureWavefrontSize64)>;3085 3086//===----------------------------------------------------------------------===//3087// HwModes3088//===----------------------------------------------------------------------===//3089 3090defvar DefaultMode_Wave64 = DefaultMode;3091defvar DefaultMode_Wave32 = HwMode<[isWave32, NotNeedsAlignedVGPRs]>;3092 3093// gfx90a-gfx950. Has AGPRs, and also the align2 VGPR/AGPR requirement. Implied3094// wave64.3095def AVAlign2LoadStoreMode : HwMode<[HasMAIInsts, NeedsAlignedVGPRs]>;3096 3097// gfx1250, has alignment requirement but no AGPRs.3098def AlignedVGPRNoAGPRMode_Wave32 : HwMode<[NotHasMAIInsts, NeedsAlignedVGPRs, isWave32Strict]>;3099def AlignedVGPRNoAGPRMode_Wave64 : HwMode<[NotHasMAIInsts, NeedsAlignedVGPRs, isWave64Strict]>;3100 3101// FIXME: This should be able to only define a separate hwmode that3102// only depends on wavesize for just ValueTypes. These use different3103// HwMode namespaces. If we don't define the full set of modes used3104// for RegClassByHwMode, tablegen crashes for some reason3105def WaveSizeVT : ValueTypeByHwMode<[3106  DefaultMode_Wave64,3107  AVAlign2LoadStoreMode,3108  AlignedVGPRNoAGPRMode_Wave64,3109  DefaultMode_Wave32,3110  AlignedVGPRNoAGPRMode_Wave32], [i64, i64, i64, i32, i32]>;3111 3112 3113// Include AMDGPU TD files3114include "SISchedule.td"3115include "GCNProcessors.td"3116include "AMDGPUInstrInfo.td"3117include "SIRegisterInfo.td"3118include "AMDGPURegisterBanks.td"3119include "AMDGPUInstructions.td"3120include "SIInstrInfo.td"3121include "AMDGPUCallingConv.td"3122include "AMDGPUSearchableTables.td"3123