diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -3316,7 +3316,8 @@
     VariadicUnsignedArgument<"CtxSelectorSets">,
     VariadicUnsignedArgument<"CtxSelectors">,
     VariadicStringArgument<"ImplVendors">,
-    VariadicStringArgument<"DeviceKinds">
+    VariadicStringArgument<"DeviceKinds">,
+    VariadicExprArgument<"DeviceISAs">
   ];
   let AdditionalMembers = [{
     void printScore(raw_ostream & OS, const PrintingPolicy &Policy, unsigned I) const {
@@ -3387,6 +3388,19 @@
             }
             OS << ")";
             break;
+          case OMP_CTX_isa:
+            assert(CtxSet == OMP_CTX_SET_device &&
+                   "Expected device context selector set.");
+            OS << "isa(";
+            if (deviceISAs_size() > 0) {
+              (*deviceISAs().begin())->printPretty(OS, nullptr, Policy);
+              for (const Expr *ISA : llvm::drop_begin(deviceISAs(), 1)) {
+                OS << ", ";
+                ISA->printPretty(OS, nullptr, Policy);
+              }
+            }
+            OS << ")";
+            break;
           case OMP_CTX_unknown:
             llvm_unreachable("Unknown context selector.");
           }
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
@@ -9634,6 +9634,14 @@
 def note_omp_marked_declare_variant_here : Note<"marked as 'declare variant' here">;
 def err_omp_one_defaultmap_each_category: Error<
   "at most one defaultmap clause for each variable-category can appear on the directive">;
+def err_omp_trait_not_string : Error<
+  "trait expression must have a string type, not %0">;
+def note_omp_trait_conversion_here : Note<
+  "conversion to a string type %0 declared here">;
+def err_omp_trait_ambiguous_conversion : Error<
+  "ambiguous conversion from type %0 to a string type">;
+def err_omp_trait_not_constant_string : Error<
+  "trait expression must be a constant string expression">;
 } // end of OpenMP category
 
 let CategoryName = "Related Result Type Issue" in {
diff --git a/clang/include/clang/Basic/OpenMPKinds.def b/clang/include/clang/Basic/OpenMPKinds.def
--- a/clang/include/clang/Basic/OpenMPKinds.def
+++ b/clang/include/clang/Basic/OpenMPKinds.def
@@ -226,6 +226,7 @@
 // OpenMP context selectors.
 OPENMP_CONTEXT_SELECTOR(vendor)
 OPENMP_CONTEXT_SELECTOR(kind)
+OPENMP_CONTEXT_SELECTOR(isa)
 
 // OpenMP directives.
 OPENMP_DIRECTIVE(threadprivate)
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -61,6 +61,7 @@
 #include <vector>
 
 namespace llvm {
+  class Any;
   class APSInt;
   template <typename ValueT> struct DenseMapInfo;
   template <typename ValueT, typename ValueInfoT> class DenseSet;
@@ -9316,7 +9317,7 @@
   /// Struct to store the context selectors info for declare variant directive.
   using OMPCtxStringType = SmallString<8>;
   using OMPCtxSelectorData =
-      OpenMPCtxSelectorData<SmallVector<OMPCtxStringType, 4>, ExprResult>;
+      OpenMPCtxSelectorData<SmallVector<llvm::Any, 4>, ExprResult>;
 
   /// Checks if the variant/multiversion functions are compatible.
   bool areMultiversionVariantFunctionsCompatible(
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -21,10 +21,12 @@
 #include "clang/Basic/BitmaskEnum.h"
 #include "llvm/ADT/ArrayRef.h"
 #include "llvm/ADT/SetOperations.h"
+#include "llvm/ADT/UniqueVector.h"
 #include "llvm/Bitcode/BitcodeReader.h"
 #include "llvm/IR/DerivedTypes.h"
 #include "llvm/IR/GlobalValue.h"
 #include "llvm/IR/Value.h"
+#include "llvm/Support/ConvertUTF.h"
 #include "llvm/Support/Format.h"
 #include "llvm/Support/raw_ostream.h"
 #include <cassert>
@@ -11028,7 +11030,7 @@
 
 namespace {
 using OMPContextSelectorData =
-    OpenMPCtxSelectorData<ArrayRef<StringRef>, llvm::APSInt>;
+    OpenMPCtxSelectorData<SmallVector<SmallString<8>, 4>, llvm::APSInt>;
 using CompleteOMPContextSelectorData = SmallVector<OMPContextSelectorData, 4>;
 } // anonymous namespace
 
@@ -11135,6 +11137,19 @@
   return true;
 }
 
+/// Checks for device={isa(<isa>)} context selector.
+/// \returns true if <isa>=current isa or false, otherwise.
+template <>
+bool checkContext<OMP_CTX_SET_device, OMP_CTX_isa, CodeGenModule &>(
+    const OMPContextSelectorData &Data, CodeGenModule &CGM) {
+  for (StringRef Name : Data.Names) {
+    llvm::Triple CustomTriple(Name);
+    if (CGM.getTarget().getTriple().getArch() != CustomTriple.getArch())
+      return false;
+  }
+  return true;
+}
+
 bool matchesContext(CodeGenModule &CGM,
                     const CompleteOMPContextSelectorData &ContextData) {
   for (const OMPContextSelectorData &Data : ContextData) {
@@ -11152,6 +11167,13 @@
                                                                            CGM))
         return false;
       break;
+    case OMP_CTX_isa:
+      assert(Data.CtxSet == OMP_CTX_SET_device &&
+             "Expected device context selector set.");
+      if (!checkContext<OMP_CTX_SET_device, OMP_CTX_isa, CodeGenModule &>(Data,
+                                                                          CGM))
+        return false;
+      break;
     case OMP_CTX_unknown:
       llvm_unreachable("Unknown context selector kind.");
     }
@@ -11177,15 +11199,47 @@
     case OMP_CTX_vendor:
       assert(CtxSet == OMP_CTX_SET_implementation &&
              "Expected implementation context selector set.");
-      Data.back().Names =
-          llvm::makeArrayRef(A->implVendors_begin(), A->implVendors_end());
+      Data.back().Names.append(A->implVendors_begin(), A->implVendors_end());
       break;
     case OMP_CTX_kind:
       assert(CtxSet == OMP_CTX_SET_device &&
              "Expected device context selector set.");
-      Data.back().Names =
-          llvm::makeArrayRef(A->deviceKinds_begin(), A->deviceKinds_end());
+      Data.back().Names.append(A->deviceKinds_begin(), A->deviceKinds_end());
       break;
+    case OMP_CTX_isa: {
+      assert(CtxSet == OMP_CTX_SET_device &&
+             "Expected device context selector set.");
+      llvm::UniqueVector<std::string> ISAs;
+      for (const Expr *ISA : A->deviceISAs()) {
+        const auto *SL = cast<StringLiteral>(ISA);
+        std::string ISAName;
+        switch(SL->getKind()) {
+        case StringLiteral::Ascii:
+        case StringLiteral::UTF8:
+          ISAName = SL->getString();
+          break;
+        case StringLiteral::UTF16:
+          if (!llvm::convertUTF16ToUTF8String(
+              llvm::makeArrayRef(SL->getBytes().begin(), SL->getBytes().end()),
+              ISAName))
+            return CompleteOMPContextSelectorData();
+          break;
+        case StringLiteral::UTF32:
+          llvm_unreachable("UTF32 is unsupported.");
+        case StringLiteral::Wide: {
+          std::wstring WideString;
+          for (int I = 0, E = SL->getLength(); I < E; ++I)
+            WideString += static_cast<wchar_t>(SL->getCodeUnit(I));
+          if (!llvm::convertWideToUTF8(WideString, ISAName))
+            return CompleteOMPContextSelectorData();
+          break;
+        }
+        }
+        ISAs.insert(ISAName);
+      }
+      Data.back().Names.append(ISAs.begin(), ISAs.end());
+      break;
+    }
     case OMP_CTX_unknown:
       llvm_unreachable("Unknown context selector kind.");
     }
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
@@ -16,6 +16,7 @@
 #include "clang/Parse/Parser.h"
 #include "clang/Parse/RAIIObjectsForParser.h"
 #include "clang/Sema/Scope.h"
+#include "llvm/ADT/Any.h"
 #include "llvm/ADT/PointerIntPair.h"
 #include "llvm/ADT/UniqueVector.h"
 
