diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h --- a/clang/include/clang/AST/OpenMPClause.h +++ b/clang/include/clang/AST/OpenMPClause.h @@ -7593,6 +7593,10 @@ 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 accept anything) and (later) extensions. + StringRef RawString; }; struct OMPTraitSelector { Expr *ScoreOrCondition = nullptr; @@ -7650,6 +7654,23 @@ 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 &&DiagUnknownTrait, + const FunctionDecl *CurrentFunctionDecl); + virtual ~TargetOMPContext() = default; + + /// See llvm::omp::OMPContext::matchesISATrait + bool matchesISATrait(StringRef RawString) const override; + +private: + std::function FeatureValidityCheck; + std::function DiagUnknownTrait; + llvm::StringMap 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 --- a/clang/include/clang/Basic/DiagnosticParseKinds.td +++ b/clang/include/clang/Basic/DiagnosticParseKinds.td @@ -1275,6 +1275,11 @@ "%select{set|selector|property}0; " "%select{set|selector|property}0 skipped">, InGroup; +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; 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 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10310,6 +10310,11 @@ : Warning<"nesting `omp begin/end declare variant` is not supported yet; " "nested context ignored">, InGroup; +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; 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 --- 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" @@ -2130,10 +2131,11 @@ llvm::APSInt CondVal; if (Selector.ScoreOrCondition->isIntegerConstantExpr(CondVal, ASTCtx)) VMI.addTrait(CondVal.isNullValue() - ? TraitProperty::user_condition_false - : TraitProperty::user_condition_true); + ? TraitProperty::user_condition_false + : TraitProperty::user_condition_true, + ""); else - VMI.addTrait(TraitProperty::user_condition_false); + VMI.addTrait(TraitProperty::user_condition_false, ""); continue; } @@ -2143,11 +2145,12 @@ if (Selector.ScoreOrCondition->isIntegerConstantExpr(Score, ASTCtx)) ScorePtr = &Score; else - VMI.addTrait(TraitProperty::user_condition_false); + VMI.addTrait(TraitProperty::user_condition_false, + ""); } 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 @@ if (!FirstProperty) OS << ", "; FirstProperty = false; - OS << getOpenMPContextTraitPropertyName(Property.Kind); + OS << getOpenMPContextTraitPropertyName(Property.Kind, + Property.RawString); } } OS << ")"; @@ -2231,7 +2235,9 @@ 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 @@ Selector.Properties.push_back(OMPTraitProperty()); OMPTraitProperty &Property = Selector.Properties.back(); std::pair 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 @@ const OMPTraitInfo *TI) { return TI ? OS << *TI : OS; } + +TargetOMPContext::TargetOMPContext( + ASTContext &ASTCtx, std::function &&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 --- a/clang/lib/Parse/ParseOpenMP.cpp +++ b/clang/lib/Parse/ParseOpenMP.cpp @@ -869,7 +869,8 @@ 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 @@ {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 @@ 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 @@ } 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 @@ 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 @@ if (!RequiresProperty) { TISelector.Properties.push_back( - {getOpenMPContextTraitPropertyForSelector(TISelector.Kind)}); + {getOpenMPContextTraitPropertyForSelector(TISelector.Kind), + getOpenMPContextTraitSelectorName(TISelector.Kind)}); return; } @@ -1157,7 +1161,8 @@ 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, ""}); return; } @@ -1236,8 +1241,8 @@ 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) @@ -1819,8 +1824,15 @@ VariantMatchInfo VMI; ASTContext &ASTCtx = Actions.getASTContext(); TI.getAsVariantMatchInfo(ASTCtx, VMI); - OMPContext OMPCtx(ASTCtx.getLangOpts().OpenMPIsDevice, - ASTCtx.getTargetInfo().getTriple()); + + std::function 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 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -5891,8 +5891,15 @@ return Call; ASTContext &Context = getASTContext(); - OMPContext OMPCtx(getLangOpts().OpenMPIsDevice, - Context.getTargetInfo().getTriple()); + std::function 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 Exprs; SmallVector VMIs; @@ -5904,7 +5911,8 @@ 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 --- /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 --- a/clang/test/OpenMP/declare_variant_messages.c +++ b/clang/test/OpenMP/declare_variant_messages.c @@ -137,6 +137,18 @@ #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 --- a/llvm/include/llvm/Frontend/OpenMP/OMPContext.h +++ b/llvm/include/llvm/Frontend/OpenMP/OMPContext.h @@ -70,15 +70,20 @@ /// 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 @@ /// 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 ISATraits; SmallVector ConstructTraits; SmallDenseMap ScoreMap; }; @@ -139,6 +156,7 @@ /// 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 @@ 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 ConstructTraits; }; diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def --- a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def +++ b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def @@ -1068,7 +1068,11 @@ __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, + "") __OMP_TRAIT_SELECTOR(device, arch, true) diff --git a/llvm/lib/Frontend/OpenMP/OMPContext.cpp b/llvm/lib/Frontend/OpenMP/OMPContext.cpp --- a/llvm/lib/Frontend/OpenMP/OMPContext.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPContext.cpp @@ -175,11 +175,11 @@ 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 @@ 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 Result = HandleTrait(Property, IsActiveTrait); if (Result.hasValue()) return Result.getValue(); @@ -225,7 +233,7 @@ if (!FoundInOrder) { LLVM_DEBUG(dbgs() << "[" << DEBUG_TYPE << "] Construct property " - << getOpenMPContextTraitPropertyName(Property) + << getOpenMPContextTraitPropertyName(Property, "") << " was not nested properly.\n"); return false; } @@ -425,8 +433,12 @@ 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 @@ #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 --- a/llvm/unittests/Frontend/OpenMPContextTest.cpp +++ b/llvm/unittests/Frontend/OpenMPContextTest.cpp @@ -38,11 +38,13 @@ #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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ } { // 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 @@ 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 @@ 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 @@ VariantMatchInfo LLVMHostUserCondTrueParallelParallelParallel = LLVMHostUserCondTrueParallelParallel; LLVMHostUserCondTrueParallelParallelParallel.addTrait( - TraitProperty::construct_parallel_parallel); + TraitProperty::construct_parallel_parallel, ""); EXPECT_FALSE(isVariantApplicableInContext( LLVMHostUserCondTrueParallelParallelParallel, HostLinuxParallelParallel)); @@ -277,9 +283,9 @@ 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 @@ 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,