[clang] ee05167 - [OpenMP] Allow traits for the OpenMP context selector `isa`
Johannes Doerfert via cfe-commits
cfe-commits at lists.llvm.org
Wed Jul 29 08:24:19 PDT 2020
Author: Johannes Doerfert
Date: 2020-07-29T10:22:27-05:00
New Revision: ee05167cc42b95f70bc2ff1bd4402969f356f53b
URL: https://github.com/llvm/llvm-project/commit/ee05167cc42b95f70bc2ff1bd4402969f356f53b
DIFF: https://github.com/llvm/llvm-project/commit/ee05167cc42b95f70bc2ff1bd4402969f356f53b.diff
LOG: [OpenMP] Allow traits for the OpenMP context selector `isa`
It was unclear what `isa` was supposed to mean so we did not provide any
traits for this context selector. With this patch we will allow *any*
string or identifier. We use the target attribute and target info to
determine if the trait matches. In other words, we will check if the
provided value is a target feature that is available (at the call site).
Fixes PR46338
Reviewed By: ABataev
Differential Revision: https://reviews.llvm.org/D83281
Added:
clang/test/OpenMP/declare_variant_device_isa_codegen_1.c
Modified:
clang/include/clang/AST/OpenMPClause.h
clang/include/clang/Basic/DiagnosticParseKinds.td
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/lib/AST/OpenMPClause.cpp
clang/lib/Parse/ParseOpenMP.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/test/OpenMP/declare_variant_messages.c
llvm/include/llvm/Frontend/OpenMP/OMPContext.h
llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
llvm/lib/Frontend/OpenMP/OMPContext.cpp
llvm/unittests/Frontend/OpenMPContextTest.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h
index c649502f765b..4f94aa7074ee 100644
--- a/clang/include/clang/AST/OpenMPClause.h
+++ b/clang/include/clang/AST/OpenMPClause.h
@@ -7635,6 +7635,10 @@ class OMPClausePrinter final : public OMPClauseVisitor<OMPClausePrinter> {
struct OMPTraitProperty {
llvm::omp::TraitProperty Kind = llvm::omp::TraitProperty::invalid;
+
+ /// The raw string as we parsed it. This is needed for the `isa` trait set
+ /// (which accepts anything) and (later) extensions.
+ StringRef RawString;
};
struct OMPTraitSelector {
Expr *ScoreOrCondition = nullptr;
@@ -7692,6 +7696,23 @@ class OMPTraitInfo {
llvm::raw_ostream &operator<<(llvm::raw_ostream &OS, const OMPTraitInfo &TI);
llvm::raw_ostream &operator<<(llvm::raw_ostream &OS, const OMPTraitInfo *TI);
+/// Clang specific specialization of the OMPContext to lookup target features.
+struct TargetOMPContext final : public llvm::omp::OMPContext {
+
+ TargetOMPContext(ASTContext &ASTCtx,
+ std::function<void(StringRef)> &&DiagUnknownTrait,
+ const FunctionDecl *CurrentFunctionDecl);
+ virtual ~TargetOMPContext() = default;
+
+ /// See llvm::omp::OMPContext::matchesISATrait
+ bool matchesISATrait(StringRef RawString) const override;
+
+private:
+ std::function<bool(StringRef)> FeatureValidityCheck;
+ std::function<void(StringRef)> DiagUnknownTrait;
+ llvm::StringMap<bool> FeatureMap;
+};
+
} // namespace clang
#endif // LLVM_CLANG_AST_OPENMPCLAUSE_H
diff --git a/clang/include/clang/Basic/DiagnosticParseKinds.td b/clang/include/clang/Basic/DiagnosticParseKinds.td
index 6138b27fb87f..08b91de31993 100644
--- a/clang/include/clang/Basic/DiagnosticParseKinds.td
+++ b/clang/include/clang/Basic/DiagnosticParseKinds.td
@@ -1278,6 +1278,11 @@ def warn_omp_declare_variant_string_literal_or_identifier
"%select{set|selector|property}0; "
"%select{set|selector|property}0 skipped">,
InGroup<OpenMPClauses>;
+def warn_unknown_begin_declare_variant_isa_trait
+ : Warning<"isa trait '%0' is not known to the current target; verify the "
+ "spelling or consider restricting the context selector with the "
+ "'arch' selector further">,
+ InGroup<SourceUsesOpenMP>;
def note_omp_declare_variant_ctx_options
: Note<"context %select{set|selector|property}0 options are: %1">;
def warn_omp_declare_variant_expected
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 8093e7ed3fbe..ae693a08108c 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -10320,6 +10320,11 @@ def warn_nested_declare_variant
: Warning<"nesting `omp begin/end declare variant` is not supported yet; "
"nested context ignored">,
InGroup<SourceUsesOpenMP>;
+def warn_unknown_declare_variant_isa_trait
+ : Warning<"isa trait '%0' is not known to the current target; verify the "
+ "spelling or consider restricting the context selector with the "
+ "'arch' selector further">,
+ InGroup<SourceUsesOpenMP>;
def err_omp_non_pointer_type_array_shaping_base : Error<
"expected expression with a pointer to a complete type as a base of an array "
"shaping operation">;
diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp
index 6933c5742552..9caa691188fd 100644
--- a/clang/lib/AST/OpenMPClause.cpp
+++ b/clang/lib/AST/OpenMPClause.cpp
@@ -17,6 +17,7 @@
#include "clang/AST/DeclOpenMP.h"
#include "clang/Basic/LLVM.h"
#include "clang/Basic/OpenMPKinds.h"
+#include "clang/Basic/TargetInfo.h"
#include "llvm/ADT/SmallPtrSet.h"
#include "llvm/Support/Casting.h"
#include "llvm/Support/ErrorHandling.h"
@@ -2131,9 +2132,10 @@ void OMPTraitInfo::getAsVariantMatchInfo(ASTContext &ASTCtx,
Selector.ScoreOrCondition->getIntegerConstantExpr(ASTCtx))
VMI.addTrait(CondVal->isNullValue()
? TraitProperty::user_condition_false
- : TraitProperty::user_condition_true);
+ : TraitProperty::user_condition_true,
+ "<condition>");
else
- VMI.addTrait(TraitProperty::user_condition_false);
+ VMI.addTrait(TraitProperty::user_condition_false, "<condition>");
continue;
}
@@ -2143,11 +2145,12 @@ void OMPTraitInfo::getAsVariantMatchInfo(ASTContext &ASTCtx,
if ((Score = Selector.ScoreOrCondition->getIntegerConstantExpr(ASTCtx)))
ScorePtr = &*Score;
else
- VMI.addTrait(TraitProperty::user_condition_false);
+ VMI.addTrait(TraitProperty::user_condition_false,
+ "<non-constant-score>");
}
for (const OMPTraitProperty &Property : Selector.Properties)
- VMI.addTrait(Set.Kind, Property.Kind, ScorePtr);
+ VMI.addTrait(Set.Kind, Property.Kind, Property.RawString, ScorePtr);
if (Set.Kind != TraitSet::construct)
continue;
@@ -2204,7 +2207,8 @@ void OMPTraitInfo::print(llvm::raw_ostream &OS,
if (!FirstProperty)
OS << ", ";
FirstProperty = false;
- OS << getOpenMPContextTraitPropertyName(Property.Kind);
+ OS << getOpenMPContextTraitPropertyName(Property.Kind,
+ Property.RawString);
}
}
OS << ")";
@@ -2231,7 +2235,9 @@ std::string OMPTraitInfo::getMangledName() const {
continue;
for (const OMPTraitProperty &Property : Selector.Properties)
- OS << '$' << 'P' << getOpenMPContextTraitPropertyName(Property.Kind);
+ OS << '$' << 'P'
+ << getOpenMPContextTraitPropertyName(Property.Kind,
+ Property.RawString);
}
}
return OS.str();
@@ -2261,8 +2267,9 @@ OMPTraitInfo::OMPTraitInfo(StringRef MangledName) {
Selector.Properties.push_back(OMPTraitProperty());
OMPTraitProperty &Property = Selector.Properties.back();
std::pair<StringRef, StringRef> PropRestPair = MangledName.split('$');
- Property.Kind =
- getOpenMPContextTraitPropertyKind(Set.Kind, PropRestPair.first);
+ Property.RawString = PropRestPair.first;
+ Property.Kind = getOpenMPContextTraitPropertyKind(
+ Set.Kind, Selector.Kind, PropRestPair.first);
MangledName = PropRestPair.second;
} while (true);
} while (true);
@@ -2280,3 +2287,24 @@ llvm::raw_ostream &clang::operator<<(llvm::raw_ostream &OS,
const OMPTraitInfo *TI) {
return TI ? OS << *TI : OS;
}
+
+TargetOMPContext::TargetOMPContext(
+ ASTContext &ASTCtx, std::function<void(StringRef)> &&DiagUnknownTrait,
+ const FunctionDecl *CurrentFunctionDecl)
+ : OMPContext(ASTCtx.getLangOpts().OpenMPIsDevice,
+ ASTCtx.getTargetInfo().getTriple()),
+ FeatureValidityCheck([&](StringRef FeatureName) {
+ return ASTCtx.getTargetInfo().isValidFeatureName(FeatureName);
+ }),
+ DiagUnknownTrait(std::move(DiagUnknownTrait)) {
+ ASTCtx.getFunctionFeatureMap(FeatureMap, CurrentFunctionDecl);
+}
+
+bool TargetOMPContext::matchesISATrait(StringRef RawString) const {
+ auto It = FeatureMap.find(RawString);
+ if (It != FeatureMap.end())
+ return It->second;
+ if (!FeatureValidityCheck(RawString))
+ DiagUnknownTrait(RawString);
+ return false;
+}
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index 6a5d7c604c6b..ac1366e6b0ef 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -869,7 +869,8 @@ void Parser::parseOMPTraitPropertyKind(
return;
}
- TIProperty.Kind = getOpenMPContextTraitPropertyKind(Set, Name);
+ TIProperty.RawString = Name;
+ TIProperty.Kind = getOpenMPContextTraitPropertyKind(Set, Selector, Name);
if (TIProperty.Kind != TraitProperty::invalid) {
if (checkForDuplicates(*this, Name, NameLoc, Seen, CONTEXT_TRAIT_LVL))
TIProperty.Kind = TraitProperty::invalid;
@@ -910,7 +911,7 @@ void Parser::parseOMPTraitPropertyKind(
{TraitSet::construct, TraitSet::user, TraitSet::implementation,
TraitSet::device}) {
TraitProperty PropertyForName =
- getOpenMPContextTraitPropertyKind(PotentialSet, Name);
+ getOpenMPContextTraitPropertyKind(PotentialSet, Selector, Name);
if (PropertyForName == TraitProperty::invalid)
continue;
Diag(NameLoc, diag::note_omp_declare_variant_ctx_try)
@@ -949,8 +950,8 @@ static bool checkExtensionProperty(Parser &P, SourceLocation Loc,
for (OMPTraitProperty &SeenProp : TISelector.Properties)
if (IsMatchExtension(SeenProp)) {
P.Diag(Loc, diag::err_omp_variant_ctx_second_match_extension);
- StringRef SeenName =
- llvm::omp::getOpenMPContextTraitPropertyName(SeenProp.Kind);
+ StringRef SeenName = llvm::omp::getOpenMPContextTraitPropertyName(
+ SeenProp.Kind, SeenProp.RawString);
SourceLocation SeenLoc = Seen[SeenName];
P.Diag(SeenLoc, diag::note_omp_declare_variant_ctx_used_here)
<< CONTEXT_TRAIT_LVL << SeenName;
@@ -995,11 +996,13 @@ void Parser::parseOMPContextProperty(OMPTraitSelector &TISelector,
}
Diag(PropertyLoc, diag::warn_omp_ctx_incompatible_property_for_selector)
- << getOpenMPContextTraitPropertyName(TIProperty.Kind)
+ << getOpenMPContextTraitPropertyName(TIProperty.Kind,
+ TIProperty.RawString)
<< getOpenMPContextTraitSelectorName(TISelector.Kind)
<< getOpenMPContextTraitSetName(Set);
Diag(PropertyLoc, diag::note_omp_ctx_compatible_set_and_selector_for_property)
- << getOpenMPContextTraitPropertyName(TIProperty.Kind)
+ << getOpenMPContextTraitPropertyName(TIProperty.Kind,
+ TIProperty.RawString)
<< getOpenMPContextTraitSelectorName(
getOpenMPContextTraitSelectorForProperty(TIProperty.Kind))
<< getOpenMPContextTraitSetName(
@@ -1045,8 +1048,8 @@ void Parser::parseOMPTraitSelectorKind(
for (const auto &PotentialSet :
{TraitSet::construct, TraitSet::user, TraitSet::implementation,
TraitSet::device}) {
- TraitProperty PropertyForName =
- getOpenMPContextTraitPropertyKind(PotentialSet, Name);
+ TraitProperty PropertyForName = getOpenMPContextTraitPropertyKind(
+ PotentialSet, TraitSelector::invalid, Name);
if (PropertyForName == TraitProperty::invalid)
continue;
Diag(NameLoc, diag::note_omp_declare_variant_ctx_is_a)
@@ -1140,7 +1143,8 @@ void Parser::parseOMPContextSelector(
if (!RequiresProperty) {
TISelector.Properties.push_back(
- {getOpenMPContextTraitPropertyForSelector(TISelector.Kind)});
+ {getOpenMPContextTraitPropertyForSelector(TISelector.Kind),
+ getOpenMPContextTraitSelectorName(TISelector.Kind)});
return;
}
@@ -1157,7 +1161,8 @@ void Parser::parseOMPContextSelector(
if (!Condition.isUsable())
return FinishSelector();
TISelector.ScoreOrCondition = Condition.get();
- TISelector.Properties.push_back({TraitProperty::user_condition_unknown});
+ TISelector.Properties.push_back(
+ {TraitProperty::user_condition_unknown, "<condition>"});
return;
}
@@ -1236,8 +1241,8 @@ void Parser::parseOMPTraitSetKind(OMPTraitSet &TISet,
for (const auto &PotentialSet :
{TraitSet::construct, TraitSet::user, TraitSet::implementation,
TraitSet::device}) {
- TraitProperty PropertyForName =
- getOpenMPContextTraitPropertyKind(PotentialSet, Name);
+ TraitProperty PropertyForName = getOpenMPContextTraitPropertyKind(
+ PotentialSet, TraitSelector::invalid, Name);
if (PropertyForName == TraitProperty::invalid)
continue;
Diag(NameLoc, diag::note_omp_declare_variant_ctx_is_a)
@@ -1820,8 +1825,15 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
VariantMatchInfo VMI;
ASTContext &ASTCtx = Actions.getASTContext();
TI.getAsVariantMatchInfo(ASTCtx, VMI);
- OMPContext OMPCtx(ASTCtx.getLangOpts().OpenMPIsDevice,
- ASTCtx.getTargetInfo().getTriple());
+
+ std::function<void(StringRef)> DiagUnknownTrait = [this, Loc](
+ StringRef ISATrait) {
+ // TODO Track the selector locations in a way that is accessible here to
+ // improve the diagnostic location.
+ Diag(Loc, diag::warn_unknown_begin_declare_variant_isa_trait) << ISATrait;
+ };
+ TargetOMPContext OMPCtx(ASTCtx, std::move(DiagUnknownTrait),
+ /* CurrentFunctionDecl */ nullptr);
if (isVariantApplicableInContext(VMI, OMPCtx, /* DeviceSetOnly */ true)) {
Actions.ActOnOpenMPBeginDeclareVariant(Loc, TI);
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 0192df3bd170..4f69975c54ed 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -5895,8 +5895,15 @@ ExprResult Sema::ActOnOpenMPCall(ExprResult Call, Scope *Scope,
return Call;
ASTContext &Context = getASTContext();
- OMPContext OMPCtx(getLangOpts().OpenMPIsDevice,
- Context.getTargetInfo().getTriple());
+ std::function<void(StringRef)> DiagUnknownTrait = [this,
+ CE](StringRef ISATrait) {
+ // TODO Track the selector locations in a way that is accessible here to
+ // improve the diagnostic location.
+ Diag(CE->getBeginLoc(), diag::warn_unknown_declare_variant_isa_trait)
+ << ISATrait;
+ };
+ TargetOMPContext OMPCtx(Context, std::move(DiagUnknownTrait),
+ getCurFunctionDecl());
SmallVector<Expr *, 4> Exprs;
SmallVector<VariantMatchInfo, 4> VMIs;
@@ -5908,7 +5915,8 @@ ExprResult Sema::ActOnOpenMPCall(ExprResult Call, Scope *Scope,
VariantMatchInfo VMI;
OMPTraitInfo &TI = A->getTraitInfo();
TI.getAsVariantMatchInfo(Context, VMI);
- if (!isVariantApplicableInContext(VMI, OMPCtx, /* DeviceSetOnly */ false))
+ if (!isVariantApplicableInContext(VMI, OMPCtx,
+ /* DeviceSetOnly */ false))
continue;
VMIs.push_back(VMI);
diff --git a/clang/test/OpenMP/declare_variant_device_isa_codegen_1.c b/clang/test/OpenMP/declare_variant_device_isa_codegen_1.c
new file mode 100644
index 000000000000..baa5eb8f8830
--- /dev/null
+++ b/clang/test/OpenMP/declare_variant_device_isa_codegen_1.c
@@ -0,0 +1,49 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple %itanium_abi_triple -emit-llvm %s -o - -fopenmp-version=50 | FileCheck %s --check-prefix=GENERIC
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 | FileCheck %s --check-prefix=GENERIC
+
+// RUN: %clang_cc1 -target-feature +avx512f -verify -fopenmp -x c -triple %itanium_abi_triple -emit-llvm %s -o - -fopenmp-version=50 | FileCheck %s --check-prefix=WITHFEATURE
+// RUN: %clang_cc1 -target-feature +avx512f -fopenmp -x c++ -std=c++11 -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s
+// RUN: %clang_cc1 -target-feature +avx512f -fopenmp -x c++ -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 | FileCheck %s --check-prefix=WITHFEATURE
+
+// expected-no-diagnostics
+
+// Test taken from PR46338 (by linna su)
+
+#ifndef HEADER
+#define HEADER
+
+void base_saxpy(int, float, float *, float *);
+void avx512_saxpy(int, float, float *, float *);
+
+#pragma omp declare variant(avx512_saxpy) \
+ match(device = {isa(avx512f)})
+void base_saxpy(int n, float s, float *x, float *y) {
+#pragma omp parallel for
+ for (int i = 0; i < n; i++)
+ y[i] = s * x[i] + y[i];
+}
+
+void avx512_saxpy(int n, float s, float *x, float *y) {
+#pragma omp parallel for simd simdlen(16) aligned(x, y : 64)
+ for (int i = 0; i < n; i++)
+ y[i] = s * x[i] + y[i];
+}
+
+void caller(int n, float s, float *x, float *y) {
+ // GENERIC: define void @{{.*}}caller
+ // GENERIC: call void @{{.*}}base_saxpy
+ // WITHFEATURE: define void @{{.*}}caller
+ // WITHFEATURE: call void @{{.*}}avx512_saxpy
+ base_saxpy(n, s, x, y);
+}
+
+__attribute__((target("avx512f"))) void variant_caller(int n, float s, float *x, float *y) {
+ // GENERIC: define void @{{.*}}variant_caller
+ // GENERIC: call void @{{.*}}avx512_saxpy
+ // WITHFEATURE: define void @{{.*}}variant_caller
+ // WITHFEATURE: call void @{{.*}}avx512_saxpy
+ base_saxpy(n, s, x, y);
+}
+
+#endif
diff --git a/clang/test/OpenMP/declare_variant_messages.c b/clang/test/OpenMP/declare_variant_messages.c
index ecbf022351a7..84a56c5fd409 100644
--- a/clang/test/OpenMP/declare_variant_messages.c
+++ b/clang/test/OpenMP/declare_variant_messages.c
@@ -137,6 +137,18 @@ void marked_variant(void);
#pragma omp declare variant(marked_variant) match(xxx={}) // expected-warning {{'xxx' is not a valid context set in a `declare variant`; set ignored}} expected-warning {{variant function in '#pragma omp declare variant' is itself marked as '#pragma omp declare variant'}} expected-note {{context set options are: 'construct' 'device' 'implementation' 'user'}} expected-note {{the ignored set spans until here}}
void marked(void);
+#pragma omp declare variant(foo) match(device = {isa("foo")})
+int unknown_isa_trait(void);
+#pragma omp declare variant(foo) match(device = {isa(foo)})
+int unknown_isa_trait2(void);
+#pragma omp declare variant(foo) match(device = {kind(fpga), isa(bar)})
+int ignored_isa_trait(void);
+
+void caller() {
+ unknown_isa_trait(); // expected-warning {{isa trait 'foo' is not known to the current target; verify the spelling or consider restricting the context selector with the 'arch' selector further}}
+ unknown_isa_trait2(); // expected-warning {{isa trait 'foo' is not known to the current target; verify the spelling or consider restricting the context selector with the 'arch' selector further}}
+ ignored_isa_trait();
+}
#pragma omp declare variant // expected-error {{function declaration is expected after 'declare variant' directive}}
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPContext.h b/llvm/include/llvm/Frontend/OpenMP/OMPContext.h
index 1a42d189db44..8a4179167c89 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPContext.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPContext.h
@@ -70,15 +70,20 @@ TraitSelector getOpenMPContextTraitSelectorForProperty(TraitProperty Property);
/// Return a textual representation of the trait selector \p Kind.
StringRef getOpenMPContextTraitSelectorName(TraitSelector Kind);
-/// Parse \p Str and return the trait set it matches or
-/// TraitProperty::invalid.
-TraitProperty getOpenMPContextTraitPropertyKind(TraitSet Set, StringRef Str);
+/// Parse \p Str and return the trait property it matches in the set \p Set and
+/// selector \p Selector or TraitProperty::invalid.
+TraitProperty getOpenMPContextTraitPropertyKind(TraitSet Set,
+ TraitSelector Selector,
+ StringRef Str);
/// Return the trait property for a singleton selector \p Selector.
TraitProperty getOpenMPContextTraitPropertyForSelector(TraitSelector Selector);
-/// Return a textual representation of the trait property \p Kind.
-StringRef getOpenMPContextTraitPropertyName(TraitProperty Kind);
+/// Return a textual representation of the trait property \p Kind, which might
+/// be the raw string we parsed (\p RawString) if we do not translate the
+/// property into a (distinct) enum.
+StringRef getOpenMPContextTraitPropertyName(TraitProperty Kind,
+ StringRef RawString);
/// Return a textual representation of the trait property \p Kind with selector
/// and set name included.
@@ -112,24 +117,36 @@ bool isValidTraitPropertyForTraitSetAndSelector(TraitProperty Property,
/// scored (via the ScoresMap). In addition, the required consturct nesting is
/// decribed as well.
struct VariantMatchInfo {
- /// Add the trait \p Property to the required trait set. If \p Score is not
- /// null, it recorded as well. If \p Property is in the `construct` set it
- /// is recorded in-order in the ConstructTraits as well.
- void addTrait(TraitProperty Property, APInt *Score = nullptr) {
- addTrait(getOpenMPContextTraitSetForProperty(Property), Property, Score);
+ /// Add the trait \p Property to the required trait set. \p RawString is the
+ /// string we parsed and derived \p Property from. If \p Score is not null, it
+ /// recorded as well. If \p Property is in the `construct` set it is recorded
+ /// in-order in the ConstructTraits as well.
+ void addTrait(TraitProperty Property, StringRef RawString,
+ APInt *Score = nullptr) {
+ addTrait(getOpenMPContextTraitSetForProperty(Property), Property, RawString,
+ Score);
}
/// Add the trait \p Property which is in set \p Set to the required trait
- /// set. If \p Score is not null, it recorded as well. If \p Set is the
- /// `construct` set it is recorded in-order in the ConstructTraits as well.
- void addTrait(TraitSet Set, TraitProperty Property, APInt *Score = nullptr) {
+ /// set. \p RawString is the string we parsed and derived \p Property from. If
+ /// \p Score is not null, it recorded as well. If \p Set is the `construct`
+ /// set it is recorded in-order in the ConstructTraits as well.
+ void addTrait(TraitSet Set, TraitProperty Property, StringRef RawString,
+ APInt *Score = nullptr) {
if (Score)
ScoreMap[Property] = *Score;
+
+ // Special handling for `device={isa(...)}` as we do not match the enum but
+ // the raw string.
+ if (Property == TraitProperty::device_isa___ANY)
+ ISATraits.push_back(RawString);
+
RequiredTraits.set(unsigned(Property));
if (Set == TraitSet::construct)
ConstructTraits.push_back(Property);
}
BitVector RequiredTraits = BitVector(unsigned(TraitProperty::Last) + 1);
+ SmallVector<StringRef, 8> ISATraits;
SmallVector<TraitProperty, 8> ConstructTraits;
SmallDenseMap<TraitProperty, APInt> ScoreMap;
};
@@ -139,6 +156,7 @@ struct VariantMatchInfo {
/// in OpenMP constructs at the location.
struct OMPContext {
OMPContext(bool IsDeviceCompilation, Triple TargetTriple);
+ virtual ~OMPContext() = default;
void addTrait(TraitProperty Property) {
addTrait(getOpenMPContextTraitSetForProperty(Property), Property);
@@ -149,6 +167,11 @@ struct OMPContext {
ConstructTraits.push_back(Property);
}
+ /// Hook for users to check if an ISA trait matches. The trait is described as
+ /// the string that got parsed and it depends on the target and context if
+ /// this matches or not.
+ virtual bool matchesISATrait(StringRef) const { return false; }
+
BitVector ActiveTraits = BitVector(unsigned(TraitProperty::Last) + 1);
SmallVector<TraitProperty, 8> ConstructTraits;
};
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
index 7771dcd72d6a..3fc87dc34cd3 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
@@ -1071,7 +1071,11 @@ __OMP_TRAIT_PROPERTY(device, kind, any)
__OMP_TRAIT_SELECTOR(device, isa, true)
-// TODO: What do we want for ISA?
+// We use "__ANY" as a placeholder in the isa property to denote the
+// conceptual "any", not the literal `any` used in kind. The string we
+// we use is not important except that it will show up in diagnostics.
+OMP_TRAIT_PROPERTY(device_isa___ANY, device, device_isa,
+ "<any, entirely target dependent>")
__OMP_TRAIT_SELECTOR(device, arch, true)
diff --git a/llvm/lib/Frontend/OpenMP/OMPContext.cpp b/llvm/lib/Frontend/OpenMP/OMPContext.cpp
index c44e858ab5ed..56a6e2b08bd9 100644
--- a/llvm/lib/Frontend/OpenMP/OMPContext.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPContext.cpp
@@ -175,11 +175,11 @@ static int isVariantApplicableInContextHelper(
LLVM_DEBUG({
if (MK == MK_ALL)
dbgs() << "[" << DEBUG_TYPE << "] Property "
- << getOpenMPContextTraitPropertyName(Property)
+ << getOpenMPContextTraitPropertyName(Property, "")
<< " was not in the OpenMP context but match kind is all.\n";
if (MK == MK_NONE)
dbgs() << "[" << DEBUG_TYPE << "] Property "
- << getOpenMPContextTraitPropertyName(Property)
+ << getOpenMPContextTraitPropertyName(Property, "")
<< " was in the OpenMP context but match kind is none.\n";
});
return false;
@@ -198,6 +198,14 @@ static int isVariantApplicableInContextHelper(
continue;
bool IsActiveTrait = Ctx.ActiveTraits.test(unsigned(Property));
+
+ // We overwrite the isa trait as it is actually up to the OMPContext hook to
+ // check the raw string(s).
+ if (Property == TraitProperty::device_isa___ANY)
+ IsActiveTrait = llvm::all_of(VMI.ISATraits, [&](StringRef RawString) {
+ return Ctx.matchesISATrait(RawString);
+ });
+
Optional<bool> Result = HandleTrait(Property, IsActiveTrait);
if (Result.hasValue())
return Result.getValue();
@@ -225,7 +233,7 @@ static int isVariantApplicableInContextHelper(
if (!FoundInOrder) {
LLVM_DEBUG(dbgs() << "[" << DEBUG_TYPE << "] Construct property "
- << getOpenMPContextTraitPropertyName(Property)
+ << getOpenMPContextTraitPropertyName(Property, "")
<< " was not nested properly.\n");
return false;
}
@@ -425,8 +433,12 @@ StringRef llvm::omp::getOpenMPContextTraitSelectorName(TraitSelector Kind) {
llvm_unreachable("Unknown trait selector!");
}
-TraitProperty llvm::omp::getOpenMPContextTraitPropertyKind(TraitSet Set,
- StringRef S) {
+TraitProperty llvm::omp::getOpenMPContextTraitPropertyKind(
+ TraitSet Set, TraitSelector Selector, StringRef S) {
+ // Special handling for `device={isa(...)}` as we accept anything here. It is
+ // up to the target to decide if the feature is available.
+ if (Set == TraitSet::device && Selector == TraitSelector::device_isa)
+ return TraitProperty::device_isa___ANY;
#define OMP_TRAIT_PROPERTY(Enum, TraitSetEnum, TraitSelectorEnum, Str) \
if (Set == TraitSet::TraitSetEnum && Str == S) \
return TraitProperty::Enum;
@@ -444,7 +456,10 @@ llvm::omp::getOpenMPContextTraitPropertyForSelector(TraitSelector Selector) {
#include "llvm/Frontend/OpenMP/OMPKinds.def"
.Default(TraitProperty::invalid);
}
-StringRef llvm::omp::getOpenMPContextTraitPropertyName(TraitProperty Kind) {
+StringRef llvm::omp::getOpenMPContextTraitPropertyName(TraitProperty Kind,
+ StringRef RawString) {
+ if (Kind == TraitProperty::device_isa___ANY)
+ return RawString;
switch (Kind) {
#define OMP_TRAIT_PROPERTY(Enum, TraitSetEnum, TraitSelectorEnum, Str) \
case TraitProperty::Enum: \
diff --git a/llvm/unittests/Frontend/OpenMPContextTest.cpp b/llvm/unittests/Frontend/OpenMPContextTest.cpp
index eb505be042cb..5938559aacbc 100644
--- a/llvm/unittests/Frontend/OpenMPContextTest.cpp
+++ b/llvm/unittests/Frontend/OpenMPContextTest.cpp
@@ -38,11 +38,13 @@ TEST_F(OpenMPContextTest, RoundTripAndAssociation) {
#define OMP_TRAIT_PROPERTY(Enum, TraitSetEnum, TraitSelectorEnum, Str) \
EXPECT_EQ(TraitProperty::Enum, \
getOpenMPContextTraitPropertyKind( \
- TraitSet::TraitSetEnum, \
- getOpenMPContextTraitPropertyName(TraitProperty::Enum))); \
+ TraitSet::TraitSetEnum, TraitSelector::TraitSelectorEnum, \
+ getOpenMPContextTraitPropertyName(TraitProperty::Enum, Str))); \
EXPECT_EQ(Str, getOpenMPContextTraitPropertyName( \
- getOpenMPContextTraitPropertyKind(TraitSet::TraitSetEnum, \
- Str))); \
+ getOpenMPContextTraitPropertyKind( \
+ TraitSet::TraitSetEnum, \
+ TraitSelector::TraitSelectorEnum, Str), \
+ Str)); \
EXPECT_EQ(TraitSet::TraitSetEnum, \
getOpenMPContextTraitSetForProperty(TraitProperty::Enum)); \
EXPECT_EQ(TraitSelector::TraitSelectorEnum, \
@@ -77,31 +79,31 @@ TEST_F(OpenMPContextTest, ApplicabilityNonConstruct) {
EXPECT_TRUE(isVariantApplicableInContext(Empty, DeviceNVPTX));
VariantMatchInfo UserCondFalse;
- UserCondFalse.addTrait(TraitProperty::user_condition_false);
+ UserCondFalse.addTrait(TraitProperty::user_condition_false, "");
EXPECT_FALSE(isVariantApplicableInContext(UserCondFalse, HostLinux));
EXPECT_FALSE(isVariantApplicableInContext(UserCondFalse, DeviceLinux));
EXPECT_FALSE(isVariantApplicableInContext(UserCondFalse, HostNVPTX));
EXPECT_FALSE(isVariantApplicableInContext(UserCondFalse, DeviceNVPTX));
VariantMatchInfo DeviceArchArm;
- DeviceArchArm.addTrait(TraitProperty::device_arch_arm);
+ DeviceArchArm.addTrait(TraitProperty::device_arch_arm, "");
EXPECT_FALSE(isVariantApplicableInContext(DeviceArchArm, HostLinux));
EXPECT_FALSE(isVariantApplicableInContext(DeviceArchArm, DeviceLinux));
EXPECT_FALSE(isVariantApplicableInContext(DeviceArchArm, HostNVPTX));
EXPECT_FALSE(isVariantApplicableInContext(DeviceArchArm, DeviceNVPTX));
VariantMatchInfo LLVMHostUserCondTrue;
- LLVMHostUserCondTrue.addTrait(TraitProperty::implementation_vendor_llvm);
- LLVMHostUserCondTrue.addTrait(TraitProperty::device_kind_host);
- LLVMHostUserCondTrue.addTrait(TraitProperty::device_kind_any);
- LLVMHostUserCondTrue.addTrait(TraitProperty::user_condition_true);
+ LLVMHostUserCondTrue.addTrait(TraitProperty::implementation_vendor_llvm, "");
+ LLVMHostUserCondTrue.addTrait(TraitProperty::device_kind_host, "");
+ LLVMHostUserCondTrue.addTrait(TraitProperty::device_kind_any, "");
+ LLVMHostUserCondTrue.addTrait(TraitProperty::user_condition_true, "");
EXPECT_TRUE(isVariantApplicableInContext(LLVMHostUserCondTrue, HostLinux));
EXPECT_FALSE(isVariantApplicableInContext(LLVMHostUserCondTrue, DeviceLinux));
EXPECT_TRUE(isVariantApplicableInContext(LLVMHostUserCondTrue, HostNVPTX));
EXPECT_FALSE(isVariantApplicableInContext(LLVMHostUserCondTrue, DeviceNVPTX));
VariantMatchInfo LLVMHostUserCondTrueCPU = LLVMHostUserCondTrue;
- LLVMHostUserCondTrueCPU.addTrait(TraitProperty::device_kind_cpu);
+ LLVMHostUserCondTrueCPU.addTrait(TraitProperty::device_kind_cpu, "");
EXPECT_TRUE(isVariantApplicableInContext(LLVMHostUserCondTrueCPU, HostLinux));
EXPECT_FALSE(
isVariantApplicableInContext(LLVMHostUserCondTrueCPU, DeviceLinux));
@@ -111,14 +113,14 @@ TEST_F(OpenMPContextTest, ApplicabilityNonConstruct) {
isVariantApplicableInContext(LLVMHostUserCondTrueCPU, DeviceNVPTX));
VariantMatchInfo GPU;
- GPU.addTrait(TraitProperty::device_kind_gpu);
+ GPU.addTrait(TraitProperty::device_kind_gpu, "");
EXPECT_FALSE(isVariantApplicableInContext(GPU, HostLinux));
EXPECT_FALSE(isVariantApplicableInContext(GPU, DeviceLinux));
EXPECT_TRUE(isVariantApplicableInContext(GPU, HostNVPTX));
EXPECT_TRUE(isVariantApplicableInContext(GPU, DeviceNVPTX));
VariantMatchInfo NoHost;
- NoHost.addTrait(TraitProperty::device_kind_nohost);
+ NoHost.addTrait(TraitProperty::device_kind_nohost, "");
EXPECT_FALSE(isVariantApplicableInContext(NoHost, HostLinux));
EXPECT_TRUE(isVariantApplicableInContext(NoHost, DeviceLinux));
EXPECT_FALSE(isVariantApplicableInContext(NoHost, HostNVPTX));
@@ -154,7 +156,7 @@ TEST_F(OpenMPContextTest, ApplicabilityAllTraits) {
isVariantApplicableInContext(Empty, DeviceNVPTXTargetTeamsParallel));
VariantMatchInfo UserCondFalse;
- UserCondFalse.addTrait(TraitProperty::user_condition_false);
+ UserCondFalse.addTrait(TraitProperty::user_condition_false, "");
EXPECT_FALSE(
isVariantApplicableInContext(UserCondFalse, HostLinuxParallelParallel));
EXPECT_FALSE(
@@ -164,7 +166,7 @@ TEST_F(OpenMPContextTest, ApplicabilityAllTraits) {
DeviceNVPTXTargetTeamsParallel));
VariantMatchInfo DeviceArchArm;
- DeviceArchArm.addTrait(TraitProperty::device_arch_arm);
+ DeviceArchArm.addTrait(TraitProperty::device_arch_arm, "");
EXPECT_FALSE(
isVariantApplicableInContext(DeviceArchArm, HostLinuxParallelParallel));
EXPECT_FALSE(
@@ -175,10 +177,12 @@ TEST_F(OpenMPContextTest, ApplicabilityAllTraits) {
APInt Score(32, 1000);
VariantMatchInfo LLVMHostUserCondTrue;
- LLVMHostUserCondTrue.addTrait(TraitProperty::implementation_vendor_llvm);
- LLVMHostUserCondTrue.addTrait(TraitProperty::device_kind_host);
- LLVMHostUserCondTrue.addTrait(TraitProperty::device_kind_any);
- LLVMHostUserCondTrue.addTrait(TraitProperty::user_condition_true, &Score);
+ LLVMHostUserCondTrue.addTrait(TraitProperty::implementation_vendor_llvm,
+ "");
+ LLVMHostUserCondTrue.addTrait(TraitProperty::device_kind_host, "");
+ LLVMHostUserCondTrue.addTrait(TraitProperty::device_kind_any, "");
+ LLVMHostUserCondTrue.addTrait(TraitProperty::user_condition_true, "",
+ &Score);
EXPECT_TRUE(isVariantApplicableInContext(LLVMHostUserCondTrue,
HostLinuxParallelParallel));
EXPECT_FALSE(isVariantApplicableInContext(LLVMHostUserCondTrue,
@@ -189,7 +193,7 @@ TEST_F(OpenMPContextTest, ApplicabilityAllTraits) {
DeviceNVPTXTargetTeamsParallel));
VariantMatchInfo LLVMHostUserCondTrueCPU = LLVMHostUserCondTrue;
- LLVMHostUserCondTrueCPU.addTrait(TraitProperty::device_kind_cpu);
+ LLVMHostUserCondTrueCPU.addTrait(TraitProperty::device_kind_cpu, "");
EXPECT_TRUE(isVariantApplicableInContext(LLVMHostUserCondTrueCPU,
HostLinuxParallelParallel));
EXPECT_FALSE(isVariantApplicableInContext(LLVMHostUserCondTrueCPU,
@@ -200,7 +204,7 @@ TEST_F(OpenMPContextTest, ApplicabilityAllTraits) {
DeviceNVPTXTargetTeamsParallel));
VariantMatchInfo GPU;
- GPU.addTrait(TraitProperty::device_kind_gpu);
+ GPU.addTrait(TraitProperty::device_kind_gpu, "");
EXPECT_FALSE(isVariantApplicableInContext(GPU, HostLinuxParallelParallel));
EXPECT_FALSE(isVariantApplicableInContext(GPU, DeviceLinuxTargetParallel));
EXPECT_TRUE(isVariantApplicableInContext(GPU, HostNVPTXFor));
@@ -208,7 +212,7 @@ TEST_F(OpenMPContextTest, ApplicabilityAllTraits) {
isVariantApplicableInContext(GPU, DeviceNVPTXTargetTeamsParallel));
VariantMatchInfo NoHost;
- NoHost.addTrait(TraitProperty::device_kind_nohost);
+ NoHost.addTrait(TraitProperty::device_kind_nohost, "");
EXPECT_FALSE(
isVariantApplicableInContext(NoHost, HostLinuxParallelParallel));
EXPECT_TRUE(
@@ -219,8 +223,9 @@ TEST_F(OpenMPContextTest, ApplicabilityAllTraits) {
}
{ // variants with all sets
VariantMatchInfo DeviceArchArmParallel;
- DeviceArchArmParallel.addTrait(TraitProperty::construct_parallel_parallel);
- DeviceArchArmParallel.addTrait(TraitProperty::device_arch_arm);
+ DeviceArchArmParallel.addTrait(TraitProperty::construct_parallel_parallel,
+ "");
+ DeviceArchArmParallel.addTrait(TraitProperty::device_arch_arm, "");
EXPECT_FALSE(isVariantApplicableInContext(DeviceArchArmParallel,
HostLinuxParallelParallel));
EXPECT_FALSE(isVariantApplicableInContext(DeviceArchArmParallel,
@@ -232,12 +237,13 @@ TEST_F(OpenMPContextTest, ApplicabilityAllTraits) {
VariantMatchInfo LLVMHostUserCondTrueParallel;
LLVMHostUserCondTrueParallel.addTrait(
- TraitProperty::implementation_vendor_llvm);
- LLVMHostUserCondTrueParallel.addTrait(TraitProperty::device_kind_host);
- LLVMHostUserCondTrueParallel.addTrait(TraitProperty::device_kind_any);
- LLVMHostUserCondTrueParallel.addTrait(TraitProperty::user_condition_true);
+ TraitProperty::implementation_vendor_llvm, "");
+ LLVMHostUserCondTrueParallel.addTrait(TraitProperty::device_kind_host, "");
+ LLVMHostUserCondTrueParallel.addTrait(TraitProperty::device_kind_any, "");
+ LLVMHostUserCondTrueParallel.addTrait(TraitProperty::user_condition_true,
+ "");
LLVMHostUserCondTrueParallel.addTrait(
- TraitProperty::construct_parallel_parallel);
+ TraitProperty::construct_parallel_parallel, "");
EXPECT_TRUE(isVariantApplicableInContext(LLVMHostUserCondTrueParallel,
HostLinuxParallelParallel));
EXPECT_FALSE(isVariantApplicableInContext(LLVMHostUserCondTrueParallel,
@@ -250,7 +256,7 @@ TEST_F(OpenMPContextTest, ApplicabilityAllTraits) {
VariantMatchInfo LLVMHostUserCondTrueParallelParallel =
LLVMHostUserCondTrueParallel;
LLVMHostUserCondTrueParallelParallel.addTrait(
- TraitProperty::construct_parallel_parallel);
+ TraitProperty::construct_parallel_parallel, "");
EXPECT_TRUE(isVariantApplicableInContext(
LLVMHostUserCondTrueParallelParallel, HostLinuxParallelParallel));
EXPECT_FALSE(isVariantApplicableInContext(
@@ -263,7 +269,7 @@ TEST_F(OpenMPContextTest, ApplicabilityAllTraits) {
VariantMatchInfo LLVMHostUserCondTrueParallelParallelParallel =
LLVMHostUserCondTrueParallelParallel;
LLVMHostUserCondTrueParallelParallelParallel.addTrait(
- TraitProperty::construct_parallel_parallel);
+ TraitProperty::construct_parallel_parallel, "");
EXPECT_FALSE(isVariantApplicableInContext(
LLVMHostUserCondTrueParallelParallelParallel,
HostLinuxParallelParallel));
@@ -277,9 +283,9 @@ TEST_F(OpenMPContextTest, ApplicabilityAllTraits) {
DeviceNVPTXTargetTeamsParallel));
VariantMatchInfo GPUTargetTeams;
- GPUTargetTeams.addTrait(TraitProperty::construct_target_target);
- GPUTargetTeams.addTrait(TraitProperty::construct_teams_teams);
- GPUTargetTeams.addTrait(TraitProperty::device_kind_gpu);
+ GPUTargetTeams.addTrait(TraitProperty::construct_target_target, "");
+ GPUTargetTeams.addTrait(TraitProperty::construct_teams_teams, "");
+ GPUTargetTeams.addTrait(TraitProperty::device_kind_gpu, "");
EXPECT_FALSE(isVariantApplicableInContext(GPUTargetTeams,
HostLinuxParallelParallel));
EXPECT_FALSE(isVariantApplicableInContext(GPUTargetTeams,
@@ -289,9 +295,9 @@ TEST_F(OpenMPContextTest, ApplicabilityAllTraits) {
DeviceNVPTXTargetTeamsParallel));
VariantMatchInfo GPUTargetParallel;
- GPUTargetParallel.addTrait(TraitProperty::construct_target_target);
- GPUTargetParallel.addTrait(TraitProperty::construct_parallel_parallel);
- GPUTargetParallel.addTrait(TraitProperty::device_kind_gpu);
+ GPUTargetParallel.addTrait(TraitProperty::construct_target_target, "");
+ GPUTargetParallel.addTrait(TraitProperty::construct_parallel_parallel, "");
+ GPUTargetParallel.addTrait(TraitProperty::device_kind_gpu, "");
EXPECT_FALSE(isVariantApplicableInContext(GPUTargetParallel,
HostLinuxParallelParallel));
EXPECT_FALSE(isVariantApplicableInContext(GPUTargetParallel,
More information about the cfe-commits
mailing list