This is an archive of the discontinued LLVM Phabricator instance.

Differential D50265

AMDGPU: Add feature vi-insts
ClosedPublic

Authored by arsenm on Aug 3 2018, 10:59 AM.

Details

Reviewers
t-tye
kzhuravl
rampitec
Summary

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?

Diff Detail

Event Timeline

arsenm created this revision.Aug 3 2018, 10:59 AM
Herald added subscribers: tpr, dstuttard, yaxunl and 2 others. · View Herald TranscriptAug 3 2018, 10:59 AM
kzhuravl accepted this revision.Aug 3 2018, 11:05 AM

LGTM

This revision is now accepted and ready to land.Aug 3 2018, 11:05 AM
arsenm closed this revision.Aug 7 2018, 12:41 AM

r339104

Revision Contents

PathSize
lib/
Target/
AMDGPU/
10 lines
1 line
1 line

Diff 159049

lib/Target/AMDGPU/AMDGPU.td

Show First 20 Lines • Show All 134 Lines • ▼ Show 20 Lines
>; >;
def FeatureCIInsts : SubtargetFeature<"ci-insts", def FeatureCIInsts : SubtargetFeature<"ci-insts",
"CIInsts", "CIInsts",
"true", "true",
"Additional instructions for CI+" "Additional instructions for CI+"
>; >;
def FeatureVIInsts : SubtargetFeature<"vi-insts",
"VIInsts",
"true",
"Additional instructions for VI+"
>;
def FeatureGFX9Insts : SubtargetFeature<"gfx9-insts", def FeatureGFX9Insts : SubtargetFeature<"gfx9-insts",
"GFX9Insts", "GFX9Insts",
"true", "true",
"Additional instructions for GFX9+" "Additional instructions for GFX9+"
>; >;
def FeatureSMemRealTime : SubtargetFeature<"s-memrealtime", def FeatureSMemRealTime : SubtargetFeature<"s-memrealtime",
"HasSMemRealTime", "HasSMemRealTime",
▲ Show 20 Lines • Show All 265 Lines • ▼ Show 20 Linesdef FeatureSeaIslands : GCNSubtargetFeatureGeneration<"SEA_ISLANDS",
[FeatureFP64, FeatureLocalMemorySize65536, FeatureMIMG_R128, [FeatureFP64, FeatureLocalMemorySize65536, FeatureMIMG_R128,
FeatureWavefrontSize64, FeatureGCN, FeatureFlatAddressSpace, FeatureWavefrontSize64, FeatureGCN, FeatureFlatAddressSpace,
FeatureCIInsts, FeatureMovrel] FeatureCIInsts, FeatureMovrel]
>; >;
def FeatureVolcanicIslands : GCNSubtargetFeatureGeneration<"VOLCANIC_ISLANDS", def FeatureVolcanicIslands : GCNSubtargetFeatureGeneration<"VOLCANIC_ISLANDS",
[FeatureFP64, FeatureLocalMemorySize65536, FeatureMIMG_R128, [FeatureFP64, FeatureLocalMemorySize65536, FeatureMIMG_R128,
FeatureWavefrontSize64, FeatureFlatAddressSpace, FeatureGCN, FeatureWavefrontSize64, FeatureFlatAddressSpace, FeatureGCN,
FeatureGCN3Encoding, FeatureCIInsts, Feature16BitInsts, FeatureGCN3Encoding, FeatureCIInsts, FeatureVIInsts, Feature16BitInsts,
FeatureSMemRealTime, FeatureVGPRIndexMode, FeatureMovrel, FeatureSMemRealTime, FeatureVGPRIndexMode, FeatureMovrel,
FeatureScalarStores, FeatureInv2PiInlineImm, FeatureScalarStores, FeatureInv2PiInlineImm,
FeatureSDWA, FeatureSDWAOutModsVOPC, FeatureSDWAMac, FeatureDPP, FeatureSDWA, FeatureSDWAOutModsVOPC, FeatureSDWAMac, FeatureDPP,
FeatureIntClamp FeatureIntClamp
] ]
>; >;
def FeatureGFX9 : GCNSubtargetFeatureGeneration<"GFX9", def FeatureGFX9 : GCNSubtargetFeatureGeneration<"GFX9",
[FeatureFP64, FeatureLocalMemorySize65536, [FeatureFP64, FeatureLocalMemorySize65536,
FeatureWavefrontSize64, FeatureFlatAddressSpace, FeatureGCN, FeatureWavefrontSize64, FeatureFlatAddressSpace, FeatureGCN,
FeatureGCN3Encoding, FeatureCIInsts, Feature16BitInsts, FeatureGCN3Encoding, FeatureCIInsts, FeatureVIInsts, Feature16BitInsts,
FeatureSMemRealTime, FeatureScalarStores, FeatureInv2PiInlineImm, FeatureSMemRealTime, FeatureScalarStores, FeatureInv2PiInlineImm,
FeatureApertureRegs, FeatureGFX9Insts, FeatureVOP3P, FeatureVGPRIndexMode, FeatureApertureRegs, FeatureGFX9Insts, FeatureVOP3P, FeatureVGPRIndexMode,
FeatureFastFMAF32, FeatureDPP, FeatureIntClamp, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp,
FeatureSDWA, FeatureSDWAOmod, FeatureSDWAScalar, FeatureSDWASdst, FeatureSDWA, FeatureSDWAOmod, FeatureSDWAScalar, FeatureSDWASdst,
FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts, FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts,
FeatureAddNoCarryInsts, FeatureScalarAtomics FeatureAddNoCarryInsts, FeatureScalarAtomics
] ]
>; >;
▲ Show 20 Lines • Show All 294 Lines • Show Last 20 Lines

