[OpenMP] Allow traits for the OpenMP context selector `isa`

[jdoerfert]

Summary

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

[JonChesterfield]

Agreed on tests. I like the mechanism - passing a string through to the backend as a way to dispatch between isa properties looks cleanly extensible. We probably do want to emit a warning when the backend claims it doesn't know anything about said string as it'll be prone to typos.

[jdoerfert]

Add warning and negative test

[jdoerfert]

ping :)

[ABataev]

clang/include/clang/AST/OpenMPClause.h

7599

No need for the default initializer here

7658

final and virtual default destructor?

7665

override?

clang/lib/Parse/ParseOpenMP.cpp

1826

Seems to me, SourceLocation is usually captured by value.

clang/lib/Sema/SemaOpenMP.cpp

5894–5902

CE can be captured by value

llvm/lib/Frontend/OpenMP/OMPContext.cpp

205

llvm::all_of?

[jdoerfert]

Addressed @ABataev comments, thx!

[ABataev]

clang/include/clang/AST/OpenMPClause.h

7663

virtual

[jdoerfert]

Add virtual

[jdoerfert]

minor fixes

[jdoerfert]

ping :)

[ABataev]

LG with a nit.

clang/lib/AST/OpenMPClause.cpp

2304–2306

It

[thakis]

This breaks check-clang on Windows: http://45.33.8.238/win/20850/step_7.txt

Please take a look, and revert while you investigate if it takes a while to fix.

[jdoerfert]

In D83281#2182507, @thakis wrote:

This breaks check-clang on Windows: http://45.33.8.238/win/20850/step_7.txt

Please take a look, and revert while you investigate if it takes a while to fix.

Should be fixed by 8723280b68b1e5ed97a699466720b36a32a9e406, apologies for the delay.

[thakis]

Looks like it's still failing with that change: http://45.33.8.238/win/20885/step_7.txt

[jdoerfert]

In D83281#2183100, @thakis wrote:

Looks like it's still failing with that change: http://45.33.8.238/win/20885/step_7.txt

dso_local is *before* the void,.. 2 min

[thakis]

That did the trick, thanks!

The bot was broken for > 5h and 43 builds. We got lucky that nothing else broke in the meantime, but red bots make it more difficult to analyze new regressions that appear when the bots are already red. Next time, please revert while you investigate / until you have time to investigate.

[mikerice]

clang/include/clang/AST/OpenMPClause.h

7599

This field doesn't seem to have serialization code. Is that expected or an oversight?

[jdoerfert]

clang/include/clang/AST/OpenMPClause.h

7599

Oversight. I assume we need it to preserve the behavior when we go the pch route.

clang/include/clang/AST/OpenMPClause.h

 
 #define OMP_CLAUSE_CLASS(Enum, Str, Class)                                     \
   void Visit##Class(Class *S);
 #include "llvm/Frontend/OpenMP/OMPKinds.def"
 };
 
 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;
   llvm::omp::TraitSelector Kind = llvm::omp::TraitSelector::invalid;
   llvm::SmallVector<OMPTraitProperty, 1> Properties;
 };
 struct OMPTraitSet {
   llvm::omp::TraitSet Kind = llvm::omp::TraitSet::invalid;
public:
   std::string getMangledName() const;
 
   /// Print a human readable representation into \p OS.
   void print(llvm::raw_ostream &OS, const PrintingPolicy &Policy) const;
 };
 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);
+  ~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

clang/include/clang/Basic/DiagnosticParseKinds.td

def err_omp_mapper_expected_declarator : Error<
   "expected declarator on 'omp declare mapper' directive">;
 def err_omp_declare_variant_wrong_clause : Error<
   "expected '%0' clause on 'omp declare variant' directive">;
 def warn_omp_declare_variant_string_literal_or_identifier
     : Warning<"expected identifier or string literal describing a context "
               "%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
     : Warning<"expected '%0' after the %1; '%0' assumed">,
       InGroup<OpenMPClauses>;
 def warn_omp_declare_variant_ctx_not_a_property
     : Warning<"'%0' is not a valid context property for the context selector "
               "'%1' and the context set '%2'; property ignored">,

clang/include/clang/Basic/DiagnosticSemaKinds.td

 def err_omp_non_lvalue_in_map_or_motion_clauses: Error<
   "expected addressable lvalue in '%0' clause">;
 def err_omp_var_expected : Error<
   "expected variable of the '%0' type%select{|, not %2}1">;
 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">;
 def err_omp_reduction_task_not_parallel_or_worksharing : Error<
   "'reduction' clause with 'task' modifier allowed only on non-simd parallel or"
   " worksharing constructs">;
 def err_omp_expected_array_alloctraits : Error<
   "expected constant sized array of 'omp_alloctrait_t' elements, not %0">;

