From 513e53421006cff7ce1235eb56404e072ccc69ab Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Tue, 7 Aug 2018 07:28:46 +0000 Subject: [PATCH] AMDGPU: Add feature vi-insts This is necessary to add a VI specific builtin, __builtin_amdgcn_s_dcache_wb. We already have an overly specific feature for one of these builtins, for s_memrealtime. I'm not sure whether it's better to add more of those, or to get rid of that and merge it with vi-insts. Alternatively, maybe this logically goes with scalar-stores? llvm-svn: 339104 --- lib/Target/AMDGPU/AMDGPU.td | 10 ++++++++-- lib/Target/AMDGPU/AMDGPUSubtarget.cpp | 1 + lib/Target/AMDGPU/AMDGPUSubtarget.h | 1 + 3 files changed, 10 insertions(+), 2 deletions(-) diff --git a/lib/Target/AMDGPU/AMDGPU.td b/lib/Target/AMDGPU/AMDGPU.td index 445b69b35eb..8945aafc4b6 100644 --- a/lib/Target/AMDGPU/AMDGPU.td +++ b/lib/Target/AMDGPU/AMDGPU.td @@ -140,6 +140,12 @@ def FeatureCIInsts : SubtargetFeature<"ci-insts", "Additional instructions for CI+" >; +def FeatureVIInsts : SubtargetFeature<"vi-insts", + "VIInsts", + "true", + "Additional instructions for VI+" +>; + def FeatureGFX9Insts : SubtargetFeature<"gfx9-insts", "GFX9Insts", "true", @@ -421,7 +427,7 @@ def FeatureSeaIslands : GCNSubtargetFeatureGeneration<"SEA_ISLANDS", def FeatureVolcanicIslands : GCNSubtargetFeatureGeneration<"VOLCANIC_ISLANDS", [FeatureFP64, FeatureLocalMemorySize65536, FeatureMIMG_R128, FeatureWavefrontSize64, FeatureFlatAddressSpace, FeatureGCN, - FeatureGCN3Encoding, FeatureCIInsts, Feature16BitInsts, + FeatureGCN3Encoding, FeatureCIInsts, FeatureVIInsts, Feature16BitInsts, FeatureSMemRealTime, FeatureVGPRIndexMode, FeatureMovrel, FeatureScalarStores, FeatureInv2PiInlineImm, FeatureSDWA, FeatureSDWAOutModsVOPC, FeatureSDWAMac, FeatureDPP, @@ -432,7 +438,7 @@ def FeatureVolcanicIslands : GCNSubtargetFeatureGeneration<"VOLCANIC_ISLANDS", def FeatureGFX9 : GCNSubtargetFeatureGeneration<"GFX9", [FeatureFP64, FeatureLocalMemorySize65536, FeatureWavefrontSize64, FeatureFlatAddressSpace, FeatureGCN, - FeatureGCN3Encoding, FeatureCIInsts, Feature16BitInsts, + FeatureGCN3Encoding, FeatureCIInsts, FeatureVIInsts, Feature16BitInsts, FeatureSMemRealTime, FeatureScalarStores, FeatureInv2PiInlineImm, FeatureApertureRegs, FeatureGFX9Insts, FeatureVOP3P, FeatureVGPRIndexMode, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp, diff --git a/lib/Target/AMDGPU/AMDGPUSubtarget.cpp b/lib/Target/AMDGPU/AMDGPUSubtarget.cpp index 98b49070fa9..3ec416a4db1 100644 --- a/lib/Target/AMDGPU/AMDGPUSubtarget.cpp +++ b/lib/Target/AMDGPU/AMDGPUSubtarget.cpp @@ -180,6 +180,7 @@ GCNSubtarget::GCNSubtarget(const Triple &TT, StringRef GPU, StringRef FS, FP64(false), GCN3Encoding(false), CIInsts(false), + VIInsts(false), GFX9Insts(false), SGPRInitBug(false), HasSMemRealTime(false), diff --git a/lib/Target/AMDGPU/AMDGPUSubtarget.h b/lib/Target/AMDGPU/AMDGPUSubtarget.h index 62310973365..bcc0a6fc048 100644 --- a/lib/Target/AMDGPU/AMDGPUSubtarget.h +++ b/lib/Target/AMDGPU/AMDGPUSubtarget.h @@ -337,6 +337,7 @@ protected: bool IsGCN; bool GCN3Encoding; bool CIInsts; + bool VIInsts; bool GFX9Insts; bool SGPRInitBug; bool HasSMemRealTime;