[OPENMP50]Add device/isa context selector support.

[ABataev]

Summary

Added basic support for device/isa context selector.

[jdoerfert]

Thanks for working on this! One nit and one concern below.

clang/include/clang/Sema/Sema.h

9322

I would like to avoid the Any type here and I hope we can if we don't allow "strings". See also my last comment.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

11147

should we normalize the names here? at least .lower might be useful.

clang/test/OpenMP/nvptx_declare_variant_device_isa_codegen.cpp

49

I'm not sure we want these to be strings. I would have assumed isa(nvptx64), similar to vendor(arm), or kind(host) and kind(gpu)

[ABataev]

clang/test/OpenMP/nvptx_declare_variant_device_isa_codegen.cpp

49

Here is the message from Joachim Protze:

vendor and kind are special, because we define the kind-name and 
vendor-name in the context definition side document. (We could change 
that to "kind-name" or define kind-name as any of "cpu", "gpu", "fpga" ???)

For the others, according to the syntax diagram in the OpenMP Context 
section, the content of the () can only be string or const-int-exp.

- Joachim

That's why we need to represent them as strings, this is per OpenMP standard.

[ABataev]

Address comments.

[ABataev]

clang/include/clang/Sema/Sema.h

9322

The only possible solution I see here is to convert kind/vendor strings into exprs and store all data as Expr *

clang/lib/CodeGen/CGOpenMPRuntime.cpp

11147

Done + added check for possibly used vendor/os/env

[merge_guards_bot]

Build result: pass - 60340 tests passed, 0 failed and 732 were skipped.

Log files: console-log.txt, CMakeCache.txt

[jdoerfert]

clang/include/clang/Sema/Sema.h

9322

Should we define a struct to be used as variant here maybe? If we only have <5 different types we might not want to go to Any or Expr directly.

clang/test/OpenMP/nvptx_declare_variant_device_isa_codegen.cpp

49

While you say it, I wanted to mention that we moved the discussion on the openmp-lang list.

[ABataev]

clang/include/clang/Sema/Sema.h

9322

What is the difference between using Any and a struct?

[jdoerfert]

clang/include/clang/Sema/Sema.h

9322

Any is generic (anything).
Having a variant (std::variant or self build struct with 2 constructors) specifies what the components can be.

[ABataev]

clang/include/clang/Sema/Sema.h

9322

It is not quite so. Anyway, I'm going to switch to Expr* since it is the only possible way to represent traits as identifiers/string literals.

clang/include/clang/Basic/Attr.td

def OMPDeclareVariant : InheritableAttr {
   let InheritEvenIfAlreadyPresent = 1;
   let Documentation = [OMPDeclareVariantDocs];
   let Args = [
     ExprArgument<"VariantFuncRef">,
     VariadicExprArgument<"Scores">,
     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 {
       if (const Expr *E = *std::next(scores_begin(), I)) {
         OS << "score(";
         E->printPretty(OS, nullptr, Policy);
         OS << "):";
       }
void printPrettyPragma(raw_ostream & OS, const PrintingPolicy &Policy)
             OS << "kind(";
             if (deviceKinds_size() > 0) {
               OS << *deviceKinds().begin();
               for (StringRef KindName : llvm::drop_begin(deviceKinds(), 1))
                 OS << ", " << KindName;
             }
             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.");
           }
         }
         OS << "}";
       }
       OS << ")";
     }

clang/include/clang/Basic/DiagnosticSemaKinds.td