clang/lib/AST/OpenMPClause.cpp

 
 #include "clang/AST/OpenMPClause.h"
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/Attr.h"
 #include "clang/AST/Decl.h"
 #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"
 #include <algorithm>
 #include <cassert>
 
 using namespace clang;
 using namespace llvm;
for (const OMPTraitSelector &Selector : Set.Selectors) {
         assert(Selector.Properties.size() == 1 &&
                Selector.Properties.front().Kind ==
                    TraitProperty::user_condition_unknown &&
                "Ill-formed user condition, expected unknown trait property!");
 
         llvm::APSInt CondVal;
         if (Selector.ScoreOrCondition->isIntegerConstantExpr(CondVal, 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;
       }
 
       llvm::APSInt Score;
       llvm::APInt *ScorePtr = nullptr;
       if (Selector.ScoreOrCondition) {
         if (Selector.ScoreOrCondition->isIntegerConstantExpr(Score, 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;
 
       // TODO: This might not hold once we implement SIMD properly.
       assert(Selector.Properties.size() == 1 &&
              Selector.Properties.front().Kind ==
                  getOpenMPContextTraitPropertyForSelector(
for (const OMPTraitSelector &Selector : Set.Selectors) {
           OS << "): ";
         }
 
         bool FirstProperty = true;
         for (const OMPTraitProperty &Property : Selector.Properties) {
           if (!FirstProperty)
             OS << ", ";
           FirstProperty = false;
-          OS << getOpenMPContextTraitPropertyName(Property.Kind);
+          OS << getOpenMPContextTraitPropertyName(Property.Kind,
+                                                  Property.RawString);
         }
       }
       OS << ")";
     }
     OS << "}";
   }
 }
 
for (const OMPTraitSelector &Selector : Set.Selectors) {
           Selector.Kind, Set.Kind, AllowsTraitScore, RequiresProperty);
       OS << '$' << 's' << unsigned(Selector.Kind);
 
       if (!RequiresProperty ||
           Selector.Kind == TraitSelector::user_condition)
         continue;
 
       for (const OMPTraitProperty &Property : Selector.Properties)
-        OS << '$' << 'P' << getOpenMPContextTraitPropertyName(Property.Kind);
+        OS << '$' << 'P'
+           << getOpenMPContextTraitPropertyName(Property.Kind,
+                                                Property.RawString);
     }
   }
   return OS.str();
 }
 
 OMPTraitInfo::OMPTraitInfo(StringRef MangledName) {
   unsigned long U;
   do {
do {
       OMPTraitSelector &Selector = Set.Selectors.back();
       Selector.Kind = TraitSelector(U);
       do {
         if (!MangledName.consume_front("$P"))
           break;
         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);
   } while (true);
 }
 
 llvm::raw_ostream &clang::operator<<(llvm::raw_ostream &OS,
                                      const OMPTraitInfo &TI) {
   LangOptions LO;
   PrintingPolicy Policy(LO);
   TI.print(OS, Policy);
   return OS;
 }
 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;
+}

clang/lib/Parse/ParseOpenMP.cpp

void Parser::parseOMPTraitPropertyKind(
   StringRef Name =
       getNameFromIdOrString(*this, Tok, CONTEXT_TRAIT_LVL);
   if (Name.empty()) {
     Diag(Tok.getLocation(), diag::note_omp_declare_variant_ctx_options)
         << CONTEXT_TRAIT_LVL << listOpenMPContextTraitProperties(Set, Selector);
     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;
     return;
   }
 
   // It follows diagnosis and helping notes.
   // FIXME: We should move the diagnosis string generation into libFrontend.
Diag(NameLoc, diag::note_omp_declare_variant_ctx_try)
                getOpenMPContextTraitSetForSelector(SelectorForName))
         << Name << (RequiresProperty ? "(<property-name>)" : "");
     return;
   }
   for (const auto &PotentialSet :
        {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)
         << getOpenMPContextTraitSetName(
                getOpenMPContextTraitSetForProperty(PropertyForName))
         << getOpenMPContextTraitSelectorName(
                getOpenMPContextTraitSelectorForProperty(PropertyForName))
         << ("(" + Name + ")").str();
return (TP.Kind ==
             TP.Kind ==
                 llvm::omp::TraitProperty::implementation_extension_match_none);
   };
 
   if (IsMatchExtension(TIProperty)) {
     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;
         return false;
       }
     return true;
   }
 
