[llvm] 99a9b94 - AMDGPU: Move GCN subtarget features to GCNFeatures.td (#230996)
via llvm-commits
llvm-commits at lists.llvm.org
Sun Oct 11 11:09:02 PDT 2026
Author: Matt Arsenault
Date: 2026-10-11T20:08:54+02:00
New Revision: 99a9b94b3f4658f2cc4d8506f820bc9cdd8e5425
URL: https://github.com/llvm/llvm-project/commit/99a9b94b3f4658f2cc4d8506f820bc9cdd8e5425
DIFF: https://github.com/llvm/llvm-project/commit/99a9b94b3f4658f2cc4d8506f820bc9cdd8e5425.diff
LOG: AMDGPU: Move GCN subtarget features to GCNFeatures.td (#230996)
Move the GCN SubtargetFeature definitions, the FeatureISAVersion lists
and AMDGPUFrontendVisibleFeatures out of AMDGPU.td, so they can be used
without parsing the instruction, register and intrinsic definitions.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
---
<sub>Stack created with <a
href="https://github.com/github/gh-stack">GitHub Stacks CLI</a> • <a
href="https://gh.io/stacks-feedback">Give Feedback 💬</a></sub>
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
Added:
llvm/lib/Target/AMDGPU/GCNFeatures.td
Modified:
llvm/lib/Target/AMDGPU/AMDGPU.td
Removed:
################################################################################
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td
index 73464a89544cc..71927e471680f 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.td
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.td
@@ -21,2777 +21,7 @@ def p5 : PtrValueType<i32, 5>;
def p6 : PtrValueType<i32, 6>;
def p15 : PtrValueType<i32, 15>;
-//===-----------------------------------------------------------------------===//
-// AMDGPU Subtarget Feature (device properties)
-//===----------------------------------------------------------------------===//
-
-// Multiclass to define a SubtargetFeature along with optional predicates.
-// Parameters:
-// - FeatureString: The feature string used in the SubtargetFeature.
-// - Description: The description of the feature.
-// - GenPredicate: If 1 (default), generates a Has#NAME predicate.
-// - GenAssemblerPredicate: If 1 (default), the predicate includes AssemblerPredicate.
-// - Deps: List of dependent SubtargetFeatures (default empty).
-//
-// Usage:
-// defm MadMixInsts : AMDGPUSubtargetFeature<"mad-mix-insts", "description">;
-// This generates:
-// - FeatureMadMixInsts : SubtargetFeature<"mad-mix-insts", "HasMadMixInsts", "true", "description">
-// - HasMadMixInsts : Predicate<"Subtarget->hasMadMixInsts()">,
-// AssemblerPredicate<(any_of FeatureMadMixInsts)>
-//
-// With GenAssemblerPredicate=0:
-// defm Foo : AMDGPUSubtargetFeature<"foo", "desc", 1, 0>;
-// This generates:
-// - FeatureFoo : SubtargetFeature<...>
-// - HasFoo : Predicate<"Subtarget->hasFoo()"> (no AssemblerPredicate)
-//
-// With dependencies:
-// defm Bar : AMDGPUSubtargetFeature<"bar", "desc", 1, 1, [FeatureFoo]>;
-// This generates:
-// - FeatureBar : SubtargetFeature<"bar", "HasBar", "true", "desc", [FeatureFoo]>
-// - HasBar : Predicate + AssemblerPredicate
-multiclass AMDGPUSubtargetFeature<string FeatureString,
- string Description,
- bit GenPredicate = 1,
- bit GenAssemblerPredicate = 1,
- list<SubtargetFeature> Deps = [],
- InlineBehavior IB = InlineDefault> {
- def Feature#NAME : SubtargetFeature<FeatureString,
- "Has"#NAME,
- "true",
- Description,
- Deps,
- IB
- >;
-
- if GenPredicate then
- if GenAssemblerPredicate then
- def Has#NAME
- : Predicate<"Subtarget->has"#NAME#"()">,
- AssemblerPredicate<(any_of !cast<SubtargetFeature>("Feature"#NAME))>;
- else
- def Has#NAME : Predicate<"Subtarget->has"#NAME#"()">;
-}
-
-// Like AMDGPUSubtargetFeature, but a generic target only needs one covered GPU
-// to support the feature.
-multiclass AMDGPUGenericAnySubtargetFeature<string FeatureString,
- string Description,
- bit GenPredicate = 1,
- bit GenAssemblerPredicate = 1,
- list<SubtargetFeature> Deps = [],
- InlineBehavior IB = InlineDefault> {
- def Feature#NAME : AMDGPUGenericAnyFeature<FeatureString,
- "Has"#NAME,
- "true",
- Description,
- Deps,
- IB
- >;
-
- if GenPredicate then
- if GenAssemblerPredicate then
- def Has#NAME
- : Predicate<"Subtarget->has"#NAME#"()">,
- AssemblerPredicate<(any_of !cast<SubtargetFeature>("Feature"#NAME))>;
- else
- def Has#NAME : Predicate<"Subtarget->has"#NAME#"()">;
-}
-
-defm FastFMAF32 : AMDGPUSubtargetFeature<"fast-fmaf",
- "Assuming f32 fma is at least as fast as mul + add",
- /*GenPredicate=*/0, /*GenAssemblerPredicate=*/1, [], InlineIgnore
->;
-
-defm FastDenormalF32 : AMDGPUSubtargetFeature<"fast-denormal-f32",
- "Enabling denormals does not cause f32 instructions to run at f64 rates",
- /*GenPredicate=*/0
->;
-
-defm MIMG_R128 : AMDGPUSubtargetFeature<"mimg-r128",
- "Support 128-bit texture resources",
- /*GenPredicate=*/0
->;
-
-defm HalfRate64Ops : AMDGPUSubtargetFeature<"half-rate-64-ops",
- "Most fp64 instructions are half rate instead of quarter",
- /*GenPredicate=*/0, /*GenAssemblerPredicate=*/1, [], InlineIgnore
->;
-
-defm FullRate64Ops : AMDGPUSubtargetFeature<"full-rate-64-ops",
- "Most fp64 instructions are full rate",
- /*GenPredicate=*/0
->;
-
-defm FlatAddressSpace : AMDGPUSubtargetFeature<"flat-address-space",
- "Support flat address space"
->;
-
-defm FlatInstOffsets : AMDGPUSubtargetFeature<"flat-inst-offsets",
- "Flat instructions have immediate offset addressing mode"
->;
-
-defm FlatGlobalInsts : AMDGPUSubtargetFeature<"flat-global-insts",
- "Have global_* flat memory instructions",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/1,
- [FeatureFlatAddressSpace]
->;
-
-defm FlatScratchInsts : AMDGPUSubtargetFeature<"flat-scratch-insts",
- "Have scratch_* flat memory instructions",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/1,
- [FeatureFlatAddressSpace]
->;
-
-defm ScalarFlatScratchInsts : AMDGPUSubtargetFeature<"scalar-flat-scratch-insts",
- "Have s_scratch_* flat memory instructions"
->;
-
-def FeatureEnableFlatScratch : SubtargetFeature<"enable-flat-scratch",
- "EnableFlatScratch",
- "true",
- "Use scratch_* flat memory instructions to access scratch"
->;
-
-defm FlatGVSMode : AMDGPUSubtargetFeature<"flat-gvs-mode",
- "Have GVS addressing mode with flat_* instructions",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/1,
- [FeatureFlatAddressSpace]
->;
-
-defm AddNoCarryInsts : AMDGPUSubtargetFeature<"add-no-carry-insts",
- "Have VALU add/sub instructions without carry out"
->;
-
-defm UnalignedBufferAccess : AMDGPUSubtargetFeature<"unaligned-buffer-access",
- "Hardware supports unaligned global loads and stores"
->;
-
-defm TrapHandler: AMDGPUSubtargetFeature<"trap-handler",
- "Trap handler support",
- /*GenPredicate=*/0, /*GenAssemblerPredicate=*/1, [], InlineIgnore
->;
-
-defm UnalignedScratchAccess : AMDGPUSubtargetFeature<"unaligned-scratch-access",
- "Support unaligned scratch loads and stores",
- /*GenPredicate=*/1, /*GenAssemblerPredicate=*/1, [], InlineIgnore
->;
-
-defm UnalignedDSAccess : AMDGPUSubtargetFeature<"unaligned-ds-access",
- "Hardware supports unaligned local and region loads and stores"
->;
-
-defm DX10ClampAndIEEEMode : AMDGPUSubtargetFeature<"dx10-clamp-and-ieee-mode",
- "Target has DX10_CLAMP and IEEE_MODE kernel descriptor bits"
->;
-
-defm ApertureRegs : AMDGPUSubtargetFeature<"aperture-regs",
- "Has Memory Aperture Base and Size Registers",
- /*GenPredicate=*/0
->;
-
-defm GetDoorbellID : AMDGPUSubtargetFeature<"get-doorbell-id",
- "Supports reading the doorbell ID through S_GETREG",
- /*GenPredicate=*/0, /*GenAssemblerPredicate=*/false, [], InlineIgnore
->;
-
-defm AGPRAlloc : AMDGPUSubtargetFeature<"agpr-alloc",
- "Supports allocation of AGPRs",
- /*GenPredicate=*/0, /*GenAssemblerPredicate=*/false, [], InlineIgnore
->;
-
-defm MadMixInsts : AMDGPUSubtargetFeature<"mad-mix-insts",
- "Has v_mad_mix_f32, v_mad_mixlo_f16, v_mad_mixhi_f16 instructions"
->;
-
-defm FmaMixInsts : AMDGPUSubtargetFeature<"fma-mix-insts",
- "Has v_fma_mix_f32, v_fma_mixlo_f16, v_fma_mixhi_f16 instructions"
->;
-
-defm FmaMixBF16Insts : AMDGPUSubtargetFeature<"fma-mix-bf16-insts",
- "Has v_fma_mix_f32_bf16, v_fma_mixlo_bf16, v_fma_mixhi_bf16 instructions"
->;
-
-defm FmaLegacy32Insts : AMDGPUSubtargetFeature<"fma-legacy32-insts",
- "Has single-precision legacy FMA instructions"
->;
-
-defm IEEEMinimumMaximumInsts : AMDGPUSubtargetFeature<"ieee-minimum-maximum-insts",
- "Has v_minimum/maximum_f16/f32/f64, v_minimummaximum/maximumminimum_f16/f32,"
- "v_pk_minmax/maxmin_num_f16, v_min/max_num_f16/f32/f64 and"
- "v_pk_min/max_num_f16 instructions"
->;
-
-defm SALUMinimumMaximumInsts : AMDGPUSubtargetFeature<"salu-minimum-maximum-insts",
- "Has s_minimum/maximum_f16/f32 instructions"
->;
-
-defm Minimum3Maximum3F32 : AMDGPUSubtargetFeature<"minimum3-maximum3-f32",
- "Has v_minimum3_f32 and v_maximum3_f32 instructions"
->;
-
-defm Minimum3Maximum3F16 : AMDGPUSubtargetFeature<"minimum3-maximum3-f16",
- "Has v_minimum3_f16 and v_maximum3_f16 instructions"
->;
-
-defm Min3Max3PKF16 : AMDGPUSubtargetFeature<"min3-max3-pkf16",
- "Has v_pk_min3_num_f16 and v_pk_max3_num_f16 instructions"
->;
-
-defm Minimum3Maximum3PKF16 : AMDGPUSubtargetFeature<"minimum3-maximum3-pkf16",
- "Has v_pk_minimum3_f16 and v_pk_maximum3_f16 instructions"
->;
-
-def FeatureSupportsXNACK : SubtargetFeature<"xnack-support",
- "SupportsXNACK",
- "true",
- "Hardware supports XNACK",
- [], InlineIgnore
->;
-
-defm XNACKOnOffModes : AMDGPUSubtargetFeature<"xnack-on-off-modes",
- "Target supports XNACK on/off modes",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/0,
- [FeatureSupportsXNACK],
- InlineIgnore
->;
-
-// Whether the hardware supports WGP (work-group processor) execution
-// mode, as opposed to the per-kernel CU-mode selection
-// (FeatureCuMode).
-defm SupportsWGP : AMDGPUSubtargetFeature<"supports-wgp",
- "Hardware supports WGP mode",
- /*GenPredicate=*/false, /*GenAssemblerPredicate=*/false,
- [], InlineIgnore
->;
-
-def FeatureCuMode : SubtargetFeature<"cumode",
- "EnableCuMode",
- "true",
- "Enable CU wavefront execution mode"
->;
-
-def FeaturePreciseMemory
- : SubtargetFeature<"precise-memory", "EnablePreciseMemory",
- "true", "Enable precise memory mode">;
-
-defm SGPRInitBug : AMDGPUSubtargetFeature<"sgpr-init-bug",
- "VI SGPR initialization bug requiring a fixed SGPR allocation size",
- /*GenPredicate=*/1, /*GenAssemblerPredicate=*/1, [], InlineIgnore
->;
-
-defm UserSGPRInit16Bug : AMDGPUGenericAnySubtargetFeature<"user-sgpr-init16-bug",
- "Bug requiring at least 16 user+system SGPRs to be enabled",
- /*GenPredicate=*/0
->;
-
-defm LDSMisalignedBug : AMDGPUSubtargetFeature<"lds-misaligned-bug",
- "Some GFX10 bug with multi-dword LDS and flat access that is not naturally aligned in WGP mode",
- /*GenPredicate=*/0
->;
-
-// Set where a work-group can address only half of the physical LDS block:
-// gfx6 has 64 KiB physical and 32 KiB addressable, while gfx10/11/12 have
-// 128 KiB physical and 64 KiB addressable.
-defm HalfAddressablePhysicalLocalMemory : AMDGPUSubtargetFeature<"half-addressable-physical-local-memory",
- "A work-group can address only half of the physical LDS block",
- /*GenPredicate=*/0
->;
-
-defm MFMAInlineLiteralBug : AMDGPUSubtargetFeature<"mfma-inline-literal-bug",
- "MFMA cannot use inline literal as SrcC",
- /*GenPredicate=*/0
->;
-
-defm VcmpxPermlaneHazard : AMDGPUSubtargetFeature<"vcmpx-permlane-hazard",
- "TODO: describe me",
- /*GenPredicate=*/0
->;
-
-defm VMEMtoScalarWriteHazard : AMDGPUSubtargetFeature<"vmem-to-scalar-write-hazard",
- "VMEM instruction followed by scalar writing to EXEC mask, M0 or SGPR leads to incorrect execution.",
- /*GenPredicate=*/0
->;
-
-defm SMEMtoVectorWriteHazard : AMDGPUSubtargetFeature<"smem-to-vector-write-hazard",
- "s_load_dword followed by v_cmp page faults",
- /*GenPredicate=*/0
->;
-
-defm InstFwdPrefetchBug : AMDGPUSubtargetFeature<"inst-fwd-prefetch-bug",
- "S_INST_PREFETCH instruction causes shader to hang",
- /*GenPredicate=*/0
->;
-
-defm VmemPrefInsts : AMDGPUSubtargetFeature<"vmem-pref-insts",
- "Has flat_prefect_b8 and global_prefetch_b8 instructions"
->;
-
-defm SmemPrefetchInsts : AMDGPUSubtargetFeature<"smem-prefetch-insts",
- "Has s_prefetch_inst, s_prefetch_inst_pc_rel, s_prefetch_data and s_prefetch_data_pc_rel instructions"
->;
-
-defm SafeSmemPrefetch : AMDGPUSubtargetFeature<"safe-smem-prefetch",
- "SMEM prefetches do not fail on illegal address",
- /*GenPredicate=*/0
->;
-
-defm SafeCUPrefetch : AMDGPUSubtargetFeature<"safe-cu-prefetch",
- "VMEM CU scope prefetches do not fail on illegal address",
- /*GenPredicate=*/0
->;
-
-defm VcmpxExecWARHazard : AMDGPUSubtargetFeature<"vcmpx-exec-war-hazard",
- "V_CMPX WAR hazard on EXEC (V_CMPX issue ONLY)",
- /*GenPredicate=*/0
->;
-
-defm LdsBranchVmemWARHazard : AMDGPUSubtargetFeature<"lds-branch-vmem-war-hazard",
- "Switching between LDS and VMEM-tex not waiting VM_VSRC=0",
- /*GenPredicate=*/0
->;
-
-class FeatureMaxHardClauseLength<int size> : SubtargetFeature<
- "max-hard-clause-length-"#size,
- "MaxHardClauseLength",
- !cast<string>(size),
- "Maximum number of instructions in an explicit S_CLAUSE is "#size
->;
-
-/// Work around a hardware bug on some chips that can be triggered
-/// under certain circumstances when clauses are longer than 32 operations.
-def FeatureMaxHardClauseLength32 : FeatureMaxHardClauseLength<32>;
-/// While the S_CLAUSE instruction permits encoding clause lengths up to 64,
-/// hardware documentation for gfx10+ indicates that 63 is the maximum
-/// permitted clause length.
-def FeatureMaxHardClauseLength63 : FeatureMaxHardClauseLength<63>;
-
-defm NSAtoVMEMBug : AMDGPUSubtargetFeature<"nsa-to-vmem-bug",
- "MIMG-NSA followed by VMEM fail if EXEC_LO or EXEC_HI equals zero",
- /*GenPredicate=*/0
->;
-
-defm NSAClauseBug : AMDGPUSubtargetFeature<"nsa-clause-bug",
- "MIMG-NSA in a hard clause has unpredictable results on GFX10.1",
- /*GenPredicate=*/0
->;
-
-defm FlatSegmentOffsetBug : AMDGPUSubtargetFeature<"flat-segment-offset-bug",
- "GFX10 bug where inst_offset is ignored when flat instructions access global memory",
- /*GenPredicate=*/0
->;
-
-defm NegativeScratchOffsetBug : AMDGPUSubtargetFeature<"negative-scratch-offset-bug",
- "Negative immediate offsets in scratch instructions with an SGPR offset page fault on GFX9"
->;
-
-defm NegativeUnalignedScratchOffsetBug : AMDGPUSubtargetFeature<"negative-unaligned-scratch-offset-bug",
- "Scratch instructions with a VGPR offset and a negative immediate offset that"
- "is not a multiple of 4 read wrong memory on GFX10",
- /*GenPredicate=*/0
->;
-
-defm Offset3fBug : AMDGPUSubtargetFeature<"offset-3f-bug",
- "Branch offset of 3f hardware bug",
- /*GenPredicate=*/0
->;
-
-defm ImageStoreD16Bug : AMDGPUSubtargetFeature<"image-store-d16-bug",
- "Image Store D16 hardware bug",
- /*GenPredicate=*/0
->;
-
-defm ImageGather4D16Bug : AMDGPUSubtargetFeature<"image-gather4-d16-bug",
- "Image Gather4 D16 hardware bug",
- /*GenPredicate=*/0
->;
-
-defm MADIntraFwdBug : AMDGPUGenericAnySubtargetFeature<"mad-intra-fwd-bug",
- "MAD_U64/I64 intra instruction forwarding bug",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/0
->;
-
-defm MSAALoadDstSelBug : AMDGPUGenericAnySubtargetFeature<"msaa-load-dst-sel-bug",
- "MSAA loads not honoring dst_sel bug",
- /*GenPredicate=*/0
->;
-
-defm PrivEnabledTrap2NopBug : AMDGPUGenericAnySubtargetFeature<"priv-enabled-trap2-nop-bug",
- "Hardware that runs with PRIV=1 interpreting 's_trap 2' as a nop bug",
- /*GenPredicate=*/0
->;
-
-// A generic target may use an LDS bank count from one of its covered GPUs.
-class SubtargetFeatureLDSBankCount <int Value> : AMDGPUGenericAnyFeature <
- "ldsbankcount"#Value,
- "LDSBankCount",
- !cast<string>(Value),
- "The number of LDS banks per compute unit."
->;
-
-def FeatureLDSBankCount16 : SubtargetFeatureLDSBankCount<16>;
-def FeatureLDSBankCount32 : SubtargetFeatureLDSBankCount<32>;
-def FeatureLDSBankCount64 : SubtargetFeatureLDSBankCount<64>;
-
-class SubtargetFeatureMaxWavesPerEU <int Value> : SubtargetFeature <
- "max-waves-per-eu-"#Value,
- "MaxWavesPerEU",
- !cast<string>(Value),
- "Maximum number of waves an execution unit can hold."
->;
-
-def FeatureMaxWavesPerEU8 : SubtargetFeatureMaxWavesPerEU<8>;
-def FeatureMaxWavesPerEU10 : SubtargetFeatureMaxWavesPerEU<10>;
-def FeatureMaxWavesPerEU16 : SubtargetFeatureMaxWavesPerEU<16>;
-def FeatureMaxWavesPerEU20 : SubtargetFeatureMaxWavesPerEU<20>;
-
-class SubtargetFeatureInstCacheLineSize <int Value> : SubtargetFeature <
- "instcachelinesize"#Value,
- "InstCacheLineSize",
- !cast<string>(Value),
- "Instruction cache line size in bytes."
->;
-
-def FeatureInstCacheLineSize64 : SubtargetFeatureInstCacheLineSize<64>;
-def FeatureInstCacheLineSize128 : SubtargetFeatureInstCacheLineSize<128>;
-
-class SubtargetFeatureDataCacheLineSize <int Value> : SubtargetFeature <
- "datacachelinesize"#Value,
- "DataCacheLineSize",
- !cast<string>(Value),
- "Data (VMEM) cache line size in bytes."
->;
-
-def FeatureDataCacheLineSize128 : SubtargetFeatureDataCacheLineSize<128>;
-
-defm GCN3Encoding : AMDGPUSubtargetFeature<"gcn3-encoding",
- "Encoding format for VI",
- /*GenPredicate=*/0
->;
-
-defm CIInsts : AMDGPUSubtargetFeature<"ci-insts",
- "Additional instructions for CI+",
- /*GenPredicate=*/0
->;
-
-defm GFX8Insts : AMDGPUSubtargetFeature<"gfx8-insts",
- "Additional instructions for GFX8+",
- /*GenPredicate=*/0
->;
-
-defm GFX9Insts : AMDGPUSubtargetFeature<"gfx9-insts",
- "Additional instructions for GFX9+",
- /*GenPredicate=*/0
->;
-
-def FeatureRequiresAlignedVGPRs : SubtargetFeature<"vgpr-align2",
- "RequiresAlignVGPR",
- "true",
- "VGPR and AGPR tuple operands require even alignment"
->;
-
-defm GFX90AInsts : AMDGPUSubtargetFeature<"gfx90a-insts",
- "Additional instructions for GFX90A+",
- /*GenPredicate=*/0
->;
-
-defm TgSplitSupport : AMDGPUSubtargetFeature<"tgsplit-support",
- "Hardware support for threadgroup split execution mode"
->;
-
-defm GFX940Insts : AMDGPUSubtargetFeature<"gfx940-insts",
- "Additional instructions for GFX940+",
- /*GenPredicate=*/0
->;
-
-defm BufferInvInst : AMDGPUSubtargetFeature<"buffer-inv-inst",
- "Has buffer_inv instruction"
->;
-
-defm Permlane16Insts : AMDGPUSubtargetFeature<"permlane16-insts",
- "Has v_permlane16_b32/v_permlanex16_b32 instructions"
->;
-
-defm Permlane16Swap : AMDGPUSubtargetFeature<"permlane16-swap",
- "Has v_permlane16_swap_b32 instructions"
->;
-
-defm Permlane32Swap : AMDGPUSubtargetFeature<"permlane32-swap",
- "Has v_permlane32_swap_b32 instructions"
->;
-
-defm FP8ConversionScaleInsts : AMDGPUSubtargetFeature<"fp8-cvt-scale-insts",
- "Has fp8 conversion scale instructions"
->;
-
-defm BF8ConversionScaleInsts : AMDGPUSubtargetFeature<"bf8-cvt-scale-insts",
- "Has bf8 conversion scale instructions"
->;
-
-defm FP4ConversionScaleInsts : AMDGPUSubtargetFeature<"fp4-cvt-scale-insts",
- "Has fp4 conversion scale instructions"
->;
-
-defm FP6BF6ConversionScaleInsts : AMDGPUSubtargetFeature<"fp6bf6-cvt-scale-insts",
- "Has fp6 and bf6 conversion scale instructions"
->;
-
-defm F16BF16ToFP6BF6ConversionScaleInsts : AMDGPUSubtargetFeature<"f16bf16-to-fp6bf6-cvt-scale-insts",
- "Has f16bf16 to fp6bf6 conversion scale instructions"
->;
-
-defm F32ToFP6BF6ConversionScaleInsts : AMDGPUSubtargetFeature<"f32-to-fp6bf6-cvt-scale-insts",
- "Has f32 to fp6bf6 conversion scale instructions"
->;
-
-defm F32ToF16BF16ConversionSRInsts : AMDGPUSubtargetFeature<"f32-to-f16bf16-cvt-sr-insts",
- "Has f32 to f16bf16 conversion scale instructions"
->;
-
-defm FP6BF6ToF16BF16F32ConversionScaleInsts : AMDGPUSubtargetFeature<"fp6bf6-to-f16bf16f32-cvt-scale-insts",
- "Has fp6bf6 to f16bf16f32 conversion scale instructions"
->;
-
-defm Block16ConversionScaleInsts : AMDGPUSubtargetFeature<"block16-cvt-scale-insts",
- "Has block16 conversion scale instructions"
->;
-
-defm AshrPkInsts : AMDGPUSubtargetFeature<"ashr-pk-insts",
- "Has Arithmetic Shift Pack instructions"
->;
-
-defm ExclusiveScanInsts : AMDGPUSubtargetFeature<"exclusive-scan-insts",
- "Has exclusive scan instructions"
->;
-
-defm CvtPkF16F32Inst : AMDGPUSubtargetFeature<"cvt-pk-f16-f32-inst",
- "Has cvt_pk_f16_f32 instruction"
->;
-
-defm CvtSrPkBF16F32Inst : AMDGPUSubtargetFeature<"cvt-sr-pk-bf16-f32-inst",
- "Has cvt_sr_pk_bf16_f32 instruction"
->;
-
-defm McastLoadInsts : AMDGPUSubtargetFeature<"mcast-load-insts",
- "Has multicast load instructions"
->;
-
-defm AsyncLoadToLDSInsts : AMDGPUSubtargetFeature<"async-load-to-lds-insts",
- "Has global async load to LDS instructions"
->;
-
-defm AsyncStoreFromLDSInsts : AMDGPUSubtargetFeature<"async-store-from-lds-insts",
- "Has global async store from LDS instructions"
->;
-
-defm SWakeupImm : AMDGPUSubtargetFeature<"s-wakeup-imm",
- "s_wakeup takes an immediate operand"
->;
-
-defm SBarrierLeaveImm : AMDGPUSubtargetFeature<"s-barrier-leave-imm",
- "s_barrier_leave takes an immediate operand"
->;
-
-defm MTBUFInsts : AMDGPUSubtargetFeature<"mtbuf-insts",
- "Has memory typed buffer instructions."
->;
-
-defm FormattedMUBUFInsts : AMDGPUSubtargetFeature<"formatted-mubuf-insts",
- "Has formatted memory untyped buffer instructions."
->;
-
-defm VMovB64Inst : AMDGPUSubtargetFeature<"v-mov-b64-inst",
- "Has v_mov_b64 instruction"
->;
-
-defm VMulU64Inst : AMDGPUSubtargetFeature<"v-mul-u64-inst",
- "Has v_mul_u64 instruction."
->;
-
-defm GFX950Insts : AMDGPUSubtargetFeature<"gfx950-insts",
- "Additional instructions for GFX950+",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/1,
- [FeaturePermlane16Swap,
- FeaturePermlane32Swap,
- FeatureAshrPkInsts,
- FeatureFP8ConversionScaleInsts,
- FeatureBF8ConversionScaleInsts,
- FeatureFP4ConversionScaleInsts,
- FeatureFP6BF6ConversionScaleInsts,
- FeatureF16BF16ToFP6BF6ConversionScaleInsts,
- FeatureF32ToF16BF16ConversionSRInsts,
- FeatureCvtPkF16F32Inst,
- FeatureMinimum3Maximum3F32,
- FeatureMinimum3Maximum3PKF16,
- ]
->;
-
-defm GFX10Insts : AMDGPUSubtargetFeature<"gfx10-insts",
- "Additional instructions for GFX10+",
- /*GenPredicate=*/0
->;
-
-defm GFX11Insts : AMDGPUSubtargetFeature<"gfx11-insts",
- "Additional instructions for GFX11+",
- /*GenPredicate=*/0
->;
-
-defm GFX12Insts : AMDGPUSubtargetFeature<"gfx12-insts",
- "Additional instructions for GFX12+",
- /*GenPredicate=*/0
->;
-
-defm GFX1250Insts : AMDGPUSubtargetFeature<"gfx1250-insts",
- "Additional instructions for GFX1250+",
- /*GenPredicate=*/0
->;
-
-defm GFX13Insts : AMDGPUSubtargetFeature<"gfx13-insts",
- "Additional instructions for GFX13+",
- /*GenPredicate=*/0,
- /*GenAssemblerPredicate=*/0,
- [FeatureSWakeupImm,
- FeatureSBarrierLeaveImm,
- ]
->;
-
-defm GFX10_3Insts : AMDGPUSubtargetFeature<"gfx10-3-insts",
- "Additional instructions for GFX10.3",
- /*GenPredicate=*/0,
- /*GenAssemblerPredicate=*/1,
- [FeatureFmaLegacy32Insts]
->;
-
-defm GFX11_7Insts : AMDGPUSubtargetFeature<"gfx11-7-insts",
- "Additional instructions for GFX11.7",
- /*GenPredicate=*/0
->;
-
-defm GFX7GFX8GFX9Insts : AMDGPUSubtargetFeature<"gfx7-gfx8-gfx9-insts",
- "Instructions shared in GFX7, GFX8, GFX9",
- /*GenPredicate=*/0
->;
-
-defm SMemRealTime : AMDGPUSubtargetFeature<"s-memrealtime",
- "Has s_memrealtime instruction"
->;
-
-defm Inv2PiInlineImm : AMDGPUSubtargetFeature<"inv-2pi-inline-imm",
- "Has 1 / (2 * pi) as inline immediate",
- /*GenPredicate=*/0
->;
-
-defm 16BitInsts : AMDGPUSubtargetFeature<"16-bit-insts",
- "Has i16/f16 instructions"
->;
-
-defm True16BitInsts : AMDGPUSubtargetFeature<"true16",
- "True 16-bit operand instructions"
->;
-
-def FeatureRealTrue16Insts : SubtargetFeature<"real-true16",
- "EnableRealTrue16Insts",
- "true",
- "Use true 16-bit registers"
->;
-
-defm D16Writes32BitVgpr : AMDGPUSubtargetFeature<"d16-write-vgpr32",
- "D16 instructions potentially have 32-bit data dependencies"
->;
-
-defm BF16TransInsts : AMDGPUSubtargetFeature<"bf16-trans-insts",
- "Has bf16 transcendental instructions"
->;
-
-defm BF16ConversionInsts : AMDGPUSubtargetFeature<"bf16-cvt-insts",
- "Has bf16 conversion instructions"
->;
-
-defm BF16PackedInsts : AMDGPUSubtargetFeature<"bf16-pk-insts",
- "Has bf16 packed instructions (fma, add, mul, max, min)"
->;
-
-defm BF16InlineConstFromUpperFP32 : AMDGPUSubtargetFeature<"bf16-from-upper-fp32",
- "bf16 inline constants are actually selected from upper 16-bits of a fp32"
->;
-
-defm VOP3PInsts : AMDGPUSubtargetFeature<"vop3p",
- "Has VOP3P packed instructions"
->;
-
-defm Movrel : AMDGPUSubtargetFeature<"movrel",
- "Has v_movrel*_b32 instructions"
->;
-
-defm VGPRIndexMode : AMDGPUSubtargetFeature<"vgpr-index-mode",
- "Has VGPR mode register indexing"
->;
-
-defm ScalarDwordx3Loads : AMDGPUSubtargetFeature<"scalar-dwordx3-loads",
- "Has 96-bit scalar load instructions"
->;
-
-defm ScalarStores : AMDGPUSubtargetFeature<"scalar-stores",
- "Has store scalar memory instructions"
->;
-
-defm ScalarAtomics : AMDGPUSubtargetFeature<"scalar-atomics",
- "Has atomic scalar memory instructions"
->;
-
-defm SDWA : AMDGPUSubtargetFeature<"sdwa",
- "Support SDWA (Sub-DWORD Addressing) extension",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/0
->;
-
-defm SDWAOmod : AMDGPUSubtargetFeature<"sdwa-omod",
- "Support OMod with SDWA (Sub-DWORD Addressing) extension",
- /*GenPredicate=*/0
->;
-
-defm SDWAScalar : AMDGPUSubtargetFeature<"sdwa-scalar",
- "Support scalar register with SDWA (Sub-DWORD Addressing) extension",
- /*GenPredicate=*/0
->;
-
-defm SDWASdst : AMDGPUSubtargetFeature<"sdwa-sdst",
- "Support scalar dst for VOPC with SDWA (Sub-DWORD Addressing) extension",
- /*GenPredicate=*/0
->;
-
-defm SDWAMac : AMDGPUSubtargetFeature<"sdwa-mav",
- "Support v_mac_f32/f16 with SDWA (Sub-DWORD Addressing) extension",
- /*GenPredicate=*/0
->;
-
-defm SDWAOutModsVOPC : AMDGPUSubtargetFeature<"sdwa-out-mods-vopc",
- "Support clamp for VOPC with SDWA (Sub-DWORD Addressing) extension",
- /*GenPredicate=*/0
->;
-
-defm DPP : AMDGPUSubtargetFeature<"dpp",
- "Support DPP (Data Parallel Primitives) extension",
- /*GenPredicate=*/0
->;
-
-// DPP8 allows arbitrary cross-lane swizzling within groups of 8 lanes.
-def FeatureDPP8 : SubtargetFeature<"dpp8",
- "HasDPP8",
- "true",
- "Support DPP8 (Data Parallel Primitives) extension"
->;
-
-defm DPALU_DPP : AMDGPUSubtargetFeature<"dpp-64bit",
- "Support DPP (Data Parallel Primitives) extension in DP ALU"
->;
-
-defm DPPSrc1SGPR : AMDGPUSubtargetFeature<"dpp-src1-sgpr",
- "Support SGPR for Src1 of DPP instructions",
- /*GenPredicate=*/0
->;
-
-defm DPPWavefrontShifts : AMDGPUSubtargetFeature<"dpp-wavefront-shifts",
- "Support DPP wave_shl/wave_rol/wave_shr/wave_ror"
->;
-
-defm DPPBroadcasts : AMDGPUSubtargetFeature<"dpp-row-bcast",
- "Support DPP row_bcast15/row_bcast31"
->;
-
-defm PackedFP32Ops : AMDGPUSubtargetFeature<"packed-fp32-ops",
- "Support packed fp32 instructions"
->;
-
-defm PackedFP32SingleSGPROps : AMDGPUSubtargetFeature<"packed-fp32-single-sgpr-ops",
- "Support packed fp32 instructions with SGPR operands propagating low to high 32 bit"
->;
-
-defm PackedFP64SingleSGPROps : AMDGPUSubtargetFeature<"packed-fp64-single-sgpr-ops",
- "Support packed fp64 instructions with SGPR operands propagating low to high 64 bit"
->;
-
-defm PackedU64SingleSGPROps : AMDGPUSubtargetFeature<"packed-u64-single-sgpr-ops",
- "Support packed uint64 instructions with SGPR operands propagating low to high 64 bit"
->;
-
-defm R128A16 : AMDGPUSubtargetFeature<"r128-a16",
- "Support gfx9-style A16 for 16-bit coordinates/gradients/lod/clamp/mip image "
- "operands, where a16 is aliased with r128"
->;
-
-defm A16 : AMDGPUSubtargetFeature<"a16",
- "Support A16 for 16-bit coordinates/gradients/lod/clamp/mip image operands"
->;
-
-defm G16 : AMDGPUSubtargetFeature<"g16",
- "Support G16 for 16-bit gradient image operands"
->;
-
-defm NSAEncoding : AMDGPUSubtargetFeature<"nsa-encoding",
- "Support NSA encoding for image instructions",
- /*GenPredicate=*/0
->;
-
-defm PartialNSAEncoding : AMDGPUSubtargetFeature<"partial-nsa-encoding",
- "Support partial NSA encoding for image instructions",
- /*GenPredicate=*/0
->;
-
-defm ImageInsts : AMDGPUSubtargetFeature<"image-insts",
- "Support image instructions"
->;
-
-defm ExtendedImageInsts : AMDGPUSubtargetFeature<"extended-image-insts",
- "Support mips != 0, lod != 0, gather4, and get_lod"
->;
-
-defm BVHRayTracingInsts : AMDGPUSubtargetFeature<"bvh-ray-tracing-insts",
- "Has BVH ray tracing instructions"
->;
-
-defm MSAALoadInsts : AMDGPUSubtargetFeature<"msaa-load-insts",
- "Has image_msaa_load instructions",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/1,
- /*Deps=*/[FeatureImageInsts]
->;
-
-defm GFX10_BEncoding : AMDGPUSubtargetFeature<"gfx10_b-encoding",
- "Encoding format GFX10_B",
- /*GenPredicate=*/0
->;
-
-defm IntClamp : AMDGPUSubtargetFeature<"int-clamp-insts",
- "Support clamp for integer destination"
->;
-
-defm UnpackedD16VMem : AMDGPUSubtargetFeature<"unpacked-d16-vmem",
- "Has unpacked d16 vmem instructions"
->;
-
-defm DLInsts : AMDGPUSubtargetFeature<"dl-insts",
- "Has v_fmac_f32 and v_xnor_b32 instructions"
->;
-
-defm FmacF64Inst : AMDGPUSubtargetFeature<"fmacf64-inst",
- "Has v_fmac_f64 instruction"
->;
-
-defm Dot1Insts : AMDGPUSubtargetFeature<"dot1-insts",
- "Has v_dot4_i32_i8 and v_dot8_i32_i4 instructions"
->;
-
-defm Dot2Insts : AMDGPUSubtargetFeature<"dot2-insts",
- "Has v_dot2_i32_i16, v_dot2_u32_u16 instructions"
->;
-
-defm Dot3Insts : AMDGPUSubtargetFeature<"dot3-insts",
- "Has v_dot8c_i32_i4 instruction"
->;
-
-defm Dot4Insts : AMDGPUSubtargetFeature<"dot4-insts",
- "Has v_dot2c_i32_i16 instruction"
->;
-
-defm Dot5Insts : AMDGPUSubtargetFeature<"dot5-insts",
- "Has v_dot2c_f32_f16 instruction"
->;
-
-defm Dot6Insts : AMDGPUSubtargetFeature<"dot6-insts",
- "Has v_dot4c_i32_i8 instruction"
->;
-
-defm Dot7Insts : AMDGPUSubtargetFeature<"dot7-insts",
- "Has v_dot4_u32_u8, v_dot8_u32_u4 instructions"
->;
-
-defm Dot8Insts : AMDGPUSubtargetFeature<"dot8-insts",
- "Has v_dot4_i32_iu8, v_dot8_i32_iu4 instructions"
->;
-
-defm Dot9Insts : AMDGPUSubtargetFeature<"dot9-insts",
- "Has v_dot2_f16_f16, v_dot2_bf16_bf16 instructions"
->;
-
-defm Dot10Insts : AMDGPUSubtargetFeature<"dot10-insts",
- "Has v_dot2_f32_f16 instruction"
->;
-
-defm Dot11Insts : AMDGPUSubtargetFeature<"dot11-insts",
- "Has v_dot4_f32_fp8_fp8, v_dot4_f32_fp8_bf8, v_dot4_f32_bf8_fp8, v_dot4_f32_bf8_bf8 instructions"
->;
-
-defm Dot12Insts : AMDGPUSubtargetFeature<"dot12-insts",
- "Has v_dot2_f32_bf16 instructions"
->;
-
-defm Dot13Insts : AMDGPUSubtargetFeature<"dot13-insts",
- "Has v_dot2c_f32_bf16 instructions"
->;
-
-defm MAIInsts : AMDGPUSubtargetFeature<"mai-insts",
- "Has mAI instructions"
->;
-
-defm FP8Insts : AMDGPUSubtargetFeature<"fp8-insts",
- "Has fp8 and bf8 instructions"
->;
-
-defm FP8ConversionInsts : AMDGPUSubtargetFeature<"fp8-conversion-insts",
- "Has fp8 and bf8 conversion instructions"
->;
-
-// OCP fp8 semantics (E4M3FN, E5M2). Distinct from the FNUZ formats used by
-// gfx942.
-defm OCPFP8ConversionInsts : AMDGPUSubtargetFeature<"ocp-fp8-conversion-insts",
- "Has OCP fp8 (E4M3FN, E5M2) conversion instructions"
->;
-
-defm FP8F16ConversionInsts : AMDGPUSubtargetFeature<"fp8-f16-conversion-insts",
- "Has fp8 and bf8 conversion to f16 instructions"
->;
-
-// The opposite direction of FP8F16ConversionInsts above.
-defm F16FP8ConversionInsts : AMDGPUSubtargetFeature<"f16-fp8-conversion-insts",
- "Has f16 conversion to fp8 and bf8 instructions"
->;
-
-defm FP8E5M3Insts : AMDGPUSubtargetFeature<"fp8e5m3-insts",
- "Has fp8 e5m3 format support"
->;
-
-defm CvtFP8VOP1Bug : AMDGPUSubtargetFeature<"cvt-fp8-vop1-bug",
- "FP8/BF8 VOP1 form of conversion to F32 is unreliable",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/0,
- [FeatureFP8ConversionInsts]
->;
-
-// The fp8/bf8 conversions take the byte index in one of two ways, and a
-// subtarget uses the same one in both directions.
-defm CvtFP8SDWASrcSel : AMDGPUSubtargetFeature<"cvt-fp8-sdwa-src-sel",
- "FP8/BF8 conversions take the byte index from the SDWA src0_sel or op_sel field",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/1,
- [FeatureFP8ConversionInsts]
->;
-
-defm CvtFP8ByteSel : AMDGPUSubtargetFeature<"cvt-fp8-byte-sel",
- "FP8/BF8 conversions take the byte index from a byte_sel operand",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/1,
- [FeatureFP8ConversionInsts]
->;
-
-defm WMMA256bInsts : AMDGPUSubtargetFeature<"wmma-256b-insts",
- "Has WMMA instructions where A and B matrices have duplicated data"
->;
-
-defm WMMA128bInsts : AMDGPUSubtargetFeature<"wmma-128b-insts",
- "Has WMMA instructions where A and B matrices do not have duplicated data"
->;
-
-defm WMMAN16Insts : AMDGPUSubtargetFeature<"wmma-n16-insts",
- "Has WMMA instructions with N = 16"
->;
-
-defm WMMAF4Insts : AMDGPUSubtargetFeature<"wmma-f4-insts",
- "Has WMMA_*_F4 instructions"
->;
-
-defm WMMA16x16x128FP8Insts : AMDGPUSubtargetFeature<"wmma-16x16x128fp8-insts",
- "Has 16x16x128 FP8/BF8 WMMA instructions"
->;
-
-defm SWMMACGfx1200Insts : AMDGPUSubtargetFeature<"swmmac-gfx1200-insts",
- "Has GFX1200 SWMMAC instructions"
->;
-
-defm SWMMACGfx1250Insts : AMDGPUSubtargetFeature<"swmmac-gfx1250-insts",
- "Has GFX1250 SWMMAC instructions"
->;
-
-defm PkFmacF16Inst : AMDGPUSubtargetFeature<"pk-fmac-f16-inst",
- "Has v_pk_fmac_f16 instruction"
->;
-
-defm CubeInsts : AMDGPUSubtargetFeature<"cube-insts",
- "Has v_cube* instructions"
->;
-
-defm IndexedResources : AMDGPUSubtargetFeature<"indexed-resources",
- "Has indexed resources"
->;
-
-defm LerpInst : AMDGPUSubtargetFeature<"lerp-inst",
- "Has v_lerp_u8 instruction"
->;
-
-defm SadInsts : AMDGPUSubtargetFeature<"sad-insts",
- "Has v_sad* instructions"
->;
-
-defm QsadInsts : AMDGPUSubtargetFeature<"qsad-insts",
- "Has v_qsad* instructions"
->;
-
-defm MsadInsts : AMDGPUSubtargetFeature<"msad-insts",
- "Has v_msad_u8 instruction"
->;
-
-defm MqsadPkInsts : AMDGPUSubtargetFeature<"mqsad-pk-insts",
- "Has v_mqsad_pk_u16_u8 instruction"
->;
-
-defm MqsadInsts : AMDGPUSubtargetFeature<"mqsad-insts",
- "Has v_mqsad_u32_u8 instruction"
->;
-
-defm CvtNormInsts : AMDGPUSubtargetFeature<"cvt-norm-insts",
- "Has v_cvt_norm* instructions"
->;
-
-defm CvtPkNormVOP2Insts : AMDGPUSubtargetFeature<"cvt-pknorm-vop2-insts",
- "Has v_cvt_pk_norm_*f32 instructions/Has v_cvt_pk_norm_*_f16 instructions"
->;
-
-defm CvtPkNormVOP3Insts : AMDGPUSubtargetFeature<"cvt-pknorm-vop3-insts",
- "Has v_cvt_pk_norm_*f32 instructions/Has v_cvt_pk_norm_*_f16 instructions"
->;
-
-defm AtomicDsPkAdd16Insts : AMDGPUSubtargetFeature<"atomic-ds-pk-add-16-insts",
- "Has ds_pk_add_bf16, ds_pk_add_f16, ds_pk_add_rtn_bf16, "
- "ds_pk_add_rtn_f16 instructions"
->;
-
-defm AtomicDsCondSubClampInsts : AMDGPUSubtargetFeature<"atomic-ds-cond-sub-clamp-insts",
- "Has ds_cond_sub_u32, ds_cond_sub_rtn_u32, ds_sub_clamp_u32 and "
- "ds_sub_clamp_rtn_u32 instructions"
->;
-
-defm LdsAtomicAddF64 : AMDGPUSubtargetFeature<"lds-atomic-add-f64",
- "Has ds_add_f64 and ds_add_rtn_f64 instructions"
->;
-
-defm AtomicFlatPkAdd16Insts : AMDGPUSubtargetFeature<"atomic-flat-pk-add-16-insts",
- "Has flat_atomic_pk_add_f16 and flat_atomic_pk_add_bf16 instructions"
->;
-
-defm AtomicCondSubClampFlatInsts : AMDGPUSubtargetFeature<"atomic-cond-sub-clamp-flat-insts",
- "Has flat_atomic_cond_sub_u32 and flat_atomic_sub_clamp_u32 instructions"
->;
-
-defm AtomicFaddRtnInsts : AMDGPUSubtargetFeature<"atomic-fadd-rtn-insts",
- "Has buffer_atomic_add_f32 and global_atomic_add_f32 instructions that "
- "return original value",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/1,
- [FeatureFlatGlobalInsts]
->;
-
-defm AtomicFMinFMaxF32GlobalInsts : AMDGPUSubtargetFeature<"atomic-fmin-fmax-global-f32",
- "Has global/buffer instructions for atomicrmw fmin/fmax for float"
->;
-
-defm AtomicFMinFMaxF64GlobalInsts : AMDGPUSubtargetFeature<"atomic-fmin-fmax-global-f64",
- "Has global/buffer instructions for atomicrmw fmin/fmax for float"
->;
-
-defm AtomicFMinFMaxF32FlatInsts : AMDGPUSubtargetFeature<"atomic-fmin-fmax-flat-f32",
- "Has flat memory instructions for atomicrmw fmin/fmax for float",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/1,
- [FeatureFlatAddressSpace]
->;
-
-defm AtomicFMinFMaxF64FlatInsts : AMDGPUSubtargetFeature<"atomic-fmin-fmax-flat-f64",
- "Has flat memory instructions for atomicrmw fmin/fmax for double",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/1,
- [FeatureFlatAddressSpace]
->;
-
-defm AtomicFaddNoRtnInsts : AMDGPUSubtargetFeature<"atomic-fadd-no-rtn-insts",
- "Has buffer_atomic_add_f32 and global_atomic_add_f32 instructions that "
- "don't return original value",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/1,
- [FeatureFlatGlobalInsts]
->;
-
-def FeatureAtomicBufferGlobalPkAddF16NoRtnInsts
- : SubtargetFeature<"atomic-buffer-global-pk-add-f16-no-rtn-insts",
- "HasAtomicBufferGlobalPkAddF16NoRtnInsts",
- "true",
- "Has buffer_atomic_pk_add_f16 and global_atomic_pk_add_f16 instructions that "
- "don't return original value",
- [FeatureFlatGlobalInsts]
->;
-
-defm AtomicBufferGlobalPkAddF16Insts : AMDGPUSubtargetFeature<"atomic-buffer-global-pk-add-f16-insts",
- "Has buffer_atomic_pk_add_f16 and global_atomic_pk_add_f16 instructions that "
- "can return original value",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/1,
- [FeatureFlatGlobalInsts]
->;
-
-defm AtomicGlobalPkAddBF16Inst : AMDGPUSubtargetFeature<"atomic-global-pk-add-bf16-inst",
- "Has global_atomic_pk_add_bf16 instruction",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/1,
- [FeatureFlatGlobalInsts]
->;
-
-defm AtomicBufferPkAddBF16Inst : AMDGPUSubtargetFeature<"atomic-buffer-pk-add-bf16-inst",
- "Has buffer_atomic_pk_add_bf16 instruction"
->;
-
-defm AtomicCSubNoRtnInsts : AMDGPUSubtargetFeature<"atomic-csub-no-rtn-insts",
- "Has buffer_atomic_csub and global_atomic_csub instructions that don't "
- "return original value",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/0
->;
-
-defm FlatAtomicFaddF32Inst : AMDGPUSubtargetFeature<"flat-atomic-fadd-f32-inst",
- "Has flat_atomic_add_f32 instruction",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/1,
- [FeatureFlatAddressSpace]
->;
-
-defm GFX1251GEMMInsts : AMDGPUSubtargetFeature<"gfx1251-gemm-insts",
- "Has additional gfx1251 DGEMM instructions"
->;
-
-defm GFX125xLowestRateWMMA : AMDGPUGenericAnySubtargetFeature<"gfx125x-lowest-rate-wmma",
- "Has the lowest rate wmma in the gfx125x family, mainly for the gfx12-5-generic"
->;
-
-defm FlatBufferGlobalAtomicFaddF64Inst : AMDGPUSubtargetFeature<"flat-buffer-global-fadd-f64-inst",
- "Has flat, buffer, and global instructions for f64 atomic fadd"
->;
-
-def FeatureMemoryAtomicFAddF32DenormalSupport
- : SubtargetFeature<"memory-atomic-fadd-f32-denormal-support",
- "HasMemoryAtomicFaddF32DenormalSupport",
- "true",
- "global/flat/buffer atomic fadd for float supports denormal handling"
->;
-
-defm AgentScopeFineGrainedRemoteMemoryAtomics : AMDGPUSubtargetFeature<
- "agent-scope-fine-grained-remote-memory-atomics",
- "Agent (device) scoped atomic operations, excluding those directly "
- "supported by PCIe (i.e. integer atomic add, exchange, and "
- "compare-and-swap), are functional for allocations in host or peer "
- "device memory.",
- /*GenPredicate=*/0
->;
-
-defm EmulatedSystemScopeAtomics : AMDGPUSubtargetFeature<
- "emulated-system-scope-atomics",
- "System scope atomics unsupported by the PCI-e are emulated in HW via CAS "
- "loop and functional.",
- /*GenPredicate=*/0
->;
-
-defm DefaultComponentZero : AMDGPUSubtargetFeature<"default-component-zero",
- "BUFFER/IMAGE store instructions set unspecified components to zero (before GFX12)"
->;
-
-defm DefaultComponentBroadcast : AMDGPUSubtargetFeature<"default-component-broadcast",
- "BUFFER/IMAGE store instructions set unspecified components to x component (GFX12)"
->;
-
-// The default assumption needs to be ecc is enabled, but no directly
-// exposed operations depend on it, so it can be safely inlined.
-def FeatureSupportsSRAMECC : SubtargetFeature<"sramecc-support",
- "SupportsSRAMECC",
- "true",
- "Hardware supports SRAMECC",
- [], InlineIgnore
->;
-
-defm SRAMECCOnOffModes : AMDGPUSubtargetFeature<"sramecc-on-off-modes",
- "Target supports SRAMECC on/off modes",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/0,
- [FeatureSupportsSRAMECC],
- InlineIgnore
->;
-
-defm NoSdstCMPX : AMDGPUSubtargetFeature<"no-sdst-cmpx",
- "V_CMPX does not write VCC/SGPR in addition to EXEC"
->;
-
-defm Vscnt : AMDGPUSubtargetFeature<"vscnt",
- "Has separate store vscnt counter",
- /*GenPredicate=*/0
->;
-
-defm Asynccnt : AMDGPUSubtargetFeature<"asynccnt",
- "Has separate asynccnt counter",
- /*GenPredicate=*/0
->;
-
-defm GetWaveIdInst : AMDGPUSubtargetFeature<"get-wave-id-inst",
- "Has s_get_waveid_in_workgroup instruction"
->;
-
-defm PopsExitingWaveID : AMDGPUSubtargetFeature<"pops-exiting-wave-id",
- "Has POPS_EXITING_WAVE_ID hardware register"
->;
-
-defm SMemTimeInst : AMDGPUSubtargetFeature<"s-memtime-inst",
- "Has s_memtime instruction"
->;
-
-defm ShaderCyclesRegister : AMDGPUSubtargetFeature<"shader-cycles-register",
- "Has SHADER_CYCLES hardware register"
->;
-
-defm ShaderCyclesHiLoRegisters : AMDGPUSubtargetFeature<"shader-cycles-hi-lo-registers",
- "Has SHADER_CYCLES_HI/LO hardware registers",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/0
->;
-
-defm MadMacF32Insts : AMDGPUSubtargetFeature<"mad-mac-f32-insts",
- "Has v_mad_f32/v_mac_f32/v_madak_f32/v_madmk_f32 instructions"
->;
-
-defm DsSrc2Insts : AMDGPUSubtargetFeature<"ds-src2-insts",
- "Has ds_*_src2 instructions"
->;
-
-defm VOP3Literal : AMDGPUSubtargetFeature<"vop3-literal",
- "Can use one literal in VOP3",
- /*GenPredicate=*/0
->;
-
-defm NoDataDepHazard : AMDGPUSubtargetFeature<"no-data-dep-hazard",
- "Does not need SW waitstates",
- /*GenPredicate=*/0
->;
-
-// 1536 physical VGPRs for wave32 and 768 VGPRs for wave64
-// with allocation granularity 24 for wave32 and 12 for wave64
-defm 1536VGPRs : AMDGPUSubtargetFeature<"1536-physical-vgprs",
- "Has 1536 physical VGPRs per SIMD",
- /*GenPredicate=*/0
->;
-
-defm VOPDInsts : AMDGPUSubtargetFeature<"vopd",
- "Has VOPD dual issue wave32 instructions",
- /*GenPredicate=*/0
->;
-
-defm VOPDInterlockHazard : AMDGPUSubtargetFeature<"vopd-interlock-hazard",
- "VOPD destination buffer forwarding drops the interlock when SRC0 or SRC1 "
- "X/Y are distinct VGPRs with the same register-number parity",
- /*GenPredicate=*/0
->;
-
-defm VALUTransUseHazard : AMDGPUGenericAnySubtargetFeature<"valu-trans-use-hazard",
- "Hazard when TRANS instructions are closely followed by a use of the result",
- /*GenPredicate=*/0
->;
-
-defm TransCoexecutionHazard : AMDGPUSubtargetFeature<"trans-coexecution-hazard",
- "Hazard when TRANS co-executes with other VALU instruction",
- /*GenPredicate=*/0
->;
-
-defm GFX1250VALUBlockingCycles : AMDGPUSubtargetFeature<"gfx1250-valu-blocking-cycles",
- "VALU pipe occupancy is modeled with GFX1250BlockingCycles",
- /*GenPredicate=*/0
->;
-
-defm VPermPk16Hazard : AMDGPUSubtargetFeature<"vperm-pk16-hazard",
- "V_PERM_PK16 must be immediately followed by a safe VALU (or an inserted V_NOP)",
- /*GenPredicate=*/0
->;
-
-defm WMMACoexecutionHazards : AMDGPUSubtargetFeature<"wmma-coexecution-hazards",
- "Hazards when WMMA co-executes with other WMMA/VALU instruction",
- /*GenPredicate=*/0
->;
-
-defm VOPD3F64OPYSrcHazard : AMDGPUSubtargetFeature<"vopd3-f64-opy-src-hazard",
- "Hazard when an OPY VGPR source of a VOPD3 with a 64-bit OPX names the last "
- "VGPR allocated to the wave",
- /*GenPredicate=*/0
->;
-
-defm SALUFloatInsts : AMDGPUSubtargetFeature<"salu-float",
- "Has SALU floating point instructions"
->;
-
-defm PseudoScalarTrans : AMDGPUSubtargetFeature<"pseudo-scalar-trans",
- "Has Pseudo Scalar Transcendental instructions"
->;
-
-defm NoF16PseudoScalarTransInlineConstants : AMDGPUSubtargetFeature<"no-f16-pseudo-scalar-trans-inline-constants",
- "Inline constants are not supported for F16 pseudo scalar transcendentals"
->;
-
-defm RestrictedSOffset : AMDGPUSubtargetFeature<"restricted-soffset",
- "Has restricted SOffset (immediate not supported)."
->;
-
-defm RequiredExportPriority : AMDGPUGenericAnySubtargetFeature<"required-export-priority",
- "Export priority must be explicitly manipulated on GFX11.5",
- /*GenPredicate=*/0
->;
-
-defm VmemWriteVgprInOrder : AMDGPUSubtargetFeature<"vmem-write-vgpr-in-order",
- "VMEM instructions of the same type write VGPR results in order",
- /*GenPredicate=*/0
->;
-
-defm BitOp3Insts : AMDGPUSubtargetFeature<"bitop3-insts",
- "Has v_bitop3_b32/v_bitop3_b16 instructions"
->;
-
-defm TanhInsts : AMDGPUSubtargetFeature<"tanh-insts",
- "Has v_tanh_f32/f16 instructions"
->;
-
-defm TensorCvtLutInsts : AMDGPUSubtargetFeature<"tensor-cvt-lut-insts",
- "Has v_perm_pk16* instructions"
->;
-
-defm TDMInsts : AMDGPUSubtargetFeature<"tdm-insts",
- "Has tensor_load_to_lds/tensor_store_from_lds instructions"
->;
-
-defm TransposeLoadF4F6Insts : AMDGPUSubtargetFeature<"transpose-load-f4f6-insts",
- "Has ds_load_tr4/tr6 and global_load_tr4/tr6 instructions"
->;
-
-defm PrngInst : AMDGPUSubtargetFeature<"prng-inst",
- "Has v_prng_b32 instruction"
->;
-
-defm BVHDualAndBVH8Insts : AMDGPUSubtargetFeature<"bvh-dual-bvh-8-insts",
- "Has image_bvh_dual_intersect_ray and image_bvh8_intersect_ray instructions"
->;
-
-defm PointSampleAccel : AMDGPUSubtargetFeature<"point-sample-accel",
- "Has point sample acceleration feature",
- /*GenPredicate=*/0
->;
-
-defm 64BitLiterals : AMDGPUSubtargetFeature<"64-bit-literals",
- "Can use 64-bit literals with single DWORD instructions"
->;
-
-defm 1024AddressableVGPRs : AMDGPUSubtargetFeature<"1024-addressable-vgprs",
- "Has 1024 addressable VGPRs"
->;
-
-defm SetregVGPRMSBFixup : AMDGPUGenericAnySubtargetFeature<"setreg-vgpr-msb-fixup",
- "S_SETREG to MODE clobbers VGPR MSB bits, requires fixup",
- /*GenPredicate=*/0
->;
-
-defm WaitXcnt : AMDGPUSubtargetFeature<"wait-xcnt",
- "Has s_wait_xcnt instruction"
->;
-
-defm SetPrioIncWgInst : AMDGPUSubtargetFeature<"setprio-inc-wg-inst",
- "Has s_setprio_inc_wg instruction."
->;
-
-defm SWakeupBarrier : AMDGPUSubtargetFeature<"s-wakeup-barrier-inst",
- "Has s_wakeup_barrier instruction."
->;
-
-defm FlatSignedOffset : AMDGPUSubtargetFeature<"flat-signed-offset",
- "Immediate offset of FLAT instructions are always signed"
->;
-
-defm RequiresInitialUnclausedVmem : AMDGPUSubtargetFeature<"requires-initial-unclause-vmem",
- "Hardware entrypoints require first VMEM instruction to be unclaused",
- GenPredicate=false
->;
-
-defm SlowMaxMinMulI64Insts : AMDGPUSubtargetFeature<"slow-max-min-mul-i64-insts",
- "64-bit integer max/min/mul VALU instructions are slower than expanding to an equivalent code sequence",
- GenPredicate=false
->;
-
-//===------------------------------------------------------------===//
-// Subtarget Features (options and debugging)
-//===------------------------------------------------------------===//
-
-// Ugly hack to accomodate assembling modules with mixed
-// wavesizes. Ideally we would have a mapping symbol in assembly which
-// would keep track of which sections of code should be treated as
-// wave32 and wave64. Instead what users do is assemble with both
-// wavesizes enabled. We translate this into this special mode so this
-// only influences assembler behavior and nothing else.
-defm AssemblerPermissiveWavesize : AMDGPUSubtargetFeature<"assembler-permissive-wavesize",
- "Allow parsing wave32 and wave64 variants of instructions",
- /*GenPredicate=*/0
->;
-
-class FeatureMaxPrivateElementSize<int size> : SubtargetFeature<
- "max-private-element-size-"#size,
- "MaxPrivateElementSize",
- !cast<string>(size),
- "Maximum private access size may be "#size
->;
-
-def FeatureMaxPrivateElementSize4 : FeatureMaxPrivateElementSize<4>;
-def FeatureMaxPrivateElementSize8 : FeatureMaxPrivateElementSize<8>;
-def FeatureMaxPrivateElementSize16 : FeatureMaxPrivateElementSize<16>;
-
-def FeatureDumpCodeLower : SubtargetFeature <"dumpcode",
- "DumpCode",
- "true",
- "Dump MachineInstrs in the CodeEmitter"
->;
-
-// XXX - This should probably be removed once enabled by default
-def FeatureEnableLoadStoreOpt : SubtargetFeature <"load-store-opt",
- "EnableLoadStoreOpt",
- "true",
- "Enable SI load/store optimizer pass",
- [], InlineIgnore
->;
-
-// Performance debugging feature. Allow using DS instruction immediate
-// offsets even if the base pointer can't be proven to be base. On SI,
-// base pointer values that won't give the same result as a 16-bit add
-// are not safe to fold, but this will override the conservative test
-// for the base pointer.
-def FeatureEnableUnsafeDSOffsetFolding : SubtargetFeature <
- "unsafe-ds-offset-folding",
- "EnableUnsafeDSOffsetFolding",
- "true",
- "Force using DS instruction immediate offsets on SI",
- [], InlineIgnore
->;
-
-def FeatureEnableSIScheduler : SubtargetFeature<"si-scheduler",
- "EnableSIScheduler",
- "true",
- "Enable SI Machine Scheduler",
- [], InlineIgnore
->;
-
-def FeatureEnableDS128 : SubtargetFeature<"enable-ds128",
- "EnableDS128",
- "true",
- "Use ds_{read|write}_b128"
->;
-
-// Sparse texture support requires that all result registers are zeroed when
-// PRTStrictNull is set to true. This feature is turned on for all architectures
-// but is enabled as a feature in case there are situations where PRTStrictNull
-// is disabled by the driver.
-def FeatureEnablePRTStrictNull : SubtargetFeature<"enable-prt-strict-null",
- "EnablePRTStrictNull",
- "true",
- "Enable zeroing of result registers for sparse texture fetches"
->;
-
-// Unless +-flat-for-global is specified, turn on FlatForGlobal for
-// all OS-es on VI and newer hardware to avoid assertion failures due
-// to missing ADDR64 variants of MUBUF instructions.
-// FIXME: moveToVALU should be able to handle converting addr64 MUBUF
-// instructions.
-
-def FeatureUseFlatForGlobal : SubtargetFeature<"flat-for-global",
- "UseFlatForGlobal",
- "true",
- "Force to generate flat instruction for global",
- [], InlineIgnore
->;
-
-defm AutoWaitcntBeforeBarrier : AMDGPUSubtargetFeature <"auto-waitcnt-before-barrier",
- "Hardware automatically inserts waitcnt before barrier",
- /*GenPredicate=*/0, /*GenAssemblerPredicate=*/1, [], InlineIgnore
->;
-
-defm BackOffBarrier : AMDGPUSubtargetFeature <"back-off-barrier",
- "Hardware supports backing off s_barrier if an exception occurs",
- /*GenPredicate=*/0
->;
-
-defm TrigReducedRange : AMDGPUSubtargetFeature<"trig-reduced-range",
- "Requires use of fract on arguments to trig instructions",
- /*GenPredicate=*/0
->;
-
-defm KernargPreload : AMDGPUSubtargetFeature <"kernarg-preload",
- "Hardware supports preloading of kernel arguments in user SGPRs.",
- /*GenPredicate=*/0
->;
-
-// Alignment enforcement is controlled by a configuration register:
-// SH_MEM_CONFIG.alignment_mode
-defm UnalignedAccessMode : AMDGPUSubtargetFeature<"unaligned-access-mode",
- "Enable unaligned global, local and region loads and stores if the hardware"
- " supports it",
- /*GenPredicate=*/1, /*GenAssemblerPredicate=*/1, [], InlineIgnore
->;
-
-defm PackedTID : AMDGPUSubtargetFeature<"packed-tid",
- "Workitem IDs are packed into v0 at kernel launch",
- /*GenPredicate=*/0
->;
-
-defm ArchitectedFlatScratch : AMDGPUSubtargetFeature<"architected-flat-scratch",
- "Flat Scratch register is a readonly SPI initialized architected register",
- /*GenPredicate=*/0
->;
-
-defm ArchitectedSGPRs : AMDGPUSubtargetFeature<"architected-sgprs",
- "Enable the architected SGPRs",
- /*GenPredicate=*/0
->;
-
-defm GDS : AMDGPUSubtargetFeature<"gds",
- "Has Global Data Share",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/0
->;
-
-defm GWS : AMDGPUSubtargetFeature<"gws",
- "Has Global Wave Sync",
- /*GenPredicate=*/1,
- /*GenAssemblerPredicate=*/0
->;
-
-// Generic targets require code object v6 independently of their covered GPUs,
-// so this feature does not use the any-covered-GPU classification.
-def FeatureRequiresCOV6 : SubtargetFeature<"requires-cov6",
- "RequiresCOV6",
- "true",
- "Target Requires Code Object V6"
->;
-
-defm XF32Insts : AMDGPUSubtargetFeature<"xf32-insts",
- "Has instructions that support xf32 format, such as "
- "v_mfma_f32_16x16x8_xf32 and v_mfma_f32_32x32x4_xf32"
->;
-
-defm GloballyAddressableScratch : AMDGPUSubtargetFeature<"globally-addressable-scratch",
- "FLAT instructions can access scratch memory for any thread in any wave",
- /*GenPredicate=*/0
->;
-
-// Enable the use of SCRATCH_STORE/LOAD_BLOCK instructions for saving and
-// restoring the callee-saved registers.
-def FeatureUseBlockVGPROpsForCSR : SubtargetFeature<"block-vgpr-csr",
- "UseBlockVGPROpsForCSR",
- "true",
- "Use block load/store for VGPR callee saved registers"
->;
-
-defm LshlAddU64Inst : AMDGPUSubtargetFeature<"lshl-add-u64-inst",
- "Has v_lshl_add_u64 instruction"
->;
-
-defm AddSubU64Insts : AMDGPUSubtargetFeature<"add-sub-u64-insts",
- "Has v_add_u64 and v_sub_u64 instructions"
->;
-
-defm MadU32Inst : AMDGPUSubtargetFeature<"mad-u32-inst",
- "Has v_mad_u32 instruction"
->;
-
-defm MadNC64_32Insts : AMDGPUSubtargetFeature<"mad-nc-64-32-insts",
- "Has v_mad_nc_{u64_u32|i64_i32} instructions"
->;
-
-defm AddMinMaxInsts : AMDGPUSubtargetFeature<"add-min-max-insts",
- "Has v_add_{min|max}_{i|u}32 instructions"
->;
-
-defm MinMaxI64Insts : AMDGPUSubtargetFeature<"min-max-i64-insts",
- "Has v_{min|max}_{i|u}64 instructions"
->;
-
-defm PkAddMinMaxInsts : AMDGPUSubtargetFeature<"pk-add-min-max-insts",
- "Has v_pk_add_{min|max}_{i|u}16 instructions"
->;
-
-defm VMemToLDSLoad : AMDGPUSubtargetFeature<"vmem-to-lds-load-insts",
- "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."
->;
-
-// Manual predicate for hasAsyncMark() which combines HasVMemToLDSLoad and GFX1250Plus
-def HasAsyncMark : Predicate<"Subtarget->hasAsyncMark()">;
-
-defm LdsBarrierArriveAtomic : AMDGPUSubtargetFeature<"lds-barrier-arrive-atomic",
- "Has LDS barrier-arrive atomic instructions"
->;
-
-// The width of the num_records field of the buffer resource (V#).
-class FeatureNumRecordsBufferResource<int width> : SubtargetFeature<
- !cast<string>(width)#"-bit-num-records-buffer-resource",
- "BufferResourceNumRecordsWidth",
- !cast<string>(width),
- "The buffer resource (V#) has "#width#"-bit num_records"
->;
-
-def Feature32BitNumRecordsBufferResource : FeatureNumRecordsBufferResource<32>;
-def Feature45BitNumRecordsBufferResource : FeatureNumRecordsBufferResource<45>;
-
-defm Clusters : AMDGPUSubtargetFeature<"clusters",
- "Has clusters of workgroups support",
- /*GenPredicate=*/0
->;
-
-def FeatureWaitsBeforeSystemScopeStores : SubtargetFeature<
- "waits-before-system-scope-stores",
- "RequiresWaitsBeforeSystemScopeStores",
- "true",
- "Target requires waits for loads and atomics before system scope stores"
->;
-
-def FeatureUseAddPC64Inst : SubtargetFeature<"use-add-pc64-inst",
- "UseAddPC64Inst",
- "true",
- "Use s_add_pc_i64 instruction."
->;
-
-defm GFX1250_STRICT : AMDGPUSubtargetFeature<"gfx1250-strict",
- "gfx1250 strict mode",
- /*GenPredicate=*/0
->;
-
-defm VOP3PX2IncrementsVaVdstTwice : AMDGPUSubtargetFeature<"vop3px2-increments-va-vdst-twice",
- "VOP3PX2 encoded instructions increment the VaVdst counter twice", GenPredicate=false
->;
-
-defm NeedsTDMDrain : AMDGPUSubtargetFeature<"needs-tdm-drain",
- "Needs to keep tensor count low",
- /*GenPredicate=*/0
->;
-
-defm NeedsAligned2addrDS : AMDGPUSubtargetFeature<"needs-aligned-2addr-lds",
- "Needs 2addr LDS instructions to be aligned to payload size",
- /*GenPredicate=*/1
->;
-
-defm NoSleepForever : AMDGPUSubtargetFeature<"no-sleep-forever",
- "s_monitor_sleep should not use sleep_forever",
- /*GenPredicate=*/1
->;
-
-defm WaveMatchInsts : AMDGPUSubtargetFeature<"wave-match-insts",
- "Has v_wave_match_b32 instruction"
->;
-
-//===----------------------------------------------------------------------===//
-
-class GCNSubtargetFeatureGeneration <string Value,
- string FeatureName,
- list<SubtargetFeature> Implies> :
- SubtargetFeatureGeneration <Value, FeatureName, "GCNSubtarget", Implies>;
-
-def FeatureSouthernIslands : GCNSubtargetFeatureGeneration<"SOUTHERN_ISLANDS",
- "southern-islands",
- [FeatureFP64, FeatureAddressableLocalMemorySize32768,
- FeatureHalfAddressablePhysicalLocalMemory,
- FeatureLDSAllocGranularity256,
- FeatureLDSEncodingGranularity256,
- FeatureMIMG_R128,
- FeatureWavefrontSize64, FeatureSupportsWave64, FeatureSMemTimeInst,
- FeatureMadMacF32Insts,
- FeatureDsSrc2Insts, FeatureLDSBankCount32, FeatureMovrel,
- FeatureTrigReducedRange, FeatureExtendedImageInsts, FeatureImageInsts,
- FeatureInstCacheLineSize64,
- FeatureGDS, FeatureGWS, FeatureDefaultComponentZero,
- FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF64GlobalInsts,
- FeatureVmemWriteVgprInOrder, FeatureCubeInsts, FeatureLerpInst,
- FeatureSadInsts, FeatureMsadInsts, FeatureMqsadPkInsts,
- FeatureCvtPkNormVOP2Insts, FeatureDX10ClampAndIEEEMode, FeatureMTBUFInsts,
- FeatureFormattedMUBUFInsts, Feature32BitNumRecordsBufferResource,
- FeatureMaxWavesPerEU10
- ]
->;
-
-def FeatureSeaIslands : GCNSubtargetFeatureGeneration<"SEA_ISLANDS",
- "sea-islands",
- [FeatureFP64, FeatureAddressableLocalMemorySize65536,
- FeatureLDSAllocGranularity512,
- FeatureLDSEncodingGranularity512,
- FeatureMIMG_R128,
- FeatureWavefrontSize64, FeatureSupportsWave64, FeatureFlatAddressSpace,
- FeatureCIInsts, FeatureMovrel, FeatureTrigReducedRange,
- FeatureGFX7GFX8GFX9Insts, FeatureSMemTimeInst, FeatureMadMacF32Insts,
- FeatureDsSrc2Insts, FeatureExtendedImageInsts, FeatureUnalignedBufferAccess,
- FeatureImageInsts, FeatureGDS, FeatureGWS, FeatureDefaultComponentZero,
- FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF64GlobalInsts,
- FeatureAtomicFMinFMaxF32FlatInsts, FeatureAtomicFMinFMaxF64FlatInsts,
- FeatureVmemWriteVgprInOrder, FeatureCubeInsts, FeatureLerpInst,
- FeatureSadInsts, FeatureQsadInsts, FeatureMsadInsts, FeatureMqsadPkInsts,
- FeatureMqsadInsts, FeatureCvtPkNormVOP2Insts,
- FeatureDX10ClampAndIEEEMode, FeatureInstCacheLineSize64,
- FeatureMTBUFInsts, FeatureFormattedMUBUFInsts,
- Feature32BitNumRecordsBufferResource,
- FeatureMaxWavesPerEU10
- ]
->;
-
-def FeatureVolcanicIslands : GCNSubtargetFeatureGeneration<"VOLCANIC_ISLANDS",
- "volcanic-islands",
- [FeatureFP64, FeatureAddressableLocalMemorySize65536,
- FeatureLDSAllocGranularity512,
- FeatureLDSEncodingGranularity512,
- FeatureMIMG_R128,
- FeatureWavefrontSize64, FeatureSupportsWave64, FeatureFlatAddressSpace,
- FeatureGCN3Encoding, FeatureCIInsts, Feature16BitInsts,
- FeatureSMemRealTime, FeatureVGPRIndexMode, FeatureMovrel,
- FeatureScalarStores, FeatureInv2PiInlineImm,
- FeatureSDWA, FeatureSDWAOutModsVOPC, FeatureSDWAMac, FeatureDPP,
- FeatureDPPWavefrontShifts, FeatureDPPBroadcasts,
- FeatureIntClamp, FeatureTrigReducedRange, FeatureGFX8Insts,
- FeatureGFX7GFX8GFX9Insts, FeatureSMemTimeInst, FeatureMadMacF32Insts,
- FeatureDsSrc2Insts, FeatureExtendedImageInsts, FeatureFastDenormalF32,
- FeatureUnalignedBufferAccess, FeatureImageInsts, FeatureGDS, FeatureGWS,
- FeatureDefaultComponentZero, FeatureVmemWriteVgprInOrder, FeatureCubeInsts,
- FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts, FeatureMsadInsts,
- FeatureMqsadPkInsts, FeatureMqsadInsts,
- FeatureCvtPkNormVOP2Insts, FeatureDX10ClampAndIEEEMode,
- FeatureInstCacheLineSize64, FeatureMTBUFInsts, FeatureFormattedMUBUFInsts,
- Feature32BitNumRecordsBufferResource,
- FeatureMaxWavesPerEU10
- ]
->;
-
-def FeatureGFX9 : GCNSubtargetFeatureGeneration<"GFX9",
- "gfx9",
- [FeatureFP64,
- FeatureWavefrontSize64, FeatureSupportsWave64, FeatureFlatAddressSpace,
- FeatureGCN3Encoding, FeatureCIInsts, Feature16BitInsts,
- FeatureSMemRealTime, FeatureScalarStores, FeatureInv2PiInlineImm,
- FeatureApertureRegs, FeatureGetDoorbellID,
- FeatureGFX9Insts, FeatureVOP3PInsts, FeatureVGPRIndexMode,
- FeatureFastFMAF32, FeatureDPP, FeatureDPPWavefrontShifts, FeatureDPPBroadcasts,
- FeatureIntClamp, FeatureSDWA, FeatureSDWAOmod, FeatureSDWAScalar, FeatureSDWASdst,
- FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts,
- FeatureAddNoCarryInsts, FeatureGFX8Insts, FeatureGFX7GFX8GFX9Insts,
- FeatureScalarFlatScratchInsts, FeatureScalarAtomics, FeatureR128A16,
- FeatureA16, FeatureSMemTimeInst, FeatureFastDenormalF32, FeatureXNACKOnOffModes,
- FeatureUnalignedBufferAccess, FeatureUnalignedScratchAccess,
- FeatureUnalignedDSAccess, FeatureNegativeScratchOffsetBug, FeatureGWS,
- FeatureDefaultComponentZero,FeatureVmemWriteVgprInOrder, FeatureVMemToLDSLoad,
- FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts,
- FeatureMsadInsts, FeatureMqsadPkInsts, FeatureMqsadInsts,
- FeatureCvtNormInsts, FeatureCvtPkNormVOP2Insts,
- FeatureCvtPkNormVOP3Insts, FeatureDX10ClampAndIEEEMode,
- FeatureInstCacheLineSize64, FeatureMTBUFInsts, FeatureFormattedMUBUFInsts,
- Feature32BitNumRecordsBufferResource
- ]
->;
-
-def FeatureGFX10 : GCNSubtargetFeatureGeneration<"GFX10",
- "gfx10",
- [FeatureFP64, FeatureAddressableLocalMemorySize65536,
- FeatureLDSEncodingGranularity512,
- FeatureHalfAddressablePhysicalLocalMemory, FeatureMIMG_R128,
- FeatureSupportsWave32, FeatureSupportsWave64, FeatureSupportsWGP,
- FeatureFlatAddressSpace,
- FeatureCIInsts, Feature16BitInsts,
- FeatureSMemRealTime, FeatureInv2PiInlineImm,
- FeatureApertureRegs, FeatureGetDoorbellID,
- FeatureGFX9Insts, FeatureGFX10Insts, FeatureVOP3PInsts,
- FeatureMovrel, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp,
- FeatureSDWA, FeatureSDWAOmod, FeatureSDWAScalar, FeatureSDWASdst,
- FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts,
- FeatureAddNoCarryInsts, FeatureFmaMixInsts, FeatureGFX8Insts,
- FeatureNoSdstCMPX, FeatureVscnt,
- FeatureVOP3Literal, FeatureDPP8, FeatureExtendedImageInsts,
- FeaturePermlane16Insts,
- FeatureNoDataDepHazard, FeaturePkFmacF16Inst,
- FeatureA16, FeatureSMemTimeInst, FeatureFastDenormalF32, FeatureG16,
- FeatureUnalignedBufferAccess, FeatureUnalignedScratchAccess,
- FeatureUnalignedDSAccess, FeatureImageInsts, FeatureGDS, FeatureGWS,
- FeatureDefaultComponentZero, FeatureMaxHardClauseLength63,
- FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF64GlobalInsts,
- FeatureAtomicFMinFMaxF32FlatInsts, FeatureAtomicFMinFMaxF64FlatInsts,
- FeatureVmemWriteVgprInOrder, FeatureVMemToLDSLoad, FeatureCubeInsts,
- FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts, FeatureMsadInsts,
- FeatureMqsadPkInsts, FeatureMqsadInsts,
- FeatureCvtNormInsts, FeatureCvtPkNormVOP2Insts,
- FeatureCvtPkNormVOP3Insts, FeatureDX10ClampAndIEEEMode, FeatureFlatOffsetBits12,
- FeatureInstCacheLineSize64, FeatureMTBUFInsts, FeatureFormattedMUBUFInsts,
- Feature32BitNumRecordsBufferResource
- ]
->;
-
-def FeatureGFX11 : GCNSubtargetFeatureGeneration<"GFX11",
- "gfx11",
- [FeatureFP64, FeatureAddressableLocalMemorySize65536,
- FeatureLDSAllocGranularity1024,
- FeatureLDSEncodingGranularity512,
- FeatureHalfAddressablePhysicalLocalMemory, FeatureMIMG_R128,
- FeatureSupportsWave32, FeatureSupportsWave64, FeatureSupportsWGP,
- FeatureFlatAddressSpace, Feature16BitInsts,
- FeatureInv2PiInlineImm, FeatureApertureRegs, FeatureGetDoorbellID,
- FeatureCIInsts, FeatureGFX8Insts, FeatureGFX9Insts, FeatureGFX10Insts,
- FeatureBVHRayTracingInsts, FeatureMSAALoadInsts,
- FeatureGFX10_BEncoding, FeatureGFX10_3Insts,
- FeatureGFX11Insts, FeatureVOP3PInsts, FeatureVOPDInsts, FeatureTrue16BitInsts,
- FeatureMovrel, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp,
- FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts,
- FeatureAddNoCarryInsts, FeatureFmaMixInsts,
- FeatureNoSdstCMPX, FeatureVscnt,
- FeatureVOP3Literal, FeatureDPP8, FeatureExtendedImageInsts,
- FeaturePermlane16Insts,
- FeatureNoDataDepHazard, FeaturePkFmacF16Inst,
- FeatureA16, FeatureFastDenormalF32, FeatureG16,
- FeatureUnalignedBufferAccess, FeatureUnalignedScratchAccess,
- FeatureUnalignedDSAccess, FeatureGDS, FeatureGWS,
- FeatureDefaultComponentZero, FeatureMaxHardClauseLength32,
- FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF32FlatInsts,
- FeatureVmemWriteVgprInOrder, FeatureCubeInsts, FeatureLerpInst,
- FeatureSadInsts, FeatureQsadInsts, FeatureMsadInsts, FeatureMqsadPkInsts,
- FeatureMqsadInsts, FeatureCvtNormInsts,
- FeatureCvtPkNormVOP2Insts, FeatureCvtPkNormVOP3Insts,
- FeatureInstCacheLineSize128, FeatureMTBUFInsts, FeatureFormattedMUBUFInsts,
- Feature32BitNumRecordsBufferResource,
- FeatureMaxWavesPerEU16
- ]
->;
-
-def FeatureGFX12 : GCNSubtargetFeatureGeneration<"GFX12",
- "gfx12",
- [FeatureFP64, FeatureMIMG_R128,
- FeatureSupportsWave32,
- FeatureFlatAddressSpace, Feature16BitInsts,
- FeatureInv2PiInlineImm, FeatureApertureRegs, FeatureGetDoorbellID,
- FeatureCIInsts, FeatureGFX8Insts, FeatureGFX9Insts, FeatureGFX10Insts,
- FeatureGFX10_BEncoding, FeatureGFX10_3Insts,
- FeatureGFX11Insts, FeatureGFX12Insts, FeatureVOP3PInsts, FeatureVOPDInsts,
- FeatureMovrel, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp,
- FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts,
- FeatureAddNoCarryInsts, FeatureFmaMixInsts,
- FeatureNoSdstCMPX, FeatureVscnt,
- FeatureVOP3Literal, FeatureDPP8,
- FeaturePermlane16Insts,
- FeatureNoDataDepHazard, FeaturePkFmacF16Inst,
- FeatureA16, FeatureFastDenormalF32, FeatureG16,
- FeatureUnalignedBufferAccess, FeatureUnalignedScratchAccess,
- FeatureUnalignedDSAccess, FeatureTrue16BitInsts,
- FeatureDefaultComponentBroadcast, FeatureMaxHardClauseLength32,
- FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF32FlatInsts,
- FeatureAtomicDsCondSubClampInsts, FeatureAtomicCondSubClampFlatInsts,
- FeatureIEEEMinimumMaximumInsts, FeatureSALUMinimumMaximumInsts,
- FeatureMinimum3Maximum3F32, FeatureMinimum3Maximum3F16,
- FeatureAgentScopeFineGrainedRemoteMemoryAtomics, FeatureFlatOffsetBits24,
- FeatureFlatSignedOffset, FeatureInstCacheLineSize128,
- FeatureMaxWavesPerEU16
- ]
->;
-
-def FeatureGFX13 : GCNSubtargetFeatureGeneration<"GFX13",
- "gfx13",
- [FeatureFP64, FeatureMIMG_R128,
- FeatureSupportsWave32, FeatureSupportsWave64, FeatureSupportsWGP,
- FeatureFlatAddressSpace, Feature16BitInsts,
- FeatureInv2PiInlineImm, FeatureApertureRegs, FeatureGetDoorbellID,
- FeatureCIInsts, FeatureGFX8Insts, FeatureGFX9Insts, FeatureGFX10Insts,
- FeatureMSAALoadInsts, FeatureGFX10_BEncoding, FeatureGFX10_3Insts,
- FeatureGFX11Insts, FeatureGFX12Insts, FeatureGFX13Insts, FeatureVOP3PInsts,
- FeatureVOPDInsts, FeatureMovrel, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp,
- FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts,
- FeatureAddNoCarryInsts, FeatureFmaMixInsts,
- FeatureNoSdstCMPX, FeatureVscnt,
- FeatureVOP3Literal, FeatureDPP8,
- FeaturePermlane16Insts,
- FeatureNoDataDepHazard, FeaturePkFmacF16Inst,
- FeatureA16, FeatureFastDenormalF32, FeatureG16,
- FeatureUnalignedBufferAccess, FeatureUnalignedScratchAccess, FeatureImageInsts,
- FeatureExtendedImageInsts,
- FeatureUnalignedDSAccess, FeatureTrue16BitInsts,
- FeatureDefaultComponentBroadcast, FeatureMaxHardClauseLength32,
- FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF32FlatInsts,
- FeatureAtomicDsCondSubClampInsts, FeatureAtomicCondSubClampFlatInsts,
- FeatureIEEEMinimumMaximumInsts, FeatureSALUMinimumMaximumInsts,
- FeatureMinimum3Maximum3F32, FeatureMinimum3Maximum3F16,
- FeatureAgentScopeFineGrainedRemoteMemoryAtomics, FeatureFlatOffsetBits24,
- FeatureFlatSignedOffset, FeatureInstCacheLineSize128,
- FeatureMTBUFInsts, FeatureFormattedMUBUFInsts,
- Feature45BitNumRecordsBufferResource,
- FeatureMaxWavesPerEU16
- ]
->;
-//===----------------------------------------------------------------------===//
-
-class FeatureSet<list<SubtargetFeature> Features_> {
- list<SubtargetFeature> Features = Features_;
-}
-
-def FeatureISAVersion6_0_0 : FeatureSet<[FeatureSouthernIslands,
- FeatureFastFMAF32,
- FeatureHalfRate64Ops,
- FeatureLDSBankCount32]>;
-
-def FeatureISAVersion6_0_1 : FeatureSet<
- [FeatureSouthernIslands,
- FeatureLDSBankCount32]>;
-
-def FeatureISAVersion6_0_2 : FeatureSet<
- [FeatureSouthernIslands,
- FeatureLDSBankCount32]>;
-
-def FeatureISAVersion7_0_0 : FeatureSet<
- [FeatureSeaIslands,
- FeatureLDSBankCount32]>;
-
-def FeatureISAVersion7_0_1 : FeatureSet<
- [FeatureSeaIslands,
- FeatureHalfRate64Ops,
- FeatureLDSBankCount32,
- FeatureFastFMAF32]>;
-
-def FeatureISAVersion7_0_2 : FeatureSet<
- [FeatureSeaIslands,
- FeatureLDSBankCount16,
- FeatureFastFMAF32]>;
-
-def FeatureISAVersion7_0_3 : FeatureSet<
- [FeatureSeaIslands,
- FeatureLDSBankCount16]>;
-
-def FeatureISAVersion7_0_4 : FeatureSet<
- [FeatureSeaIslands,
- FeatureLDSBankCount32]>;
-
-def FeatureISAVersion7_0_5 : FeatureSet<
- [FeatureSeaIslands,
- FeatureLDSBankCount16]>;
-
-def FeatureISAVersion8_0_Common : FeatureSet<
- [FeatureVolcanicIslands,
- FeatureLDSBankCount32,
- FeatureUnpackedD16VMem]>;
-
-def FeatureISAVersion8_0_1 : FeatureSet<
- !listconcat(FeatureISAVersion8_0_Common.Features,
- [FeatureFastFMAF32,
- FeatureHalfRate64Ops,
- FeatureXNACKOnOffModes])>;
-
-def FeatureISAVersion8_0_2 : FeatureSet<
- !listconcat(FeatureISAVersion8_0_Common.Features,
- [FeatureSGPRInitBug])>;
-
-def FeatureISAVersion8_0_3 : FeatureSet<
- !listconcat(FeatureISAVersion8_0_Common.Features,
- [])>;
-
-def FeatureISAVersion8_0_5 : FeatureSet<
- !listconcat(FeatureISAVersion8_0_Common.Features,
- [FeatureSGPRInitBug])>;
-
-def FeatureISAVersion8_1_0 : FeatureSet<
- [FeatureVolcanicIslands,
- FeatureLDSBankCount16,
- FeatureXNACKOnOffModes,
- FeatureImageStoreD16Bug,
- FeatureImageGather4D16Bug]>;
-
-def FeatureISAVersion9_0_Common : FeatureSet<
- [FeatureGFX9,
- FeatureAddressableLocalMemorySize65536,
- FeatureLDSAllocGranularity512,
- FeatureLDSEncodingGranularity512,
- FeatureLDSBankCount32,
- FeatureImageInsts,
- FeatureMadMacF32Insts]>;
-
-def FeatureISAVersion9_0_Consumer_Common : FeatureSet<
- !listconcat(FeatureISAVersion9_0_Common.Features,
- [FeatureImageGather4D16Bug,
- FeatureDsSrc2Insts,
- FeatureExtendedImageInsts,
- FeatureGDS,
- FeatureCubeInsts,
- FeaturePopsExitingWaveID,
- FeatureMaxWavesPerEU10])>;
-
-def FeatureISAVersion9_Generic : FeatureSet<
- !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,
- [FeatureRequiresCOV6])>;
-
-def FeatureISAVersion9_0_MI_Common : FeatureSet<
- !listconcat(FeatureISAVersion9_0_Common.Features,
- [FeatureCubeInsts,
- FeatureFmaMixInsts,
- FeatureDLInsts,
- FeatureDot1Insts,
- FeatureDot2Insts,
- FeatureDot3Insts,
- FeatureDot4Insts,
- FeatureDot5Insts,
- FeatureDot6Insts,
- FeatureDot7Insts,
- FeatureDot10Insts,
- FeatureMAIInsts,
- FeaturePkFmacF16Inst,
- FeatureAtomicFaddNoRtnInsts,
- FeatureSRAMECCOnOffModes])>;
-
-def FeatureISAVersion9_0_0 : FeatureSet<
- !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,
- [FeatureMadMixInsts])>;
-
-def FeatureISAVersion9_0_2 : FeatureSet<
- !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,
- [FeatureMadMixInsts])>;
-
-def FeatureISAVersion9_0_4 : FeatureSet<
- !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,
- [FeatureFmaMixInsts])>;
-
-def FeatureISAVersion9_0_6 : FeatureSet<
- !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,
- [FeatureHalfRate64Ops,
- FeatureFmaMixInsts,
- FeatureDLInsts,
- FeatureDot1Insts,
- FeatureDot2Insts,
- FeatureDot7Insts,
- FeatureDot10Insts,
- FeatureSRAMECCOnOffModes])>;
-
-def FeatureISAVersion9_0_8 : FeatureSet<
- !listconcat(FeatureISAVersion9_0_MI_Common.Features,
- [FeatureGDS,
- FeatureHalfRate64Ops,
- FeatureDsSrc2Insts,
- FeatureExtendedImageInsts,
- FeatureAtomicBufferGlobalPkAddF16NoRtnInsts,
- FeatureMFMAInlineLiteralBug,
- FeatureImageGather4D16Bug,
- FeatureMaxWavesPerEU10])>;
-
-def FeatureISAVersion9_0_9 : FeatureSet<
- !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,
- [FeatureMadMixInsts])>;
-
-def FeatureISAVersion9_0_A : FeatureSet<
- !listconcat(FeatureISAVersion9_0_MI_Common.Features,
- [FeatureGFX90AInsts,
- FeatureAGPRAlloc,
- FeatureTgSplitSupport,
- FeatureRequiresAlignedVGPRs,
- FeatureFmacF64Inst,
- FeatureDPALU_DPP,
- FeaturePackedFP32Ops,
- FeatureAtomicFaddRtnInsts,
- FeatureAtomicBufferGlobalPkAddF16Insts,
- FeaturePackedTID,
- FeatureFullRate64Ops,
- FeatureBackOffBarrier,
- FeatureKernargPreload,
- FeatureAtomicFMinFMaxF64GlobalInsts,
- FeatureAtomicFMinFMaxF64FlatInsts,
- FeatureFlatBufferGlobalAtomicFaddF64Inst,
- FeatureLdsAtomicAddF64,
- FeatureMaxWavesPerEU8
- ])>;
-
-def FeatureISAVersion9_0_C : FeatureSet<
- !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,
- [FeatureMadMixInsts])>;
-
-def FeatureISAVersion9_4_Common : FeatureSet<
- [FeatureGFX9,
- FeatureGFX90AInsts,
- FeatureAGPRAlloc,
- FeatureTgSplitSupport,
- FeatureGFX940Insts,
- FeatureBufferInvInst,
- FeatureRequiresAlignedVGPRs,
- FeatureFmaMixInsts,
- FeatureDLInsts,
- FeatureFmacF64Inst,
- FeatureDot1Insts,
- FeatureDot2Insts,
- FeatureDot3Insts,
- FeatureDot4Insts,
- FeatureDot5Insts,
- FeatureDot6Insts,
- FeatureDot7Insts,
- FeatureDot10Insts,
- FeatureAtomicDsPkAdd16Insts,
- FeatureAtomicFlatPkAdd16Insts,
- FeatureDPALU_DPP,
- FeaturePackedFP32Ops,
- FeatureMAIInsts,
- FeaturePkFmacF16Inst,
- FeatureAtomicFaddRtnInsts,
- FeatureAtomicFaddNoRtnInsts,
- FeatureAtomicBufferGlobalPkAddF16Insts,
- FeatureAtomicGlobalPkAddBF16Inst,
- FeatureFlatAtomicFaddF32Inst,
- FeatureSRAMECCOnOffModes,
- FeaturePackedTID,
- FeatureArchitectedFlatScratch,
- FeatureFullRate64Ops,
- FeatureBackOffBarrier,
- FeatureKernargPreload,
- FeatureAtomicFMinFMaxF64GlobalInsts,
- FeatureAtomicFMinFMaxF64FlatInsts,
- FeatureAgentScopeFineGrainedRemoteMemoryAtomics,
- FeatureMemoryAtomicFAddF32DenormalSupport,
- FeatureFlatBufferGlobalAtomicFaddF64Inst,
- FeatureLdsAtomicAddF64,
- FeatureLshlAddU64Inst,
- FeatureVMovB64Inst,
- FeatureMaxWavesPerEU8
- ]>;
-
-def FeatureISAVersion9_5_Common : FeatureSet<
- !listconcat(FeatureISAVersion9_4_Common.Features,
- [FeatureAddressableLocalMemorySize163840,
- FeatureLDSAllocGranularity1280,
- FeatureLDSEncodingGranularity1280,
- FeatureLDSBankCount64,
- FeatureFP8Insts,
- FeatureFP8ConversionInsts,
- FeatureCvtFP8SDWASrcSel,
- FeatureOCPFP8ConversionInsts,
- FeatureGFX950Insts,
- FeaturePrngInst,
- FeatureBF16ConversionInsts,
- FeatureBitOp3Insts,
- FeatureFP8ConversionScaleInsts,
- FeatureBF8ConversionScaleInsts,
- FeatureFP4ConversionScaleInsts,
- FeatureFP6BF6ConversionScaleInsts,
- FeatureDot12Insts,
- FeatureDot13Insts,
- FeatureAtomicBufferPkAddBF16Inst
- ])>;
-
-def FeatureISAVersion9_4_2 : FeatureSet<
- !listconcat(FeatureISAVersion9_4_Common.Features,
- [
- FeatureAddressableLocalMemorySize65536,
- FeatureLDSAllocGranularity512,
- FeatureLDSEncodingGranularity512,
- FeatureLDSBankCount32,
- FeatureFP8Insts,
- FeatureFP8ConversionInsts,
- FeatureCvtFP8SDWASrcSel,
- FeatureCvtFP8VOP1Bug,
- FeatureXF32Insts
- ])>;
-
-def FeatureISAVersion9_4_Generic : FeatureSet<
- !listconcat(FeatureISAVersion9_4_Common.Features,
- [FeatureAddressableLocalMemorySize65536,
- FeatureLDSAllocGranularity1280,
- FeatureLDSEncodingGranularity1280,
- FeatureLDSBankCount32,
- FeatureRequiresCOV6])>;
-
-def FeatureISAVersion9_5_0 : FeatureSet<FeatureISAVersion9_5_Common.Features>;
-
-def FeatureISAVersion10_Common : FeatureSet<
- [FeatureGFX10,
- FeatureLDSBankCount32,
- FeatureDLInsts,
- FeatureNSAEncoding,
- FeatureBackOffBarrier,
- FeatureMsadInsts,
- FeatureMqsadPkInsts,
- FeatureMqsadInsts,
- FeaturePopsExitingWaveID]>;
-
-def FeatureISAVersion10_1_Common : FeatureSet<
- !listconcat(FeatureISAVersion10_Common.Features,
- [FeatureLDSAllocGranularity512,
- FeatureScalarStores,
- FeatureScalarAtomics,
- FeatureScalarFlatScratchInsts,
- FeatureGetWaveIdInst,
- FeatureMadMacF32Insts,
- FeatureDsSrc2Insts,
- FeatureLDSMisalignedBug,
- FeatureXNACKOnOffModes,
- // gfx101x bugs
- FeatureVcmpxPermlaneHazard,
- FeatureVMEMtoScalarWriteHazard,
- FeatureSMEMtoVectorWriteHazard,
- FeatureInstFwdPrefetchBug,
- FeatureVcmpxExecWARHazard,
- FeatureLdsBranchVmemWARHazard,
- FeatureNSAtoVMEMBug,
- FeatureNSAClauseBug,
- FeatureOffset3fBug,
- FeatureFlatSegmentOffsetBug,
- FeatureNegativeUnalignedScratchOffsetBug,
- FeatureMaxWavesPerEU20])>;
-
-def FeatureISAVersion10_1_Generic : FeatureSet<
- !listconcat(FeatureISAVersion10_1_Common.Features,
- [FeatureRequiresCOV6])>;
-
-def FeatureISAVersion10_1_0 : FeatureSet<
- !listconcat(FeatureISAVersion10_1_Common.Features,
- [])>;
-
-def FeatureISAVersion10_1_1 : FeatureSet<
- !listconcat(FeatureISAVersion10_1_Common.Features,
- [FeatureDot1Insts,
- FeatureDot2Insts,
- FeatureDot5Insts,
- FeatureDot6Insts,
- FeatureDot7Insts,
- FeatureDot10Insts])>;
-
-def FeatureISAVersion10_1_2 : FeatureSet<
- !listconcat(FeatureISAVersion10_1_Common.Features,
- [FeatureDot1Insts,
- FeatureDot2Insts,
- FeatureDot5Insts,
- FeatureDot6Insts,
- FeatureDot7Insts,
- FeatureDot10Insts])>;
-
-def FeatureISAVersion10_1_3 : FeatureSet<
- !listconcat(FeatureISAVersion10_1_Common.Features,
- [FeatureBVHRayTracingInsts,
- FeatureMSAALoadInsts])>;
-
-def FeatureISAVersion10_3_0 : FeatureSet<
- !listconcat(FeatureISAVersion10_Common.Features,
- [FeatureLDSAllocGranularity1024,
- FeatureBVHRayTracingInsts,
- FeatureMSAALoadInsts,
- FeatureGFX10_BEncoding,
- FeatureGFX10_3Insts,
- FeatureDot1Insts,
- FeatureDot2Insts,
- FeatureDot5Insts,
- FeatureDot6Insts,
- FeatureDot7Insts,
- FeatureDot10Insts,
- FeatureShaderCyclesRegister,
- FeatureMaxWavesPerEU16])>;
-
-def FeatureISAVersion10_3_Generic: FeatureSet<
- !listconcat(FeatureISAVersion10_3_0.Features,
- [FeatureRequiresCOV6])>;
-
-def FeatureISAVersion11_Common : FeatureSet<
- [FeatureGFX11,
- FeatureBackOffBarrier,
- FeatureLDSBankCount32,
- FeatureDLInsts,
- FeatureDot7Insts,
- FeatureDot8Insts,
- FeatureDot9Insts,
- FeatureDot10Insts,
- FeatureDot12Insts,
- FeatureNSAEncoding,
- FeaturePartialNSAEncoding,
- FeatureShaderCyclesRegister,
- FeatureArchitectedFlatScratch,
- FeatureAtomicFaddRtnInsts,
- FeatureAtomicFaddNoRtnInsts,
- FeatureFlatAtomicFaddF32Inst,
- FeatureImageInsts,
- FeaturePackedTID,
- FeatureVcmpxPermlaneHazard,
- FeatureMemoryAtomicFAddF32DenormalSupport,
- FeatureRealTrue16Insts,
- FeatureD16Writes32BitVgpr,
- FeatureMsadInsts,
- FeatureMqsadPkInsts,
- FeatureMqsadInsts,
-]>;
-
-// There are few workarounds that need to be added to all targets. This
-// pessimizes codegen a bit on the generic GFX11 target. This generic target
-// does not include GFX11_7 due to incompatible changes.
-def FeatureISAVersion11_Generic: FeatureSet<
- !listconcat(FeatureISAVersion11_Common.Features,
- [FeatureMSAALoadDstSelBug,
- FeatureVALUTransUseHazard,
- FeatureVOPDInterlockHazard,
- FeatureUserSGPRInit16Bug,
- FeatureMADIntraFwdBug,
- FeaturePrivEnabledTrap2NopBug,
- FeatureRequiresCOV6,
- FeatureRequiredExportPriority,
- FeatureDot5Insts,
- FeatureWMMA256bInsts,
- FeatureDX10ClampAndIEEEMode])>;
-
-def FeatureISAVersion11_0_Common : FeatureSet<
- !listconcat(FeatureISAVersion11_Common.Features,
- [FeatureMSAALoadDstSelBug,
- FeatureVALUTransUseHazard,
- FeatureVOPDInterlockHazard,
- FeatureMADIntraFwdBug,
- FeaturePrivEnabledTrap2NopBug,
- FeatureDot5Insts,
- FeatureWMMA256bInsts,
- FeatureDX10ClampAndIEEEMode])>;
-
-def FeatureISAVersion11_0_0 : FeatureSet<
- !listconcat(FeatureISAVersion11_0_Common.Features,
- [Feature1536VGPRs,
- FeatureUserSGPRInit16Bug])>;
-
-def FeatureISAVersion11_0_1 : FeatureSet<
- !listconcat(FeatureISAVersion11_0_Common.Features,
- [Feature1536VGPRs])>;
-
-def FeatureISAVersion11_0_2 : FeatureSet<
- !listconcat(FeatureISAVersion11_0_Common.Features,
- [FeatureUserSGPRInit16Bug])>;
-
-def FeatureISAVersion11_0_3 : FeatureSet<
- !listconcat(FeatureISAVersion11_0_Common.Features,
- [])>;
-
-def FeatureISAVersion11_5_Common : FeatureSet<
- !listconcat(FeatureISAVersion11_Common.Features,
- [FeatureVOPDInterlockHazard,
- FeatureSALUFloatInsts,
- FeatureDPPSrc1SGPR,
- FeatureRequiredExportPriority,
- FeatureDot5Insts,
- FeatureWMMA256bInsts,
- FeatureDX10ClampAndIEEEMode])>;
-
-def FeatureISAVersion11_5_0 : FeatureSet<
- !listconcat(FeatureISAVersion11_5_Common.Features,
- [FeaturePointSampleAccel])>;
-
-def FeatureISAVersion11_5_1 : FeatureSet<
- !listconcat(FeatureISAVersion11_5_Common.Features,
- [Feature1536VGPRs,
- FeaturePointSampleAccel])>;
-
-def FeatureISAVersion11_5_2 : FeatureSet<
- !listconcat(FeatureISAVersion11_5_Common.Features,
- [FeaturePointSampleAccel])>;
-
-def FeatureISAVersion11_7_Common : FeatureSet<
- !listconcat(FeatureISAVersion11_Common.Features,
- [FeatureGFX11_7Insts,
- FeatureSALUFloatInsts,
- FeatureDPPSrc1SGPR,
- FeatureFP8ConversionInsts,
- FeatureCvtFP8ByteSel,
- FeatureOCPFP8ConversionInsts,
- FeatureDot11Insts,
- FeatureWMMA128bInsts,
- FeatureSWMMACGfx1200Insts,
- FeatureIEEEMinimumMaximumInsts,
- FeatureMinimum3Maximum3F32,
- FeatureMinimum3Maximum3F16,
- FeatureNoF16PseudoScalarTransInlineConstants])>;
-
-def FeatureISAVersion11_7_Generic: FeatureSet<
- !listconcat(FeatureISAVersion11_7_Common.Features,
- [FeatureRequiresCOV6])>;
-
-def FeatureISAVersion12 : FeatureSet<
- [FeatureGFX12,
- Feature32BitNumRecordsBufferResource,
- FeatureSupportsWave64, FeatureSupportsWGP,
- FeatureBackOffBarrier,
- FeatureAddressableLocalMemorySize65536,
- FeatureLDSAllocGranularity1024,
- FeatureLDSEncodingGranularity512,
- FeatureHalfAddressablePhysicalLocalMemory,
- FeatureLDSBankCount32,
- FeatureDLInsts,
- FeatureDot7Insts,
- FeatureDot8Insts,
- FeatureDot9Insts,
- FeatureDot10Insts,
- FeatureDot11Insts,
- FeatureDot12Insts,
- FeatureNSAEncoding,
- FeaturePartialNSAEncoding,
- FeatureShaderCyclesHiLoRegisters,
- FeatureArchitectedFlatScratch,
- FeatureArchitectedSGPRs,
- FeatureAtomicFaddRtnInsts,
- FeatureAtomicFaddNoRtnInsts,
- FeatureAtomicDsPkAdd16Insts,
- FeatureAtomicFlatPkAdd16Insts,
- FeatureAtomicBufferGlobalPkAddF16Insts,
- FeatureAtomicGlobalPkAddBF16Inst,
- FeatureAtomicBufferPkAddBF16Inst,
- FeatureFlatAtomicFaddF32Inst,
- FeatureImageInsts,
- FeatureExtendedImageInsts,
- FeatureBVHRayTracingInsts,
- FeatureMSAALoadInsts,
- FeatureFP8ConversionInsts,
- FeatureCvtFP8ByteSel,
- FeatureOCPFP8ConversionInsts,
- FeatureWMMA128bInsts,
- FeatureSWMMACGfx1200Insts,
- FeatureIEEEMinimumMaximumInsts,
- FeaturePackedTID,
- FeatureVcmpxPermlaneHazard,
- FeatureSALUFloatInsts,
- FeaturePseudoScalarTrans,
- FeatureRestrictedSOffset,
- FeatureScalarDwordx3Loads,
- FeatureDPPSrc1SGPR,
- FeatureMaxHardClauseLength32,
- Feature1536VGPRs,
- FeatureMemoryAtomicFAddF32DenormalSupport,
- FeatureBVHDualAndBVH8Insts,
- FeatureWaitsBeforeSystemScopeStores,
- FeatureD16Writes32BitVgpr,
- FeatureCubeInsts,
- FeatureLerpInst,
- FeatureSadInsts,
- FeatureQsadInsts,
- FeatureMsadInsts,
- FeatureMqsadPkInsts,
- FeatureMqsadInsts,
- FeatureCvtNormInsts,
- FeatureCvtPkNormVOP2Insts,
- FeatureCvtPkNormVOP3Insts,
- FeatureSmemPrefetchInsts,
- FeatureNoF16PseudoScalarTransInlineConstants,
- FeatureRealTrue16Insts,
- FeatureMTBUFInsts,
- FeatureFormattedMUBUFInsts,
- FeatureDataCacheLineSize128
- ]>;
-
-def FeatureISAVersion12_50_Common : FeatureSet<
- [FeatureGFX12,
- FeatureGFX1250Insts,
- FeatureBackOffBarrier,
- FeatureRequiresAlignedVGPRs,
- FeatureCuMode,
- Feature1024AddressableVGPRs,
- Feature64BitLiterals,
- FeatureDLInsts,
- FeatureFmacF64Inst,
- FeaturePackedFP32SingleSGPROps,
- FeatureDot7Insts,
- FeatureDot8Insts,
- FeatureWavefrontSize32,
- FeatureShaderCyclesHiLoRegisters,
- FeatureArchitectedFlatScratch,
- FeatureArchitectedSGPRs,
- FeatureFlatGVSMode,
- FeatureAtomicFaddRtnInsts,
- FeatureAtomicFaddNoRtnInsts,
- FeatureAtomicDsPkAdd16Insts,
- FeatureAtomicFlatPkAdd16Insts,
- FeatureAtomicBufferGlobalPkAddF16Insts,
- FeatureAtomicGlobalPkAddBF16Inst,
- FeatureAtomicBufferPkAddBF16Inst,
- FeatureFlatAtomicFaddF32Inst,
- FeatureFP8ConversionInsts,
- FeatureCvtFP8ByteSel,
- FeatureOCPFP8ConversionInsts,
- FeatureFP8F16ConversionInsts,
- FeatureF16FP8ConversionInsts,
- FeaturePackedTID,
- FeatureVcmpxPermlaneHazard,
- FeatureSALUFloatInsts,
- FeaturePseudoScalarTrans,
- FeatureRestrictedSOffset,
- FeatureScalarDwordx3Loads,
- FeatureDPPSrc1SGPR,
- FeatureBitOp3Insts,
- FeatureTanhInsts,
- FeatureTensorCvtLutInsts,
- FeatureTDMInsts,
- FeatureVPermPk16Hazard,
- FeatureTransposeLoadF4F6Insts,
- FeatureBF16TransInsts,
- FeatureBF16ConversionInsts,
- FeatureBF16PackedInsts,
- FeatureBF16InlineConstFromUpperFP32,
- FeatureCvtPkF16F32Inst,
- FeatureCvtSrPkBF16F32Inst,
- FeatureFmaMixBF16Insts,
- FeatureMin3Max3PKF16,
- FeatureMinimum3Maximum3PKF16,
- FeaturePrngInst,
- FeaturePermlane16Swap,
- FeatureAshrPkInsts,
- FeatureSRAMECCOnOffModes,
- FeatureMaxHardClauseLength63,
- FeatureWaitXcnt,
- FeatureAsynccnt,
- FeatureAtomicFMinFMaxF64GlobalInsts,
- FeatureAtomicFMinFMaxF64FlatInsts,
- FeatureFlatBufferGlobalAtomicFaddF64Inst,
- FeatureMemoryAtomicFAddF32DenormalSupport,
- FeatureEmulatedSystemScopeAtomics,
- FeatureGloballyAddressableScratch,
- FeatureKernargPreload,
- FeatureVmemPrefInsts,
- FeatureLshlAddU64Inst,
- FeatureAddSubU64Insts,
- FeatureMadU32Inst,
- FeatureMadNC64_32Insts,
- FeatureAddMinMaxInsts,
- FeatureMinMaxI64Insts,
- FeaturePkAddMinMaxInsts,
- FeatureLdsBarrierArriveAtomic,
- FeatureLdsAtomicAddF64,
- FeatureSetPrioIncWgInst,
- FeatureSWakeupBarrier,
- Feature45BitNumRecordsBufferResource,
- FeatureClusters,
- FeatureD16Writes32BitVgpr,
- FeatureMcastLoadInsts,
- FeatureAsyncLoadToLDSInsts,
- FeatureAsyncStoreFromLDSInsts,
- FeatureSmemPrefetchInsts,
- FeatureNoF16PseudoScalarTransInlineConstants,
- FeatureRealTrue16Insts,
- FeatureVMovB64Inst,
- FeatureVMulU64Inst,
-]>;
-
-def FeatureISAVersion12_50_STRICT : FeatureSet<
- !listconcat(FeatureISAVersion12_50_Common.Features,
- [FeatureVOPD3F64OPYSrcHazard,
- FeatureGFX1250_STRICT,
- FeatureAddressableLocalMemorySize327680,
- FeatureLDSAllocGranularity2048,
- FeatureLDSEncodingGranularity2048,
- FeatureLDSBankCount64,
- FeatureWMMAN16Insts,
- FeatureVOP3PX2IncrementsVaVdstTwice,
- FeatureSetregVGPRMSBFixup,
- FeatureCubeInsts,
- FeatureLerpInst,
- FeatureSadInsts,
- FeatureQsadInsts,
- FeatureMsadInsts,
- FeatureMqsadPkInsts,
- FeatureMqsadInsts,
- FeatureCvtNormInsts,
- FeatureCvtPkNormVOP2Insts,
- FeatureCvtPkNormVOP3Insts,
- FeatureSWMMACGfx1200Insts,
- FeatureSWMMACGfx1250Insts,
- FeatureTransCoexecutionHazard,
- FeatureGFX1250VALUBlockingCycles,
- FeatureWMMACoexecutionHazards,
- FeatureDataCacheLineSize128,
- FeatureRequiresInitialUnclausedVmem,
- FeatureNeedsTDMDrain,
- FeatureNeedsAligned2addrDS,
- FeatureNoSleepForever,
- FeatureSlowMaxMinMulI64Insts,
- FeatureSupportsXNACK,
- ])>;
-
-def FeatureISAVersion12_50 : FeatureSet<
- !listconcat(FeatureISAVersion12_50_Common.Features,
- [FeatureVOPD3F64OPYSrcHazard,
- FeatureAddressableLocalMemorySize327680,
- FeatureLDSAllocGranularity2048,
- FeatureLDSEncodingGranularity2048,
- FeatureLDSBankCount64,
- FeatureWMMAN16Insts,
- FeatureWMMAF4Insts,
- FeatureWMMA16x16x128FP8Insts,
- FeatureVOP3PX2IncrementsVaVdstTwice,
- FeatureSetregVGPRMSBFixup,
- FeatureCubeInsts,
- FeatureLerpInst,
- FeatureSadInsts,
- FeatureQsadInsts,
- FeatureMsadInsts,
- FeatureMqsadPkInsts,
- FeatureMqsadInsts,
- FeatureCvtNormInsts,
- FeatureCvtPkNormVOP2Insts,
- FeatureCvtPkNormVOP3Insts,
- FeatureBlock16ConversionScaleInsts,
- FeatureFP8E5M3Insts,
- FeatureSWMMACGfx1200Insts,
- FeatureSWMMACGfx1250Insts,
- FeatureTransCoexecutionHazard,
- FeatureGFX1250VALUBlockingCycles,
- FeatureWMMACoexecutionHazards,
- FeatureDataCacheLineSize128,
- FeatureRequiresInitialUnclausedVmem,
- FeatureNeedsTDMDrain,
- FeatureNoSleepForever,
- FeatureSlowMaxMinMulI64Insts,
- FeatureSupportsXNACK,
- ])>;
-
-def FeatureISAVersion12_51 : FeatureSet<
- !listconcat(FeatureISAVersion12_50_Common.Features,
- [FeatureVOPD3F64OPYSrcHazard,
- FeatureAddressableLocalMemorySize327680,
- FeatureLDSAllocGranularity2048,
- FeatureLDSEncodingGranularity2048,
- FeatureLDSBankCount64,
- FeatureWMMAN16Insts,
- FeatureWMMAF4Insts,
- FeatureWMMA16x16x128FP8Insts,
- FeatureFullRate64Ops,
- FeatureVOP3PX2IncrementsVaVdstTwice,
- FeatureDPALU_DPP,
- FeaturePackedFP64SingleSGPROps,
- FeaturePackedU64SingleSGPROps,
- FeatureGFX1251GEMMInsts,
- FeatureGFX125xLowestRateWMMA,
- FeatureCubeInsts,
- FeatureLerpInst,
- FeatureSadInsts,
- FeatureQsadInsts,
- FeatureMsadInsts,
- FeatureMqsadPkInsts,
- FeatureMqsadInsts,
- FeatureCvtNormInsts,
- FeatureCvtPkNormVOP2Insts,
- FeatureCvtPkNormVOP3Insts,
- FeatureBlock16ConversionScaleInsts,
- FeatureFP8E5M3Insts,
- FeatureSWMMACGfx1200Insts,
- FeatureSWMMACGfx1250Insts,
- FeatureUseAddPC64Inst,
- FeatureTransCoexecutionHazard,
- FeatureWMMACoexecutionHazards,
- FeatureDataCacheLineSize128,
- FeatureRequiresInitialUnclausedVmem,
- FeatureNeedsTDMDrain,
- FeatureSupportsXNACK,
- ])>;
-
-def FeatureISAVersion12_Generic: FeatureSet<
- !listconcat(FeatureISAVersion12.Features,
- [FeatureRequiresCOV6])>;
-
-def FeatureISAVersion12_5_Generic: FeatureSet<
- !listconcat(FeatureISAVersion12_50_Common.Features,
- [FeatureVOPD3F64OPYSrcHazard,
- FeatureAddressableLocalMemorySize327680,
- FeatureLDSAllocGranularity2048,
- FeatureLDSEncodingGranularity2048,
- FeatureLDSBankCount64,
- FeatureBlock16ConversionScaleInsts,
- FeatureFP8E5M3Insts,
- FeatureSetregVGPRMSBFixup,
- FeatureRequiresCOV6,
- FeatureGFX125xLowestRateWMMA,
- FeatureTransCoexecutionHazard,
- FeatureGFX1250VALUBlockingCycles,
- FeatureWMMACoexecutionHazards,
- FeatureDataCacheLineSize128,
- FeatureRequiresInitialUnclausedVmem,
- FeatureNeedsTDMDrain,
- FeatureSupportsXNACK,
- ])>;
-
-def FeatureISAVersion13 : FeatureSet<
- [FeatureGFX13,
- FeatureGFX1250Insts,
- FeatureWaveMatchInsts,
- FeatureExclusiveScanInsts,
- FeatureAddressableLocalMemorySize196608,
- FeatureLDSAllocGranularity1024,
- FeatureLDSEncodingGranularity1024,
- Feature64BitLiterals,
- FeatureLDSBankCount32,
- FeatureDLInsts,
- FeatureFmacF64Inst,
- FeatureDot7Insts,
- FeatureDot8Insts,
- FeatureNSAEncoding,
- FeaturePartialNSAEncoding,
- FeatureShaderCyclesRegister,
- FeatureArchitectedFlatScratch,
- FeatureArchitectedSGPRs,
- FeatureFlatGVSMode,
- FeatureAtomicFaddRtnInsts,
- FeatureAtomicFaddNoRtnInsts,
- FeatureAtomicDsPkAdd16Insts,
- FeatureAtomicFlatPkAdd16Insts,
- FeatureAtomicBufferGlobalPkAddF16Insts,
- FeatureAtomicGlobalPkAddBF16Inst,
- FeatureAtomicBufferPkAddBF16Inst,
- FeatureFlatAtomicFaddF32Inst,
- FeatureFP8ConversionInsts,
- FeatureCvtFP8ByteSel,
- FeatureOCPFP8ConversionInsts,
- FeatureFP8F16ConversionInsts,
- FeatureF16FP8ConversionInsts,
- FeaturePackedTID,
- FeatureVcmpxPermlaneHazard,
- FeatureSALUFloatInsts,
- FeaturePseudoScalarTrans,
- FeatureRestrictedSOffset,
- FeatureScalarDwordx3Loads,
- FeatureDPPSrc1SGPR,
- FeatureBitOp3Insts,
- FeatureTanhInsts,
- FeatureTensorCvtLutInsts,
- FeatureBF16TransInsts,
- FeatureBF16ConversionInsts,
- FeatureBF16PackedInsts,
- FeatureBF16InlineConstFromUpperFP32,
- FeaturePrngInst,
- FeaturePermlane16Swap,
- FeatureAshrPkInsts,
- FeatureAtomicFMinFMaxF64GlobalInsts,
- FeatureAtomicFMinFMaxF64FlatInsts,
- FeatureFmaMixBF16Insts,
- FeatureGloballyAddressableScratch,
- FeatureCvtPkF16F32Inst,
- FeatureCvtSrPkBF16F32Inst,
- FeatureF16BF16ToFP6BF6ConversionScaleInsts,
- FeatureF32ToFP6BF6ConversionScaleInsts,
- FeatureFP6BF6ToF16BF16F32ConversionScaleInsts,
- FeatureIEEEMinimumMaximumInsts,
- FeatureSWakeupBarrier,
- FeatureClusters,
- FeatureCubeInsts,
- FeatureLerpInst,
- FeatureSadInsts,
- FeatureQsadInsts,
- FeatureMsadInsts,
- FeatureMqsadPkInsts,
- FeatureMqsadInsts,
- FeatureCvtNormInsts,
- FeatureCvtPkNormVOP2Insts,
- FeatureCvtPkNormVOP3Insts,
- FeatureBlock16ConversionScaleInsts,
- FeatureSmemPrefetchInsts,
- FeatureAsyncLoadToLDSInsts,
- FeatureIndexedResources
-]>;
-
-def FeatureISAVersion13_Generic: FeatureSet<
- !listconcat(FeatureISAVersion13.Features,
- [FeatureRequiresCOV6])>;
+include "GCNFeatures.td"
//===----------------------------------------------------------------------===//
@@ -3361,67 +591,6 @@ def WaveSizeVT : ValueTypeByHwMode<[
AlignedVGPRNoAGPRMode_Wave32], [i64, i64, i64, i32, i32]>;
-// The subset of SubtargetFeatures which should be exposed outside of
-// the backend.
-def AMDGPUFrontendVisibleFeatures {
- list<SubtargetFeature> Features = [
- Feature16BitInsts, FeatureAddMinMaxInsts, FeatureAshrPkInsts,
- FeatureAsyncLoadToLDSInsts, FeatureAsyncStoreFromLDSInsts, FeatureAsynccnt,
- FeatureAtomicBufferGlobalPkAddF16Insts, FeatureAtomicBufferPkAddBF16Inst, FeatureAtomicDsPkAdd16Insts,
- FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF64GlobalInsts, FeatureAtomicFaddNoRtnInsts,
- FeatureAtomicFaddRtnInsts, FeatureAtomicFlatPkAdd16Insts, FeatureAtomicGlobalPkAddBF16Inst,
- FeatureBF16ConversionInsts,
- FeatureBF16PackedInsts, FeatureBF16TransInsts, FeatureBF8ConversionScaleInsts,
- FeatureBVHRayTracingInsts, FeatureBitOp3Insts, FeatureBufferInvInst,
- FeatureCIInsts,
- FeatureClusters, FeatureCubeInsts, FeatureCvtPkNormVOP2Insts,
- FeatureCvtSrPkBF16F32Inst, FeatureDLInsts, FeatureDPP,
- FeatureDot10Insts, FeatureDot11Insts, FeatureDot12Insts,
- FeatureDot13Insts, FeatureDot1Insts, FeatureDot2Insts,
- FeatureDot3Insts, FeatureDot4Insts, FeatureDot5Insts,
- FeatureDot6Insts, FeatureDot7Insts, FeatureDot8Insts,
- FeatureDot9Insts, FeatureExclusiveScanInsts, FeatureExtendedImageInsts, FeatureF16BF16ToFP6BF6ConversionScaleInsts,
- FeatureF32ToF16BF16ConversionSRInsts, FeatureF32ToFP6BF6ConversionScaleInsts, FeatureFP4ConversionScaleInsts,
- FeatureFP6BF6ConversionScaleInsts, FeatureFP8ConversionInsts, FeatureFP8ConversionScaleInsts,
- FeatureBlock16ConversionScaleInsts, FeatureFP8E5M3Insts, FeatureFP8Insts,
- FeatureFP6BF6ToF16BF16F32ConversionScaleInsts,
- FeatureFlatBufferGlobalAtomicFaddF64Inst, FeatureFlatGlobalInsts,
- FeatureGFX10Insts, FeatureGFX10_3Insts, FeatureGFX11Insts,
- FeatureGFX1250Insts, FeatureGFX1251GEMMInsts, FeatureGFX12Insts,
- FeatureGFX13Insts, FeatureGFX8Insts, FeatureGFX90AInsts,
- FeatureGFX940Insts, FeatureGFX950Insts, FeatureGFX9Insts,
- FeatureGWS, FeatureImageInsts, FeatureLdsBarrierArriveAtomic,
- FeatureLerpInst, FeatureMAIInsts, FeatureMcastLoadInsts,
- FeatureMqsadInsts, FeatureMqsadPkInsts, FeatureMsadInsts,
- FeatureOCPFP8ConversionInsts, FeaturePermlane16Insts, FeaturePermlane16Swap,
- FeaturePermlane32Swap, FeaturePkAddMinMaxInsts, FeaturePrngInst,
- FeatureQsadInsts, FeatureSMemRealTime, FeatureSMemTimeInst,
- FeatureSWMMACGfx1200Insts, FeatureSWMMACGfx1250Insts, FeatureSWakeupBarrier,
- FeatureSadInsts, FeatureSetPrioIncWgInst, FeatureSmemPrefetchInsts,
- FeatureTDMInsts, FeatureTanhInsts, FeatureTensorCvtLutInsts,
- FeatureTransposeLoadF4F6Insts, FeatureVMemToLDSLoad, FeatureVmemPrefInsts,
- FeatureWaveMatchInsts, FeatureWMMA128bInsts, FeatureWMMA256bInsts,
- FeatureWMMAN16Insts, FeatureWMMAF4Insts, FeatureWMMA16x16x128FP8Insts, FeatureXF32Insts,
- FeatureWavefrontSize32, FeatureWavefrontSize64, FeatureSupportsWGP,
- FeatureSupportsWave32, FeatureFastFMAF32, FeatureFastDenormalF32,
- FeatureSupportsXNACK, FeatureSupportsSRAMECC, FeatureXNACKOnOffModes,
- FeatureSRAMECCOnOffModes,
- FeatureSGPRInitBug, FeatureApertureRegs, FeatureGetDoorbellID,
- FeatureAGPRAlloc, Feature1536VGPRs, Feature1024AddressableVGPRs,
- FeatureHalfAddressablePhysicalLocalMemory, FeatureBackOffBarrier,
- FeatureLDSAllocGranularity256,
- FeatureLDSAllocGranularity512,
- FeatureLDSAllocGranularity1024,
- FeatureLDSAllocGranularity1280,
- FeatureLDSAllocGranularity2048,
- FeatureLDSEncodingGranularity256,
- FeatureLDSEncodingGranularity512,
- FeatureLDSEncodingGranularity1024,
- FeatureLDSEncodingGranularity1280,
- FeatureLDSEncodingGranularity2048,
- ];
-}
-
// Include AMDGPU TD files
include "SISchedule.td"
include "GCNProcessors.td"
diff --git a/llvm/lib/Target/AMDGPU/GCNFeatures.td b/llvm/lib/Target/AMDGPU/GCNFeatures.td
new file mode 100644
index 0000000000000..24926023d4696
--- /dev/null
+++ b/llvm/lib/Target/AMDGPU/GCNFeatures.td
@@ -0,0 +1,2840 @@
+//===-- GCNFeatures.td - GCN Feature Definitions -----------*- tablegen -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+//===-----------------------------------------------------------------------===//
+// AMDGPU Subtarget Feature (device properties)
+//===----------------------------------------------------------------------===//
+
+// Multiclass to define a SubtargetFeature along with optional predicates.
+// Parameters:
+// - FeatureString: The feature string used in the SubtargetFeature.
+// - Description: The description of the feature.
+// - GenPredicate: If 1 (default), generates a Has#NAME predicate.
+// - GenAssemblerPredicate: If 1 (default), the predicate includes AssemblerPredicate.
+// - Deps: List of dependent SubtargetFeatures (default empty).
+//
+// Usage:
+// defm MadMixInsts : AMDGPUSubtargetFeature<"mad-mix-insts", "description">;
+// This generates:
+// - FeatureMadMixInsts : SubtargetFeature<"mad-mix-insts", "HasMadMixInsts", "true", "description">
+// - HasMadMixInsts : Predicate<"Subtarget->hasMadMixInsts()">,
+// AssemblerPredicate<(any_of FeatureMadMixInsts)>
+//
+// With GenAssemblerPredicate=0:
+// defm Foo : AMDGPUSubtargetFeature<"foo", "desc", 1, 0>;
+// This generates:
+// - FeatureFoo : SubtargetFeature<...>
+// - HasFoo : Predicate<"Subtarget->hasFoo()"> (no AssemblerPredicate)
+//
+// With dependencies:
+// defm Bar : AMDGPUSubtargetFeature<"bar", "desc", 1, 1, [FeatureFoo]>;
+// This generates:
+// - FeatureBar : SubtargetFeature<"bar", "HasBar", "true", "desc", [FeatureFoo]>
+// - HasBar : Predicate + AssemblerPredicate
+multiclass AMDGPUSubtargetFeature<string FeatureString,
+ string Description,
+ bit GenPredicate = 1,
+ bit GenAssemblerPredicate = 1,
+ list<SubtargetFeature> Deps = [],
+ InlineBehavior IB = InlineDefault> {
+ def Feature#NAME : SubtargetFeature<FeatureString,
+ "Has"#NAME,
+ "true",
+ Description,
+ Deps,
+ IB
+ >;
+
+ if GenPredicate then
+ if GenAssemblerPredicate then
+ def Has#NAME
+ : Predicate<"Subtarget->has"#NAME#"()">,
+ AssemblerPredicate<(any_of !cast<SubtargetFeature>("Feature"#NAME))>;
+ else
+ def Has#NAME : Predicate<"Subtarget->has"#NAME#"()">;
+}
+
+// Like AMDGPUSubtargetFeature, but a generic target only needs one covered GPU
+// to support the feature.
+multiclass AMDGPUGenericAnySubtargetFeature<string FeatureString,
+ string Description,
+ bit GenPredicate = 1,
+ bit GenAssemblerPredicate = 1,
+ list<SubtargetFeature> Deps = [],
+ InlineBehavior IB = InlineDefault> {
+ def Feature#NAME : AMDGPUGenericAnyFeature<FeatureString,
+ "Has"#NAME,
+ "true",
+ Description,
+ Deps,
+ IB
+ >;
+
+ if GenPredicate then
+ if GenAssemblerPredicate then
+ def Has#NAME
+ : Predicate<"Subtarget->has"#NAME#"()">,
+ AssemblerPredicate<(any_of !cast<SubtargetFeature>("Feature"#NAME))>;
+ else
+ def Has#NAME : Predicate<"Subtarget->has"#NAME#"()">;
+}
+
+defm FastFMAF32 : AMDGPUSubtargetFeature<"fast-fmaf",
+ "Assuming f32 fma is at least as fast as mul + add",
+ /*GenPredicate=*/0, /*GenAssemblerPredicate=*/1, [], InlineIgnore
+>;
+
+defm FastDenormalF32 : AMDGPUSubtargetFeature<"fast-denormal-f32",
+ "Enabling denormals does not cause f32 instructions to run at f64 rates",
+ /*GenPredicate=*/0
+>;
+
+defm MIMG_R128 : AMDGPUSubtargetFeature<"mimg-r128",
+ "Support 128-bit texture resources",
+ /*GenPredicate=*/0
+>;
+
+defm HalfRate64Ops : AMDGPUSubtargetFeature<"half-rate-64-ops",
+ "Most fp64 instructions are half rate instead of quarter",
+ /*GenPredicate=*/0, /*GenAssemblerPredicate=*/1, [], InlineIgnore
+>;
+
+defm FullRate64Ops : AMDGPUSubtargetFeature<"full-rate-64-ops",
+ "Most fp64 instructions are full rate",
+ /*GenPredicate=*/0
+>;
+
+defm FlatAddressSpace : AMDGPUSubtargetFeature<"flat-address-space",
+ "Support flat address space"
+>;
+
+defm FlatInstOffsets : AMDGPUSubtargetFeature<"flat-inst-offsets",
+ "Flat instructions have immediate offset addressing mode"
+>;
+
+defm FlatGlobalInsts : AMDGPUSubtargetFeature<"flat-global-insts",
+ "Have global_* flat memory instructions",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/1,
+ [FeatureFlatAddressSpace]
+>;
+
+defm FlatScratchInsts : AMDGPUSubtargetFeature<"flat-scratch-insts",
+ "Have scratch_* flat memory instructions",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/1,
+ [FeatureFlatAddressSpace]
+>;
+
+defm ScalarFlatScratchInsts : AMDGPUSubtargetFeature<"scalar-flat-scratch-insts",
+ "Have s_scratch_* flat memory instructions"
+>;
+
+def FeatureEnableFlatScratch : SubtargetFeature<"enable-flat-scratch",
+ "EnableFlatScratch",
+ "true",
+ "Use scratch_* flat memory instructions to access scratch"
+>;
+
+defm FlatGVSMode : AMDGPUSubtargetFeature<"flat-gvs-mode",
+ "Have GVS addressing mode with flat_* instructions",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/1,
+ [FeatureFlatAddressSpace]
+>;
+
+defm AddNoCarryInsts : AMDGPUSubtargetFeature<"add-no-carry-insts",
+ "Have VALU add/sub instructions without carry out"
+>;
+
+defm UnalignedBufferAccess : AMDGPUSubtargetFeature<"unaligned-buffer-access",
+ "Hardware supports unaligned global loads and stores"
+>;
+
+defm TrapHandler: AMDGPUSubtargetFeature<"trap-handler",
+ "Trap handler support",
+ /*GenPredicate=*/0, /*GenAssemblerPredicate=*/1, [], InlineIgnore
+>;
+
+defm UnalignedScratchAccess : AMDGPUSubtargetFeature<"unaligned-scratch-access",
+ "Support unaligned scratch loads and stores",
+ /*GenPredicate=*/1, /*GenAssemblerPredicate=*/1, [], InlineIgnore
+>;
+
+defm UnalignedDSAccess : AMDGPUSubtargetFeature<"unaligned-ds-access",
+ "Hardware supports unaligned local and region loads and stores"
+>;
+
+defm DX10ClampAndIEEEMode : AMDGPUSubtargetFeature<"dx10-clamp-and-ieee-mode",
+ "Target has DX10_CLAMP and IEEE_MODE kernel descriptor bits"
+>;
+
+defm ApertureRegs : AMDGPUSubtargetFeature<"aperture-regs",
+ "Has Memory Aperture Base and Size Registers",
+ /*GenPredicate=*/0
+>;
+
+defm GetDoorbellID : AMDGPUSubtargetFeature<"get-doorbell-id",
+ "Supports reading the doorbell ID through S_GETREG",
+ /*GenPredicate=*/0, /*GenAssemblerPredicate=*/false, [], InlineIgnore
+>;
+
+defm AGPRAlloc : AMDGPUSubtargetFeature<"agpr-alloc",
+ "Supports allocation of AGPRs",
+ /*GenPredicate=*/0, /*GenAssemblerPredicate=*/false, [], InlineIgnore
+>;
+
+defm MadMixInsts : AMDGPUSubtargetFeature<"mad-mix-insts",
+ "Has v_mad_mix_f32, v_mad_mixlo_f16, v_mad_mixhi_f16 instructions"
+>;
+
+defm FmaMixInsts : AMDGPUSubtargetFeature<"fma-mix-insts",
+ "Has v_fma_mix_f32, v_fma_mixlo_f16, v_fma_mixhi_f16 instructions"
+>;
+
+defm FmaMixBF16Insts : AMDGPUSubtargetFeature<"fma-mix-bf16-insts",
+ "Has v_fma_mix_f32_bf16, v_fma_mixlo_bf16, v_fma_mixhi_bf16 instructions"
+>;
+
+defm FmaLegacy32Insts : AMDGPUSubtargetFeature<"fma-legacy32-insts",
+ "Has single-precision legacy FMA instructions"
+>;
+
+defm IEEEMinimumMaximumInsts : AMDGPUSubtargetFeature<"ieee-minimum-maximum-insts",
+ "Has v_minimum/maximum_f16/f32/f64, v_minimummaximum/maximumminimum_f16/f32,"
+ "v_pk_minmax/maxmin_num_f16, v_min/max_num_f16/f32/f64 and"
+ "v_pk_min/max_num_f16 instructions"
+>;
+
+defm SALUMinimumMaximumInsts : AMDGPUSubtargetFeature<"salu-minimum-maximum-insts",
+ "Has s_minimum/maximum_f16/f32 instructions"
+>;
+
+defm Minimum3Maximum3F32 : AMDGPUSubtargetFeature<"minimum3-maximum3-f32",
+ "Has v_minimum3_f32 and v_maximum3_f32 instructions"
+>;
+
+defm Minimum3Maximum3F16 : AMDGPUSubtargetFeature<"minimum3-maximum3-f16",
+ "Has v_minimum3_f16 and v_maximum3_f16 instructions"
+>;
+
+defm Min3Max3PKF16 : AMDGPUSubtargetFeature<"min3-max3-pkf16",
+ "Has v_pk_min3_num_f16 and v_pk_max3_num_f16 instructions"
+>;
+
+defm Minimum3Maximum3PKF16 : AMDGPUSubtargetFeature<"minimum3-maximum3-pkf16",
+ "Has v_pk_minimum3_f16 and v_pk_maximum3_f16 instructions"
+>;
+
+def FeatureSupportsXNACK : SubtargetFeature<"xnack-support",
+ "SupportsXNACK",
+ "true",
+ "Hardware supports XNACK",
+ [], InlineIgnore
+>;
+
+defm XNACKOnOffModes : AMDGPUSubtargetFeature<"xnack-on-off-modes",
+ "Target supports XNACK on/off modes",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/0,
+ [FeatureSupportsXNACK],
+ InlineIgnore
+>;
+
+// Whether the hardware supports WGP (work-group processor) execution
+// mode, as opposed to the per-kernel CU-mode selection
+// (FeatureCuMode).
+defm SupportsWGP : AMDGPUSubtargetFeature<"supports-wgp",
+ "Hardware supports WGP mode",
+ /*GenPredicate=*/false, /*GenAssemblerPredicate=*/false,
+ [], InlineIgnore
+>;
+
+def FeatureCuMode : SubtargetFeature<"cumode",
+ "EnableCuMode",
+ "true",
+ "Enable CU wavefront execution mode"
+>;
+
+def FeaturePreciseMemory
+ : SubtargetFeature<"precise-memory", "EnablePreciseMemory",
+ "true", "Enable precise memory mode">;
+
+defm SGPRInitBug : AMDGPUSubtargetFeature<"sgpr-init-bug",
+ "VI SGPR initialization bug requiring a fixed SGPR allocation size",
+ /*GenPredicate=*/1, /*GenAssemblerPredicate=*/1, [], InlineIgnore
+>;
+
+defm UserSGPRInit16Bug : AMDGPUGenericAnySubtargetFeature<"user-sgpr-init16-bug",
+ "Bug requiring at least 16 user+system SGPRs to be enabled",
+ /*GenPredicate=*/0
+>;
+
+defm LDSMisalignedBug : AMDGPUSubtargetFeature<"lds-misaligned-bug",
+ "Some GFX10 bug with multi-dword LDS and flat access that is not naturally aligned in WGP mode",
+ /*GenPredicate=*/0
+>;
+
+// Set where a work-group can address only half of the physical LDS block:
+// gfx6 has 64 KiB physical and 32 KiB addressable, while gfx10/11/12 have
+// 128 KiB physical and 64 KiB addressable.
+defm HalfAddressablePhysicalLocalMemory : AMDGPUSubtargetFeature<"half-addressable-physical-local-memory",
+ "A work-group can address only half of the physical LDS block",
+ /*GenPredicate=*/0
+>;
+
+defm MFMAInlineLiteralBug : AMDGPUSubtargetFeature<"mfma-inline-literal-bug",
+ "MFMA cannot use inline literal as SrcC",
+ /*GenPredicate=*/0
+>;
+
+defm VcmpxPermlaneHazard : AMDGPUSubtargetFeature<"vcmpx-permlane-hazard",
+ "TODO: describe me",
+ /*GenPredicate=*/0
+>;
+
+defm VMEMtoScalarWriteHazard : AMDGPUSubtargetFeature<"vmem-to-scalar-write-hazard",
+ "VMEM instruction followed by scalar writing to EXEC mask, M0 or SGPR leads to incorrect execution.",
+ /*GenPredicate=*/0
+>;
+
+defm SMEMtoVectorWriteHazard : AMDGPUSubtargetFeature<"smem-to-vector-write-hazard",
+ "s_load_dword followed by v_cmp page faults",
+ /*GenPredicate=*/0
+>;
+
+defm InstFwdPrefetchBug : AMDGPUSubtargetFeature<"inst-fwd-prefetch-bug",
+ "S_INST_PREFETCH instruction causes shader to hang",
+ /*GenPredicate=*/0
+>;
+
+defm VmemPrefInsts : AMDGPUSubtargetFeature<"vmem-pref-insts",
+ "Has flat_prefect_b8 and global_prefetch_b8 instructions"
+>;
+
+defm SmemPrefetchInsts : AMDGPUSubtargetFeature<"smem-prefetch-insts",
+ "Has s_prefetch_inst, s_prefetch_inst_pc_rel, s_prefetch_data and s_prefetch_data_pc_rel instructions"
+>;
+
+defm SafeSmemPrefetch : AMDGPUSubtargetFeature<"safe-smem-prefetch",
+ "SMEM prefetches do not fail on illegal address",
+ /*GenPredicate=*/0
+>;
+
+defm SafeCUPrefetch : AMDGPUSubtargetFeature<"safe-cu-prefetch",
+ "VMEM CU scope prefetches do not fail on illegal address",
+ /*GenPredicate=*/0
+>;
+
+defm VcmpxExecWARHazard : AMDGPUSubtargetFeature<"vcmpx-exec-war-hazard",
+ "V_CMPX WAR hazard on EXEC (V_CMPX issue ONLY)",
+ /*GenPredicate=*/0
+>;
+
+defm LdsBranchVmemWARHazard : AMDGPUSubtargetFeature<"lds-branch-vmem-war-hazard",
+ "Switching between LDS and VMEM-tex not waiting VM_VSRC=0",
+ /*GenPredicate=*/0
+>;
+
+class FeatureMaxHardClauseLength<int size> : SubtargetFeature<
+ "max-hard-clause-length-"#size,
+ "MaxHardClauseLength",
+ !cast<string>(size),
+ "Maximum number of instructions in an explicit S_CLAUSE is "#size
+>;
+
+/// Work around a hardware bug on some chips that can be triggered
+/// under certain circumstances when clauses are longer than 32 operations.
+def FeatureMaxHardClauseLength32 : FeatureMaxHardClauseLength<32>;
+/// While the S_CLAUSE instruction permits encoding clause lengths up to 64,
+/// hardware documentation for gfx10+ indicates that 63 is the maximum
+/// permitted clause length.
+def FeatureMaxHardClauseLength63 : FeatureMaxHardClauseLength<63>;
+
+defm NSAtoVMEMBug : AMDGPUSubtargetFeature<"nsa-to-vmem-bug",
+ "MIMG-NSA followed by VMEM fail if EXEC_LO or EXEC_HI equals zero",
+ /*GenPredicate=*/0
+>;
+
+defm NSAClauseBug : AMDGPUSubtargetFeature<"nsa-clause-bug",
+ "MIMG-NSA in a hard clause has unpredictable results on GFX10.1",
+ /*GenPredicate=*/0
+>;
+
+defm FlatSegmentOffsetBug : AMDGPUSubtargetFeature<"flat-segment-offset-bug",
+ "GFX10 bug where inst_offset is ignored when flat instructions access global memory",
+ /*GenPredicate=*/0
+>;
+
+defm NegativeScratchOffsetBug : AMDGPUSubtargetFeature<"negative-scratch-offset-bug",
+ "Negative immediate offsets in scratch instructions with an SGPR offset page fault on GFX9"
+>;
+
+defm NegativeUnalignedScratchOffsetBug : AMDGPUSubtargetFeature<"negative-unaligned-scratch-offset-bug",
+ "Scratch instructions with a VGPR offset and a negative immediate offset that"
+ "is not a multiple of 4 read wrong memory on GFX10",
+ /*GenPredicate=*/0
+>;
+
+defm Offset3fBug : AMDGPUSubtargetFeature<"offset-3f-bug",
+ "Branch offset of 3f hardware bug",
+ /*GenPredicate=*/0
+>;
+
+defm ImageStoreD16Bug : AMDGPUSubtargetFeature<"image-store-d16-bug",
+ "Image Store D16 hardware bug",
+ /*GenPredicate=*/0
+>;
+
+defm ImageGather4D16Bug : AMDGPUSubtargetFeature<"image-gather4-d16-bug",
+ "Image Gather4 D16 hardware bug",
+ /*GenPredicate=*/0
+>;
+
+defm MADIntraFwdBug : AMDGPUGenericAnySubtargetFeature<"mad-intra-fwd-bug",
+ "MAD_U64/I64 intra instruction forwarding bug",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/0
+>;
+
+defm MSAALoadDstSelBug : AMDGPUGenericAnySubtargetFeature<"msaa-load-dst-sel-bug",
+ "MSAA loads not honoring dst_sel bug",
+ /*GenPredicate=*/0
+>;
+
+defm PrivEnabledTrap2NopBug : AMDGPUGenericAnySubtargetFeature<"priv-enabled-trap2-nop-bug",
+ "Hardware that runs with PRIV=1 interpreting 's_trap 2' as a nop bug",
+ /*GenPredicate=*/0
+>;
+
+// A generic target may use an LDS bank count from one of its covered GPUs.
+class SubtargetFeatureLDSBankCount <int Value> : AMDGPUGenericAnyFeature <
+ "ldsbankcount"#Value,
+ "LDSBankCount",
+ !cast<string>(Value),
+ "The number of LDS banks per compute unit."
+>;
+
+def FeatureLDSBankCount16 : SubtargetFeatureLDSBankCount<16>;
+def FeatureLDSBankCount32 : SubtargetFeatureLDSBankCount<32>;
+def FeatureLDSBankCount64 : SubtargetFeatureLDSBankCount<64>;
+
+class SubtargetFeatureMaxWavesPerEU <int Value> : SubtargetFeature <
+ "max-waves-per-eu-"#Value,
+ "MaxWavesPerEU",
+ !cast<string>(Value),
+ "Maximum number of waves an execution unit can hold."
+>;
+
+def FeatureMaxWavesPerEU8 : SubtargetFeatureMaxWavesPerEU<8>;
+def FeatureMaxWavesPerEU10 : SubtargetFeatureMaxWavesPerEU<10>;
+def FeatureMaxWavesPerEU16 : SubtargetFeatureMaxWavesPerEU<16>;
+def FeatureMaxWavesPerEU20 : SubtargetFeatureMaxWavesPerEU<20>;
+
+class SubtargetFeatureInstCacheLineSize <int Value> : SubtargetFeature <
+ "instcachelinesize"#Value,
+ "InstCacheLineSize",
+ !cast<string>(Value),
+ "Instruction cache line size in bytes."
+>;
+
+def FeatureInstCacheLineSize64 : SubtargetFeatureInstCacheLineSize<64>;
+def FeatureInstCacheLineSize128 : SubtargetFeatureInstCacheLineSize<128>;
+
+class SubtargetFeatureDataCacheLineSize <int Value> : SubtargetFeature <
+ "datacachelinesize"#Value,
+ "DataCacheLineSize",
+ !cast<string>(Value),
+ "Data (VMEM) cache line size in bytes."
+>;
+
+def FeatureDataCacheLineSize128 : SubtargetFeatureDataCacheLineSize<128>;
+
+defm GCN3Encoding : AMDGPUSubtargetFeature<"gcn3-encoding",
+ "Encoding format for VI",
+ /*GenPredicate=*/0
+>;
+
+defm CIInsts : AMDGPUSubtargetFeature<"ci-insts",
+ "Additional instructions for CI+",
+ /*GenPredicate=*/0
+>;
+
+defm GFX8Insts : AMDGPUSubtargetFeature<"gfx8-insts",
+ "Additional instructions for GFX8+",
+ /*GenPredicate=*/0
+>;
+
+defm GFX9Insts : AMDGPUSubtargetFeature<"gfx9-insts",
+ "Additional instructions for GFX9+",
+ /*GenPredicate=*/0
+>;
+
+def FeatureRequiresAlignedVGPRs : SubtargetFeature<"vgpr-align2",
+ "RequiresAlignVGPR",
+ "true",
+ "VGPR and AGPR tuple operands require even alignment"
+>;
+
+defm GFX90AInsts : AMDGPUSubtargetFeature<"gfx90a-insts",
+ "Additional instructions for GFX90A+",
+ /*GenPredicate=*/0
+>;
+
+defm TgSplitSupport : AMDGPUSubtargetFeature<"tgsplit-support",
+ "Hardware support for threadgroup split execution mode"
+>;
+
+defm GFX940Insts : AMDGPUSubtargetFeature<"gfx940-insts",
+ "Additional instructions for GFX940+",
+ /*GenPredicate=*/0
+>;
+
+defm BufferInvInst : AMDGPUSubtargetFeature<"buffer-inv-inst",
+ "Has buffer_inv instruction"
+>;
+
+defm Permlane16Insts : AMDGPUSubtargetFeature<"permlane16-insts",
+ "Has v_permlane16_b32/v_permlanex16_b32 instructions"
+>;
+
+defm Permlane16Swap : AMDGPUSubtargetFeature<"permlane16-swap",
+ "Has v_permlane16_swap_b32 instructions"
+>;
+
+defm Permlane32Swap : AMDGPUSubtargetFeature<"permlane32-swap",
+ "Has v_permlane32_swap_b32 instructions"
+>;
+
+defm FP8ConversionScaleInsts : AMDGPUSubtargetFeature<"fp8-cvt-scale-insts",
+ "Has fp8 conversion scale instructions"
+>;
+
+defm BF8ConversionScaleInsts : AMDGPUSubtargetFeature<"bf8-cvt-scale-insts",
+ "Has bf8 conversion scale instructions"
+>;
+
+defm FP4ConversionScaleInsts : AMDGPUSubtargetFeature<"fp4-cvt-scale-insts",
+ "Has fp4 conversion scale instructions"
+>;
+
+defm FP6BF6ConversionScaleInsts : AMDGPUSubtargetFeature<"fp6bf6-cvt-scale-insts",
+ "Has fp6 and bf6 conversion scale instructions"
+>;
+
+defm F16BF16ToFP6BF6ConversionScaleInsts : AMDGPUSubtargetFeature<"f16bf16-to-fp6bf6-cvt-scale-insts",
+ "Has f16bf16 to fp6bf6 conversion scale instructions"
+>;
+
+defm F32ToFP6BF6ConversionScaleInsts : AMDGPUSubtargetFeature<"f32-to-fp6bf6-cvt-scale-insts",
+ "Has f32 to fp6bf6 conversion scale instructions"
+>;
+
+defm F32ToF16BF16ConversionSRInsts : AMDGPUSubtargetFeature<"f32-to-f16bf16-cvt-sr-insts",
+ "Has f32 to f16bf16 conversion scale instructions"
+>;
+
+defm FP6BF6ToF16BF16F32ConversionScaleInsts : AMDGPUSubtargetFeature<"fp6bf6-to-f16bf16f32-cvt-scale-insts",
+ "Has fp6bf6 to f16bf16f32 conversion scale instructions"
+>;
+
+defm Block16ConversionScaleInsts : AMDGPUSubtargetFeature<"block16-cvt-scale-insts",
+ "Has block16 conversion scale instructions"
+>;
+
+defm AshrPkInsts : AMDGPUSubtargetFeature<"ashr-pk-insts",
+ "Has Arithmetic Shift Pack instructions"
+>;
+
+defm ExclusiveScanInsts : AMDGPUSubtargetFeature<"exclusive-scan-insts",
+ "Has exclusive scan instructions"
+>;
+
+defm CvtPkF16F32Inst : AMDGPUSubtargetFeature<"cvt-pk-f16-f32-inst",
+ "Has cvt_pk_f16_f32 instruction"
+>;
+
+defm CvtSrPkBF16F32Inst : AMDGPUSubtargetFeature<"cvt-sr-pk-bf16-f32-inst",
+ "Has cvt_sr_pk_bf16_f32 instruction"
+>;
+
+defm McastLoadInsts : AMDGPUSubtargetFeature<"mcast-load-insts",
+ "Has multicast load instructions"
+>;
+
+defm AsyncLoadToLDSInsts : AMDGPUSubtargetFeature<"async-load-to-lds-insts",
+ "Has global async load to LDS instructions"
+>;
+
+defm AsyncStoreFromLDSInsts : AMDGPUSubtargetFeature<"async-store-from-lds-insts",
+ "Has global async store from LDS instructions"
+>;
+
+defm SWakeupImm : AMDGPUSubtargetFeature<"s-wakeup-imm",
+ "s_wakeup takes an immediate operand"
+>;
+
+defm SBarrierLeaveImm : AMDGPUSubtargetFeature<"s-barrier-leave-imm",
+ "s_barrier_leave takes an immediate operand"
+>;
+
+defm MTBUFInsts : AMDGPUSubtargetFeature<"mtbuf-insts",
+ "Has memory typed buffer instructions."
+>;
+
+defm FormattedMUBUFInsts : AMDGPUSubtargetFeature<"formatted-mubuf-insts",
+ "Has formatted memory untyped buffer instructions."
+>;
+
+defm VMovB64Inst : AMDGPUSubtargetFeature<"v-mov-b64-inst",
+ "Has v_mov_b64 instruction"
+>;
+
+defm VMulU64Inst : AMDGPUSubtargetFeature<"v-mul-u64-inst",
+ "Has v_mul_u64 instruction."
+>;
+
+defm GFX950Insts : AMDGPUSubtargetFeature<"gfx950-insts",
+ "Additional instructions for GFX950+",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/1,
+ [FeaturePermlane16Swap,
+ FeaturePermlane32Swap,
+ FeatureAshrPkInsts,
+ FeatureFP8ConversionScaleInsts,
+ FeatureBF8ConversionScaleInsts,
+ FeatureFP4ConversionScaleInsts,
+ FeatureFP6BF6ConversionScaleInsts,
+ FeatureF16BF16ToFP6BF6ConversionScaleInsts,
+ FeatureF32ToF16BF16ConversionSRInsts,
+ FeatureCvtPkF16F32Inst,
+ FeatureMinimum3Maximum3F32,
+ FeatureMinimum3Maximum3PKF16,
+ ]
+>;
+
+defm GFX10Insts : AMDGPUSubtargetFeature<"gfx10-insts",
+ "Additional instructions for GFX10+",
+ /*GenPredicate=*/0
+>;
+
+defm GFX11Insts : AMDGPUSubtargetFeature<"gfx11-insts",
+ "Additional instructions for GFX11+",
+ /*GenPredicate=*/0
+>;
+
+defm GFX12Insts : AMDGPUSubtargetFeature<"gfx12-insts",
+ "Additional instructions for GFX12+",
+ /*GenPredicate=*/0
+>;
+
+defm GFX1250Insts : AMDGPUSubtargetFeature<"gfx1250-insts",
+ "Additional instructions for GFX1250+",
+ /*GenPredicate=*/0
+>;
+
+defm GFX13Insts : AMDGPUSubtargetFeature<"gfx13-insts",
+ "Additional instructions for GFX13+",
+ /*GenPredicate=*/0,
+ /*GenAssemblerPredicate=*/0,
+ [FeatureSWakeupImm,
+ FeatureSBarrierLeaveImm,
+ ]
+>;
+
+defm GFX10_3Insts : AMDGPUSubtargetFeature<"gfx10-3-insts",
+ "Additional instructions for GFX10.3",
+ /*GenPredicate=*/0,
+ /*GenAssemblerPredicate=*/1,
+ [FeatureFmaLegacy32Insts]
+>;
+
+defm GFX11_7Insts : AMDGPUSubtargetFeature<"gfx11-7-insts",
+ "Additional instructions for GFX11.7",
+ /*GenPredicate=*/0
+>;
+
+defm GFX7GFX8GFX9Insts : AMDGPUSubtargetFeature<"gfx7-gfx8-gfx9-insts",
+ "Instructions shared in GFX7, GFX8, GFX9",
+ /*GenPredicate=*/0
+>;
+
+defm SMemRealTime : AMDGPUSubtargetFeature<"s-memrealtime",
+ "Has s_memrealtime instruction"
+>;
+
+defm Inv2PiInlineImm : AMDGPUSubtargetFeature<"inv-2pi-inline-imm",
+ "Has 1 / (2 * pi) as inline immediate",
+ /*GenPredicate=*/0
+>;
+
+defm 16BitInsts : AMDGPUSubtargetFeature<"16-bit-insts",
+ "Has i16/f16 instructions"
+>;
+
+defm True16BitInsts : AMDGPUSubtargetFeature<"true16",
+ "True 16-bit operand instructions"
+>;
+
+def FeatureRealTrue16Insts : SubtargetFeature<"real-true16",
+ "EnableRealTrue16Insts",
+ "true",
+ "Use true 16-bit registers"
+>;
+
+defm D16Writes32BitVgpr : AMDGPUSubtargetFeature<"d16-write-vgpr32",
+ "D16 instructions potentially have 32-bit data dependencies"
+>;
+
+defm BF16TransInsts : AMDGPUSubtargetFeature<"bf16-trans-insts",
+ "Has bf16 transcendental instructions"
+>;
+
+defm BF16ConversionInsts : AMDGPUSubtargetFeature<"bf16-cvt-insts",
+ "Has bf16 conversion instructions"
+>;
+
+defm BF16PackedInsts : AMDGPUSubtargetFeature<"bf16-pk-insts",
+ "Has bf16 packed instructions (fma, add, mul, max, min)"
+>;
+
+defm BF16InlineConstFromUpperFP32 : AMDGPUSubtargetFeature<"bf16-from-upper-fp32",
+ "bf16 inline constants are actually selected from upper 16-bits of a fp32"
+>;
+
+defm VOP3PInsts : AMDGPUSubtargetFeature<"vop3p",
+ "Has VOP3P packed instructions"
+>;
+
+defm Movrel : AMDGPUSubtargetFeature<"movrel",
+ "Has v_movrel*_b32 instructions"
+>;
+
+defm VGPRIndexMode : AMDGPUSubtargetFeature<"vgpr-index-mode",
+ "Has VGPR mode register indexing"
+>;
+
+defm ScalarDwordx3Loads : AMDGPUSubtargetFeature<"scalar-dwordx3-loads",
+ "Has 96-bit scalar load instructions"
+>;
+
+defm ScalarStores : AMDGPUSubtargetFeature<"scalar-stores",
+ "Has store scalar memory instructions"
+>;
+
+defm ScalarAtomics : AMDGPUSubtargetFeature<"scalar-atomics",
+ "Has atomic scalar memory instructions"
+>;
+
+defm SDWA : AMDGPUSubtargetFeature<"sdwa",
+ "Support SDWA (Sub-DWORD Addressing) extension",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/0
+>;
+
+defm SDWAOmod : AMDGPUSubtargetFeature<"sdwa-omod",
+ "Support OMod with SDWA (Sub-DWORD Addressing) extension",
+ /*GenPredicate=*/0
+>;
+
+defm SDWAScalar : AMDGPUSubtargetFeature<"sdwa-scalar",
+ "Support scalar register with SDWA (Sub-DWORD Addressing) extension",
+ /*GenPredicate=*/0
+>;
+
+defm SDWASdst : AMDGPUSubtargetFeature<"sdwa-sdst",
+ "Support scalar dst for VOPC with SDWA (Sub-DWORD Addressing) extension",
+ /*GenPredicate=*/0
+>;
+
+defm SDWAMac : AMDGPUSubtargetFeature<"sdwa-mav",
+ "Support v_mac_f32/f16 with SDWA (Sub-DWORD Addressing) extension",
+ /*GenPredicate=*/0
+>;
+
+defm SDWAOutModsVOPC : AMDGPUSubtargetFeature<"sdwa-out-mods-vopc",
+ "Support clamp for VOPC with SDWA (Sub-DWORD Addressing) extension",
+ /*GenPredicate=*/0
+>;
+
+defm DPP : AMDGPUSubtargetFeature<"dpp",
+ "Support DPP (Data Parallel Primitives) extension",
+ /*GenPredicate=*/0
+>;
+
+// DPP8 allows arbitrary cross-lane swizzling within groups of 8 lanes.
+def FeatureDPP8 : SubtargetFeature<"dpp8",
+ "HasDPP8",
+ "true",
+ "Support DPP8 (Data Parallel Primitives) extension"
+>;
+
+defm DPALU_DPP : AMDGPUSubtargetFeature<"dpp-64bit",
+ "Support DPP (Data Parallel Primitives) extension in DP ALU"
+>;
+
+defm DPPSrc1SGPR : AMDGPUSubtargetFeature<"dpp-src1-sgpr",
+ "Support SGPR for Src1 of DPP instructions",
+ /*GenPredicate=*/0
+>;
+
+defm DPPWavefrontShifts : AMDGPUSubtargetFeature<"dpp-wavefront-shifts",
+ "Support DPP wave_shl/wave_rol/wave_shr/wave_ror"
+>;
+
+defm DPPBroadcasts : AMDGPUSubtargetFeature<"dpp-row-bcast",
+ "Support DPP row_bcast15/row_bcast31"
+>;
+
+defm PackedFP32Ops : AMDGPUSubtargetFeature<"packed-fp32-ops",
+ "Support packed fp32 instructions"
+>;
+
+defm PackedFP32SingleSGPROps : AMDGPUSubtargetFeature<"packed-fp32-single-sgpr-ops",
+ "Support packed fp32 instructions with SGPR operands propagating low to high 32 bit"
+>;
+
+defm PackedFP64SingleSGPROps : AMDGPUSubtargetFeature<"packed-fp64-single-sgpr-ops",
+ "Support packed fp64 instructions with SGPR operands propagating low to high 64 bit"
+>;
+
+defm PackedU64SingleSGPROps : AMDGPUSubtargetFeature<"packed-u64-single-sgpr-ops",
+ "Support packed uint64 instructions with SGPR operands propagating low to high 64 bit"
+>;
+
+defm R128A16 : AMDGPUSubtargetFeature<"r128-a16",
+ "Support gfx9-style A16 for 16-bit coordinates/gradients/lod/clamp/mip image "
+ "operands, where a16 is aliased with r128"
+>;
+
+defm A16 : AMDGPUSubtargetFeature<"a16",
+ "Support A16 for 16-bit coordinates/gradients/lod/clamp/mip image operands"
+>;
+
+defm G16 : AMDGPUSubtargetFeature<"g16",
+ "Support G16 for 16-bit gradient image operands"
+>;
+
+defm NSAEncoding : AMDGPUSubtargetFeature<"nsa-encoding",
+ "Support NSA encoding for image instructions",
+ /*GenPredicate=*/0
+>;
+
+defm PartialNSAEncoding : AMDGPUSubtargetFeature<"partial-nsa-encoding",
+ "Support partial NSA encoding for image instructions",
+ /*GenPredicate=*/0
+>;
+
+defm ImageInsts : AMDGPUSubtargetFeature<"image-insts",
+ "Support image instructions"
+>;
+
+defm ExtendedImageInsts : AMDGPUSubtargetFeature<"extended-image-insts",
+ "Support mips != 0, lod != 0, gather4, and get_lod"
+>;
+
+defm BVHRayTracingInsts : AMDGPUSubtargetFeature<"bvh-ray-tracing-insts",
+ "Has BVH ray tracing instructions"
+>;
+
+defm MSAALoadInsts : AMDGPUSubtargetFeature<"msaa-load-insts",
+ "Has image_msaa_load instructions",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/1,
+ /*Deps=*/[FeatureImageInsts]
+>;
+
+defm GFX10_BEncoding : AMDGPUSubtargetFeature<"gfx10_b-encoding",
+ "Encoding format GFX10_B",
+ /*GenPredicate=*/0
+>;
+
+defm IntClamp : AMDGPUSubtargetFeature<"int-clamp-insts",
+ "Support clamp for integer destination"
+>;
+
+defm UnpackedD16VMem : AMDGPUSubtargetFeature<"unpacked-d16-vmem",
+ "Has unpacked d16 vmem instructions"
+>;
+
+defm DLInsts : AMDGPUSubtargetFeature<"dl-insts",
+ "Has v_fmac_f32 and v_xnor_b32 instructions"
+>;
+
+defm FmacF64Inst : AMDGPUSubtargetFeature<"fmacf64-inst",
+ "Has v_fmac_f64 instruction"
+>;
+
+defm Dot1Insts : AMDGPUSubtargetFeature<"dot1-insts",
+ "Has v_dot4_i32_i8 and v_dot8_i32_i4 instructions"
+>;
+
+defm Dot2Insts : AMDGPUSubtargetFeature<"dot2-insts",
+ "Has v_dot2_i32_i16, v_dot2_u32_u16 instructions"
+>;
+
+defm Dot3Insts : AMDGPUSubtargetFeature<"dot3-insts",
+ "Has v_dot8c_i32_i4 instruction"
+>;
+
+defm Dot4Insts : AMDGPUSubtargetFeature<"dot4-insts",
+ "Has v_dot2c_i32_i16 instruction"
+>;
+
+defm Dot5Insts : AMDGPUSubtargetFeature<"dot5-insts",
+ "Has v_dot2c_f32_f16 instruction"
+>;
+
+defm Dot6Insts : AMDGPUSubtargetFeature<"dot6-insts",
+ "Has v_dot4c_i32_i8 instruction"
+>;
+
+defm Dot7Insts : AMDGPUSubtargetFeature<"dot7-insts",
+ "Has v_dot4_u32_u8, v_dot8_u32_u4 instructions"
+>;
+
+defm Dot8Insts : AMDGPUSubtargetFeature<"dot8-insts",
+ "Has v_dot4_i32_iu8, v_dot8_i32_iu4 instructions"
+>;
+
+defm Dot9Insts : AMDGPUSubtargetFeature<"dot9-insts",
+ "Has v_dot2_f16_f16, v_dot2_bf16_bf16 instructions"
+>;
+
+defm Dot10Insts : AMDGPUSubtargetFeature<"dot10-insts",
+ "Has v_dot2_f32_f16 instruction"
+>;
+
+defm Dot11Insts : AMDGPUSubtargetFeature<"dot11-insts",
+ "Has v_dot4_f32_fp8_fp8, v_dot4_f32_fp8_bf8, v_dot4_f32_bf8_fp8, v_dot4_f32_bf8_bf8 instructions"
+>;
+
+defm Dot12Insts : AMDGPUSubtargetFeature<"dot12-insts",
+ "Has v_dot2_f32_bf16 instructions"
+>;
+
+defm Dot13Insts : AMDGPUSubtargetFeature<"dot13-insts",
+ "Has v_dot2c_f32_bf16 instructions"
+>;
+
+defm MAIInsts : AMDGPUSubtargetFeature<"mai-insts",
+ "Has mAI instructions"
+>;
+
+defm FP8Insts : AMDGPUSubtargetFeature<"fp8-insts",
+ "Has fp8 and bf8 instructions"
+>;
+
+defm FP8ConversionInsts : AMDGPUSubtargetFeature<"fp8-conversion-insts",
+ "Has fp8 and bf8 conversion instructions"
+>;
+
+// OCP fp8 semantics (E4M3FN, E5M2). Distinct from the FNUZ formats used by
+// gfx942.
+defm OCPFP8ConversionInsts : AMDGPUSubtargetFeature<"ocp-fp8-conversion-insts",
+ "Has OCP fp8 (E4M3FN, E5M2) conversion instructions"
+>;
+
+defm FP8F16ConversionInsts : AMDGPUSubtargetFeature<"fp8-f16-conversion-insts",
+ "Has fp8 and bf8 conversion to f16 instructions"
+>;
+
+// The opposite direction of FP8F16ConversionInsts above.
+defm F16FP8ConversionInsts : AMDGPUSubtargetFeature<"f16-fp8-conversion-insts",
+ "Has f16 conversion to fp8 and bf8 instructions"
+>;
+
+defm FP8E5M3Insts : AMDGPUSubtargetFeature<"fp8e5m3-insts",
+ "Has fp8 e5m3 format support"
+>;
+
+defm CvtFP8VOP1Bug : AMDGPUSubtargetFeature<"cvt-fp8-vop1-bug",
+ "FP8/BF8 VOP1 form of conversion to F32 is unreliable",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/0,
+ [FeatureFP8ConversionInsts]
+>;
+
+// The fp8/bf8 conversions take the byte index in one of two ways, and a
+// subtarget uses the same one in both directions.
+defm CvtFP8SDWASrcSel : AMDGPUSubtargetFeature<"cvt-fp8-sdwa-src-sel",
+ "FP8/BF8 conversions take the byte index from the SDWA src0_sel or op_sel field",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/1,
+ [FeatureFP8ConversionInsts]
+>;
+
+defm CvtFP8ByteSel : AMDGPUSubtargetFeature<"cvt-fp8-byte-sel",
+ "FP8/BF8 conversions take the byte index from a byte_sel operand",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/1,
+ [FeatureFP8ConversionInsts]
+>;
+
+defm WMMA256bInsts : AMDGPUSubtargetFeature<"wmma-256b-insts",
+ "Has WMMA instructions where A and B matrices have duplicated data"
+>;
+
+defm WMMA128bInsts : AMDGPUSubtargetFeature<"wmma-128b-insts",
+ "Has WMMA instructions where A and B matrices do not have duplicated data"
+>;
+
+defm WMMAN16Insts : AMDGPUSubtargetFeature<"wmma-n16-insts",
+ "Has WMMA instructions with N = 16"
+>;
+
+defm WMMAF4Insts : AMDGPUSubtargetFeature<"wmma-f4-insts",
+ "Has WMMA_*_F4 instructions"
+>;
+
+defm WMMA16x16x128FP8Insts : AMDGPUSubtargetFeature<"wmma-16x16x128fp8-insts",
+ "Has 16x16x128 FP8/BF8 WMMA instructions"
+>;
+
+defm SWMMACGfx1200Insts : AMDGPUSubtargetFeature<"swmmac-gfx1200-insts",
+ "Has GFX1200 SWMMAC instructions"
+>;
+
+defm SWMMACGfx1250Insts : AMDGPUSubtargetFeature<"swmmac-gfx1250-insts",
+ "Has GFX1250 SWMMAC instructions"
+>;
+
+defm PkFmacF16Inst : AMDGPUSubtargetFeature<"pk-fmac-f16-inst",
+ "Has v_pk_fmac_f16 instruction"
+>;
+
+defm CubeInsts : AMDGPUSubtargetFeature<"cube-insts",
+ "Has v_cube* instructions"
+>;
+
+defm IndexedResources : AMDGPUSubtargetFeature<"indexed-resources",
+ "Has indexed resources"
+>;
+
+defm LerpInst : AMDGPUSubtargetFeature<"lerp-inst",
+ "Has v_lerp_u8 instruction"
+>;
+
+defm SadInsts : AMDGPUSubtargetFeature<"sad-insts",
+ "Has v_sad* instructions"
+>;
+
+defm QsadInsts : AMDGPUSubtargetFeature<"qsad-insts",
+ "Has v_qsad* instructions"
+>;
+
+defm MsadInsts : AMDGPUSubtargetFeature<"msad-insts",
+ "Has v_msad_u8 instruction"
+>;
+
+defm MqsadPkInsts : AMDGPUSubtargetFeature<"mqsad-pk-insts",
+ "Has v_mqsad_pk_u16_u8 instruction"
+>;
+
+defm MqsadInsts : AMDGPUSubtargetFeature<"mqsad-insts",
+ "Has v_mqsad_u32_u8 instruction"
+>;
+
+defm CvtNormInsts : AMDGPUSubtargetFeature<"cvt-norm-insts",
+ "Has v_cvt_norm* instructions"
+>;
+
+defm CvtPkNormVOP2Insts : AMDGPUSubtargetFeature<"cvt-pknorm-vop2-insts",
+ "Has v_cvt_pk_norm_*f32 instructions/Has v_cvt_pk_norm_*_f16 instructions"
+>;
+
+defm CvtPkNormVOP3Insts : AMDGPUSubtargetFeature<"cvt-pknorm-vop3-insts",
+ "Has v_cvt_pk_norm_*f32 instructions/Has v_cvt_pk_norm_*_f16 instructions"
+>;
+
+defm AtomicDsPkAdd16Insts : AMDGPUSubtargetFeature<"atomic-ds-pk-add-16-insts",
+ "Has ds_pk_add_bf16, ds_pk_add_f16, ds_pk_add_rtn_bf16, "
+ "ds_pk_add_rtn_f16 instructions"
+>;
+
+defm AtomicDsCondSubClampInsts : AMDGPUSubtargetFeature<"atomic-ds-cond-sub-clamp-insts",
+ "Has ds_cond_sub_u32, ds_cond_sub_rtn_u32, ds_sub_clamp_u32 and "
+ "ds_sub_clamp_rtn_u32 instructions"
+>;
+
+defm LdsAtomicAddF64 : AMDGPUSubtargetFeature<"lds-atomic-add-f64",
+ "Has ds_add_f64 and ds_add_rtn_f64 instructions"
+>;
+
+defm AtomicFlatPkAdd16Insts : AMDGPUSubtargetFeature<"atomic-flat-pk-add-16-insts",
+ "Has flat_atomic_pk_add_f16 and flat_atomic_pk_add_bf16 instructions"
+>;
+
+defm AtomicCondSubClampFlatInsts : AMDGPUSubtargetFeature<"atomic-cond-sub-clamp-flat-insts",
+ "Has flat_atomic_cond_sub_u32 and flat_atomic_sub_clamp_u32 instructions"
+>;
+
+defm AtomicFaddRtnInsts : AMDGPUSubtargetFeature<"atomic-fadd-rtn-insts",
+ "Has buffer_atomic_add_f32 and global_atomic_add_f32 instructions that "
+ "return original value",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/1,
+ [FeatureFlatGlobalInsts]
+>;
+
+defm AtomicFMinFMaxF32GlobalInsts : AMDGPUSubtargetFeature<"atomic-fmin-fmax-global-f32",
+ "Has global/buffer instructions for atomicrmw fmin/fmax for float"
+>;
+
+defm AtomicFMinFMaxF64GlobalInsts : AMDGPUSubtargetFeature<"atomic-fmin-fmax-global-f64",
+ "Has global/buffer instructions for atomicrmw fmin/fmax for float"
+>;
+
+defm AtomicFMinFMaxF32FlatInsts : AMDGPUSubtargetFeature<"atomic-fmin-fmax-flat-f32",
+ "Has flat memory instructions for atomicrmw fmin/fmax for float",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/1,
+ [FeatureFlatAddressSpace]
+>;
+
+defm AtomicFMinFMaxF64FlatInsts : AMDGPUSubtargetFeature<"atomic-fmin-fmax-flat-f64",
+ "Has flat memory instructions for atomicrmw fmin/fmax for double",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/1,
+ [FeatureFlatAddressSpace]
+>;
+
+defm AtomicFaddNoRtnInsts : AMDGPUSubtargetFeature<"atomic-fadd-no-rtn-insts",
+ "Has buffer_atomic_add_f32 and global_atomic_add_f32 instructions that "
+ "don't return original value",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/1,
+ [FeatureFlatGlobalInsts]
+>;
+
+def FeatureAtomicBufferGlobalPkAddF16NoRtnInsts
+ : SubtargetFeature<"atomic-buffer-global-pk-add-f16-no-rtn-insts",
+ "HasAtomicBufferGlobalPkAddF16NoRtnInsts",
+ "true",
+ "Has buffer_atomic_pk_add_f16 and global_atomic_pk_add_f16 instructions that "
+ "don't return original value",
+ [FeatureFlatGlobalInsts]
+>;
+
+defm AtomicBufferGlobalPkAddF16Insts : AMDGPUSubtargetFeature<"atomic-buffer-global-pk-add-f16-insts",
+ "Has buffer_atomic_pk_add_f16 and global_atomic_pk_add_f16 instructions that "
+ "can return original value",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/1,
+ [FeatureFlatGlobalInsts]
+>;
+
+defm AtomicGlobalPkAddBF16Inst : AMDGPUSubtargetFeature<"atomic-global-pk-add-bf16-inst",
+ "Has global_atomic_pk_add_bf16 instruction",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/1,
+ [FeatureFlatGlobalInsts]
+>;
+
+defm AtomicBufferPkAddBF16Inst : AMDGPUSubtargetFeature<"atomic-buffer-pk-add-bf16-inst",
+ "Has buffer_atomic_pk_add_bf16 instruction"
+>;
+
+defm AtomicCSubNoRtnInsts : AMDGPUSubtargetFeature<"atomic-csub-no-rtn-insts",
+ "Has buffer_atomic_csub and global_atomic_csub instructions that don't "
+ "return original value",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/0
+>;
+
+defm FlatAtomicFaddF32Inst : AMDGPUSubtargetFeature<"flat-atomic-fadd-f32-inst",
+ "Has flat_atomic_add_f32 instruction",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/1,
+ [FeatureFlatAddressSpace]
+>;
+
+defm GFX1251GEMMInsts : AMDGPUSubtargetFeature<"gfx1251-gemm-insts",
+ "Has additional gfx1251 DGEMM instructions"
+>;
+
+defm GFX125xLowestRateWMMA : AMDGPUGenericAnySubtargetFeature<"gfx125x-lowest-rate-wmma",
+ "Has the lowest rate wmma in the gfx125x family, mainly for the gfx12-5-generic"
+>;
+
+defm FlatBufferGlobalAtomicFaddF64Inst : AMDGPUSubtargetFeature<"flat-buffer-global-fadd-f64-inst",
+ "Has flat, buffer, and global instructions for f64 atomic fadd"
+>;
+
+def FeatureMemoryAtomicFAddF32DenormalSupport
+ : SubtargetFeature<"memory-atomic-fadd-f32-denormal-support",
+ "HasMemoryAtomicFaddF32DenormalSupport",
+ "true",
+ "global/flat/buffer atomic fadd for float supports denormal handling"
+>;
+
+defm AgentScopeFineGrainedRemoteMemoryAtomics : AMDGPUSubtargetFeature<
+ "agent-scope-fine-grained-remote-memory-atomics",
+ "Agent (device) scoped atomic operations, excluding those directly "
+ "supported by PCIe (i.e. integer atomic add, exchange, and "
+ "compare-and-swap), are functional for allocations in host or peer "
+ "device memory.",
+ /*GenPredicate=*/0
+>;
+
+defm EmulatedSystemScopeAtomics : AMDGPUSubtargetFeature<
+ "emulated-system-scope-atomics",
+ "System scope atomics unsupported by the PCI-e are emulated in HW via CAS "
+ "loop and functional.",
+ /*GenPredicate=*/0
+>;
+
+defm DefaultComponentZero : AMDGPUSubtargetFeature<"default-component-zero",
+ "BUFFER/IMAGE store instructions set unspecified components to zero (before GFX12)"
+>;
+
+defm DefaultComponentBroadcast : AMDGPUSubtargetFeature<"default-component-broadcast",
+ "BUFFER/IMAGE store instructions set unspecified components to x component (GFX12)"
+>;
+
+// The default assumption needs to be ecc is enabled, but no directly
+// exposed operations depend on it, so it can be safely inlined.
+def FeatureSupportsSRAMECC : SubtargetFeature<"sramecc-support",
+ "SupportsSRAMECC",
+ "true",
+ "Hardware supports SRAMECC",
+ [], InlineIgnore
+>;
+
+defm SRAMECCOnOffModes : AMDGPUSubtargetFeature<"sramecc-on-off-modes",
+ "Target supports SRAMECC on/off modes",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/0,
+ [FeatureSupportsSRAMECC],
+ InlineIgnore
+>;
+
+defm NoSdstCMPX : AMDGPUSubtargetFeature<"no-sdst-cmpx",
+ "V_CMPX does not write VCC/SGPR in addition to EXEC"
+>;
+
+defm Vscnt : AMDGPUSubtargetFeature<"vscnt",
+ "Has separate store vscnt counter",
+ /*GenPredicate=*/0
+>;
+
+defm Asynccnt : AMDGPUSubtargetFeature<"asynccnt",
+ "Has separate asynccnt counter",
+ /*GenPredicate=*/0
+>;
+
+defm GetWaveIdInst : AMDGPUSubtargetFeature<"get-wave-id-inst",
+ "Has s_get_waveid_in_workgroup instruction"
+>;
+
+defm PopsExitingWaveID : AMDGPUSubtargetFeature<"pops-exiting-wave-id",
+ "Has POPS_EXITING_WAVE_ID hardware register"
+>;
+
+defm SMemTimeInst : AMDGPUSubtargetFeature<"s-memtime-inst",
+ "Has s_memtime instruction"
+>;
+
+defm ShaderCyclesRegister : AMDGPUSubtargetFeature<"shader-cycles-register",
+ "Has SHADER_CYCLES hardware register"
+>;
+
+defm ShaderCyclesHiLoRegisters : AMDGPUSubtargetFeature<"shader-cycles-hi-lo-registers",
+ "Has SHADER_CYCLES_HI/LO hardware registers",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/0
+>;
+
+defm MadMacF32Insts : AMDGPUSubtargetFeature<"mad-mac-f32-insts",
+ "Has v_mad_f32/v_mac_f32/v_madak_f32/v_madmk_f32 instructions"
+>;
+
+defm DsSrc2Insts : AMDGPUSubtargetFeature<"ds-src2-insts",
+ "Has ds_*_src2 instructions"
+>;
+
+defm VOP3Literal : AMDGPUSubtargetFeature<"vop3-literal",
+ "Can use one literal in VOP3",
+ /*GenPredicate=*/0
+>;
+
+defm NoDataDepHazard : AMDGPUSubtargetFeature<"no-data-dep-hazard",
+ "Does not need SW waitstates",
+ /*GenPredicate=*/0
+>;
+
+// 1536 physical VGPRs for wave32 and 768 VGPRs for wave64
+// with allocation granularity 24 for wave32 and 12 for wave64
+defm 1536VGPRs : AMDGPUSubtargetFeature<"1536-physical-vgprs",
+ "Has 1536 physical VGPRs per SIMD",
+ /*GenPredicate=*/0
+>;
+
+defm VOPDInsts : AMDGPUSubtargetFeature<"vopd",
+ "Has VOPD dual issue wave32 instructions",
+ /*GenPredicate=*/0
+>;
+
+defm VOPDInterlockHazard : AMDGPUSubtargetFeature<"vopd-interlock-hazard",
+ "VOPD destination buffer forwarding drops the interlock when SRC0 or SRC1 "
+ "X/Y are distinct VGPRs with the same register-number parity",
+ /*GenPredicate=*/0
+>;
+
+defm VALUTransUseHazard : AMDGPUGenericAnySubtargetFeature<"valu-trans-use-hazard",
+ "Hazard when TRANS instructions are closely followed by a use of the result",
+ /*GenPredicate=*/0
+>;
+
+defm TransCoexecutionHazard : AMDGPUSubtargetFeature<"trans-coexecution-hazard",
+ "Hazard when TRANS co-executes with other VALU instruction",
+ /*GenPredicate=*/0
+>;
+
+defm GFX1250VALUBlockingCycles : AMDGPUSubtargetFeature<"gfx1250-valu-blocking-cycles",
+ "VALU pipe occupancy is modeled with GFX1250BlockingCycles",
+ /*GenPredicate=*/0
+>;
+
+defm VPermPk16Hazard : AMDGPUSubtargetFeature<"vperm-pk16-hazard",
+ "V_PERM_PK16 must be immediately followed by a safe VALU (or an inserted V_NOP)",
+ /*GenPredicate=*/0
+>;
+
+defm WMMACoexecutionHazards : AMDGPUSubtargetFeature<"wmma-coexecution-hazards",
+ "Hazards when WMMA co-executes with other WMMA/VALU instruction",
+ /*GenPredicate=*/0
+>;
+
+defm VOPD3F64OPYSrcHazard : AMDGPUSubtargetFeature<"vopd3-f64-opy-src-hazard",
+ "Hazard when an OPY VGPR source of a VOPD3 with a 64-bit OPX names the last "
+ "VGPR allocated to the wave",
+ /*GenPredicate=*/0
+>;
+
+defm SALUFloatInsts : AMDGPUSubtargetFeature<"salu-float",
+ "Has SALU floating point instructions"
+>;
+
+defm PseudoScalarTrans : AMDGPUSubtargetFeature<"pseudo-scalar-trans",
+ "Has Pseudo Scalar Transcendental instructions"
+>;
+
+defm NoF16PseudoScalarTransInlineConstants : AMDGPUSubtargetFeature<"no-f16-pseudo-scalar-trans-inline-constants",
+ "Inline constants are not supported for F16 pseudo scalar transcendentals"
+>;
+
+defm RestrictedSOffset : AMDGPUSubtargetFeature<"restricted-soffset",
+ "Has restricted SOffset (immediate not supported)."
+>;
+
+defm RequiredExportPriority : AMDGPUGenericAnySubtargetFeature<"required-export-priority",
+ "Export priority must be explicitly manipulated on GFX11.5",
+ /*GenPredicate=*/0
+>;
+
+defm VmemWriteVgprInOrder : AMDGPUSubtargetFeature<"vmem-write-vgpr-in-order",
+ "VMEM instructions of the same type write VGPR results in order",
+ /*GenPredicate=*/0
+>;
+
+defm BitOp3Insts : AMDGPUSubtargetFeature<"bitop3-insts",
+ "Has v_bitop3_b32/v_bitop3_b16 instructions"
+>;
+
+defm TanhInsts : AMDGPUSubtargetFeature<"tanh-insts",
+ "Has v_tanh_f32/f16 instructions"
+>;
+
+defm TensorCvtLutInsts : AMDGPUSubtargetFeature<"tensor-cvt-lut-insts",
+ "Has v_perm_pk16* instructions"
+>;
+
+defm TDMInsts : AMDGPUSubtargetFeature<"tdm-insts",
+ "Has tensor_load_to_lds/tensor_store_from_lds instructions"
+>;
+
+defm TransposeLoadF4F6Insts : AMDGPUSubtargetFeature<"transpose-load-f4f6-insts",
+ "Has ds_load_tr4/tr6 and global_load_tr4/tr6 instructions"
+>;
+
+defm PrngInst : AMDGPUSubtargetFeature<"prng-inst",
+ "Has v_prng_b32 instruction"
+>;
+
+defm BVHDualAndBVH8Insts : AMDGPUSubtargetFeature<"bvh-dual-bvh-8-insts",
+ "Has image_bvh_dual_intersect_ray and image_bvh8_intersect_ray instructions"
+>;
+
+defm PointSampleAccel : AMDGPUSubtargetFeature<"point-sample-accel",
+ "Has point sample acceleration feature",
+ /*GenPredicate=*/0
+>;
+
+defm 64BitLiterals : AMDGPUSubtargetFeature<"64-bit-literals",
+ "Can use 64-bit literals with single DWORD instructions"
+>;
+
+defm 1024AddressableVGPRs : AMDGPUSubtargetFeature<"1024-addressable-vgprs",
+ "Has 1024 addressable VGPRs"
+>;
+
+defm SetregVGPRMSBFixup : AMDGPUGenericAnySubtargetFeature<"setreg-vgpr-msb-fixup",
+ "S_SETREG to MODE clobbers VGPR MSB bits, requires fixup",
+ /*GenPredicate=*/0
+>;
+
+defm WaitXcnt : AMDGPUSubtargetFeature<"wait-xcnt",
+ "Has s_wait_xcnt instruction"
+>;
+
+defm SetPrioIncWgInst : AMDGPUSubtargetFeature<"setprio-inc-wg-inst",
+ "Has s_setprio_inc_wg instruction."
+>;
+
+defm SWakeupBarrier : AMDGPUSubtargetFeature<"s-wakeup-barrier-inst",
+ "Has s_wakeup_barrier instruction."
+>;
+
+defm FlatSignedOffset : AMDGPUSubtargetFeature<"flat-signed-offset",
+ "Immediate offset of FLAT instructions are always signed"
+>;
+
+defm RequiresInitialUnclausedVmem : AMDGPUSubtargetFeature<"requires-initial-unclause-vmem",
+ "Hardware entrypoints require first VMEM instruction to be unclaused",
+ GenPredicate=false
+>;
+
+defm SlowMaxMinMulI64Insts : AMDGPUSubtargetFeature<"slow-max-min-mul-i64-insts",
+ "64-bit integer max/min/mul VALU instructions are slower than expanding to an equivalent code sequence",
+ GenPredicate=false
+>;
+
+//===------------------------------------------------------------===//
+// Subtarget Features (options and debugging)
+//===------------------------------------------------------------===//
+
+// Ugly hack to accomodate assembling modules with mixed
+// wavesizes. Ideally we would have a mapping symbol in assembly which
+// would keep track of which sections of code should be treated as
+// wave32 and wave64. Instead what users do is assemble with both
+// wavesizes enabled. We translate this into this special mode so this
+// only influences assembler behavior and nothing else.
+defm AssemblerPermissiveWavesize : AMDGPUSubtargetFeature<"assembler-permissive-wavesize",
+ "Allow parsing wave32 and wave64 variants of instructions",
+ /*GenPredicate=*/0
+>;
+
+class FeatureMaxPrivateElementSize<int size> : SubtargetFeature<
+ "max-private-element-size-"#size,
+ "MaxPrivateElementSize",
+ !cast<string>(size),
+ "Maximum private access size may be "#size
+>;
+
+def FeatureMaxPrivateElementSize4 : FeatureMaxPrivateElementSize<4>;
+def FeatureMaxPrivateElementSize8 : FeatureMaxPrivateElementSize<8>;
+def FeatureMaxPrivateElementSize16 : FeatureMaxPrivateElementSize<16>;
+
+def FeatureDumpCodeLower : SubtargetFeature <"dumpcode",
+ "DumpCode",
+ "true",
+ "Dump MachineInstrs in the CodeEmitter"
+>;
+
+// XXX - This should probably be removed once enabled by default
+def FeatureEnableLoadStoreOpt : SubtargetFeature <"load-store-opt",
+ "EnableLoadStoreOpt",
+ "true",
+ "Enable SI load/store optimizer pass",
+ [], InlineIgnore
+>;
+
+// Performance debugging feature. Allow using DS instruction immediate
+// offsets even if the base pointer can't be proven to be base. On SI,
+// base pointer values that won't give the same result as a 16-bit add
+// are not safe to fold, but this will override the conservative test
+// for the base pointer.
+def FeatureEnableUnsafeDSOffsetFolding : SubtargetFeature <
+ "unsafe-ds-offset-folding",
+ "EnableUnsafeDSOffsetFolding",
+ "true",
+ "Force using DS instruction immediate offsets on SI",
+ [], InlineIgnore
+>;
+
+def FeatureEnableSIScheduler : SubtargetFeature<"si-scheduler",
+ "EnableSIScheduler",
+ "true",
+ "Enable SI Machine Scheduler",
+ [], InlineIgnore
+>;
+
+def FeatureEnableDS128 : SubtargetFeature<"enable-ds128",
+ "EnableDS128",
+ "true",
+ "Use ds_{read|write}_b128"
+>;
+
+// Sparse texture support requires that all result registers are zeroed when
+// PRTStrictNull is set to true. This feature is turned on for all architectures
+// but is enabled as a feature in case there are situations where PRTStrictNull
+// is disabled by the driver.
+def FeatureEnablePRTStrictNull : SubtargetFeature<"enable-prt-strict-null",
+ "EnablePRTStrictNull",
+ "true",
+ "Enable zeroing of result registers for sparse texture fetches"
+>;
+
+// Unless +-flat-for-global is specified, turn on FlatForGlobal for
+// all OS-es on VI and newer hardware to avoid assertion failures due
+// to missing ADDR64 variants of MUBUF instructions.
+// FIXME: moveToVALU should be able to handle converting addr64 MUBUF
+// instructions.
+
+def FeatureUseFlatForGlobal : SubtargetFeature<"flat-for-global",
+ "UseFlatForGlobal",
+ "true",
+ "Force to generate flat instruction for global",
+ [], InlineIgnore
+>;
+
+defm AutoWaitcntBeforeBarrier : AMDGPUSubtargetFeature <"auto-waitcnt-before-barrier",
+ "Hardware automatically inserts waitcnt before barrier",
+ /*GenPredicate=*/0, /*GenAssemblerPredicate=*/1, [], InlineIgnore
+>;
+
+defm BackOffBarrier : AMDGPUSubtargetFeature <"back-off-barrier",
+ "Hardware supports backing off s_barrier if an exception occurs",
+ /*GenPredicate=*/0
+>;
+
+defm TrigReducedRange : AMDGPUSubtargetFeature<"trig-reduced-range",
+ "Requires use of fract on arguments to trig instructions",
+ /*GenPredicate=*/0
+>;
+
+defm KernargPreload : AMDGPUSubtargetFeature <"kernarg-preload",
+ "Hardware supports preloading of kernel arguments in user SGPRs.",
+ /*GenPredicate=*/0
+>;
+
+// Alignment enforcement is controlled by a configuration register:
+// SH_MEM_CONFIG.alignment_mode
+defm UnalignedAccessMode : AMDGPUSubtargetFeature<"unaligned-access-mode",
+ "Enable unaligned global, local and region loads and stores if the hardware"
+ " supports it",
+ /*GenPredicate=*/1, /*GenAssemblerPredicate=*/1, [], InlineIgnore
+>;
+
+defm PackedTID : AMDGPUSubtargetFeature<"packed-tid",
+ "Workitem IDs are packed into v0 at kernel launch",
+ /*GenPredicate=*/0
+>;
+
+defm ArchitectedFlatScratch : AMDGPUSubtargetFeature<"architected-flat-scratch",
+ "Flat Scratch register is a readonly SPI initialized architected register",
+ /*GenPredicate=*/0
+>;
+
+defm ArchitectedSGPRs : AMDGPUSubtargetFeature<"architected-sgprs",
+ "Enable the architected SGPRs",
+ /*GenPredicate=*/0
+>;
+
+defm GDS : AMDGPUSubtargetFeature<"gds",
+ "Has Global Data Share",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/0
+>;
+
+defm GWS : AMDGPUSubtargetFeature<"gws",
+ "Has Global Wave Sync",
+ /*GenPredicate=*/1,
+ /*GenAssemblerPredicate=*/0
+>;
+
+// Generic targets require code object v6 independently of their covered GPUs,
+// so this feature does not use the any-covered-GPU classification.
+def FeatureRequiresCOV6 : SubtargetFeature<"requires-cov6",
+ "RequiresCOV6",
+ "true",
+ "Target Requires Code Object V6"
+>;
+
+defm XF32Insts : AMDGPUSubtargetFeature<"xf32-insts",
+ "Has instructions that support xf32 format, such as "
+ "v_mfma_f32_16x16x8_xf32 and v_mfma_f32_32x32x4_xf32"
+>;
+
+defm GloballyAddressableScratch : AMDGPUSubtargetFeature<"globally-addressable-scratch",
+ "FLAT instructions can access scratch memory for any thread in any wave",
+ /*GenPredicate=*/0
+>;
+
+// Enable the use of SCRATCH_STORE/LOAD_BLOCK instructions for saving and
+// restoring the callee-saved registers.
+def FeatureUseBlockVGPROpsForCSR : SubtargetFeature<"block-vgpr-csr",
+ "UseBlockVGPROpsForCSR",
+ "true",
+ "Use block load/store for VGPR callee saved registers"
+>;
+
+defm LshlAddU64Inst : AMDGPUSubtargetFeature<"lshl-add-u64-inst",
+ "Has v_lshl_add_u64 instruction"
+>;
+
+defm AddSubU64Insts : AMDGPUSubtargetFeature<"add-sub-u64-insts",
+ "Has v_add_u64 and v_sub_u64 instructions"
+>;
+
+defm MadU32Inst : AMDGPUSubtargetFeature<"mad-u32-inst",
+ "Has v_mad_u32 instruction"
+>;
+
+defm MadNC64_32Insts : AMDGPUSubtargetFeature<"mad-nc-64-32-insts",
+ "Has v_mad_nc_{u64_u32|i64_i32} instructions"
+>;
+
+defm AddMinMaxInsts : AMDGPUSubtargetFeature<"add-min-max-insts",
+ "Has v_add_{min|max}_{i|u}32 instructions"
+>;
+
+defm MinMaxI64Insts : AMDGPUSubtargetFeature<"min-max-i64-insts",
+ "Has v_{min|max}_{i|u}64 instructions"
+>;
+
+defm PkAddMinMaxInsts : AMDGPUSubtargetFeature<"pk-add-min-max-insts",
+ "Has v_pk_add_{min|max}_{i|u}16 instructions"
+>;
+
+defm VMemToLDSLoad : AMDGPUSubtargetFeature<"vmem-to-lds-load-insts",
+ "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."
+>;
+
+// Manual predicate for hasAsyncMark() which combines HasVMemToLDSLoad and GFX1250Plus
+def HasAsyncMark : Predicate<"Subtarget->hasAsyncMark()">;
+
+defm LdsBarrierArriveAtomic : AMDGPUSubtargetFeature<"lds-barrier-arrive-atomic",
+ "Has LDS barrier-arrive atomic instructions"
+>;
+
+// The width of the num_records field of the buffer resource (V#).
+class FeatureNumRecordsBufferResource<int width> : SubtargetFeature<
+ !cast<string>(width)#"-bit-num-records-buffer-resource",
+ "BufferResourceNumRecordsWidth",
+ !cast<string>(width),
+ "The buffer resource (V#) has "#width#"-bit num_records"
+>;
+
+def Feature32BitNumRecordsBufferResource : FeatureNumRecordsBufferResource<32>;
+def Feature45BitNumRecordsBufferResource : FeatureNumRecordsBufferResource<45>;
+
+defm Clusters : AMDGPUSubtargetFeature<"clusters",
+ "Has clusters of workgroups support",
+ /*GenPredicate=*/0
+>;
+
+def FeatureWaitsBeforeSystemScopeStores : SubtargetFeature<
+ "waits-before-system-scope-stores",
+ "RequiresWaitsBeforeSystemScopeStores",
+ "true",
+ "Target requires waits for loads and atomics before system scope stores"
+>;
+
+def FeatureUseAddPC64Inst : SubtargetFeature<"use-add-pc64-inst",
+ "UseAddPC64Inst",
+ "true",
+ "Use s_add_pc_i64 instruction."
+>;
+
+defm GFX1250_STRICT : AMDGPUSubtargetFeature<"gfx1250-strict",
+ "gfx1250 strict mode",
+ /*GenPredicate=*/0
+>;
+
+defm VOP3PX2IncrementsVaVdstTwice : AMDGPUSubtargetFeature<"vop3px2-increments-va-vdst-twice",
+ "VOP3PX2 encoded instructions increment the VaVdst counter twice", GenPredicate=false
+>;
+
+defm NeedsTDMDrain : AMDGPUSubtargetFeature<"needs-tdm-drain",
+ "Needs to keep tensor count low",
+ /*GenPredicate=*/0
+>;
+
+defm NeedsAligned2addrDS : AMDGPUSubtargetFeature<"needs-aligned-2addr-lds",
+ "Needs 2addr LDS instructions to be aligned to payload size",
+ /*GenPredicate=*/1
+>;
+
+defm NoSleepForever : AMDGPUSubtargetFeature<"no-sleep-forever",
+ "s_monitor_sleep should not use sleep_forever",
+ /*GenPredicate=*/1
+>;
+
+defm WaveMatchInsts : AMDGPUSubtargetFeature<"wave-match-insts",
+ "Has v_wave_match_b32 instruction"
+>;
+
+//===----------------------------------------------------------------------===//
+
+class GCNSubtargetFeatureGeneration <string Value,
+ string FeatureName,
+ list<SubtargetFeature> Implies> :
+ SubtargetFeatureGeneration <Value, FeatureName, "GCNSubtarget", Implies>;
+
+def FeatureSouthernIslands : GCNSubtargetFeatureGeneration<"SOUTHERN_ISLANDS",
+ "southern-islands",
+ [FeatureFP64, FeatureAddressableLocalMemorySize32768,
+ FeatureHalfAddressablePhysicalLocalMemory,
+ FeatureLDSAllocGranularity256,
+ FeatureLDSEncodingGranularity256,
+ FeatureMIMG_R128,
+ FeatureWavefrontSize64, FeatureSupportsWave64, FeatureSMemTimeInst,
+ FeatureMadMacF32Insts,
+ FeatureDsSrc2Insts, FeatureLDSBankCount32, FeatureMovrel,
+ FeatureTrigReducedRange, FeatureExtendedImageInsts, FeatureImageInsts,
+ FeatureInstCacheLineSize64,
+ FeatureGDS, FeatureGWS, FeatureDefaultComponentZero,
+ FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF64GlobalInsts,
+ FeatureVmemWriteVgprInOrder, FeatureCubeInsts, FeatureLerpInst,
+ FeatureSadInsts, FeatureMsadInsts, FeatureMqsadPkInsts,
+ FeatureCvtPkNormVOP2Insts, FeatureDX10ClampAndIEEEMode, FeatureMTBUFInsts,
+ FeatureFormattedMUBUFInsts, Feature32BitNumRecordsBufferResource,
+ FeatureMaxWavesPerEU10
+ ]
+>;
+
+def FeatureSeaIslands : GCNSubtargetFeatureGeneration<"SEA_ISLANDS",
+ "sea-islands",
+ [FeatureFP64, FeatureAddressableLocalMemorySize65536,
+ FeatureLDSAllocGranularity512,
+ FeatureLDSEncodingGranularity512,
+ FeatureMIMG_R128,
+ FeatureWavefrontSize64, FeatureSupportsWave64, FeatureFlatAddressSpace,
+ FeatureCIInsts, FeatureMovrel, FeatureTrigReducedRange,
+ FeatureGFX7GFX8GFX9Insts, FeatureSMemTimeInst, FeatureMadMacF32Insts,
+ FeatureDsSrc2Insts, FeatureExtendedImageInsts, FeatureUnalignedBufferAccess,
+ FeatureImageInsts, FeatureGDS, FeatureGWS, FeatureDefaultComponentZero,
+ FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF64GlobalInsts,
+ FeatureAtomicFMinFMaxF32FlatInsts, FeatureAtomicFMinFMaxF64FlatInsts,
+ FeatureVmemWriteVgprInOrder, FeatureCubeInsts, FeatureLerpInst,
+ FeatureSadInsts, FeatureQsadInsts, FeatureMsadInsts, FeatureMqsadPkInsts,
+ FeatureMqsadInsts, FeatureCvtPkNormVOP2Insts,
+ FeatureDX10ClampAndIEEEMode, FeatureInstCacheLineSize64,
+ FeatureMTBUFInsts, FeatureFormattedMUBUFInsts,
+ Feature32BitNumRecordsBufferResource,
+ FeatureMaxWavesPerEU10
+ ]
+>;
+
+def FeatureVolcanicIslands : GCNSubtargetFeatureGeneration<"VOLCANIC_ISLANDS",
+ "volcanic-islands",
+ [FeatureFP64, FeatureAddressableLocalMemorySize65536,
+ FeatureLDSAllocGranularity512,
+ FeatureLDSEncodingGranularity512,
+ FeatureMIMG_R128,
+ FeatureWavefrontSize64, FeatureSupportsWave64, FeatureFlatAddressSpace,
+ FeatureGCN3Encoding, FeatureCIInsts, Feature16BitInsts,
+ FeatureSMemRealTime, FeatureVGPRIndexMode, FeatureMovrel,
+ FeatureScalarStores, FeatureInv2PiInlineImm,
+ FeatureSDWA, FeatureSDWAOutModsVOPC, FeatureSDWAMac, FeatureDPP,
+ FeatureDPPWavefrontShifts, FeatureDPPBroadcasts,
+ FeatureIntClamp, FeatureTrigReducedRange, FeatureGFX8Insts,
+ FeatureGFX7GFX8GFX9Insts, FeatureSMemTimeInst, FeatureMadMacF32Insts,
+ FeatureDsSrc2Insts, FeatureExtendedImageInsts, FeatureFastDenormalF32,
+ FeatureUnalignedBufferAccess, FeatureImageInsts, FeatureGDS, FeatureGWS,
+ FeatureDefaultComponentZero, FeatureVmemWriteVgprInOrder, FeatureCubeInsts,
+ FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts, FeatureMsadInsts,
+ FeatureMqsadPkInsts, FeatureMqsadInsts,
+ FeatureCvtPkNormVOP2Insts, FeatureDX10ClampAndIEEEMode,
+ FeatureInstCacheLineSize64, FeatureMTBUFInsts, FeatureFormattedMUBUFInsts,
+ Feature32BitNumRecordsBufferResource,
+ FeatureMaxWavesPerEU10
+ ]
+>;
+
+def FeatureGFX9 : GCNSubtargetFeatureGeneration<"GFX9",
+ "gfx9",
+ [FeatureFP64,
+ FeatureWavefrontSize64, FeatureSupportsWave64, FeatureFlatAddressSpace,
+ FeatureGCN3Encoding, FeatureCIInsts, Feature16BitInsts,
+ FeatureSMemRealTime, FeatureScalarStores, FeatureInv2PiInlineImm,
+ FeatureApertureRegs, FeatureGetDoorbellID,
+ FeatureGFX9Insts, FeatureVOP3PInsts, FeatureVGPRIndexMode,
+ FeatureFastFMAF32, FeatureDPP, FeatureDPPWavefrontShifts, FeatureDPPBroadcasts,
+ FeatureIntClamp, FeatureSDWA, FeatureSDWAOmod, FeatureSDWAScalar, FeatureSDWASdst,
+ FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts,
+ FeatureAddNoCarryInsts, FeatureGFX8Insts, FeatureGFX7GFX8GFX9Insts,
+ FeatureScalarFlatScratchInsts, FeatureScalarAtomics, FeatureR128A16,
+ FeatureA16, FeatureSMemTimeInst, FeatureFastDenormalF32, FeatureXNACKOnOffModes,
+ FeatureUnalignedBufferAccess, FeatureUnalignedScratchAccess,
+ FeatureUnalignedDSAccess, FeatureNegativeScratchOffsetBug, FeatureGWS,
+ FeatureDefaultComponentZero,FeatureVmemWriteVgprInOrder, FeatureVMemToLDSLoad,
+ FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts,
+ FeatureMsadInsts, FeatureMqsadPkInsts, FeatureMqsadInsts,
+ FeatureCvtNormInsts, FeatureCvtPkNormVOP2Insts,
+ FeatureCvtPkNormVOP3Insts, FeatureDX10ClampAndIEEEMode,
+ FeatureInstCacheLineSize64, FeatureMTBUFInsts, FeatureFormattedMUBUFInsts,
+ Feature32BitNumRecordsBufferResource
+ ]
+>;
+
+def FeatureGFX10 : GCNSubtargetFeatureGeneration<"GFX10",
+ "gfx10",
+ [FeatureFP64, FeatureAddressableLocalMemorySize65536,
+ FeatureLDSEncodingGranularity512,
+ FeatureHalfAddressablePhysicalLocalMemory, FeatureMIMG_R128,
+ FeatureSupportsWave32, FeatureSupportsWave64, FeatureSupportsWGP,
+ FeatureFlatAddressSpace,
+ FeatureCIInsts, Feature16BitInsts,
+ FeatureSMemRealTime, FeatureInv2PiInlineImm,
+ FeatureApertureRegs, FeatureGetDoorbellID,
+ FeatureGFX9Insts, FeatureGFX10Insts, FeatureVOP3PInsts,
+ FeatureMovrel, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp,
+ FeatureSDWA, FeatureSDWAOmod, FeatureSDWAScalar, FeatureSDWASdst,
+ FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts,
+ FeatureAddNoCarryInsts, FeatureFmaMixInsts, FeatureGFX8Insts,
+ FeatureNoSdstCMPX, FeatureVscnt,
+ FeatureVOP3Literal, FeatureDPP8, FeatureExtendedImageInsts,
+ FeaturePermlane16Insts,
+ FeatureNoDataDepHazard, FeaturePkFmacF16Inst,
+ FeatureA16, FeatureSMemTimeInst, FeatureFastDenormalF32, FeatureG16,
+ FeatureUnalignedBufferAccess, FeatureUnalignedScratchAccess,
+ FeatureUnalignedDSAccess, FeatureImageInsts, FeatureGDS, FeatureGWS,
+ FeatureDefaultComponentZero, FeatureMaxHardClauseLength63,
+ FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF64GlobalInsts,
+ FeatureAtomicFMinFMaxF32FlatInsts, FeatureAtomicFMinFMaxF64FlatInsts,
+ FeatureVmemWriteVgprInOrder, FeatureVMemToLDSLoad, FeatureCubeInsts,
+ FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts, FeatureMsadInsts,
+ FeatureMqsadPkInsts, FeatureMqsadInsts,
+ FeatureCvtNormInsts, FeatureCvtPkNormVOP2Insts,
+ FeatureCvtPkNormVOP3Insts, FeatureDX10ClampAndIEEEMode, FeatureFlatOffsetBits12,
+ FeatureInstCacheLineSize64, FeatureMTBUFInsts, FeatureFormattedMUBUFInsts,
+ Feature32BitNumRecordsBufferResource
+ ]
+>;
+
+def FeatureGFX11 : GCNSubtargetFeatureGeneration<"GFX11",
+ "gfx11",
+ [FeatureFP64, FeatureAddressableLocalMemorySize65536,
+ FeatureLDSAllocGranularity1024,
+ FeatureLDSEncodingGranularity512,
+ FeatureHalfAddressablePhysicalLocalMemory, FeatureMIMG_R128,
+ FeatureSupportsWave32, FeatureSupportsWave64, FeatureSupportsWGP,
+ FeatureFlatAddressSpace, Feature16BitInsts,
+ FeatureInv2PiInlineImm, FeatureApertureRegs, FeatureGetDoorbellID,
+ FeatureCIInsts, FeatureGFX8Insts, FeatureGFX9Insts, FeatureGFX10Insts,
+ FeatureBVHRayTracingInsts, FeatureMSAALoadInsts,
+ FeatureGFX10_BEncoding, FeatureGFX10_3Insts,
+ FeatureGFX11Insts, FeatureVOP3PInsts, FeatureVOPDInsts, FeatureTrue16BitInsts,
+ FeatureMovrel, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp,
+ FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts,
+ FeatureAddNoCarryInsts, FeatureFmaMixInsts,
+ FeatureNoSdstCMPX, FeatureVscnt,
+ FeatureVOP3Literal, FeatureDPP8, FeatureExtendedImageInsts,
+ FeaturePermlane16Insts,
+ FeatureNoDataDepHazard, FeaturePkFmacF16Inst,
+ FeatureA16, FeatureFastDenormalF32, FeatureG16,
+ FeatureUnalignedBufferAccess, FeatureUnalignedScratchAccess,
+ FeatureUnalignedDSAccess, FeatureGDS, FeatureGWS,
+ FeatureDefaultComponentZero, FeatureMaxHardClauseLength32,
+ FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF32FlatInsts,
+ FeatureVmemWriteVgprInOrder, FeatureCubeInsts, FeatureLerpInst,
+ FeatureSadInsts, FeatureQsadInsts, FeatureMsadInsts, FeatureMqsadPkInsts,
+ FeatureMqsadInsts, FeatureCvtNormInsts,
+ FeatureCvtPkNormVOP2Insts, FeatureCvtPkNormVOP3Insts,
+ FeatureInstCacheLineSize128, FeatureMTBUFInsts, FeatureFormattedMUBUFInsts,
+ Feature32BitNumRecordsBufferResource,
+ FeatureMaxWavesPerEU16
+ ]
+>;
+
+def FeatureGFX12 : GCNSubtargetFeatureGeneration<"GFX12",
+ "gfx12",
+ [FeatureFP64, FeatureMIMG_R128,
+ FeatureSupportsWave32,
+ FeatureFlatAddressSpace, Feature16BitInsts,
+ FeatureInv2PiInlineImm, FeatureApertureRegs, FeatureGetDoorbellID,
+ FeatureCIInsts, FeatureGFX8Insts, FeatureGFX9Insts, FeatureGFX10Insts,
+ FeatureGFX10_BEncoding, FeatureGFX10_3Insts,
+ FeatureGFX11Insts, FeatureGFX12Insts, FeatureVOP3PInsts, FeatureVOPDInsts,
+ FeatureMovrel, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp,
+ FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts,
+ FeatureAddNoCarryInsts, FeatureFmaMixInsts,
+ FeatureNoSdstCMPX, FeatureVscnt,
+ FeatureVOP3Literal, FeatureDPP8,
+ FeaturePermlane16Insts,
+ FeatureNoDataDepHazard, FeaturePkFmacF16Inst,
+ FeatureA16, FeatureFastDenormalF32, FeatureG16,
+ FeatureUnalignedBufferAccess, FeatureUnalignedScratchAccess,
+ FeatureUnalignedDSAccess, FeatureTrue16BitInsts,
+ FeatureDefaultComponentBroadcast, FeatureMaxHardClauseLength32,
+ FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF32FlatInsts,
+ FeatureAtomicDsCondSubClampInsts, FeatureAtomicCondSubClampFlatInsts,
+ FeatureIEEEMinimumMaximumInsts, FeatureSALUMinimumMaximumInsts,
+ FeatureMinimum3Maximum3F32, FeatureMinimum3Maximum3F16,
+ FeatureAgentScopeFineGrainedRemoteMemoryAtomics, FeatureFlatOffsetBits24,
+ FeatureFlatSignedOffset, FeatureInstCacheLineSize128,
+ FeatureMaxWavesPerEU16
+ ]
+>;
+
+def FeatureGFX13 : GCNSubtargetFeatureGeneration<"GFX13",
+ "gfx13",
+ [FeatureFP64, FeatureMIMG_R128,
+ FeatureSupportsWave32, FeatureSupportsWave64, FeatureSupportsWGP,
+ FeatureFlatAddressSpace, Feature16BitInsts,
+ FeatureInv2PiInlineImm, FeatureApertureRegs, FeatureGetDoorbellID,
+ FeatureCIInsts, FeatureGFX8Insts, FeatureGFX9Insts, FeatureGFX10Insts,
+ FeatureMSAALoadInsts, FeatureGFX10_BEncoding, FeatureGFX10_3Insts,
+ FeatureGFX11Insts, FeatureGFX12Insts, FeatureGFX13Insts, FeatureVOP3PInsts,
+ FeatureVOPDInsts, FeatureMovrel, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp,
+ FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts,
+ FeatureAddNoCarryInsts, FeatureFmaMixInsts,
+ FeatureNoSdstCMPX, FeatureVscnt,
+ FeatureVOP3Literal, FeatureDPP8,
+ FeaturePermlane16Insts,
+ FeatureNoDataDepHazard, FeaturePkFmacF16Inst,
+ FeatureA16, FeatureFastDenormalF32, FeatureG16,
+ FeatureUnalignedBufferAccess, FeatureUnalignedScratchAccess, FeatureImageInsts,
+ FeatureExtendedImageInsts,
+ FeatureUnalignedDSAccess, FeatureTrue16BitInsts,
+ FeatureDefaultComponentBroadcast, FeatureMaxHardClauseLength32,
+ FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF32FlatInsts,
+ FeatureAtomicDsCondSubClampInsts, FeatureAtomicCondSubClampFlatInsts,
+ FeatureIEEEMinimumMaximumInsts, FeatureSALUMinimumMaximumInsts,
+ FeatureMinimum3Maximum3F32, FeatureMinimum3Maximum3F16,
+ FeatureAgentScopeFineGrainedRemoteMemoryAtomics, FeatureFlatOffsetBits24,
+ FeatureFlatSignedOffset, FeatureInstCacheLineSize128,
+ FeatureMTBUFInsts, FeatureFormattedMUBUFInsts,
+ Feature45BitNumRecordsBufferResource,
+ FeatureMaxWavesPerEU16
+ ]
+>;
+//===----------------------------------------------------------------------===//
+
+class FeatureSet<list<SubtargetFeature> Features_> {
+ list<SubtargetFeature> Features = Features_;
+}
+
+def FeatureISAVersion6_0_0 : FeatureSet<[FeatureSouthernIslands,
+ FeatureFastFMAF32,
+ FeatureHalfRate64Ops,
+ FeatureLDSBankCount32]>;
+
+def FeatureISAVersion6_0_1 : FeatureSet<
+ [FeatureSouthernIslands,
+ FeatureLDSBankCount32]>;
+
+def FeatureISAVersion6_0_2 : FeatureSet<
+ [FeatureSouthernIslands,
+ FeatureLDSBankCount32]>;
+
+def FeatureISAVersion7_0_0 : FeatureSet<
+ [FeatureSeaIslands,
+ FeatureLDSBankCount32]>;
+
+def FeatureISAVersion7_0_1 : FeatureSet<
+ [FeatureSeaIslands,
+ FeatureHalfRate64Ops,
+ FeatureLDSBankCount32,
+ FeatureFastFMAF32]>;
+
+def FeatureISAVersion7_0_2 : FeatureSet<
+ [FeatureSeaIslands,
+ FeatureLDSBankCount16,
+ FeatureFastFMAF32]>;
+
+def FeatureISAVersion7_0_3 : FeatureSet<
+ [FeatureSeaIslands,
+ FeatureLDSBankCount16]>;
+
+def FeatureISAVersion7_0_4 : FeatureSet<
+ [FeatureSeaIslands,
+ FeatureLDSBankCount32]>;
+
+def FeatureISAVersion7_0_5 : FeatureSet<
+ [FeatureSeaIslands,
+ FeatureLDSBankCount16]>;
+
+def FeatureISAVersion8_0_Common : FeatureSet<
+ [FeatureVolcanicIslands,
+ FeatureLDSBankCount32,
+ FeatureUnpackedD16VMem]>;
+
+def FeatureISAVersion8_0_1 : FeatureSet<
+ !listconcat(FeatureISAVersion8_0_Common.Features,
+ [FeatureFastFMAF32,
+ FeatureHalfRate64Ops,
+ FeatureXNACKOnOffModes])>;
+
+def FeatureISAVersion8_0_2 : FeatureSet<
+ !listconcat(FeatureISAVersion8_0_Common.Features,
+ [FeatureSGPRInitBug])>;
+
+def FeatureISAVersion8_0_3 : FeatureSet<
+ !listconcat(FeatureISAVersion8_0_Common.Features,
+ [])>;
+
+def FeatureISAVersion8_0_5 : FeatureSet<
+ !listconcat(FeatureISAVersion8_0_Common.Features,
+ [FeatureSGPRInitBug])>;
+
+def FeatureISAVersion8_1_0 : FeatureSet<
+ [FeatureVolcanicIslands,
+ FeatureLDSBankCount16,
+ FeatureXNACKOnOffModes,
+ FeatureImageStoreD16Bug,
+ FeatureImageGather4D16Bug]>;
+
+def FeatureISAVersion9_0_Common : FeatureSet<
+ [FeatureGFX9,
+ FeatureAddressableLocalMemorySize65536,
+ FeatureLDSAllocGranularity512,
+ FeatureLDSEncodingGranularity512,
+ FeatureLDSBankCount32,
+ FeatureImageInsts,
+ FeatureMadMacF32Insts]>;
+
+def FeatureISAVersion9_0_Consumer_Common : FeatureSet<
+ !listconcat(FeatureISAVersion9_0_Common.Features,
+ [FeatureImageGather4D16Bug,
+ FeatureDsSrc2Insts,
+ FeatureExtendedImageInsts,
+ FeatureGDS,
+ FeatureCubeInsts,
+ FeaturePopsExitingWaveID,
+ FeatureMaxWavesPerEU10])>;
+
+def FeatureISAVersion9_Generic : FeatureSet<
+ !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,
+ [FeatureRequiresCOV6])>;
+
+def FeatureISAVersion9_0_MI_Common : FeatureSet<
+ !listconcat(FeatureISAVersion9_0_Common.Features,
+ [FeatureCubeInsts,
+ FeatureFmaMixInsts,
+ FeatureDLInsts,
+ FeatureDot1Insts,
+ FeatureDot2Insts,
+ FeatureDot3Insts,
+ FeatureDot4Insts,
+ FeatureDot5Insts,
+ FeatureDot6Insts,
+ FeatureDot7Insts,
+ FeatureDot10Insts,
+ FeatureMAIInsts,
+ FeaturePkFmacF16Inst,
+ FeatureAtomicFaddNoRtnInsts,
+ FeatureSRAMECCOnOffModes])>;
+
+def FeatureISAVersion9_0_0 : FeatureSet<
+ !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,
+ [FeatureMadMixInsts])>;
+
+def FeatureISAVersion9_0_2 : FeatureSet<
+ !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,
+ [FeatureMadMixInsts])>;
+
+def FeatureISAVersion9_0_4 : FeatureSet<
+ !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,
+ [FeatureFmaMixInsts])>;
+
+def FeatureISAVersion9_0_6 : FeatureSet<
+ !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,
+ [FeatureHalfRate64Ops,
+ FeatureFmaMixInsts,
+ FeatureDLInsts,
+ FeatureDot1Insts,
+ FeatureDot2Insts,
+ FeatureDot7Insts,
+ FeatureDot10Insts,
+ FeatureSRAMECCOnOffModes])>;
+
+def FeatureISAVersion9_0_8 : FeatureSet<
+ !listconcat(FeatureISAVersion9_0_MI_Common.Features,
+ [FeatureGDS,
+ FeatureHalfRate64Ops,
+ FeatureDsSrc2Insts,
+ FeatureExtendedImageInsts,
+ FeatureAtomicBufferGlobalPkAddF16NoRtnInsts,
+ FeatureMFMAInlineLiteralBug,
+ FeatureImageGather4D16Bug,
+ FeatureMaxWavesPerEU10])>;
+
+def FeatureISAVersion9_0_9 : FeatureSet<
+ !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,
+ [FeatureMadMixInsts])>;
+
+def FeatureISAVersion9_0_A : FeatureSet<
+ !listconcat(FeatureISAVersion9_0_MI_Common.Features,
+ [FeatureGFX90AInsts,
+ FeatureAGPRAlloc,
+ FeatureTgSplitSupport,
+ FeatureRequiresAlignedVGPRs,
+ FeatureFmacF64Inst,
+ FeatureDPALU_DPP,
+ FeaturePackedFP32Ops,
+ FeatureAtomicFaddRtnInsts,
+ FeatureAtomicBufferGlobalPkAddF16Insts,
+ FeaturePackedTID,
+ FeatureFullRate64Ops,
+ FeatureBackOffBarrier,
+ FeatureKernargPreload,
+ FeatureAtomicFMinFMaxF64GlobalInsts,
+ FeatureAtomicFMinFMaxF64FlatInsts,
+ FeatureFlatBufferGlobalAtomicFaddF64Inst,
+ FeatureLdsAtomicAddF64,
+ FeatureMaxWavesPerEU8
+ ])>;
+
+def FeatureISAVersion9_0_C : FeatureSet<
+ !listconcat(FeatureISAVersion9_0_Consumer_Common.Features,
+ [FeatureMadMixInsts])>;
+
+def FeatureISAVersion9_4_Common : FeatureSet<
+ [FeatureGFX9,
+ FeatureGFX90AInsts,
+ FeatureAGPRAlloc,
+ FeatureTgSplitSupport,
+ FeatureGFX940Insts,
+ FeatureBufferInvInst,
+ FeatureRequiresAlignedVGPRs,
+ FeatureFmaMixInsts,
+ FeatureDLInsts,
+ FeatureFmacF64Inst,
+ FeatureDot1Insts,
+ FeatureDot2Insts,
+ FeatureDot3Insts,
+ FeatureDot4Insts,
+ FeatureDot5Insts,
+ FeatureDot6Insts,
+ FeatureDot7Insts,
+ FeatureDot10Insts,
+ FeatureAtomicDsPkAdd16Insts,
+ FeatureAtomicFlatPkAdd16Insts,
+ FeatureDPALU_DPP,
+ FeaturePackedFP32Ops,
+ FeatureMAIInsts,
+ FeaturePkFmacF16Inst,
+ FeatureAtomicFaddRtnInsts,
+ FeatureAtomicFaddNoRtnInsts,
+ FeatureAtomicBufferGlobalPkAddF16Insts,
+ FeatureAtomicGlobalPkAddBF16Inst,
+ FeatureFlatAtomicFaddF32Inst,
+ FeatureSRAMECCOnOffModes,
+ FeaturePackedTID,
+ FeatureArchitectedFlatScratch,
+ FeatureFullRate64Ops,
+ FeatureBackOffBarrier,
+ FeatureKernargPreload,
+ FeatureAtomicFMinFMaxF64GlobalInsts,
+ FeatureAtomicFMinFMaxF64FlatInsts,
+ FeatureAgentScopeFineGrainedRemoteMemoryAtomics,
+ FeatureMemoryAtomicFAddF32DenormalSupport,
+ FeatureFlatBufferGlobalAtomicFaddF64Inst,
+ FeatureLdsAtomicAddF64,
+ FeatureLshlAddU64Inst,
+ FeatureVMovB64Inst,
+ FeatureMaxWavesPerEU8
+ ]>;
+
+def FeatureISAVersion9_5_Common : FeatureSet<
+ !listconcat(FeatureISAVersion9_4_Common.Features,
+ [FeatureAddressableLocalMemorySize163840,
+ FeatureLDSAllocGranularity1280,
+ FeatureLDSEncodingGranularity1280,
+ FeatureLDSBankCount64,
+ FeatureFP8Insts,
+ FeatureFP8ConversionInsts,
+ FeatureCvtFP8SDWASrcSel,
+ FeatureOCPFP8ConversionInsts,
+ FeatureGFX950Insts,
+ FeaturePrngInst,
+ FeatureBF16ConversionInsts,
+ FeatureBitOp3Insts,
+ FeatureFP8ConversionScaleInsts,
+ FeatureBF8ConversionScaleInsts,
+ FeatureFP4ConversionScaleInsts,
+ FeatureFP6BF6ConversionScaleInsts,
+ FeatureDot12Insts,
+ FeatureDot13Insts,
+ FeatureAtomicBufferPkAddBF16Inst
+ ])>;
+
+def FeatureISAVersion9_4_2 : FeatureSet<
+ !listconcat(FeatureISAVersion9_4_Common.Features,
+ [
+ FeatureAddressableLocalMemorySize65536,
+ FeatureLDSAllocGranularity512,
+ FeatureLDSEncodingGranularity512,
+ FeatureLDSBankCount32,
+ FeatureFP8Insts,
+ FeatureFP8ConversionInsts,
+ FeatureCvtFP8SDWASrcSel,
+ FeatureCvtFP8VOP1Bug,
+ FeatureXF32Insts
+ ])>;
+
+def FeatureISAVersion9_4_Generic : FeatureSet<
+ !listconcat(FeatureISAVersion9_4_Common.Features,
+ [FeatureAddressableLocalMemorySize65536,
+ FeatureLDSAllocGranularity1280,
+ FeatureLDSEncodingGranularity1280,
+ FeatureLDSBankCount32,
+ FeatureRequiresCOV6])>;
+
+def FeatureISAVersion9_5_0 : FeatureSet<FeatureISAVersion9_5_Common.Features>;
+
+def FeatureISAVersion10_Common : FeatureSet<
+ [FeatureGFX10,
+ FeatureLDSBankCount32,
+ FeatureDLInsts,
+ FeatureNSAEncoding,
+ FeatureBackOffBarrier,
+ FeatureMsadInsts,
+ FeatureMqsadPkInsts,
+ FeatureMqsadInsts,
+ FeaturePopsExitingWaveID]>;
+
+def FeatureISAVersion10_1_Common : FeatureSet<
+ !listconcat(FeatureISAVersion10_Common.Features,
+ [FeatureLDSAllocGranularity512,
+ FeatureScalarStores,
+ FeatureScalarAtomics,
+ FeatureScalarFlatScratchInsts,
+ FeatureGetWaveIdInst,
+ FeatureMadMacF32Insts,
+ FeatureDsSrc2Insts,
+ FeatureLDSMisalignedBug,
+ FeatureXNACKOnOffModes,
+ // gfx101x bugs
+ FeatureVcmpxPermlaneHazard,
+ FeatureVMEMtoScalarWriteHazard,
+ FeatureSMEMtoVectorWriteHazard,
+ FeatureInstFwdPrefetchBug,
+ FeatureVcmpxExecWARHazard,
+ FeatureLdsBranchVmemWARHazard,
+ FeatureNSAtoVMEMBug,
+ FeatureNSAClauseBug,
+ FeatureOffset3fBug,
+ FeatureFlatSegmentOffsetBug,
+ FeatureNegativeUnalignedScratchOffsetBug,
+ FeatureMaxWavesPerEU20])>;
+
+def FeatureISAVersion10_1_Generic : FeatureSet<
+ !listconcat(FeatureISAVersion10_1_Common.Features,
+ [FeatureRequiresCOV6])>;
+
+def FeatureISAVersion10_1_0 : FeatureSet<
+ !listconcat(FeatureISAVersion10_1_Common.Features,
+ [])>;
+
+def FeatureISAVersion10_1_1 : FeatureSet<
+ !listconcat(FeatureISAVersion10_1_Common.Features,
+ [FeatureDot1Insts,
+ FeatureDot2Insts,
+ FeatureDot5Insts,
+ FeatureDot6Insts,
+ FeatureDot7Insts,
+ FeatureDot10Insts])>;
+
+def FeatureISAVersion10_1_2 : FeatureSet<
+ !listconcat(FeatureISAVersion10_1_Common.Features,
+ [FeatureDot1Insts,
+ FeatureDot2Insts,
+ FeatureDot5Insts,
+ FeatureDot6Insts,
+ FeatureDot7Insts,
+ FeatureDot10Insts])>;
+
+def FeatureISAVersion10_1_3 : FeatureSet<
+ !listconcat(FeatureISAVersion10_1_Common.Features,
+ [FeatureBVHRayTracingInsts,
+ FeatureMSAALoadInsts])>;
+
+def FeatureISAVersion10_3_0 : FeatureSet<
+ !listconcat(FeatureISAVersion10_Common.Features,
+ [FeatureLDSAllocGranularity1024,
+ FeatureBVHRayTracingInsts,
+ FeatureMSAALoadInsts,
+ FeatureGFX10_BEncoding,
+ FeatureGFX10_3Insts,
+ FeatureDot1Insts,
+ FeatureDot2Insts,
+ FeatureDot5Insts,
+ FeatureDot6Insts,
+ FeatureDot7Insts,
+ FeatureDot10Insts,
+ FeatureShaderCyclesRegister,
+ FeatureMaxWavesPerEU16])>;
+
+def FeatureISAVersion10_3_Generic: FeatureSet<
+ !listconcat(FeatureISAVersion10_3_0.Features,
+ [FeatureRequiresCOV6])>;
+
+def FeatureISAVersion11_Common : FeatureSet<
+ [FeatureGFX11,
+ FeatureBackOffBarrier,
+ FeatureLDSBankCount32,
+ FeatureDLInsts,
+ FeatureDot7Insts,
+ FeatureDot8Insts,
+ FeatureDot9Insts,
+ FeatureDot10Insts,
+ FeatureDot12Insts,
+ FeatureNSAEncoding,
+ FeaturePartialNSAEncoding,
+ FeatureShaderCyclesRegister,
+ FeatureArchitectedFlatScratch,
+ FeatureAtomicFaddRtnInsts,
+ FeatureAtomicFaddNoRtnInsts,
+ FeatureFlatAtomicFaddF32Inst,
+ FeatureImageInsts,
+ FeaturePackedTID,
+ FeatureVcmpxPermlaneHazard,
+ FeatureMemoryAtomicFAddF32DenormalSupport,
+ FeatureRealTrue16Insts,
+ FeatureD16Writes32BitVgpr,
+ FeatureMsadInsts,
+ FeatureMqsadPkInsts,
+ FeatureMqsadInsts,
+]>;
+
+// There are few workarounds that need to be added to all targets. This
+// pessimizes codegen a bit on the generic GFX11 target. This generic target
+// does not include GFX11_7 due to incompatible changes.
+def FeatureISAVersion11_Generic: FeatureSet<
+ !listconcat(FeatureISAVersion11_Common.Features,
+ [FeatureMSAALoadDstSelBug,
+ FeatureVALUTransUseHazard,
+ FeatureVOPDInterlockHazard,
+ FeatureUserSGPRInit16Bug,
+ FeatureMADIntraFwdBug,
+ FeaturePrivEnabledTrap2NopBug,
+ FeatureRequiresCOV6,
+ FeatureRequiredExportPriority,
+ FeatureDot5Insts,
+ FeatureWMMA256bInsts,
+ FeatureDX10ClampAndIEEEMode])>;
+
+def FeatureISAVersion11_0_Common : FeatureSet<
+ !listconcat(FeatureISAVersion11_Common.Features,
+ [FeatureMSAALoadDstSelBug,
+ FeatureVALUTransUseHazard,
+ FeatureVOPDInterlockHazard,
+ FeatureMADIntraFwdBug,
+ FeaturePrivEnabledTrap2NopBug,
+ FeatureDot5Insts,
+ FeatureWMMA256bInsts,
+ FeatureDX10ClampAndIEEEMode])>;
+
+def FeatureISAVersion11_0_0 : FeatureSet<
+ !listconcat(FeatureISAVersion11_0_Common.Features,
+ [Feature1536VGPRs,
+ FeatureUserSGPRInit16Bug])>;
+
+def FeatureISAVersion11_0_1 : FeatureSet<
+ !listconcat(FeatureISAVersion11_0_Common.Features,
+ [Feature1536VGPRs])>;
+
+def FeatureISAVersion11_0_2 : FeatureSet<
+ !listconcat(FeatureISAVersion11_0_Common.Features,
+ [FeatureUserSGPRInit16Bug])>;
+
+def FeatureISAVersion11_0_3 : FeatureSet<
+ !listconcat(FeatureISAVersion11_0_Common.Features,
+ [])>;
+
+def FeatureISAVersion11_5_Common : FeatureSet<
+ !listconcat(FeatureISAVersion11_Common.Features,
+ [FeatureVOPDInterlockHazard,
+ FeatureSALUFloatInsts,
+ FeatureDPPSrc1SGPR,
+ FeatureRequiredExportPriority,
+ FeatureDot5Insts,
+ FeatureWMMA256bInsts,
+ FeatureDX10ClampAndIEEEMode])>;
+
+def FeatureISAVersion11_5_0 : FeatureSet<
+ !listconcat(FeatureISAVersion11_5_Common.Features,
+ [FeaturePointSampleAccel])>;
+
+def FeatureISAVersion11_5_1 : FeatureSet<
+ !listconcat(FeatureISAVersion11_5_Common.Features,
+ [Feature1536VGPRs,
+ FeaturePointSampleAccel])>;
+
+def FeatureISAVersion11_5_2 : FeatureSet<
+ !listconcat(FeatureISAVersion11_5_Common.Features,
+ [FeaturePointSampleAccel])>;
+
+def FeatureISAVersion11_7_Common : FeatureSet<
+ !listconcat(FeatureISAVersion11_Common.Features,
+ [FeatureGFX11_7Insts,
+ FeatureSALUFloatInsts,
+ FeatureDPPSrc1SGPR,
+ FeatureFP8ConversionInsts,
+ FeatureCvtFP8ByteSel,
+ FeatureOCPFP8ConversionInsts,
+ FeatureDot11Insts,
+ FeatureWMMA128bInsts,
+ FeatureSWMMACGfx1200Insts,
+ FeatureIEEEMinimumMaximumInsts,
+ FeatureMinimum3Maximum3F32,
+ FeatureMinimum3Maximum3F16,
+ FeatureNoF16PseudoScalarTransInlineConstants])>;
+
+def FeatureISAVersion11_7_Generic: FeatureSet<
+ !listconcat(FeatureISAVersion11_7_Common.Features,
+ [FeatureRequiresCOV6])>;
+
+def FeatureISAVersion12 : FeatureSet<
+ [FeatureGFX12,
+ Feature32BitNumRecordsBufferResource,
+ FeatureSupportsWave64, FeatureSupportsWGP,
+ FeatureBackOffBarrier,
+ FeatureAddressableLocalMemorySize65536,
+ FeatureLDSAllocGranularity1024,
+ FeatureLDSEncodingGranularity512,
+ FeatureHalfAddressablePhysicalLocalMemory,
+ FeatureLDSBankCount32,
+ FeatureDLInsts,
+ FeatureDot7Insts,
+ FeatureDot8Insts,
+ FeatureDot9Insts,
+ FeatureDot10Insts,
+ FeatureDot11Insts,
+ FeatureDot12Insts,
+ FeatureNSAEncoding,
+ FeaturePartialNSAEncoding,
+ FeatureShaderCyclesHiLoRegisters,
+ FeatureArchitectedFlatScratch,
+ FeatureArchitectedSGPRs,
+ FeatureAtomicFaddRtnInsts,
+ FeatureAtomicFaddNoRtnInsts,
+ FeatureAtomicDsPkAdd16Insts,
+ FeatureAtomicFlatPkAdd16Insts,
+ FeatureAtomicBufferGlobalPkAddF16Insts,
+ FeatureAtomicGlobalPkAddBF16Inst,
+ FeatureAtomicBufferPkAddBF16Inst,
+ FeatureFlatAtomicFaddF32Inst,
+ FeatureImageInsts,
+ FeatureExtendedImageInsts,
+ FeatureBVHRayTracingInsts,
+ FeatureMSAALoadInsts,
+ FeatureFP8ConversionInsts,
+ FeatureCvtFP8ByteSel,
+ FeatureOCPFP8ConversionInsts,
+ FeatureWMMA128bInsts,
+ FeatureSWMMACGfx1200Insts,
+ FeatureIEEEMinimumMaximumInsts,
+ FeaturePackedTID,
+ FeatureVcmpxPermlaneHazard,
+ FeatureSALUFloatInsts,
+ FeaturePseudoScalarTrans,
+ FeatureRestrictedSOffset,
+ FeatureScalarDwordx3Loads,
+ FeatureDPPSrc1SGPR,
+ FeatureMaxHardClauseLength32,
+ Feature1536VGPRs,
+ FeatureMemoryAtomicFAddF32DenormalSupport,
+ FeatureBVHDualAndBVH8Insts,
+ FeatureWaitsBeforeSystemScopeStores,
+ FeatureD16Writes32BitVgpr,
+ FeatureCubeInsts,
+ FeatureLerpInst,
+ FeatureSadInsts,
+ FeatureQsadInsts,
+ FeatureMsadInsts,
+ FeatureMqsadPkInsts,
+ FeatureMqsadInsts,
+ FeatureCvtNormInsts,
+ FeatureCvtPkNormVOP2Insts,
+ FeatureCvtPkNormVOP3Insts,
+ FeatureSmemPrefetchInsts,
+ FeatureNoF16PseudoScalarTransInlineConstants,
+ FeatureRealTrue16Insts,
+ FeatureMTBUFInsts,
+ FeatureFormattedMUBUFInsts,
+ FeatureDataCacheLineSize128
+ ]>;
+
+def FeatureISAVersion12_50_Common : FeatureSet<
+ [FeatureGFX12,
+ FeatureGFX1250Insts,
+ FeatureBackOffBarrier,
+ FeatureRequiresAlignedVGPRs,
+ FeatureCuMode,
+ Feature1024AddressableVGPRs,
+ Feature64BitLiterals,
+ FeatureDLInsts,
+ FeatureFmacF64Inst,
+ FeaturePackedFP32SingleSGPROps,
+ FeatureDot7Insts,
+ FeatureDot8Insts,
+ FeatureWavefrontSize32,
+ FeatureShaderCyclesHiLoRegisters,
+ FeatureArchitectedFlatScratch,
+ FeatureArchitectedSGPRs,
+ FeatureFlatGVSMode,
+ FeatureAtomicFaddRtnInsts,
+ FeatureAtomicFaddNoRtnInsts,
+ FeatureAtomicDsPkAdd16Insts,
+ FeatureAtomicFlatPkAdd16Insts,
+ FeatureAtomicBufferGlobalPkAddF16Insts,
+ FeatureAtomicGlobalPkAddBF16Inst,
+ FeatureAtomicBufferPkAddBF16Inst,
+ FeatureFlatAtomicFaddF32Inst,
+ FeatureFP8ConversionInsts,
+ FeatureCvtFP8ByteSel,
+ FeatureOCPFP8ConversionInsts,
+ FeatureFP8F16ConversionInsts,
+ FeatureF16FP8ConversionInsts,
+ FeaturePackedTID,
+ FeatureVcmpxPermlaneHazard,
+ FeatureSALUFloatInsts,
+ FeaturePseudoScalarTrans,
+ FeatureRestrictedSOffset,
+ FeatureScalarDwordx3Loads,
+ FeatureDPPSrc1SGPR,
+ FeatureBitOp3Insts,
+ FeatureTanhInsts,
+ FeatureTensorCvtLutInsts,
+ FeatureTDMInsts,
+ FeatureVPermPk16Hazard,
+ FeatureTransposeLoadF4F6Insts,
+ FeatureBF16TransInsts,
+ FeatureBF16ConversionInsts,
+ FeatureBF16PackedInsts,
+ FeatureBF16InlineConstFromUpperFP32,
+ FeatureCvtPkF16F32Inst,
+ FeatureCvtSrPkBF16F32Inst,
+ FeatureFmaMixBF16Insts,
+ FeatureMin3Max3PKF16,
+ FeatureMinimum3Maximum3PKF16,
+ FeaturePrngInst,
+ FeaturePermlane16Swap,
+ FeatureAshrPkInsts,
+ FeatureSRAMECCOnOffModes,
+ FeatureMaxHardClauseLength63,
+ FeatureWaitXcnt,
+ FeatureAsynccnt,
+ FeatureAtomicFMinFMaxF64GlobalInsts,
+ FeatureAtomicFMinFMaxF64FlatInsts,
+ FeatureFlatBufferGlobalAtomicFaddF64Inst,
+ FeatureMemoryAtomicFAddF32DenormalSupport,
+ FeatureEmulatedSystemScopeAtomics,
+ FeatureGloballyAddressableScratch,
+ FeatureKernargPreload,
+ FeatureVmemPrefInsts,
+ FeatureLshlAddU64Inst,
+ FeatureAddSubU64Insts,
+ FeatureMadU32Inst,
+ FeatureMadNC64_32Insts,
+ FeatureAddMinMaxInsts,
+ FeatureMinMaxI64Insts,
+ FeaturePkAddMinMaxInsts,
+ FeatureLdsBarrierArriveAtomic,
+ FeatureLdsAtomicAddF64,
+ FeatureSetPrioIncWgInst,
+ FeatureSWakeupBarrier,
+ Feature45BitNumRecordsBufferResource,
+ FeatureClusters,
+ FeatureD16Writes32BitVgpr,
+ FeatureMcastLoadInsts,
+ FeatureAsyncLoadToLDSInsts,
+ FeatureAsyncStoreFromLDSInsts,
+ FeatureSmemPrefetchInsts,
+ FeatureNoF16PseudoScalarTransInlineConstants,
+ FeatureRealTrue16Insts,
+ FeatureVMovB64Inst,
+ FeatureVMulU64Inst,
+]>;
+
+def FeatureISAVersion12_50_STRICT : FeatureSet<
+ !listconcat(FeatureISAVersion12_50_Common.Features,
+ [FeatureVOPD3F64OPYSrcHazard,
+ FeatureGFX1250_STRICT,
+ FeatureAddressableLocalMemorySize327680,
+ FeatureLDSAllocGranularity2048,
+ FeatureLDSEncodingGranularity2048,
+ FeatureLDSBankCount64,
+ FeatureWMMAN16Insts,
+ FeatureVOP3PX2IncrementsVaVdstTwice,
+ FeatureSetregVGPRMSBFixup,
+ FeatureCubeInsts,
+ FeatureLerpInst,
+ FeatureSadInsts,
+ FeatureQsadInsts,
+ FeatureMsadInsts,
+ FeatureMqsadPkInsts,
+ FeatureMqsadInsts,
+ FeatureCvtNormInsts,
+ FeatureCvtPkNormVOP2Insts,
+ FeatureCvtPkNormVOP3Insts,
+ FeatureSWMMACGfx1200Insts,
+ FeatureSWMMACGfx1250Insts,
+ FeatureTransCoexecutionHazard,
+ FeatureGFX1250VALUBlockingCycles,
+ FeatureWMMACoexecutionHazards,
+ FeatureDataCacheLineSize128,
+ FeatureRequiresInitialUnclausedVmem,
+ FeatureNeedsTDMDrain,
+ FeatureNeedsAligned2addrDS,
+ FeatureNoSleepForever,
+ FeatureSlowMaxMinMulI64Insts,
+ FeatureSupportsXNACK,
+ ])>;
+
+def FeatureISAVersion12_50 : FeatureSet<
+ !listconcat(FeatureISAVersion12_50_Common.Features,
+ [FeatureVOPD3F64OPYSrcHazard,
+ FeatureAddressableLocalMemorySize327680,
+ FeatureLDSAllocGranularity2048,
+ FeatureLDSEncodingGranularity2048,
+ FeatureLDSBankCount64,
+ FeatureWMMAN16Insts,
+ FeatureWMMAF4Insts,
+ FeatureWMMA16x16x128FP8Insts,
+ FeatureVOP3PX2IncrementsVaVdstTwice,
+ FeatureSetregVGPRMSBFixup,
+ FeatureCubeInsts,
+ FeatureLerpInst,
+ FeatureSadInsts,
+ FeatureQsadInsts,
+ FeatureMsadInsts,
+ FeatureMqsadPkInsts,
+ FeatureMqsadInsts,
+ FeatureCvtNormInsts,
+ FeatureCvtPkNormVOP2Insts,
+ FeatureCvtPkNormVOP3Insts,
+ FeatureBlock16ConversionScaleInsts,
+ FeatureFP8E5M3Insts,
+ FeatureSWMMACGfx1200Insts,
+ FeatureSWMMACGfx1250Insts,
+ FeatureTransCoexecutionHazard,
+ FeatureGFX1250VALUBlockingCycles,
+ FeatureWMMACoexecutionHazards,
+ FeatureDataCacheLineSize128,
+ FeatureRequiresInitialUnclausedVmem,
+ FeatureNeedsTDMDrain,
+ FeatureNoSleepForever,
+ FeatureSlowMaxMinMulI64Insts,
+ FeatureSupportsXNACK,
+ ])>;
+
+def FeatureISAVersion12_51 : FeatureSet<
+ !listconcat(FeatureISAVersion12_50_Common.Features,
+ [FeatureVOPD3F64OPYSrcHazard,
+ FeatureAddressableLocalMemorySize327680,
+ FeatureLDSAllocGranularity2048,
+ FeatureLDSEncodingGranularity2048,
+ FeatureLDSBankCount64,
+ FeatureWMMAN16Insts,
+ FeatureWMMAF4Insts,
+ FeatureWMMA16x16x128FP8Insts,
+ FeatureFullRate64Ops,
+ FeatureVOP3PX2IncrementsVaVdstTwice,
+ FeatureDPALU_DPP,
+ FeaturePackedFP64SingleSGPROps,
+ FeaturePackedU64SingleSGPROps,
+ FeatureGFX1251GEMMInsts,
+ FeatureGFX125xLowestRateWMMA,
+ FeatureCubeInsts,
+ FeatureLerpInst,
+ FeatureSadInsts,
+ FeatureQsadInsts,
+ FeatureMsadInsts,
+ FeatureMqsadPkInsts,
+ FeatureMqsadInsts,
+ FeatureCvtNormInsts,
+ FeatureCvtPkNormVOP2Insts,
+ FeatureCvtPkNormVOP3Insts,
+ FeatureBlock16ConversionScaleInsts,
+ FeatureFP8E5M3Insts,
+ FeatureSWMMACGfx1200Insts,
+ FeatureSWMMACGfx1250Insts,
+ FeatureUseAddPC64Inst,
+ FeatureTransCoexecutionHazard,
+ FeatureWMMACoexecutionHazards,
+ FeatureDataCacheLineSize128,
+ FeatureRequiresInitialUnclausedVmem,
+ FeatureNeedsTDMDrain,
+ FeatureSupportsXNACK,
+ ])>;
+
+def FeatureISAVersion12_Generic: FeatureSet<
+ !listconcat(FeatureISAVersion12.Features,
+ [FeatureRequiresCOV6])>;
+
+def FeatureISAVersion12_5_Generic: FeatureSet<
+ !listconcat(FeatureISAVersion12_50_Common.Features,
+ [FeatureVOPD3F64OPYSrcHazard,
+ FeatureAddressableLocalMemorySize327680,
+ FeatureLDSAllocGranularity2048,
+ FeatureLDSEncodingGranularity2048,
+ FeatureLDSBankCount64,
+ FeatureBlock16ConversionScaleInsts,
+ FeatureFP8E5M3Insts,
+ FeatureSetregVGPRMSBFixup,
+ FeatureRequiresCOV6,
+ FeatureGFX125xLowestRateWMMA,
+ FeatureTransCoexecutionHazard,
+ FeatureGFX1250VALUBlockingCycles,
+ FeatureWMMACoexecutionHazards,
+ FeatureDataCacheLineSize128,
+ FeatureRequiresInitialUnclausedVmem,
+ FeatureNeedsTDMDrain,
+ FeatureSupportsXNACK,
+ ])>;
+
+def FeatureISAVersion13 : FeatureSet<
+ [FeatureGFX13,
+ FeatureGFX1250Insts,
+ FeatureWaveMatchInsts,
+ FeatureExclusiveScanInsts,
+ FeatureAddressableLocalMemorySize196608,
+ FeatureLDSAllocGranularity1024,
+ FeatureLDSEncodingGranularity1024,
+ Feature64BitLiterals,
+ FeatureLDSBankCount32,
+ FeatureDLInsts,
+ FeatureFmacF64Inst,
+ FeatureDot7Insts,
+ FeatureDot8Insts,
+ FeatureNSAEncoding,
+ FeaturePartialNSAEncoding,
+ FeatureShaderCyclesRegister,
+ FeatureArchitectedFlatScratch,
+ FeatureArchitectedSGPRs,
+ FeatureFlatGVSMode,
+ FeatureAtomicFaddRtnInsts,
+ FeatureAtomicFaddNoRtnInsts,
+ FeatureAtomicDsPkAdd16Insts,
+ FeatureAtomicFlatPkAdd16Insts,
+ FeatureAtomicBufferGlobalPkAddF16Insts,
+ FeatureAtomicGlobalPkAddBF16Inst,
+ FeatureAtomicBufferPkAddBF16Inst,
+ FeatureFlatAtomicFaddF32Inst,
+ FeatureFP8ConversionInsts,
+ FeatureCvtFP8ByteSel,
+ FeatureOCPFP8ConversionInsts,
+ FeatureFP8F16ConversionInsts,
+ FeatureF16FP8ConversionInsts,
+ FeaturePackedTID,
+ FeatureVcmpxPermlaneHazard,
+ FeatureSALUFloatInsts,
+ FeaturePseudoScalarTrans,
+ FeatureRestrictedSOffset,
+ FeatureScalarDwordx3Loads,
+ FeatureDPPSrc1SGPR,
+ FeatureBitOp3Insts,
+ FeatureTanhInsts,
+ FeatureTensorCvtLutInsts,
+ FeatureBF16TransInsts,
+ FeatureBF16ConversionInsts,
+ FeatureBF16PackedInsts,
+ FeatureBF16InlineConstFromUpperFP32,
+ FeaturePrngInst,
+ FeaturePermlane16Swap,
+ FeatureAshrPkInsts,
+ FeatureAtomicFMinFMaxF64GlobalInsts,
+ FeatureAtomicFMinFMaxF64FlatInsts,
+ FeatureFmaMixBF16Insts,
+ FeatureGloballyAddressableScratch,
+ FeatureCvtPkF16F32Inst,
+ FeatureCvtSrPkBF16F32Inst,
+ FeatureF16BF16ToFP6BF6ConversionScaleInsts,
+ FeatureF32ToFP6BF6ConversionScaleInsts,
+ FeatureFP6BF6ToF16BF16F32ConversionScaleInsts,
+ FeatureIEEEMinimumMaximumInsts,
+ FeatureSWakeupBarrier,
+ FeatureClusters,
+ FeatureCubeInsts,
+ FeatureLerpInst,
+ FeatureSadInsts,
+ FeatureQsadInsts,
+ FeatureMsadInsts,
+ FeatureMqsadPkInsts,
+ FeatureMqsadInsts,
+ FeatureCvtNormInsts,
+ FeatureCvtPkNormVOP2Insts,
+ FeatureCvtPkNormVOP3Insts,
+ FeatureBlock16ConversionScaleInsts,
+ FeatureSmemPrefetchInsts,
+ FeatureAsyncLoadToLDSInsts,
+ FeatureIndexedResources
+]>;
+
+def FeatureISAVersion13_Generic: FeatureSet<
+ !listconcat(FeatureISAVersion13.Features,
+ [FeatureRequiresCOV6])>;
+
+// The subset of SubtargetFeatures which should be exposed outside of
+// the backend.
+def AMDGPUFrontendVisibleFeatures {
+ list<SubtargetFeature> Features = [
+ Feature16BitInsts, FeatureAddMinMaxInsts, FeatureAshrPkInsts,
+ FeatureAsyncLoadToLDSInsts, FeatureAsyncStoreFromLDSInsts, FeatureAsynccnt,
+ FeatureAtomicBufferGlobalPkAddF16Insts, FeatureAtomicBufferPkAddBF16Inst, FeatureAtomicDsPkAdd16Insts,
+ FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF64GlobalInsts, FeatureAtomicFaddNoRtnInsts,
+ FeatureAtomicFaddRtnInsts, FeatureAtomicFlatPkAdd16Insts, FeatureAtomicGlobalPkAddBF16Inst,
+ FeatureBF16ConversionInsts,
+ FeatureBF16PackedInsts, FeatureBF16TransInsts, FeatureBF8ConversionScaleInsts,
+ FeatureBVHRayTracingInsts, FeatureBitOp3Insts, FeatureBufferInvInst,
+ FeatureCIInsts,
+ FeatureClusters, FeatureCubeInsts, FeatureCvtPkNormVOP2Insts,
+ FeatureCvtSrPkBF16F32Inst, FeatureDLInsts, FeatureDPP,
+ FeatureDot10Insts, FeatureDot11Insts, FeatureDot12Insts,
+ FeatureDot13Insts, FeatureDot1Insts, FeatureDot2Insts,
+ FeatureDot3Insts, FeatureDot4Insts, FeatureDot5Insts,
+ FeatureDot6Insts, FeatureDot7Insts, FeatureDot8Insts,
+ FeatureDot9Insts, FeatureExclusiveScanInsts, FeatureExtendedImageInsts, FeatureF16BF16ToFP6BF6ConversionScaleInsts,
+ FeatureF32ToF16BF16ConversionSRInsts, FeatureF32ToFP6BF6ConversionScaleInsts, FeatureFP4ConversionScaleInsts,
+ FeatureFP6BF6ConversionScaleInsts, FeatureFP8ConversionInsts, FeatureFP8ConversionScaleInsts,
+ FeatureBlock16ConversionScaleInsts, FeatureFP8E5M3Insts, FeatureFP8Insts,
+ FeatureFP6BF6ToF16BF16F32ConversionScaleInsts,
+ FeatureFlatBufferGlobalAtomicFaddF64Inst, FeatureFlatGlobalInsts,
+ FeatureGFX10Insts, FeatureGFX10_3Insts, FeatureGFX11Insts,
+ FeatureGFX1250Insts, FeatureGFX1251GEMMInsts, FeatureGFX12Insts,
+ FeatureGFX13Insts, FeatureGFX8Insts, FeatureGFX90AInsts,
+ FeatureGFX940Insts, FeatureGFX950Insts, FeatureGFX9Insts,
+ FeatureGWS, FeatureImageInsts, FeatureLdsBarrierArriveAtomic,
+ FeatureLerpInst, FeatureMAIInsts, FeatureMcastLoadInsts,
+ FeatureMqsadInsts, FeatureMqsadPkInsts, FeatureMsadInsts,
+ FeatureOCPFP8ConversionInsts, FeaturePermlane16Insts, FeaturePermlane16Swap,
+ FeaturePermlane32Swap, FeaturePkAddMinMaxInsts, FeaturePrngInst,
+ FeatureQsadInsts, FeatureSMemRealTime, FeatureSMemTimeInst,
+ FeatureSWMMACGfx1200Insts, FeatureSWMMACGfx1250Insts, FeatureSWakeupBarrier,
+ FeatureSadInsts, FeatureSetPrioIncWgInst, FeatureSmemPrefetchInsts,
+ FeatureTDMInsts, FeatureTanhInsts, FeatureTensorCvtLutInsts,
+ FeatureTransposeLoadF4F6Insts, FeatureVMemToLDSLoad, FeatureVmemPrefInsts,
+ FeatureWaveMatchInsts, FeatureWMMA128bInsts, FeatureWMMA256bInsts,
+ FeatureWMMAN16Insts, FeatureWMMAF4Insts, FeatureWMMA16x16x128FP8Insts, FeatureXF32Insts,
+ FeatureWavefrontSize32, FeatureWavefrontSize64, FeatureSupportsWGP,
+ FeatureSupportsWave32, FeatureFastFMAF32, FeatureFastDenormalF32,
+ FeatureSupportsXNACK, FeatureSupportsSRAMECC, FeatureXNACKOnOffModes,
+ FeatureSRAMECCOnOffModes,
+ FeatureSGPRInitBug, FeatureApertureRegs, FeatureGetDoorbellID,
+ FeatureAGPRAlloc, Feature1536VGPRs, Feature1024AddressableVGPRs,
+ FeatureHalfAddressablePhysicalLocalMemory, FeatureBackOffBarrier,
+ FeatureLDSAllocGranularity256,
+ FeatureLDSAllocGranularity512,
+ FeatureLDSAllocGranularity1024,
+ FeatureLDSAllocGranularity1280,
+ FeatureLDSAllocGranularity2048,
+ FeatureLDSEncodingGranularity256,
+ FeatureLDSEncodingGranularity512,
+ FeatureLDSEncodingGranularity1024,
+ FeatureLDSEncodingGranularity1280,
+ FeatureLDSEncodingGranularity2048,
+ ];
+}
More information about the llvm-commits
mailing list