def err_omp_declare_variant_incompat_types : Error<
   "variant in '#pragma omp declare variant' with type %0 is incompatible with type %1"
   >;
 def warn_omp_declare_variant_marked_as_declare_variant : Warning<
   "variant function in '#pragma omp declare variant' is itself marked as '#pragma omp declare variant'"
   >, InGroup<SourceUsesOpenMP>;
 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 {
 // Objective-C related result type compatibility
 def warn_related_result_type_compatibility_class : Warning<
   "method is expected to return an instance of its class type "
   "%diff{$, but is declared to return $|"
   ", but is declared to return different type}0,1">;

clang/include/clang/Basic/OpenMPKinds.def

 
 // OpenMP context selector sets.
 OPENMP_CONTEXT_SELECTOR_SET(implementation)
 OPENMP_CONTEXT_SELECTOR_SET(device)
 
 // OpenMP context selectors.
 OPENMP_CONTEXT_SELECTOR(vendor)
 OPENMP_CONTEXT_SELECTOR(kind)
+OPENMP_CONTEXT_SELECTOR(isa)
 
 // OpenMP directives.
 OPENMP_DIRECTIVE(threadprivate)
 OPENMP_DIRECTIVE(parallel)
 OPENMP_DIRECTIVE(task)
 OPENMP_DIRECTIVE(simd)
 OPENMP_DIRECTIVE(for)
 OPENMP_DIRECTIVE(sections)

clang/include/clang/Sema/Sema.h

 #include "llvm/ADT/TinyPtrVector.h"
 #include <deque>
 #include <memory>
 #include <string>
 #include <tuple>
 #include <vector>
 
 namespace llvm {
+  class Any;
   class APSInt;
   template <typename ValueT> struct DenseMapInfo;
   template <typename ValueT, typename ValueInfoT> class DenseSet;
   class SmallBitVector;
   struct InlineAsmIdentifierInfo;
 }
 
 namespace clang {
private:
   void markOpenMPDeclareVariantFuncsReferenced(SourceLocation Loc,
                                                FunctionDecl *Func,
                                                bool MightBeOdrUse);
 
 public:
   /// 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(
       const FunctionDecl *OldFD, const FunctionDecl *NewFD,
       const PartialDiagnostic &NoProtoDiagID,
       const PartialDiagnosticAt &NoteCausedDiagIDAt,
       const PartialDiagnosticAt &NoSupportDiagIDAt,
       const PartialDiagnosticAt &DiffDiagIDAt, bool TemplatesSupported,

clang/lib/CodeGen/CGOpenMPRuntime.cpp

 #include "CGRecordLayout.h"
 #include "CodeGenFunction.h"
 #include "clang/CodeGen/ConstantInitBuilder.h"
 #include "clang/AST/Decl.h"
 #include "clang/AST/StmtOpenMP.h"
 #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>
 
 using namespace clang;
 using namespace CodeGen;
 
 namespace {
Addr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
       Addr,
       CGF.ConvertTypeForMem(CGM.getContext().getPointerType(CVD->getType())),
       CVD->getName() + ".addr");
   return Address(Addr, Align);
 }
 
 namespace {
 using OMPContextSelectorData =
-    OpenMPCtxSelectorData<ArrayRef<StringRef>, llvm::APSInt>;
+    OpenMPCtxSelectorData<SmallVector<SmallString<8>, 4>, llvm::APSInt>;
 using CompleteOMPContextSelectorData = SmallVector<OMPContextSelectorData, 4>;
 } // anonymous namespace
 
 /// Checks current context and returns true if it matches the context selector.
 template <OpenMPContextSelectorSetKind CtxSet, OpenMPContextSelectorKind Ctx,
           typename... Arguments>
 static bool checkContext(const OMPContextSelectorData &Data,
                          Arguments... Params) {
for (StringRef Name : Data.Names) {
     case llvm::Triple::renderscript32:
     case llvm::Triple::renderscript64:
       return false;
     }
   }
   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) {
+    /// Consider arch with vendor/os/env as invalid.
+    if (Name.contains('-'))
+      return false;
+    llvm::Triple CustomTriple(Name.lower());
+    if (CustomTriple.getArch() != CGM.getTarget().getTriple().getArch() ||
+        (CustomTriple.getSubArch() != llvm::Triple::NoSubArch &&
+         CustomTriple.getSubArch() != CGM.getTarget().getTriple().getSubArch()))
+      return false;
+  }
+  return true;
+}
+
 bool matchesContext(CodeGenModule &CGM,
                     const CompleteOMPContextSelectorData &ContextData) {
   for (const OMPContextSelectorData &Data : ContextData) {
     switch (Data.Ctx) {
     case OMP_CTX_vendor:
       assert(Data.CtxSet == OMP_CTX_SET_implementation &&
              "Expected implementation context selector set.");
       if (!checkContext<OMP_CTX_SET_implementation, OMP_CTX_vendor>(Data))
         return false;
       break;
     case OMP_CTX_kind:
       assert(Data.CtxSet == OMP_CTX_SET_device &&
              "Expected device context selector set.");
       if (!checkContext<OMP_CTX_SET_device, OMP_CTX_kind, CodeGenModule &>(Data,
                                                                            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.");
     }
   }
   return true;
 }
 
 static CompleteOMPContextSelectorData
for (unsigned I = 0, E = A->scores_size(); I < E; ++I) {
     Data.back().CtxSet = CtxSet;
     Data.back().Ctx = Ctx;
     const Expr *Score = *std::next(A->scores_begin(), I);
     Data.back().Score = Score->EvaluateKnownConstInt(C);
     switch (Ctx) {
     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.");
     }
   }
   return Data;
 }
 
 static bool isStrictSubset(const CompleteOMPContextSelectorData &LHS,

clang/lib/Parse/ParseOpenMP.cpp

 //===----------------------------------------------------------------------===//
 
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/StmtOpenMP.h"
 #include "clang/Parse/ParseDiagnostic.h"
 #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"
 
 using namespace clang;
 
 //===----------------------------------------------------------------------===//
 // OpenMP declarative directives.
 //===----------------------------------------------------------------------===//
case OMP_CTX_vendor: {
     } while (Tok.is(tok::identifier));
     // Parse ')'.
     (void)T.consumeClose();
     if (!Vendors.empty())
       Data.emplace_back(OMP_CTX_SET_implementation, CSKind, Score, Vendors);
     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";
     // Skip until either '}', ')', or end of directive.
     while (!P.SkipUntil(tok::r_brace, tok::r_paren,
                         tok::annot_pragma_openmp_end, Parser::StopBeforeMatch))
       ;
     return;
   }
 }
 
 /// Parse context selector for 'device' selector set:
 /// 'kind' '(' <kind> { ',' <kind> } ')'