if (isValidTraitPropertyForTraitSetAndSelector(TIProperty.Kind,
 
     // If we make it here the property, selector, set, score, condition, ... are
     // all valid (or have been corrected). Thus we can record the property.
     TISelector.Properties.push_back(TIProperty);
     return;
   }
 
   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(
              getOpenMPContextTraitSetForProperty(TIProperty.Kind));
   Diag(Tok.getLocation(), diag::note_omp_declare_variant_ctx_continue_here)
       << CONTEXT_TRAIT_LVL;
 }
 
if (SetForName != TraitSet::invalid) {
     Diag(NameLoc, diag::note_omp_declare_variant_ctx_try)
         << Name << "<selector-name>"
         << "<property-name>";
     return;
   }
   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)
         << Name << CONTEXT_TRAIT_LVL << CONTEXT_SELECTOR_LVL;
     Diag(NameLoc, diag::note_omp_declare_variant_ctx_try)
         << getOpenMPContextTraitSetName(
                getOpenMPContextTraitSetForProperty(PropertyForName))
         << getOpenMPContextTraitSelectorName(
Diag(NameLoc, diag::note_omp_declare_variant_ctx_try)
         << getOpenMPContextTraitSetName(
                getOpenMPContextTraitSetForSelector(SelectorForName))
         << Name << (RequiresProperty ? "(<property-name>)" : "");
     return;
   }
   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)
         << Name << CONTEXT_TRAIT_LVL << CONTEXT_SELECTOR_SET_LVL;
     Diag(NameLoc, diag::note_omp_declare_variant_ctx_try)
         << getOpenMPContextTraitSetName(
                getOpenMPContextTraitSetForProperty(PropertyForName))
         << getOpenMPContextTraitSelectorName(
case OMPD_begin_declare_variant: {
     // Skip last tokens.
     skipUntilPragmaOpenMPEnd(OMPD_begin_declare_variant);
 
     ParsingOpenMPDirectiveRAII NormalScope(*this, /*Value=*/false);
 
     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);
       break;
     }
 
     // Elide all the code till the matching end declare variant was found.
     unsigned Nesting = 1;

clang/lib/Sema/SemaOpenMP.cpp

ExprResult Sema::ActOnOpenMPCall(ExprResult Call, Scope *Scope,
   FunctionDecl *CalleeFnDecl = CE->getDirectCallee();
   if (!CalleeFnDecl)
     return Call;
 
   if (!CalleeFnDecl->hasAttr<OMPDeclareVariantAttr>())
     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;
   while (CalleeFnDecl) {
     for (OMPDeclareVariantAttr *A :
          CalleeFnDecl->specific_attrs<OMPDeclareVariantAttr>()) {
       Expr *VariantRef = A->getVariantFuncRef();
 
       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);
       Exprs.push_back(VariantRef);
     }
 
     CalleeFnDecl = CalleeFnDecl->getPreviousDecl();
   }

clang/test/OpenMP/declare_variant_device_isa_codegen_1.c

+// 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

clang/test/OpenMP/declare_variant_messages.c

 void not_marked(void);
 
 #pragma omp declare variant(not_marked) match(implementation={vendor(unknown)}, device={kind(cpu)}) // expected-note {{marked as 'declare variant' here}}
 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}}
 
 #pragma omp declare variant // expected-error {{function declaration is expected after 'declare variant' directive}}

llvm/include/llvm/Frontend/OpenMP/OMPContext.h

 TraitSelector getOpenMPContextTraitSelectorKind(StringRef Str);
 
 /// Return the trait selector for which \p Property is a property.
 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.
 StringRef getOpenMPContextTraitPropertyFullName(TraitProperty Kind);
 
 /// Return a string listing all trait sets.
 std::string listOpenMPContextTraitSets();
 

 bool isValidTraitPropertyForTraitSetAndSelector(TraitProperty Property,
                                                 TraitSelector Selector,
                                                 TraitSet Set);
 
 /// Variant match information describes the required traits and how they are
 /// 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;
 };
 
 /// The context for a source location is made up of active property traits,
 /// e.g., device={kind(host)}, and constructs traits which describe the nesting
 /// in OpenMP constructs at the location.
 struct OMPContext {
   OMPContext(bool IsDeviceCompilation, Triple TargetTriple);
 
   void addTrait(TraitProperty Property) {
     addTrait(getOpenMPContextTraitSetForProperty(Property), Property);
   }
   void addTrait(TraitSet Set, TraitProperty Property) {
     ActiveTraits.set(unsigned(Property));
     if (Set == TraitSet::construct)
       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;
 };
 
 /// Return true if \p VMI is applicable in \p Ctx, that is, all traits required
 /// by \p VMI are available in the OpenMP context \p Ctx. If \p DeviceSetOnly is
 /// true, only the device selector set, if present, are checked. Note that we
 /// still honor extension traits provided by the user.

llvm/include/llvm/Frontend/OpenMP/OMPKinds.def

 __OMP_TRAIT_PROPERTY(device, kind, nohost)
 __OMP_TRAIT_PROPERTY(device, kind, cpu)
 __OMP_TRAIT_PROPERTY(device, kind, gpu)
 __OMP_TRAIT_PROPERTY(device, kind, fpga)
 __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)
 
 __OMP_TRAIT_PROPERTY(device, arch, arm)
 __OMP_TRAIT_PROPERTY(device, arch, armeb)
 __OMP_TRAIT_PROPERTY(device, arch, aarch64)
 __OMP_TRAIT_PROPERTY(device, arch, aarch64_be)
 __OMP_TRAIT_PROPERTY(device, arch, aarch64_32)

