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,17 @@ 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 : public llvm::omp::OMPContext { + TargetOMPContext(ASTContext &ASTCtx, const FunctionDecl *CurrentFunctionDecl); + + /// See llvm::omp::OMPContext::matchesISATrait + virtual bool matchesISATrait(StringRef) const override; + +private: + llvm::StringMap FeatureMap; +}; + } // namespace clang #endif // LLVM_CLANG_AST_OPENMPCLAUSE_H 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,17 @@ const OMPTraitInfo *TI) { return TI ? OS << *TI : OS; } + +TargetOMPContext::TargetOMPContext(ASTContext &ASTCtx, + const FunctionDecl *CurrentFunctionDecl) + : OMPContext(ASTCtx.getLangOpts().OpenMPIsDevice, + ASTCtx.getTargetInfo().getTriple()) { + ASTCtx.getFunctionFeatureMap(FeatureMap, CurrentFunctionDecl); +} + +bool TargetOMPContext::matchesISATrait(StringRef RawString) const { + auto it = FeatureMap.find(RawString); + if (it != FeatureMap.end()) + return it->second; + 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) @@ -1236,8 +1239,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 +1822,7 @@ VariantMatchInfo VMI; ASTContext &ASTCtx = Actions.getASTContext(); TI.getAsVariantMatchInfo(ASTCtx, VMI); - OMPContext OMPCtx(ASTCtx.getLangOpts().OpenMPIsDevice, - ASTCtx.getTargetInfo().getTriple()); + TargetOMPContext OMPCtx(ASTCtx, /* 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 @@ -5853,8 +5853,7 @@ return Call; ASTContext &Context = getASTContext(); - OMPContext OMPCtx(getLangOpts().OpenMPIsDevice, - Context.getTargetInfo().getTriple()); + TargetOMPContext OMPCtx(Context, getCurFunctionDecl()); SmallVector Exprs; SmallVector VMIs; 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/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; }; @@ -149,6 +166,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 @@ -1067,7 +1067,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 = std::all_of( + VMI.ISATraits.begin(), VMI.ISATraits.end(), + [&](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, \