@@ -875,6 +876,7 @@
     break;
   }
   case OMP_CTX_kind:
+  case OMP_CTX_isa:
   case OMP_CTX_unknown:
     P.Diag(Tok.getLocation(), diag::warn_omp_declare_variant_cs_name_expected)
         << "implementation";
@@ -888,6 +890,7 @@
 
 /// Parse context selector for 'device' selector set:
 /// 'kind' '(' <kind> { ',' <kind> } ')'
+/// 'isa' '(' <ISA> { ',' <ISA> } ')'
 static void
 parseDeviceSelector(Parser &P, SourceLocation Loc,
                     llvm::StringMap<SourceLocation> &UsedCtx,
@@ -959,6 +962,30 @@
       Data.emplace_back(OMP_CTX_SET_device, CSKind, ExprResult(), Kinds);
     break;
   }
+  case OMP_CTX_isa: {
+    // Parse '('.
+    BalancedDelimiterTracker T(P, tok::l_paren, tok::annot_pragma_openmp_end);
+    (void)T.expectAndConsume(diag::err_expected_lparen_after,
+                             CtxSelectorName.data());
+    SmallVector<ExprResult, 4> ISAs;
+    SourceLocation PrevLoc;
+    do {
+      PrevLoc = Tok.getLocation();
+      // Parse <ISA>, which is constant expression of string type.
+      ExprResult ISA = P.ParseConstantExpression();
+      if (ISA.isUsable())
+        ISAs.push_back(ISA);
+      if (!P.TryConsumeToken(tok::comma) && Tok.isNot(tok::r_paren))
+        P.Diag(Tok, diag::err_expected_punc) << "instruction set architecture";
+    } while (Tok.isNot(tok::r_paren) &&
+             Tok.isNot(tok::annot_pragma_openmp_end) &&
+             PrevLoc.getRawEncoding() != Tok.getLocation().getRawEncoding());
+    // Parse ')'.
+    (void)T.consumeClose();
+    if (!ISAs.empty())
+      Data.emplace_back(OMP_CTX_SET_device, CSKind, ExprResult(), ISAs);
+    break;
+  }
   case OMP_CTX_vendor:
   case OMP_CTX_unknown:
     P.Diag(Tok.getLocation(), diag::warn_omp_declare_variant_cs_name_expected)
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
@@ -28,6 +28,7 @@
 #include "clang/Sema/Scope.h"
 #include "clang/Sema/ScopeInfo.h"
 #include "clang/Sema/SemaInternal.h"
+#include "llvm/ADT/Any.h"
 #include "llvm/ADT/PointerEmbeddedInt.h"
 using namespace clang;
 
@@ -5362,6 +5363,92 @@
   return std::make_pair(FD, cast<Expr>(DRE));
 }
 
+static ExprResult performOpenMPImplicitStringConversion(Sema &S, Expr *E) {
+  if (E->isTypeDependent() || E->isValueDependent() ||
+      E->isInstantiationDependent())
+    return E;
+  E = S.DefaultFunctionArrayLvalueConversion(E).get();
+  if (!E)
+    return ExprError();
+  class StringConverter : public Sema::ContextualImplicitConverter {
+  public:
+    StringConverter()
+        : Sema::ContextualImplicitConverter(/*Suppress=*/false,
+                                            /*SuppressConversion=*/true) {}
+
+    bool match(QualType ConvType) override {
+      if (ConvType->isArrayType() || ConvType->isPointerType()) {
+        const Type *ElTy = ConvType->getPointeeOrArrayElementType();
+        return ElTy->isWideCharType() || ElTy->isCharType() ||
+               ElTy->isChar8Type() || ElTy->isChar16Type();
+      }
+      return false;
+    }
+
+    Sema::SemaDiagnosticBuilder diagnoseNoMatch(Sema &S, SourceLocation Loc,
+                                                QualType T) override {
+      return S.Diag(Loc, diag::err_omp_trait_not_string) << T;
+    }
+
+    Sema::SemaDiagnosticBuilder diagnoseIncomplete(Sema &S, SourceLocation Loc,
+                                                   QualType T) override {
+      return S.Diag(Loc, diag::err_omp_incomplete_type) << T;
+    }
+
+    Sema::SemaDiagnosticBuilder diagnoseExplicitConv(Sema &S,
+                                                     SourceLocation Loc,
+                                                     QualType T,
+                                                     QualType ConvTy) override {
+      return S.Diag(Loc, diag::err_omp_explicit_conversion) << T << ConvTy;
+    }
+
+    Sema::SemaDiagnosticBuilder noteExplicitConv(Sema &S,
+                                                 CXXConversionDecl *Conv,
+                                                 QualType ConvTy) override {
+      return S.Diag(Conv->getLocation(), diag::note_omp_trait_conversion_here)
+             << ConvTy;
+    }
+
+    Sema::SemaDiagnosticBuilder diagnoseAmbiguous(Sema &S, SourceLocation Loc,
+                                                  QualType T) override {
+      return S.Diag(Loc, diag::err_omp_trait_ambiguous_conversion) << T;
+    }
+
+    Sema::SemaDiagnosticBuilder noteAmbiguous(Sema &S, CXXConversionDecl *Conv,
+                                              QualType ConvTy) override {
+      return S.Diag(Conv->getLocation(), diag::note_omp_trait_conversion_here)
+             << ConvTy;
+    }
+
+    Sema::SemaDiagnosticBuilder diagnoseConversion(Sema &S, SourceLocation Loc,
+                                                   QualType T,
+                                                   QualType ConvTy) override {
+      llvm_unreachable("conversion functions are permitted");
+    }
+  } Converter;
+
+  ExprResult ER =
+      S.PerformContextualImplicitConversion(E->getBeginLoc(), E, Converter);
+  if (!ER.isUsable())
+    return ExprError();
+  E = ER.get();
+  Expr::EvalResult Res;
+  if (!E->EvaluateAsRValue(Res, S.Context) || !Res.Val.isLValue()) {
+    S.Diag(E->getExprLoc(), diag::err_omp_trait_not_constant_string)
+        << E->getSourceRange();
+    return ExprError();
+  }
+  const auto *SL = dyn_cast_or_null<StringLiteral>(
+      Res.Val.getLValueBase().dyn_cast<const Expr *>());
+  if (!SL) {
+    S.Diag(E->getExprLoc(), diag::err_omp_trait_not_constant_string)
+        << E->getSourceRange();
+    return ExprError();
+  }
+  return StringLiteral::Create(S.Context, SL->getBytes(), SL->getKind(),
+                               SL->isPascal(), SL->getType(), E->getExprLoc());
+}
+
 void Sema::ActOnOpenMPDeclareVariantDirective(
     FunctionDecl *FD, Expr *VariantRef, SourceRange SR,
     ArrayRef<OMPCtxSelectorData> Data) {
@@ -5371,6 +5458,7 @@
   SmallVector<unsigned, 4> CtxSets;
   SmallVector<unsigned, 4> Ctxs;
   SmallVector<StringRef, 4> ImplVendors, DeviceKinds;
+  SmallVector<Expr *, 4> DeviceISAs;
   bool IsError = false;
   for (const OMPCtxSelectorData &D : Data) {
     OpenMPContextSelectorSetKind CtxSet = D.CtxSet;
@@ -5394,29 +5482,45 @@
       // The kind, arch, and isa selectors are given the values 2^l, 2^(l+1) and
       // 2^(l+2), respectively, where l is the number of traits in the construct
       // set.
-      // TODO: implement correct logic for isa and arch traits.
+      // TODO: implement correct logic for arch traits.
       // TODO: take the construct context set into account when it is
       // implemented.
       int L = 0; // Currently set the number of traits in construct set to 0,
                  // since the construct trait set in not supported yet.
       if (CtxSet == OMP_CTX_SET_device && Ctx == OMP_CTX_kind)
         Score = ActOnIntegerConstant(SourceLocation(), std::pow(2, L)).get();
+      else if (CtxSet == OMP_CTX_SET_device && Ctx == OMP_CTX_isa)
+        Score =
+            ActOnIntegerConstant(SourceLocation(), std::pow(2, L + 1)).get();
       else
         Score = ActOnIntegerConstant(SourceLocation(), 0).get();
     }
-    switch (Ctx) {
-    case OMP_CTX_vendor:
-      assert(CtxSet == OMP_CTX_SET_implementation &&
-             "Expected implementation context selector set.");
-      ImplVendors.append(D.Names.begin(), D.Names.end());
-      break;
-    case OMP_CTX_kind:
-      assert(CtxSet == OMP_CTX_SET_device &&
-             "Expected device context selector set.");
-      DeviceKinds.append(D.Names.begin(), D.Names.end());
-      break;
-    case OMP_CTX_unknown:
-      llvm_unreachable("Unknown context selector kind.");
+    for (const llvm::Any &Val : D.Names) {
+      switch (Ctx) {
+      case OMP_CTX_vendor:
+        assert(CtxSet == OMP_CTX_SET_implementation &&
+               "Expected implementation context selector set.");
+        ImplVendors.push_back(*llvm::any_cast<OMPCtxStringType>(&Val));
+        break;
+      case OMP_CTX_kind:
+        assert(CtxSet == OMP_CTX_SET_device &&
+               "Expected device context selector set.");
+        DeviceKinds.push_back(*llvm::any_cast<OMPCtxStringType>(&Val));
+        break;
+      case OMP_CTX_isa: {
+        assert(CtxSet == OMP_CTX_SET_device &&
+               "Expected device context selector set.");
+        auto ER = llvm::any_cast<ExprResult>(Val);
+        ER = performOpenMPImplicitStringConversion(*this, ER.get());
+        if (ER.isUsable())
+          DeviceISAs.push_back(ER.get());
+        else
+          IsError = true;
+        break;
+      }
+      case OMP_CTX_unknown:
+        llvm_unreachable("Unknown context selector kind.");
+      }
     }
     IsError = IsError || !Score;
     CtxSets.push_back(CtxSet);
@@ -5428,7 +5532,7 @@
         Context, VariantRef, CtxScores.begin(), CtxScores.size(),
         CtxSets.begin(), CtxSets.size(), Ctxs.begin(), Ctxs.size(),
         ImplVendors.begin(), ImplVendors.size(), DeviceKinds.begin(),
-        DeviceKinds.size(), SR);
+        DeviceKinds.size(), DeviceISAs.begin(), DeviceISAs.size(), SR);
     FD->addAttr(NewAttr);
   }
 }
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -23,6 +23,7 @@
 #include "clang/Sema/Lookup.h"
 #include "clang/Sema/Template.h"
 #include "clang/Sema/TemplateInstCallback.h"
