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<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
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<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
--- 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<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
--- 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,
+                       "<condition>");
         else
-          VMI.addTrait(TraitProperty::user_condition_false);
+          VMI.addTrait(TraitProperty::user_condition_false, "<condition>");
         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,
+                       "<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 @@
           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<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 @@
                                      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
--- 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,15 @@
     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
--- 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<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;
@@ -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<StringRef, 8> ISATraits;
   SmallVector<TraitProperty, 8> ConstructTraits;
   SmallDenseMap<TraitProperty, APInt> 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<TraitProperty, 8> 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,
+                   "<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
--- 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<bool> 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,                                  \