[llvm-branch-commits] [llvm] AMDGPU: Move GCN subtarget features to GCNFeatures.td (PR #230996)
Matt Arsenault via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Sun Oct 11 09:03:13 PDT 2026
https://github.com/arsenm updated https://github.com/llvm/llvm-project/pull/230996
>From 61455cde1437a61f8b198c6c9f4b4e1bfc99df3d Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Sun, 11 Oct 2026 02:33:30 -0700
Subject: [PATCH] AMDGPU: Move GCN subtarget features to GCNFeatures.td
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>
---
llvm/lib/Target/AMDGPU/AMDGPU.td | 2833 +-----------------------
llvm/lib/Target/AMDGPU/GCNFeatures.td | 2840 +++++++++++++++++++++++++
2 files changed, 2841 insertions(+), 2832 deletions(-)
create mode 100644 llvm/lib/Target/AMDGPU/GCNFeatures.td
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-branch-commits
mailing list