+#include "llvm/ADT/Any.h"
 #include "llvm/Support/TimeProfiler.h"
 
 using namespace clang;
@@ -404,30 +405,33 @@
         *std::next(Attr.ctxSelectorSets_begin(), I));
     auto Ctx = static_cast<OpenMPContextSelectorKind>(
         *std::next(Attr.ctxSelectors_begin(), I));
-    switch (CtxSet) {
-    case OMP_CTX_SET_implementation:
-      switch (Ctx) {
-      case OMP_CTX_vendor:
-        Data.emplace_back(CtxSet, Ctx, Score, Attr.implVendors());
-        break;
-      case OMP_CTX_kind:
-      case OMP_CTX_unknown:
-        llvm_unreachable("Unexpected context selector kind.");
-      }
+    SmallVector<llvm::Any, 4> Values;
+    switch (Ctx) {
+    case OMP_CTX_vendor:
+      assert(CtxSet == OMP_CTX_SET_implementation &&
+             "Expected implementation context selector set.");
+      for (StringRef Vendor : Attr.implVendors())
+        Values.push_back(Sema::OMPCtxStringType(Vendor));
       break;
-    case OMP_CTX_SET_device:
-      switch (Ctx) {
-      case OMP_CTX_kind:
-        Data.emplace_back(CtxSet, Ctx, Score, Attr.deviceKinds());
-        break;
-      case OMP_CTX_vendor:
-      case OMP_CTX_unknown:
-        llvm_unreachable("Unexpected context selector kind.");
+    case OMP_CTX_kind:
+      assert(CtxSet == OMP_CTX_SET_device &&
+             "Expected device context selector set.");
+      for (StringRef Kind : Attr.deviceKinds())
+        Values.push_back(Sema::OMPCtxStringType(Kind));
+      break;
+    case OMP_CTX_isa:
+      assert(CtxSet == OMP_CTX_SET_device &&
+             "Expected device context selector set.");
+      for (Expr *ISA : Attr.deviceISAs()) {
+        ExprResult ISARes = Subst(ISA);
+        if (ISARes.isUsable())
+          Values.push_back(ISARes);
       }
       break;
-    case OMP_CTX_SET_unknown:
-      llvm_unreachable("Unexpected context selector set kind.");
+    case OMP_CTX_unknown:
+      llvm_unreachable("Unexpected context selector kind.");
     }
+    Data.emplace_back(CtxSet, Ctx, Score, Values);
   }
   S.ActOnOpenMPDeclareVariantDirective(DeclVarData.getValue().first,
                                        DeclVarData.getValue().second,
diff --git a/clang/test/OpenMP/declare_variant_ast_print.c b/clang/test/OpenMP/declare_variant_ast_print.c
--- a/clang/test/OpenMP/declare_variant_ast_print.c
+++ b/clang/test/OpenMP/declare_variant_ast_print.c
@@ -8,10 +8,10 @@
 
 #pragma omp declare variant(foo) match(xxx={}, yyy={ccc})
 #pragma omp declare variant(foo) match(xxx={vvv})
-#pragma omp declare variant(foo) match(implementation={vendor(llvm)}, device={kind(fpga)})
+#pragma omp declare variant(foo) match(implementation={vendor(llvm)}, device={kind(fpga), isa("i586")})
 #pragma omp declare variant(foo) match(implementation={vendor(llvm), xxx})
-#pragma omp declare variant(foo) match(implementation={vendor(unknown)}, device={kind(gpu)})
-#pragma omp declare variant(foo) match(implementation={vendor(score(5): ibm, xxx, ibm)}, device={kind(cpu, nohost)})
+#pragma omp declare variant(foo) match(implementation={vendor(unknown)}, device={kind(gpu), isa("x86""_64", "arm")})
+#pragma omp declare variant(foo) match(implementation={vendor(score(5): ibm, xxx, ibm)}, device={kind(cpu, nohost), isa("ppc64le")})
 #pragma omp declare variant(foo) match(device={kind(host)})
 #pragma omp declare variant(foo) match(device={kind(nohost), xxx})
 int bar(void);
@@ -19,8 +19,8 @@
 // CHECK:      int foo();
 // CHECK-NEXT: #pragma omp declare variant(foo) match(device={kind(nohost)})
 // CHECK-NEXT: #pragma omp declare variant(foo) match(device={kind(host)})
-// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(5):ibm, xxx)},device={kind(cpu, nohost)})
-// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0):unknown)},device={kind(gpu)})
+// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(5):ibm, xxx)},device={kind(cpu, nohost),isa("ppc64le")})
+// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0):unknown)},device={kind(gpu),isa("x86_64", "arm")})
 // CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0):llvm)})
-// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0):llvm)},device={kind(fpga)})
+// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0):llvm)},device={kind(fpga),isa("i586")})
 // CHECK-NEXT: int bar();
diff --git a/clang/test/OpenMP/declare_variant_ast_print.cpp b/clang/test/OpenMP/declare_variant_ast_print.cpp
--- a/clang/test/OpenMP/declare_variant_ast_print.cpp
+++ b/clang/test/OpenMP/declare_variant_ast_print.cpp
@@ -17,20 +17,22 @@
 // CHECK-NEXT: return int();
 // CHECK-NEXT: }
 
-// CHECK:      #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(5):ibm)},device={kind(fpga)})
+// CHECK:      #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(5):ibm)},device={kind(fpga),isa("ppc64le")})
 // CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):unknown)})
-// CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):llvm)},device={kind(cpu)})
+// CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):llvm)},device={kind(cpu),isa("x86_64", "arm")})
 // CHECK-NEXT: int bar();
 #pragma omp declare variant(foofoo <int>) match(xxx = {})
 #pragma omp declare variant(foofoo <int>) match(xxx = {vvv})
-#pragma omp declare variant(foofoo <int>) match(implementation={vendor(llvm), xxx}, device={kind(cpu)})
+#pragma omp declare variant(foofoo <int>) match(implementation={vendor(llvm), xxx}, device={kind(cpu), isa("x86""_64", "arm")})
 #pragma omp declare variant(foofoo <int>) match(implementation={vendor(unknown)})
-#pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(5): ibm)}, device={kind(fpga)})
+#pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(5): ibm)}, device={kind(fpga), isa("ppc64le")})
 int bar();
 
-// CHECK:      #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(C + 5):ibm, xxx)},device={kind(cpu, host)})
+constexpr auto Arch = "x86""_64";
+
+// CHECK:      #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(C + 5):ibm, xxx)},device={kind(cpu, host),isa("ppc64le")})
 // CHECK-NEXT: #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(0):unknown)})
-// CHECK-NEXT: #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(0):llvm)},device={kind(cpu)})
+// CHECK-NEXT: #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(0):llvm)},device={kind(cpu),isa("x86_64", "arm")})
 // CHECK-NEXT: template <typename T, int C> T barbar();
 #pragma omp declare variant(foofoo <T>) match(xxx = {})
 #pragma omp declare variant(foofoo <T>) match(xxx = {vvv})
@@ -38,15 +40,15 @@
 #pragma omp declare variant(foofoo <T>) match(user = {score(<expr>) : condition(<expr>)})
 #pragma omp declare variant(foofoo <T>) match(user = {condition(<expr>)})
 #pragma omp declare variant(foofoo <T>) match(user = {condition(<expr>)})
-#pragma omp declare variant(foofoo <T>) match(implementation={vendor(llvm)},device={kind(cpu)})
+#pragma omp declare variant(foofoo <T>) match(implementation={vendor(llvm)},device={kind(cpu), isa(Arch, "arm")})
 #pragma omp declare variant(foofoo <T>) match(implementation={vendor(unknown)})
-#pragma omp declare variant(foofoo <T>) match(implementation={vendor(score(C+5): ibm, xxx, ibm)},device={kind(cpu,host)})
+#pragma omp declare variant(foofoo <T>) match(implementation={vendor(score(C+5): ibm, xxx, ibm)},device={kind(cpu,host), isa("ppc64le")})
 template <typename T, int C>
 T barbar();
 
-// CHECK:      #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(3 + 5):ibm, xxx)},device={kind(cpu, host)})
+// CHECK:      #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(3 + 5):ibm, xxx)},device={kind(cpu, host),isa("ppc64le")})
 // CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):unknown)})
-// CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):llvm)},device={kind(cpu)})
+// CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):llvm)},device={kind(cpu),isa("x86_64", "arm")})
 // CHECK-NEXT: template<> int barbar<int, 3>();
 
 // CHECK-NEXT: int baz() {
@@ -66,19 +68,19 @@
 void h_ref(C *hp, C *hp2, C *hq, C *lin) {
 }
 
-// CHECK:      #pragma omp declare variant(h_ref<C>) match(implementation={vendor(score(0):unknown)},device={kind(nohost)})
-// CHECK-NEXT: #pragma omp declare variant(h_ref<C>) match(implementation={vendor(score(0):llvm)},device={kind(gpu)})
+// CHECK:      #pragma omp declare variant(h_ref<C>) match(implementation={vendor(score(0):unknown)},device={kind(nohost),isa("x86", "aarch64")})
+// CHECK-NEXT: #pragma omp declare variant(h_ref<C>) match(implementation={vendor(score(0):llvm)},device={kind(gpu),isa("ppc64")})
 // CHECK-NEXT: template <class C> void h(C *hp, C *hp2, C *hq, C *lin) {
 // CHECK-NEXT: }
 #pragma omp declare variant(h_ref <C>) match(xxx = {})
-#pragma omp declare variant(h_ref <C>) match(implementation={vendor(llvm)}, device={kind(gpu)})
-#pragma omp declare variant(h_ref <C>) match(implementation={vendor(unknown)},device={kind(nohost)})
+#pragma omp declare variant(h_ref <C>) match(implementation={vendor(llvm)}, device={kind(gpu),isa("ppc64")})
+#pragma omp declare variant(h_ref <C>) match(implementation={vendor(unknown)},device={kind(nohost), isa("x86","aarch64")})
 template <class C>
 void h(C *hp, C *hp2, C *hq, C *lin) {
 }
 
-// CHECK:      #pragma omp declare variant(h_ref<float>) match(implementation={vendor(score(0):unknown)},device={kind(nohost)})
-// CHECK-NEXT: #pragma omp declare variant(h_ref<float>) match(implementation={vendor(score(0):llvm)},device={kind(gpu)})
+// CHECK:      #pragma omp declare variant(h_ref<float>) match(implementation={vendor(score(0):unknown)},device={kind(nohost),isa("x86", "aarch64")})
+// CHECK-NEXT: #pragma omp declare variant(h_ref<float>) match(implementation={vendor(score(0):llvm)},device={kind(gpu),isa("ppc64")})
 // CHECK-NEXT: template<> void h<float>(float *hp, float *hp2, float *hq, float *lin) {
 // CHECK-NEXT: }
 
@@ -86,7 +88,7 @@
 // CHECK-NEXT:   h((float *)hp, (float *)hp2, (float *)hq, (float *)lin);
 // CHECK-NEXT: }
 #pragma omp declare variant(h_ref <double>) match(xxx = {})
-#pragma omp declare variant(h_ref <double>) match(implementation={vendor(ibm)},device={kind(cpu,gpu)})
+#pragma omp declare variant(h_ref <double>) match(implementation={vendor(ibm)},device={kind(cpu,gpu),isa("x86", "aarch64")})
 #pragma omp declare variant(h_ref <double>) match(implementation={vendor(unknown)})
 template <>
 void h(double *hp, double *hp2, double *hq, double *lin) {
@@ -97,36 +99,36 @@
 int fn();
 // CHECK: int fn(int);
 int fn(int);
-// CHECK:      #pragma omp declare variant(fn) match(implementation={vendor(score(0):unknown)},device={kind(cpu, gpu)})
+// CHECK:      #pragma omp declare variant(fn) match(implementation={vendor(score(0):unknown)},device={kind(cpu, gpu),isa("x86", "aarch64")})
 // CHECK-NEXT: #pragma omp declare variant(fn) match(implementation={vendor(score(0):llvm)})
 // CHECK-NEXT: int overload();
 #pragma omp declare variant(fn) match(xxx = {})
 #pragma omp declare variant(fn) match(implementation={vendor(llvm)})
-#pragma omp declare variant(fn) match(implementation={vendor(unknown)},device={kind(cpu,gpu)})
+#pragma omp declare variant(fn) match(implementation={vendor(unknown)},device={kind(cpu,gpu),isa("x86", "aarch64")})
 int overload(void);
 
 // CHECK:      int fn_deduced_variant() {
 // CHECK-NEXT: return 0;
 // CHECK-NEXT: }
 auto fn_deduced_variant() { return 0; }
-// CHECK:      #pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(score(0):unknown)},device={kind(gpu, nohost)})
-// CHECK-NEXT: #pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(score(0):llvm)},device={kind(cpu, host)})
+// CHECK:      #pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(score(0):unknown)},device={kind(gpu, nohost),isa("ppc64le", "aarch64")})
+// CHECK-NEXT: #pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(score(0):llvm)},device={kind(cpu, host),isa("x86", "aarch64")})
 // CHECK-NEXT: int fn_deduced();
 #pragma omp declare variant(fn_deduced_variant) match(xxx = {})
-#pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(llvm)},device={kind(cpu,host)})
-#pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(unknown)},device={kind(gpu,nohost)})
+#pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(llvm)},device={kind(cpu,host),isa("x86", "aarch64")})
+#pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(unknown)},device={kind(gpu,nohost),isa("ppc64le", "aarch64")})
 int fn_deduced();
 
 // CHECK: int fn_deduced_variant1();
 int fn_deduced_variant1();
-// CHECK:      #pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(score(0):unknown)},device={kind(cpu, host)})
-// CHECK-NEXT: #pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(score(0):ibm)},device={kind(gpu, nohost)})
+// CHECK:      #pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(score(0):unknown)},device={kind(cpu, host),isa("ppc64le", "aarch64")})
+// CHECK-NEXT: #pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(score(0):ibm)},device={kind(gpu, nohost),isa("x86", "aarch64")})
 // CHECK-NEXT: int fn_deduced1() {
 // CHECK-NEXT: return 0;
 // CHECK-NEXT: }
 #pragma omp declare variant(fn_deduced_variant1) match(xxx = {})
-#pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(ibm)},device={kind(gpu,nohost)})
-#pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(unknown)},device={kind(cpu,host)})
+#pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(ibm)},device={kind(gpu,nohost),isa("x86", "aarch64")})
+#pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(unknown)},device={kind(cpu,host),isa("ppc64le", "aarch64")})
 auto fn_deduced1() { return 0; }
 
 // CHECK:      struct SpecialFuncs {
@@ -140,11 +142,11 @@
 // CHECK-NEXT: }
 // CHECK-NEXT: void bar(int) {
 // CHECK-NEXT: }
-// CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)},device={kind(nohost)})
-// CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::bar) match(implementation={vendor(score(0):ibm)},device={kind(cpu)})
+// CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)},device={kind(nohost),isa("ppc64le", "aarch64")})
+// CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::bar) match(implementation={vendor(score(0):ibm)},device={kind(cpu),isa("x86", "aarch64")})
 // CHECK-NEXT: void foo1() {
 // CHECK-NEXT: }
-// CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)},device={kind(cpu, host)})
+// CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)},device={kind(cpu, host),isa("x86_64")})
 // CHECK-NEXT: void xxx();
 // CHECK-NEXT: } s;
 struct SpecialFuncs {
@@ -157,14 +159,14 @@
   void bar(int) {}
 #pragma omp declare variant(SpecialFuncs::baz) match(xxx = {})
 #pragma omp declare variant(SpecialFuncs::bar) match(xxx = {})
-#pragma omp declare variant(SpecialFuncs::bar) match(implementation={vendor(ibm)},device={kind(cpu)})
-#pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(unknown)},device={kind(nohost)})
+#pragma omp declare variant(SpecialFuncs::bar) match(implementation={vendor(ibm)},device={kind(cpu),isa("x86", "aarch64")})
+#pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(unknown)},device={kind(nohost),isa("ppc64le", "aarch64")})
   void foo1() {}
-#pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(unknown)},device={kind(cpu, host)})
+#pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(unknown)},device={kind(cpu, host),isa("x86_64")})
   void xxx();
 } s;
 
-// CHECK:      #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)},device={kind(cpu, host)})
+// CHECK:      #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)},device={kind(cpu, host),isa("x86_64")})
 // CHECK-NEXT: void SpecialFuncs::xxx() {
 // CHECK-NEXT: }
 void SpecialFuncs::xxx() {}
@@ -173,11 +175,11 @@
 // CHECK-NEXT: }
 static void static_f_variant() {}
 // CHECK:      #pragma omp declare variant(static_f_variant) match(implementation={vendor(score(0):unknown)})
-// CHECK-NEXT: #pragma omp declare variant(static_f_variant) match(implementation={vendor(score(0):llvm)},device={kind(fpga)})
+// CHECK-NEXT: #pragma omp declare variant(static_f_variant) match(implementation={vendor(score(0):llvm)},device={kind(fpga),isa("aarch64")})
 // CHECK-NEXT: static void static_f() {
 // CHECK-NEXT: }
 #pragma omp declare variant(static_f_variant) match(xxx = {})
-#pragma omp declare variant(static_f_variant) match(implementation={vendor(llvm)},device={kind(fpga)})
+#pragma omp declare variant(static_f_variant) match(implementation={vendor(llvm)},device={kind(fpga),isa("aarch64")})
 #pragma omp declare variant(static_f_variant) match(implementation={vendor(unknown)})
 static void static_f() {}
 
@@ -192,19 +194,19 @@
 
 // CHECK: int fn_linkage_variant();
 // CHECK: extern "C" {
-// CHECK:     #pragma omp declare variant(fn_linkage_variant) match(implementation={vendor(score(0):xxx)},device={kind(cpu, host)})
+// CHECK:     #pragma omp declare variant(fn_linkage_variant) match(implementation={vendor(score(0):xxx)},device={kind(cpu, host),isa("x86", "aarch64")})
 // CHECK:     int fn_linkage();
 // CHECK: }
 int fn_linkage_variant();
 extern "C" {
-#pragma omp declare variant(fn_linkage_variant) match(implementation = {vendor(xxx)},device={kind(cpu,host)})
+#pragma omp declare variant(fn_linkage_variant) match(implementation = {vendor(xxx)},device={kind(cpu,host),isa("x86", "aarch64")})
 int fn_linkage();
 }
 
 // CHECK: extern "C" int fn_linkage_variant1()
-// CHECK: #pragma omp declare variant(fn_linkage_variant1) match(implementation={vendor(score(0):xxx)},device={kind(cpu, host)})
+// CHECK: #pragma omp declare variant(fn_linkage_variant1) match(implementation={vendor(score(0):xxx)},device={kind(cpu, host),isa("x86", "aarch64")})
 // CHECK: int fn_linkage1();
 extern "C" int fn_linkage_variant1();
-#pragma omp declare variant(fn_linkage_variant1) match(implementation = {vendor(xxx)},device={kind(cpu,host)})
+#pragma omp declare variant(fn_linkage_variant1) match(implementation = {vendor(xxx)},device={kind(cpu,host),isa("x86", "aarch64")})
 int fn_linkage1();
 
diff --git a/clang/test/OpenMP/declare_variant_device_isa_codegen.cpp b/clang/test/OpenMP/declare_variant_device_isa_codegen.cpp
new file mode 100644
--- /dev/null
+++ b/clang/test/OpenMP/declare_variant_device_isa_codegen.cpp
@@ -0,0 +1,152 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -DCORRECT=\"x86_64\" | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s -DCORRECT=\"x86_64\"
+// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 -DCORRECT=\"x86_64\" | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple aarch64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -DCORRECT=\"aarch64\" | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple aarch64-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s -DCORRECT=\"aarch64\"
+// RUN: %clang_cc1 -fopenmp -x c++ -triple aarch64-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 -DCORRECT=\"aarch64\" | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -DCORRECT=\"powerpc64le\" | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple ppc64le-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s -DCORRECT=\"ppc64le\"
+// RUN: %clang_cc1 -fopenmp -x c++ -triple ppc64le-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 -DCORRECT=\"ppc64le\" | FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=x86_64-unknown-linux -emit-llvm-bc %s -o %t-host.bc -DCORRECT=\"x86_64\"
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - -DCORRECT=\"x86_64\" | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -emit-pch -o %t -DCORRECT=\"x86_64\"
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -include-pch %t -o - -DCORRECT=\"x86_64\" | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -fopenmp-targets=ppc64le-unknown-linux -emit-llvm-bc %s -o %t-host.bc -DCORRECT=\"ppc64le\"
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - -DCORRECT=\"ppc64le\" | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -emit-pch -o %t -DCORRECT=\"ppc64le\"
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -include-pch %t -o - -DCORRECT=\"ppc64le\" | FileCheck %s
+
+// expected-no-diagnostics
+
+// CHECK-NOT: ret i32 {{1|4|81|84}}
+// CHECK-DAG: @_Z3barv = {{.*}}alias i32 (), i32 ()* @_Z3foov
+// CHECK-DAG: @_ZN16SpecSpecialFuncs6MethodEv = {{.*}}alias i32 (%struct.SpecSpecialFuncs*), i32 (%struct.SpecSpecialFuncs*)* @_ZN16SpecSpecialFuncs7method_Ev
+// CHECK-DAG: @_ZN16SpecSpecialFuncs6methodEv = linkonce_odr {{.*}}alias i32 (%struct.SpecSpecialFuncs*), i32 (%struct.SpecSpecialFuncs*)* @_ZN16SpecSpecialFuncs7method_Ev
+// CHECK-DAG: @_ZN12SpecialFuncs6methodEv = linkonce_odr {{.*}}alias i32 (%struct.SpecialFuncs*), i32 (%struct.SpecialFuncs*)* @_ZN12SpecialFuncs7method_Ev
+// CHECK-DAG: @_Z5prio_v = {{.*}}alias i32 (), i32 ()* @_Z5prio1v
+// CHECK-DAG: @_ZL6prio1_v = internal alias i32 (), i32 ()* @_ZL5prio2v
+// CHECK-DAG: @_Z4callv = {{.*}}alias i32 (), i32 ()* @_Z4testv
+// CHECK-DAG: @_ZL9stat_usedv = internal alias i32 (), i32 ()* @_ZL10stat_used_v
+// CHECK-DAG: @_ZN12SpecialFuncs6MethodEv = {{.*}}alias i32 (%struct.SpecialFuncs*), i32 (%struct.SpecialFuncs*)* @_ZN12SpecialFuncs7method_Ev
+// CHECK-DAG: @fn_linkage = {{.*}}alias i32 (), i32 ()* @_Z18fn_linkage_variantv
+// CHECK-DAG: @_Z11fn_linkage1v = {{.*}}alias i32 (), i32 ()* @fn_linkage_variant1
+// CHECK-DAG: declare {{.*}}i32 @_Z5bazzzv()
+// CHECK-DAG: declare {{.*}}i32 @_Z3bazv()
+// CHECK-DAG: ret i32 2
+// CHECK-DAG: ret i32 3
+// CHECK-DAG: ret i32 5
+// CHECK-DAG: ret i32 6
+// CHECK-DAG: ret i32 7
+// CHECK-DAG: ret i32 82
+// CHECK-DAG: ret i32 83
+// CHECK-DAG: ret i32 85
+// CHECK-DAG: ret i32 86
+// CHECK-DAG: ret i32 87
+// CHECK-NOT: ret i32 {{1|4|81|84}}
+
+#ifndef HEADER
+#define HEADER
+
+#pragma omp declare target
+#define WRONG "mips"
+
+int foo() { return 2; }
+
+#pragma omp declare variant(foo) match(device = {isa(CORRECT)})
+int bar() { return 1; }
+
+int bazzz();
+#pragma omp declare variant(bazzz) match(device = {isa(CORRECT)})
+int baz() { return 1; }
+
+int test();
+#pragma omp declare variant(test) match(device = {isa(CORRECT)})
+int call() { return 1; }
+
+static int stat_unused_();
+#pragma omp declare variant(stat_unused_) match(device = {isa(CORRECT)})
+static int stat_unused() { return 1; }
+
+static int stat_used_();
+#pragma omp declare variant(stat_used_) match(device = {isa(CORRECT)})
+static int stat_used() { return 1; }
+
+int main() { return bar() + baz() + call() + stat_used(); }
+
+int test() { return 3; }
+static int stat_unused_() { return 4; }
+static int stat_used_() { return 5; }
+
+struct SpecialFuncs {
+  void vd() {}
+  SpecialFuncs();
+  ~SpecialFuncs();
+
+  int method_() { return 6; }
+#pragma omp declare variant(SpecialFuncs::method_) \
+    match(device = {isa(CORRECT)})
+  int method() { return 1; }
+#pragma omp declare variant(SpecialFuncs::method_) \
+    match(device = {isa(CORRECT)})
+  int Method();
+} s;
+
+int SpecialFuncs::Method() { return 1; }
+
+struct SpecSpecialFuncs {
+  void vd() {}
+  SpecSpecialFuncs();
+  ~SpecSpecialFuncs();
+
+  int method_();
+#pragma omp declare variant(SpecSpecialFuncs::method_) \
+    match(device = {isa(CORRECT)})
+  int method() { return 1; }
+#pragma omp declare variant(SpecSpecialFuncs::method_) \
+    match(device = {isa(CORRECT)})
+  int Method();
+} s1;
+
+int SpecSpecialFuncs::method_() { return 7; }
+int SpecSpecialFuncs::Method() { return 1; }
+
+void xxx() {
+  (void)s.method();
+  (void)s1.method();
+}
+
+int prio() { return 81; }
+int prio1() { return 82; }
+
+#pragma omp declare variant(prio) match(device = {isa(CORRECT), kind(cpu)})
+#pragma omp declare variant(prio1) match(device = {isa(CORRECT)})
+int prio_() { return 1; }
+
+static int prio2() { return 83; }
+static int prio3() { return 84; }
+static int prio4() { return 84; }
+
+#pragma omp declare variant(prio4) match(device = {isa(CORRECT),kind(cpu)})
+#pragma omp declare variant(prio2) match(device = {isa(CORRECT)})
+#pragma omp declare variant(prio3) match(device = {isa(CORRECT), kind(nohost)})
+static int prio1_() { return 1; }
+
+int int_fn() { return prio1_(); }
+
+int fn_linkage_variant() { return 85; }
+extern "C" {
+#pragma omp declare variant(fn_linkage_variant) match(device = {isa(CORRECT)})
+int fn_linkage() { return 1; }
+}
+
+extern "C" int fn_linkage_variant1() { return 86; }
+#pragma omp declare variant(fn_linkage_variant1) match(device = {isa(CORRECT)})
+int fn_linkage1() { return 1; }
+
+int fn_variant2() { return 1; }
+#pragma omp declare variant(fn_variant2) match(device = {isa(WRONG)})
+int fn2() { return 87; }
+
+#pragma omp end declare target
+#endif // HEADER
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
@@ -1,6 +1,6 @@
-// RUN: %clang_cc1 -triple=x86_64-pc-win32 -verify -fopenmp -x c -std=c99 -fms-extensions -Wno-pragma-pack %s
+// RUN: %clang_cc1 -triple=x86_64-pc-win32 -verify -fopenmp -x c -std=c99 -fms-extensions -Wno-pragma-pack %s -Wno-implicit-function-declaration
 
-// RUN: %clang_cc1 -triple=x86_64-pc-win32 -verify -fopenmp-simd -x c -std=c99 -fms-extensions -Wno-pragma-pack %s
+// RUN: %clang_cc1 -triple=x86_64-pc-win32 -verify -fopenmp-simd -x c -std=c99 -fms-extensions -Wno-pragma-pack %s -Wno-implicit-function-declaration
 
 // expected-error@+1 {{expected an OpenMP directive}}
 #pragma omp declare
@@ -45,7 +45,17 @@
 #pragma omp declare variant(foo) match(device={kind(score(2 gpu)}) // expected-error 2 {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
 #pragma omp declare variant(foo) match(device={kind(score(foo()) ibm)}) // expected-error {{expected ')' or ',' after 'score'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
 #pragma omp declare variant(foo) match(device={kind(score(5): host), kind(llvm)}) // expected-error {{context trait selector 'kind' is used already in the same 'device' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'kind' used here}} expected-error {{expected ')' or ',' after 'score'}} expected-note {{to match this '('}} expected-error {{expected ')'}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}} expected-error {{unknown 'llvm' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
-#pragma omp declare variant(foo) match(device={kind(score(5): nohost), vendor(llvm)}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}} expected-error {{expected ')' or ',' after 'score'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}}
+#pragma omp declare variant(foo) match(device={kind(score(5): nohost), vendor(llvm)}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}} expected-error {{expected ')' or ',' after 'score'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
+#pragma omp declare variant(foo) match(device={xxx}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}}
+#pragma omp declare variant(foo) match(device={isa}) // expected-error {{expected '(' after 'isa'}} expected-error {{expected ')'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+#pragma omp declare variant(foo) match(device={isa(}) // expected-error 2 {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+#pragma omp declare variant(foo) match(device={isa()}) // expected-error {{expected expression}}}
+#pragma omp declare variant(foo) match(device={isa(score "i586")}) // expected-error {{use of undeclared identifier 'score'}} expected-error {{expected ')' or ',' after 'instruction set architecture'}}
+#pragma omp declare variant(foo) match(device={isa(score( "x86_64")}) // expected-error 3 {{expected ')'}} expected-note {{to match this '('}} expected-error {{trait expression must have a string type, not 'int'}} expected-error {{trait expression must be a constant string expression}} expected-error {{expected expression}}
+#pragma omp declare variant(foo) match(device={isa(score(2 "arm")}) // expected-error 4 {{expected ')'}} expected-note 2 {{to match this '('}} expected-error {{expected expression}}
+#pragma omp declare variant(foo) match(device={isa(score(foo()) "powerpc64le")}) // expected-error {{expected ')'}} expected-error {{trait expression must have a string type, not 'int'}} expected-error {{trait expression must be a constant string expression}}
+#pragma omp declare variant(foo) match(device={isa(score(5): "x86"), isa(llvm)}) // expected-error {{context trait selector 'isa' is used already in the same 'device' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'isa' used here}} expected-note {{to match this '('}} expected-error 3 {{expected ')'}} expected-error {{use of undeclared identifier 'llvm'}} expected-error {{trait expression must have a string type, not 'int'}} expected-error {{trait expression must be a constant string expression}} expected-error {{expected expression}}
+#pragma omp declare variant(foo) match(device={isa(score(5): "x86"), vendor(llvm)}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}} expected-error 3 {{expected ')'}} expected-note {{to match this '('}} expected-error {{trait expression must have a string type, not 'int'}} expected-error {{trait expression must be a constant string expression}} expected-error {{expected expression}}
 int bar(void);
 
 // expected-error@+2 {{'#pragma omp declare variant' can only be applied to functions}}
diff --git a/clang/test/OpenMP/declare_variant_messages.cpp b/clang/test/OpenMP/declare_variant_messages.cpp
--- a/clang/test/OpenMP/declare_variant_messages.cpp
+++ b/clang/test/OpenMP/declare_variant_messages.cpp
@@ -89,6 +89,16 @@
 #pragma omp declare variant(foofoo <int>) match(device={kind(score(foofoo <int>()) ibm)}) // expected-error {{expected ')' or ',' after 'score'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
 #pragma omp declare variant(foofoo <int>) match(device={kind(score(C+5): host), kind(llvm)}) // expected-error {{context trait selector 'kind' is used already in the same 'device' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'kind' used here}} expected-error {{expected ')' or ',' after 'score'}} expected-note {{to match this '('}} expected-error {{expected ')'}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}} expected-error {{unknown 'llvm' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
 #pragma omp declare variant(foofoo <int>) match(device={kind(score(C+5): nohost), vendor(llvm)}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}} expected-error {{expected ')' or ',' after 'score'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{unknown 'score' device kind trait in the 'device' context selector set, expected one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'}}
+#pragma omp declare variant(foofoo <int>) match(device={xxx}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}}
+#pragma omp declare variant(foofoo <int>) match(device={isa}) // expected-error {{expected '(' after 'isa'}}  expected-error {{expected ')'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+#pragma omp declare variant(foofoo <int>) match(device={isa(}) // expected-error 2 {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+#pragma omp declare variant(foofoo <int>) match(device={isa()}) // expected-error {{expected expression}}
+#pragma omp declare variant(foofoo <int>) match(device={isa(score "x86_64")}) // expected-error {{use of undeclared identifier 'score'}} expected-error {{expected ')' or ',' after 'instruction set architecture'}}
+#pragma omp declare variant(foofoo <int>) match(device={isa(score( "powerpc64le")}) // expected-error 3 {{expected ')'}} expected-note {{to match this '('}} expected-error {{use of undeclared identifier 'score'}} expected-error {{expected expression}}
+#pragma omp declare variant(foofoo <int>) match(device={isa(score(C "nvptx")}) // expected-error 4 {{expected ')'}} expected-note 2 {{to match this '('}} expected-error {{expected expression}}
+#pragma omp declare variant(foofoo <int>) match(device={isa(score(foofoo <int>()) "arm")}) // expected-error {{expected ')'}} expected-error {{use of undeclared identifier 'score'}}
+#pragma omp declare variant(foofoo <int>) match(device={isa(score(C+5): "i586"), isa(llvm)}) // expected-error {{context trait selector 'isa' is used already in the same 'device' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'isa' used here}} expected-error 3 {{expected ')'}} expected-error {{use of undeclared identifier 'score'}} expected-error {{use of undeclared identifier 'llvm'}} expected-error {{expected expression}} expected-note {{to match this '('}}
+#pragma omp declare variant(foofoo <int>) match(device={isa(score(C+5): "aarch64"), vendor(llvm)}) // expected-warning {{unknown context selector in 'device' context selector set of 'omp declare variant' directive, ignored}} expected-error 2 {{expected ')'}} expected-error {{use of undeclared identifier 'score'}} expected-error {{expected ')'}} expected-error {{expected expression}} expected-note {{to match this '('}}
 template <typename T, int C>
 T barbar();
 
diff --git a/clang/test/OpenMP/declare_variant_mixed_codegen.cpp b/clang/test/OpenMP/declare_variant_mixed_codegen.cpp
--- a/clang/test/OpenMP/declare_variant_mixed_codegen.cpp
+++ b/clang/test/OpenMP/declare_variant_mixed_codegen.cpp
@@ -1,6 +1,12 @@
-// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope | FileCheck %s
-// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s
-// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -DARCH=\"x86_64\"| FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s -DARCH=\"x86_64\"
+// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 -DARCH=\"x86_64\" | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple aarch64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -DARCH=\"aarch64\" | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple aarch64-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s -DARCH=\"aarch64\"
+// RUN: %clang_cc1 -fopenmp -x c++ -triple aarch64-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 -DARCH=\"aarch64\" | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope -DARCH=\"ppc64le\" | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple ppc64le-unknown-linux -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s -DARCH=\"ppc64le\"
+// RUN: %clang_cc1 -fopenmp -x c++ -triple ppc64le-unknown-linux -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 -DARCH=\"ppc64le\" | FileCheck %s
 // expected-no-diagnostics
 
 // CHECK-NOT: ret i32 {{1|4|81|84}}
@@ -34,23 +40,23 @@
 
 int foo() { return 2; }
 
-#pragma omp declare variant(foo) match(implementation = {vendor(llvm)}, device={kind(cpu)})
+#pragma omp declare variant(foo) match(implementation = {vendor(llvm)}, device={kind(cpu), isa(ARCH)})
 int bar() { return 1; }
 
 int bazzz();
-#pragma omp declare variant(bazzz) match(implementation = {vendor(llvm)}, device={kind(host)})
+#pragma omp declare variant(bazzz) match(implementation = {vendor(llvm)}, device={kind(host), isa(ARCH)})
 int baz() { return 1; }
 
 int test();
-#pragma omp declare variant(test) match(implementation = {vendor(llvm)}, device={kind(cpu)})
+#pragma omp declare variant(test) match(implementation = {vendor(llvm)}, device={kind(cpu), isa(ARCH)})
 int call() { return 1; }
 
 static int stat_unused_();
-#pragma omp declare variant(stat_unused_) match(implementation = {vendor(llvm)}, device={kind(cpu)})
+#pragma omp declare variant(stat_unused_) match(implementation = {vendor(llvm)}, device={kind(cpu), isa(ARCH)})
 static int stat_unused() { return 1; }
 
 static int stat_used_();
-#pragma omp declare variant(stat_used_) match(implementation = {vendor(llvm)}, device={kind(host)})
+#pragma omp declare variant(stat_used_) match(implementation = {vendor(llvm)}, device={kind(host), isa(ARCH)})
 static int stat_used() { return 1; }
 
 int main() { return bar() + baz() + call() + stat_used(); }
@@ -66,10 +72,10 @@
 
   int method_() { return 6; }
 #pragma omp declare variant(SpecialFuncs::method_)                             \
-    match(implementation = {vendor(llvm)}, device={kind(cpu)})
+    match(implementation = {vendor(llvm)}, device={kind(cpu), isa(ARCH)})
   int method() { return 1; }
 #pragma omp declare variant(SpecialFuncs::method_)                             \
-    match(implementation = {vendor(llvm)}, device={kind(host)})
+    match(implementation = {vendor(llvm)}, device={kind(host), isa(ARCH)})
   int Method();
 } s;
 
@@ -82,10 +88,10 @@
 
   int method_();
 #pragma omp declare variant(SpecSpecialFuncs::method_)                         \
-    match(implementation = {vendor(llvm)}, device={kind(cpu)})
+    match(implementation = {vendor(llvm)}, device={kind(cpu), isa(ARCH)})
   int method() { return 1; }
 #pragma omp declare variant(SpecSpecialFuncs::method_)                         \
-    match(implementation = {vendor(llvm)}, device={kind(host)})
+    match(implementation = {vendor(llvm)}, device={kind(host), isa(ARCH)})
   int Method();
 } s1;
 
@@ -100,29 +106,29 @@
 int prio() { return 81; }
 int prio1() { return 82; }
 
-#pragma omp declare variant(prio) match(implementation = {vendor(score(2): llvm)}, device={kind(cpu,host)})
-#pragma omp declare variant(prio1) match(implementation = {vendor(score(1): llvm)}, device={kind(cpu)})
+#pragma omp declare variant(prio) match(implementation = {vendor(score(2): llvm)}, device={kind(cpu,host), isa(ARCH)})
+#pragma omp declare variant(prio1) match(implementation = {vendor(score(1): llvm)}, device={kind(cpu), isa(ARCH)})
 int prio_() { return 1; }
 
 static int prio2() { return 83; }
 static int prio3() { return 84; }
 static int prio4() { return 84; }
 
-#pragma omp declare variant(prio4) match(implementation = {vendor(score(8): llvm)},device={kind(cpu,host)})
-#pragma omp declare variant(prio2) match(implementation = {vendor(score(5): llvm)})
-#pragma omp declare variant(prio3) match(implementation = {vendor(score(7): llvm)}, device={kind(cpu)})
+#pragma omp declare variant(prio4) match(implementation = {vendor(score(8): llvm)},device={kind(cpu,host), isa(ARCH)})
+#pragma omp declare variant(prio2) match(implementation = {vendor(score(5): llvm)}, device={isa(ARCH)})
+#pragma omp declare variant(prio3) match(implementation = {vendor(score(7): llvm)}, device={kind(cpu), isa("i586")})
 static int prio1_() { return 1; }
 
 int int_fn() { return prio1_(); }
 
 int fn_linkage_variant() { return 85; }
 extern "C" {
-#pragma omp declare variant(fn_linkage_variant) match(implementation = {vendor(llvm)}, device={kind(cpu)})
+#pragma omp declare variant(fn_linkage_variant) match(implementation = {vendor(llvm)}, device={kind(cpu), isa(ARCH)})
 int fn_linkage() { return 1; }
 }
 
 extern "C" int fn_linkage_variant1() { return 86; }
-#pragma omp declare variant(fn_linkage_variant1) match(implementation = {vendor(llvm)}, device={kind(host)})
+#pragma omp declare variant(fn_linkage_variant1) match(implementation = {vendor(llvm)}, device={kind(host), isa(ARCH)})
 int fn_linkage1() { return 1; }
 
 int fn_variant2() { return 1; }
@@ -132,6 +138,7 @@
 #pragma omp declare variant(fn_variant2) match(implementation = {vendor(llvm)}, device={kind(cpu,nohost)})
 #pragma omp declare variant(fn_variant2) match(implementation = {vendor(llvm)}, device={kind(gpu)})
 #pragma omp declare variant(fn_variant2) match(implementation = {vendor(llvm)}, device={kind(fpga)})
+#pragma omp declare variant(fn_variant2) match(implementation = {vendor(llvm)}, device={isa("mips")})
 int fn2() { return 87; }
 
 #endif // HEADER
diff --git a/clang/test/OpenMP/nvptx_declare_variant_device_isa_codegen.cpp b/clang/test/OpenMP/nvptx_declare_variant_device_isa_codegen.cpp
new file mode 100644
--- /dev/null
+++ b/clang/test/OpenMP/nvptx_declare_variant_device_isa_codegen.cpp
@@ -0,0 +1,154 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -fopenmp-version=50
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -fopenmp-version=50 | FileCheck %s --implicit-check-not='ret i32 {{1|81|84}}'
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t -fopenmp-version=50
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - -fopenmp-version=50 | FileCheck %s --implicit-check-not='ret i32 {{1|81|84}}'
+// expected-no-diagnostics
+
+// CHECK-NOT: ret i32 {{1|81|84}}
+// CHECK-DAG: define {{.*}}i32 @_Z3barv()
+// CHECK-DAG: define {{.*}}i32 @_ZN16SpecSpecialFuncs6MethodEv(%struct.SpecSpecialFuncs* %{{.+}})
+// CHECK-DAG: define {{.*}}i32 @_ZN12SpecialFuncs6MethodEv(%struct.SpecialFuncs* %{{.+}})
+// CHECK-DAG: define linkonce_odr {{.*}}i32 @_ZN16SpecSpecialFuncs6methodEv(%struct.SpecSpecialFuncs* %{{.+}})
+// CHECK-DAG: define linkonce_odr {{.*}}i32 @_ZN12SpecialFuncs6methodEv(%struct.SpecialFuncs* %{{.+}})
+// CHECK-DAG: define {{.*}}i32 @_Z5prio_v()
+// CHECK-DAG: define internal i32 @_ZL6prio1_v()
+// CHECK-DAG: define {{.*}}i32 @_Z4callv()
+// CHECK-DAG: define internal i32 @_ZL9stat_usedv()
+// CHECK-DAG: define {{.*}}i32 @fn_linkage()
+// CHECK-DAG: define {{.*}}i32 @_Z11fn_linkage1v()
+
+// CHECK-DAG: ret i32 2
+// CHECK-DAG: ret i32 3
+// CHECK-DAG: ret i32 4
+// CHECK-DAG: ret i32 5
+// CHECK-DAG: ret i32 6
+// CHECK-DAG: ret i32 7
+// CHECK-DAG: ret i32 82
+// CHECK-DAG: ret i32 83
+// CHECK-DAG: ret i32 85
+// CHECK-DAG: ret i32 86
+// CHECK-DAG: ret i32 87
+
+// Outputs for function members
+// CHECK-DAG: ret i32 6
+// CHECK-DAG: ret i32 7
+// CHECK-NOT: ret i32 {{1|81|84}}
+
+#ifndef HEADER
+#define HEADER
+
+int foo() { return 2; }
+int bazzz();
+int test();
+static int stat_unused_();
+static int stat_used_();
+
+#pragma omp declare target
+
+#pragma omp declare variant(foo) match(device = {isa("nvptx64")})
+int bar() { return 1; }
+
+#pragma omp declare variant(bazzz) match(device = {isa("nvptx64")})
+int baz() { return 1; }
+
+#pragma omp declare variant(test) match(device = {isa("nvptx64")})
+int call() { return 1; }
+
+#pragma omp declare variant(stat_unused_) match(device = {isa("nvptx64")})
+static int stat_unused() { return 1; }
+
+#pragma omp declare variant(stat_used_) match(device = {isa("nvptx64")})
+static int stat_used() { return 1; }
+
+#pragma omp end declare target
+
+int main() {
+  int res;
+#pragma omp target map(from \
+                       : res)
+  res = bar() + baz() + call();
+  return res;
+}
+
+int test() { return 3; }
+static int stat_unused_() { return 4; }
+static int stat_used_() { return 5; }
+
+#pragma omp declare target
+
+struct SpecialFuncs {
+  void vd() {}
+  SpecialFuncs();
+  ~SpecialFuncs();
+
+  int method_() { return 6; }
+#pragma omp declare variant(SpecialFuncs::method_) \
+    match(device = {isa("nvptx64")})
+  int method() { return 1; }
+#pragma omp declare variant(SpecialFuncs::method_) \
+    match(device = {isa("nvptx64")})
+  int Method();
+} s;
+
+int SpecialFuncs::Method() { return 1; }
+
+struct SpecSpecialFuncs {
+  void vd() {}
+  SpecSpecialFuncs();
+  ~SpecSpecialFuncs();
+
+  int method_();
+#pragma omp declare variant(SpecSpecialFuncs::method_) \
+    match(device = {isa("nvptx64")})
+  int method() { return 1; }
+#pragma omp declare variant(SpecSpecialFuncs::method_) \
+    match(device = {isa("nvptx64")})
+  int Method();
+} s1;
+
+#pragma omp end declare target
+
+int SpecSpecialFuncs::method_() { return 7; }
+int SpecSpecialFuncs::Method() { return 1; }
+
+int prio() { return 81; }
+int prio1() { return 82; }
+static int prio2() { return 83; }
+static int prio3() { return 84; }
+static int prio4() { return 84; }
+int fn_linkage_variant() { return 85; }
+extern "C" int fn_linkage_variant1() { return 86; }
+int fn_variant2() { return 1; }
+
+#pragma omp declare target
+
+void xxx() {
+  (void)s.method();
+  (void)s1.method();
+}
+
+#pragma omp declare variant(prio) match(device = {isa("nvptx64"), kind(gpu)})
+#pragma omp declare variant(prio1) match(device = {isa("nvptx64", "nvptx64")})
+int prio_() { return 1; }
+
+#pragma omp declare variant(prio4) match(device = {isa("nvptx64"), kind(gpu)})
+#pragma omp declare variant(prio2) match(device = {isa("nvptx64")})
+#pragma omp declare variant(prio3) match(device = {isa("nvptx64"), kind(gpu)})
+static int prio1_() { return 1; }
+
+int int_fn() { return prio1_(); }
+
+extern "C" {
+#pragma omp declare variant(fn_linkage_variant) match(device = {isa("nvptx64")})
+int fn_linkage() { return 1; }
+}
+
+#pragma omp declare variant(fn_linkage_variant1) match(device = {isa("nvptx64")})
+int fn_linkage1() { return 1; }
+
+#pragma omp declare variant(fn_variant2) match(device = {isa("amdgcn")})
+int fn2() { return 87; }
+
+#pragma omp end declare target
+
+#endif // HEADER