Skip to content

Commit 4dab4c3

Browse files
committed
Addressed review comments
1. Delegated logic to Subtarget API from Tablegen. 2. Added a comment stating sm_101 is renamed to sm_110.
1 parent 9138b8e commit 4dab4c3

File tree

2 files changed

+10
-27
lines changed

2 files changed

+10
-27
lines changed

llvm/lib/Target/NVPTX/NVPTXInstrInfo.td

Lines changed: 5 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -111,31 +111,21 @@ class AnyPred<list<Predicate> Preds>
111111
// Checks PTX version and family-specific and architecture-specific SM versions.
112112
// For example, sm_100{f/a} and any future variants in the same family will match.
113113
class PTXWithFamilySMs<int PTXVersion, list<int> SMVersions> :
114-
Predicate<"Subtarget->getPTXVersion() >= " # PTXVersion #
115-
" && Subtarget->hasFamilySpecificFeatures()" #
116-
" && (" #
117-
!interleave(!foreach(sm, SMVersions,
118-
"(Subtarget->getSmFamilyVersion() == " # !div(sm, 10) #
119-
" && Subtarget->getSmVersion() >= " # sm # ")"),
120-
" || ") #
121-
")">;
114+
Predicate<"Subtarget->hasPTXWithFamilySMs(" # PTXVersion # ", {" #
115+
!interleave(SMVersions, ", ") # "})">;
122116

123117
// Checks PTX version and architecture-specific SM versions.
124118
// For example, sm_100{a} will match.
125119
class PTXWithAccelSMs<int PTXVersion, list<int> SMVersions> :
126-
Predicate<"Subtarget->getPTXVersion() >= " # PTXVersion #
127-
" && Subtarget->hasArchAccelFeatures()" #
128-
" && (" #
129-
!interleave(!foreach(sm, SMVersions,
130-
"Subtarget->getSmVersion() == " # sm),
131-
" || ") #
132-
")">;
120+
Predicate<"Subtarget->hasPTXWithAccelSMs(" # PTXVersion # ", {" #
121+
!interleave(SMVersions, ", ") # "})">;
133122

134123
// Helper predicate to call a subtarget method.
135124
class callSubtarget<string SubtargetMethod> : Predicate<"Subtarget->" # SubtargetMethod # "()">;
136125

137126
// Composed predicate to check tcgen05.shift instructions support.
138127
def hasTcgen05ShiftSupport : AnyPred<[
128+
// sm_101 renamed to sm_110 in PTX 9.0
139129
PTXWithAccelSMs<90, [100, 110, 103]>,
140130
PTXWithAccelSMs<88, [100, 101, 103]>,
141131
PTXWithAccelSMs<86, [100, 101]>

llvm/lib/Target/NVPTX/NVPTXSubtarget.cpp

Lines changed: 5 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -77,25 +77,18 @@ bool NVPTXSubtarget::hasPTXWithFamilySMs(unsigned PTXVersion,
7777
if (!hasFamilySpecificFeatures() || getPTXVersion() < PTXVersion)
7878
return false;
7979

80-
for (unsigned SM : SMVersions) {
81-
if (getSmFamilyVersion() == SM / 10 && getSmVersion() >= SM)
82-
return true;
83-
}
84-
85-
return false;
80+
return llvm::any_of(SMVersions, [&](unsigned SM) {
81+
return getSmFamilyVersion() == SM / 10 && getSmVersion() >= SM;
82+
});
8683
}
8784

8885
bool NVPTXSubtarget::hasPTXWithAccelSMs(unsigned PTXVersion,
8986
ArrayRef<unsigned> SMVersions) const {
9087
if (!hasArchAccelFeatures() || getPTXVersion() < PTXVersion)
9188
return false;
9289

93-
for (unsigned SM : SMVersions) {
94-
if (getSmVersion() == SM)
95-
return true;
96-
}
97-
98-
return false;
90+
return llvm::any_of(SMVersions,
91+
[&](unsigned SM) { return getSmVersion() == SM; });
9992
}
10093

10194
bool NVPTXSubtarget::allowFP16Math() const {

0 commit comments

Comments
 (0)