llvm/lib/Frontend/OpenMP/OMPContext.cpp

auto HandleTrait = [MK](TraitProperty Property,
     // respectively and move on. We are not done yet!
     if ((WasFound && MK == MK_ALL) || (!WasFound && MK == MK_NONE))
       return None;
 
     // We missed a property, provide some debug output and indicate failure.
     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;
   };
 
   for (unsigned Bit : VMI.RequiredTraits.set_bits()) {
     TraitProperty Property = TraitProperty(Bit);
     if (DeviceSetOnly &&
         getOpenMPContextTraitSetForProperty(Property) != TraitSet::device)
       continue;
 
     // So far all extensions are handled elsewhere, we skip them here as they
     // are not part of the OpenMP context.
     if (getOpenMPContextTraitSelectorForProperty(Property) ==
         TraitSelector::implementation_extension)
       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();
   }
 
   if (!DeviceSetOnly) {
     // We could use isSubset here but we also want to record the match
     // locations.
for (TraitProperty Property : VMI.ConstructTraits) {
         ConstructMatches->push_back(ConstructIdx - 1);
 
       Optional<bool> Result = HandleTrait(Property, FoundInOrder);
       if (Result.hasValue())
         return Result.getValue();
 
       if (!FoundInOrder) {
         LLVM_DEBUG(dbgs() << "[" << DEBUG_TYPE << "] Construct property "
-                          << getOpenMPContextTraitPropertyName(Property)
+                          << getOpenMPContextTraitPropertyName(Property, "")
                           << " was not nested properly.\n");
         return false;
       }
 
       // TODO: Verify SIMD
     }
 
     assert(isSubset<TraitProperty>(VMI.ConstructTraits, Ctx.ConstructTraits) &&

 #define OMP_TRAIT_SELECTOR(Enum, TraitSetEnum, Str, ReqProp)                   \
   case TraitSelector::Enum:                                                    \
     return Str;
 #include "llvm/Frontend/OpenMP/OMPKinds.def"
   }
   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;
 #include "llvm/Frontend/OpenMP/OMPKinds.def"
   return TraitProperty::invalid;
 }
 TraitProperty
 llvm::omp::getOpenMPContextTraitPropertyForSelector(TraitSelector Selector) {
   return StringSwitch<TraitProperty>(
              getOpenMPContextTraitSelectorName(Selector))
 #define OMP_TRAIT_PROPERTY(Enum, TraitSetEnum, TraitSelectorEnum, Str)         \
   .Case(Str, Selector == TraitSelector::TraitSelectorEnum                      \
                  ? TraitProperty::Enum                                         \
                  : TraitProperty::invalid)
 #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:                                                    \
     return Str;
 #include "llvm/Frontend/OpenMP/OMPKinds.def"
   }
   llvm_unreachable("Unknown trait property!");
 }

llvm/unittests/Frontend/OpenMPContextTest.cpp

#define OMP_TRAIT_SELECTOR(Enum, TraitSetEnum, Str, RequiresProperty)          \
   EXPECT_EQ(TraitSelector::Enum,                                               \
             getOpenMPContextTraitSelectorKind(                                 \
                 getOpenMPContextTraitSelectorName(TraitSelector::Enum)));      \
   EXPECT_EQ(Str, getOpenMPContextTraitSelectorName(                            \
                      getOpenMPContextTraitSelectorKind(Str)));
 #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,                                  \
             getOpenMPContextTraitSelectorForProperty(TraitProperty::Enum));
 #include "llvm/Frontend/OpenMP/OMPKinds.def"
 }
 
 TEST_F(OpenMPContextTest, ValidNesting) {