+/// 'isa' '(' <ISA> { ',' <ISA> } ')'
 static void
 parseDeviceSelector(Parser &P, SourceLocation Loc,
                     llvm::StringMap<SourceLocation> &UsedCtx,
                     SmallVectorImpl<Sema::OMPCtxSelectorData> &Data) {
   const Token &Tok = P.getCurToken();
   // Parse inner context selector set name, if any.
   if (!Tok.is(tok::identifier)) {
     P.Diag(Tok.getLocation(), diag::warn_omp_declare_variant_cs_name_expected)
do {
       }
     } while (Tok.is(tok::identifier));
     // Parse ')'.
     (void)T.consumeClose();
     if (!Kinds.empty())
       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)
         << "device";
     // Skip until either '}', ')', or end of directive.
     while (!P.SkipUntil(tok::r_brace, tok::r_paren,
                         tok::annot_pragma_openmp_end, Parser::StopBeforeMatch))
       ;

clang/lib/Sema/SemaOpenMP.cpp

 #include "clang/AST/StmtVisitor.h"
 #include "clang/AST/TypeOrdering.h"
 #include "clang/Basic/OpenMPKinds.h"
 #include "clang/Sema/Initialization.h"
 #include "clang/Sema/Lookup.h"
 #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;
 
 //===----------------------------------------------------------------------===//
 // Stack of data-sharing attributes for variables
 //===----------------------------------------------------------------------===//
 
 static const Expr *checkMapClauseExpressionBase(
if (areMultiversionVariantFunctionsCompatible(
                               PDiag(diag::err_omp_declare_variant_diff)
                                   << FD->getLocation()),
           /*TemplatesSupported=*/true, /*ConstexprSupported=*/false,
           /*CLinkageMayDiffer=*/true))
     return None;
   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) {
   if (Data.empty())
     return;
   SmallVector<Expr *, 4> CtxScores;
   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;
     OpenMPContextSelectorKind Ctx = D.Ctx;
     if (CtxSet == OMP_CTX_SET_unknown || Ctx == OMP_CTX_unknown)
       return;
     Expr *Score = nullptr;
     if (D.Score.isUsable()) {
       Score = D.Score.get();
       if (!Score->isTypeDependent() && !Score->isValueDependent() &&
           !Score->isInstantiationDependent() &&
           !Score->containsUnexpandedParameterPack()) {
         Score =
             PerformOpenMPImplicitIntegerConversion(Score->getExprLoc(), Score)
                 .get();
         if (Score)
           Score = VerifyIntegerConstantExpression(Score).get();
       }
     } else {
       // OpenMP 5.0, 2.3.3 Matching and Scoring Context Selectors.
       // 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();
     }
+    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.append(D.Names.begin(), D.Names.end());
+        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.append(D.Names.begin(), D.Names.end());
+        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);
     Ctxs.push_back(Ctx);
     CtxScores.push_back(Score);
   }
   if (!IsError) {
     auto *NewAttr = OMPDeclareVariantAttr::CreateImplicit(
         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);
   }
 }
 
 void Sema::markOpenMPDeclareVariantFuncsReferenced(SourceLocation Loc,
                                                    FunctionDecl *Func,
                                                    bool MightBeOdrUse) {
   assert(LangOpts.OpenMP && "Expected OpenMP mode.");

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

 #include "clang/AST/Expr.h"
 #include "clang/AST/ExprCXX.h"
 #include "clang/AST/PrettyDeclStackTrace.h"
 #include "clang/AST/TypeLoc.h"
 #include "clang/Sema/Initialization.h"
 #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;
 
 static bool isDeclWithinFunction(const Decl *D) {
   const DeclContext *DC = D->getDeclContext();
   if (DC->isFunctionOrMethod())
     return true;
for (unsigned I = 0, E = Attr.scores_size(); I < E; ++I) {
     ExprResult Score;
     if (Expr *E = *std::next(Attr.scores_begin(), I))
       Score = Subst(E);
     // Instantiate the attribute.
     auto CtxSet = static_cast<OpenMPContextSelectorSetKind>(
         *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;
+      assert(CtxSet == OMP_CTX_SET_implementation &&
+             "Expected implementation context selector set.");
+      for (StringRef Vendor : Attr.implVendors())
+        Values.push_back(Sema::OMPCtxStringType(Vendor));
+      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_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.");
-      }
-      break;
-    case OMP_CTX_SET_unknown:
-      llvm_unreachable("Unexpected context selector set kind.");
-    }
+    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);
+      }
+    Data.emplace_back(CtxSet, Ctx, Score, Values);
   }
   S.ActOnOpenMPDeclareVariantDirective(DeclVarData.getValue().first,
                                        DeclVarData.getValue().second,
                                        Attr.getRange(), Data);
 }
 
 static void instantiateDependentAMDGPUFlatWorkGroupSizeAttr(
     Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,

clang/test/OpenMP/declare_variant_ast_print.c

 // RUN: %clang_cc1 -verify -fopenmp -x c -std=c99 -ast-print %s -o - -Wno-openmp-clauses | FileCheck %s
 
 // RUN: %clang_cc1 -verify -fopenmp-simd -x c -std=c99 -ast-print %s -o - -Wno-openmp-clauses | FileCheck %s
 
 // expected-no-diagnostics
 
 int foo(void);
 
 #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);
 
 // 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();

clang/test/OpenMP/declare_variant_ast_print.cpp

 // CHECK-NEXT: }
 template <typename T>
 T foofoo() { return T(); }
 
 // CHECK:      template<> int foofoo<int>() {
 // 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})
 #pragma omp declare variant(foofoo <T>) match(user = {score(<expr>) : condition(<expr>)})
 #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() {
 // CHECK-NEXT: return barbar<int, 3>();
 // CHECK-NEXT: }
 int baz() {
   return barbar<int, 3>();
 }
 
 // CHECK:      template <class C> void h_ref(C *hp, C *hp2, C *hq, C *lin) {
 // CHECK-NEXT: }
 // CHECK-NEXT: template<> void h_ref<double>(double *hp, double *hp2, double *hq, double *lin) {
 // CHECK-NEXT: }
 // CHECK-NEXT: template<> void h_ref<float>(float *hp, float *hp2, float *hq, float *lin) {
 // CHECK-NEXT: }
 template <class C>
 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: }
 
 // CHECK-NEXT: template<> void h<double>(double *hp, double *hp2, double *hq, double *lin) {
 // 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) {
   h((float *)hp, (float *)hp2, (float *)hq, (float *)lin);
 }
 
 // CHECK: int fn();
 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 {
 // CHECK-NEXT: void vd() {
 // CHECK-NEXT: }
 // CHECK-NEXT: SpecialFuncs();
 // CHECK-NEXT: ~SpecialFuncs() noexcept;
 // CHECK-NEXT: void baz() {
 // CHECK-NEXT: }
 // CHECK-NEXT: void bar() {
 // 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 {
   void vd() {}
   SpecialFuncs();
   ~SpecialFuncs();
 
   void baz() {}
   void bar() {}
   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() {}
 
 // CHECK:      static void static_f_variant() {
 // 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() {}
 
 // CHECK: void bazzzz() {
 // CHECK-NEXT: s.foo1();
 // CHECK-NEXT: static_f();
 // CHECK-NEXT: }
 void bazzzz() {
   s.foo1();
   static_f();
 }
 
 // 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();
 

clang/test/OpenMP/declare_variant_device_isa_codegen.cpp

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

clang/test/OpenMP/declare_variant_messages.c

-// 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
 
 int foo(void);
 
 #pragma omp declare variant // expected-error {{expected '(' after 'declare variant'}}
 #pragma omp declare variant(  // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}

 #pragma omp declare variant(foo) match(device={kind}) // expected-error {{expected '(' after 'kind'}} expected-error {{expected 'host', 'nohost', 'cpu', 'gpu', or 'fpga' in 'kind' context selector of 'device' selector set of 'omp declare variant' directive}} expected-error {{expected ')'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
 #pragma omp declare variant(foo) match(device={kind(}) // expected-error {{expected 'host', 'nohost', 'cpu', 'gpu', or 'fpga' in 'kind' context selector of 'device' selector set of 'omp declare variant' directive}} expected-error 2 {{expected ')'}} expected-note {{to match this '('}}
 #pragma omp declare variant(foo) match(device={kind()}) // expected-error {{expected 'host', 'nohost', 'cpu', 'gpu', or 'fpga' in 'kind' context selector of 'device' selector set of 'omp declare variant' directive}}
 #pragma omp declare variant(foo) match(device={kind(score cpu)}) // expected-error {{expected ')' or ',' after 'score'}} 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( ibm)}) // 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(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}}
 #pragma omp declare variant(foo) match(xxx={})
 int a;
 // expected-error@+2 {{'#pragma omp declare variant' can only be applied to functions}}
 #pragma omp declare variant(foo) match(xxx={})
 #pragma omp threadprivate(a)

clang/test/OpenMP/declare_variant_messages.cpp

 #pragma omp declare variant(foofoo <int>) match(device={kind(}) // expected-error {{expected 'host', 'nohost', 'cpu', 'gpu', or 'fpga' in 'kind' context selector of 'device' selector set of 'omp declare variant' directive}} expected-error 2 {{expected ')'}} expected-note {{to match this '('}}
 #pragma omp declare variant(foofoo <int>) match(device={kind()}) // expected-error {{expected 'host', 'nohost', 'cpu', 'gpu', or 'fpga' in 'kind' context selector of 'device' selector set of 'omp declare variant' directive}}
 #pragma omp declare variant(foofoo <int>) match(device={kind(score cpu)}) // expected-error {{expected ')' or ',' after 'score'}} 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( ibm)}) // 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(foofoo <int>) match(device={kind(score(C 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(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();
 
 // expected-error@+2 {{'#pragma omp declare variant' can only be applied to functions}}
 #pragma omp declare variant(barbar <int>) match(xxx = {})
 int a;
 // expected-error@+2 {{'#pragma omp declare variant' can only be applied to functions}}
 #pragma omp declare variant(barbar <int>) match(xxx = {})

clang/test/OpenMP/declare_variant_mixed_codegen.cpp

-// 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}}
 // 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: ret i32 87
 // CHECK-NOT: ret i32 {{1|4|81|84}}
 
 #ifndef HEADER
 #define HEADER
 
 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(); }
 
 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(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;
 
 int SpecialFuncs::Method() { return 1; }
 
 struct SpecSpecialFuncs {
   void vd() {}
   SpecSpecialFuncs();
   ~SpecSpecialFuncs();
 
   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;
 
 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(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; }
 #pragma omp declare variant(fn_variant2) match(implementation = {vendor(llvm, ibm)}, device={kind(cpu)})
 #pragma omp declare variant(fn_variant2) match(implementation = {vendor(llvm)}, device={kind(cpu,gpu)})
 #pragma omp declare variant(fn_variant2) match(implementation = {vendor(llvm)}, device={kind(nohost)})
 #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

clang/test/OpenMP/nvptx_declare_variant_device_isa_codegen.cpp

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