lib/Target/AMDGPU/AMDGPUSubtarget.h

Show First 20 Lines • Show All 331 Lines • ▼ Show 20 Linesprotected:
// Subtarget statically properties set by tablegen // Subtarget statically properties set by tablegen
bool FP64; bool FP64;
bool FMA; bool FMA;
bool MIMG_R128; bool MIMG_R128;
bool IsGCN; bool IsGCN;
bool GCN3Encoding; bool GCN3Encoding;
bool CIInsts; bool CIInsts;
bool VIInsts;
bool GFX9Insts; bool GFX9Insts;
bool SGPRInitBug; bool SGPRInitBug;
bool HasSMemRealTime; bool HasSMemRealTime;
bool HasIntClamp; bool HasIntClamp;
bool HasFmaMixInsts; bool HasFmaMixInsts;
bool HasMovrel; bool HasMovrel;
bool HasVGPRIndexMode; bool HasVGPRIndexMode;
bool HasScalarStores; bool HasScalarStores;
▲ Show 20 Lines • Show All 722 Lines • Show Last 20 Lines

lib/Target/AMDGPU/AMDGPUSubtarget.cpp

Show First 20 Lines • Show All 174 Lines • ▼ Show 20 LinesGCNSubtarget::GCNSubtarget(const Triple &TT, StringRef GPU, StringRef FS,
EnableUnsafeDSOffsetFolding(false), EnableUnsafeDSOffsetFolding(false),
EnableSIScheduler(false), EnableSIScheduler(false),
EnableDS128(false), EnableDS128(false),
DumpCode(false), DumpCode(false),
FP64(false), FP64(false),
GCN3Encoding(false), GCN3Encoding(false),
CIInsts(false), CIInsts(false),
VIInsts(false),
GFX9Insts(false), GFX9Insts(false),
SGPRInitBug(false), SGPRInitBug(false),
HasSMemRealTime(false), HasSMemRealTime(false),
HasIntClamp(false), HasIntClamp(false),
HasFmaMixInsts(false), HasFmaMixInsts(false),
HasMovrel(false), HasMovrel(false),
HasVGPRIndexMode(false), HasVGPRIndexMode(false),
HasScalarStores(false), HasScalarStores(false),
▲ Show 20 Lines • Show All 504 Lines • Show Last 20 Lines