[OPENMP50]Add device/kind context selector support.

[ABataev]

Summary

Added basic parsing/sema support for device/kind context selector.

[jdoerfert]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

11249

Nit: bool RHSIsSubsetOfLHS = isStrictSubset(RHS, LHS)
saves you a line later.

clang/test/OpenMP/declare_variant_device_kind_codegen.cpp

82

What is wrong about cpu, gpu together?

[ABataev]

clang/test/OpenMP/declare_variant_device_kind_codegen.cpp

82

When several traits sre used, the logic and is applied. So, cpu,gpu means that the target is cpu and gpu at the same time.

[ABataev]

Address comment

[ABataev]

More comments?

[jdoerfert]

LGTM.

[jdoerfert]

This comment was removed by jdoerfert.

clang/include/clang/Basic/Attr.td

def OMPDeclareVariant : InheritableAttr {
   let HasCustomParsing = 1;
   let InheritEvenIfAlreadyPresent = 1;
   let Documentation = [OMPDeclareVariantDocs];
   let Args = [
     ExprArgument<"VariantFuncRef">,
     VariadicExprArgument<"Scores">,
     VariadicUnsignedArgument<"CtxSelectorSets">,
     VariadicUnsignedArgument<"CtxSelectors">,
-    VariadicStringArgument<"ImplVendors">
+    VariadicStringArgument<"ImplVendors">,
+    VariadicStringArgument<"DeviceKinds">
   ];
   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)
             printScore(OS, Policy, I);
             if (implVendors_size() > 0) {
               OS << *implVendors(). begin();
               for (StringRef VendorName : llvm::drop_begin(implVendors(), 1))
                 OS << ", " << VendorName;
             }
             OS << ")";
             break;
+          case OMP_CTX_kind:
+            llvm_unreachable("Unexpected context selector in implementation set.");
+          case OMP_CTX_unknown:
+            llvm_unreachable("Unknown context selector.");
+          }
+          OS << "}";
+          break;
+        case OMP_CTX_SET_device:
+          OS << "device={";
+          switch (Ctx) {
+          case OMP_CTX_kind:
+            OS << "kind(";
+            if (deviceKinds_size() > 0) {
+              OS << *deviceKinds().begin();
+              for (StringRef KindName : llvm::drop_begin(deviceKinds(), 1))
+                OS << ", " << KindName;
+            }
+            OS << ")";
+            break;
+          case OMP_CTX_vendor:
+            llvm_unreachable("Unexpected context selector in device set.");
           case OMP_CTX_unknown:
             llvm_unreachable("Unknown context selector.");
           }
           OS << "}";
           break;
         case OMP_CTX_SET_unknown:
           llvm_unreachable("Unknown context selector set.");
         }
+        if (I != E - 1)
+          OS << ",";
       }
       OS << ")";
     }
   }];
 }
 
 def InternalLinkage : InheritableAttr {
   let Spellings = [Clang<"internal_linkage">];

clang/include/clang/Basic/DiagnosticParseKinds.td

 def err_omp_expected_comma_brace : Error<"expected '}' or ',' after '%0'">;
 def err_omp_declare_variant_ctx_mutiple_use : Error<
   "context trait selector '%0' is used already in the same '%1' context selector set of 'omp declare variant' directive">;
 def note_omp_declare_variant_ctx_used_here : Note<
   "previously context trait selector '%0' used here">;
 def warn_omp_more_one_device_type_clause : Warning<
   "more than one 'device_type' clause is specified">,
   InGroup<OpenMPClauses>;
+def err_omp_wrong_device_kind_trait : Error<
+  "unknown '%0' device kind trait in the 'device' context selector set, expected"
+  " one of 'host', 'nohost', 'cpu', 'gpu' or 'fpga'">;
 
 // Pragma loop support.
 def err_pragma_loop_missing_argument : Error<
   "missing argument; expected %select{an integer value|"
   "'enable'%select{|, 'full'}1%select{|, 'assume_safety'}2 or 'disable'}0">;
 def err_pragma_loop_invalid_option : Error<
   "%select{invalid|missing}0 option%select{ %1|}0; expected vectorize, "
   "vectorize_width, interleave, interleave_count, unroll, unroll_count, "

clang/include/clang/Basic/OpenMPKinds.def

 #define OPENMP_CONTEXT_SELECTOR_SET(Name)
 #endif
 #ifndef OPENMP_CONTEXT_SELECTOR
 #define OPENMP_CONTEXT_SELECTOR(Name)
 #endif
 
 // 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 directives.
 OPENMP_DIRECTIVE(threadprivate)
 OPENMP_DIRECTIVE(parallel)
 OPENMP_DIRECTIVE(task)
 OPENMP_DIRECTIVE(simd)
 OPENMP_DIRECTIVE(for)
 OPENMP_DIRECTIVE(sections)

clang/lib/CodeGen/CGOpenMPRuntime.cpp

 #include "CGOpenMPRuntime.h"
 #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/Bitcode/BitcodeReader.h"
 #include "llvm/IR/DerivedTypes.h"
 #include "llvm/IR/GlobalValue.h"
 #include "llvm/IR/Value.h"
 #include "llvm/Support/Format.h"
 #include "llvm/Support/raw_ostream.h"
 #include <cassert>
 

 
 namespace {
 using OMPContextSelectorData =
     OpenMPCtxSelectorData<StringRef, ArrayRef<StringRef>, 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>
-static bool checkContext(const OMPContextSelectorData &Data) {
+template <OpenMPContextSelectorSetKind CtxSet, OpenMPContextSelectorKind Ctx,
+          typename... Arguments>
+static bool checkContext(const OMPContextSelectorData &Data,
+                         Arguments... Params) {
   assert(Data.CtxSet != OMP_CTX_SET_unknown && Data.Ctx != OMP_CTX_unknown &&
          "Unknown context selector or context selector set.");
   return false;
 }
 
 /// Checks for implementation={vendor(<vendor>)} context selector.
 /// \returns true iff <vendor>="llvm", false otherwise.
 template <>
 bool checkContext<OMP_CTX_SET_implementation, OMP_CTX_vendor>(
     const OMPContextSelectorData &Data) {
   return llvm::all_of(Data.Names,
                       [](StringRef S) { return !S.compare_lower("llvm"); });
 }
 
-bool matchesContext(const CompleteOMPContextSelectorData &ContextData) {
+/// Checks for device={kind(<kind>)} context selector.
+/// \returns true if <kind>="host" and compilation is for host.
+/// true if <kind>="nohost" and compilation is for device.
+/// true if <kind>="cpu" and compilation is for Arm, X86 or PPC CPU.
+/// true if <kind>="gpu" and compilation is for NVPTX or AMDGCN.
+/// false otherwise.
+template <>
+bool checkContext<OMP_CTX_SET_device, OMP_CTX_kind, CodeGenModule &>(
+    const OMPContextSelectorData &Data, CodeGenModule &CGM) {
+  for (StringRef Name : Data.Names) {
+    if (!Name.compare_lower("host")) {
+      if (CGM.getLangOpts().OpenMPIsDevice)
+        return false;
+      continue;
+    }
+    if (!Name.compare_lower("nohost")) {
+      if (!CGM.getLangOpts().OpenMPIsDevice)
+        return false;
+      continue;
+    }
+    switch (CGM.getTriple().getArch()) {
+    case llvm::Triple::arm:
+    case llvm::Triple::armeb:
+    case llvm::Triple::aarch64:
+    case llvm::Triple::aarch64_be:
+    case llvm::Triple::aarch64_32:
+    case llvm::Triple::ppc:
+    case llvm::Triple::ppc64:
+    case llvm::Triple::ppc64le:
+    case llvm::Triple::x86:
+    case llvm::Triple::x86_64:
+      if (Name.compare_lower("cpu"))
+        return false;
+      break;
+    case llvm::Triple::amdgcn:
+    case llvm::Triple::nvptx:
+    case llvm::Triple::nvptx64:
+      if (Name.compare_lower("gpu"))
+        return false;
+      break;
+    case llvm::Triple::UnknownArch:
+    case llvm::Triple::arc:
+    case llvm::Triple::avr:
+    case llvm::Triple::bpfel:
+    case llvm::Triple::bpfeb:
+    case llvm::Triple::hexagon:
+    case llvm::Triple::mips:
+    case llvm::Triple::mipsel:
+    case llvm::Triple::mips64:
+    case llvm::Triple::mips64el:
+    case llvm::Triple::msp430:
+    case llvm::Triple::r600:
+    case llvm::Triple::riscv32:
+    case llvm::Triple::riscv64:
+    case llvm::Triple::sparc:
+    case llvm::Triple::sparcv9:
+    case llvm::Triple::sparcel:
+    case llvm::Triple::systemz:
+    case llvm::Triple::tce:
+    case llvm::Triple::tcele:
+    case llvm::Triple::thumb:
+    case llvm::Triple::thumbeb:
+    case llvm::Triple::xcore:
+    case llvm::Triple::le32:
+    case llvm::Triple::le64:
+    case llvm::Triple::amdil:
+    case llvm::Triple::amdil64:
+    case llvm::Triple::hsail:
+    case llvm::Triple::hsail64:
+    case llvm::Triple::spir:
+    case llvm::Triple::spir64:
+    case llvm::Triple::kalimba:
+    case llvm::Triple::shave:
+    case llvm::Triple::lanai:
+    case llvm::Triple::wasm32:
+    case llvm::Triple::wasm64:
+    case llvm::Triple::renderscript32:
+    case llvm::Triple::renderscript64:
+      return false;
+    }
+  }
+  return true;
+}
+
+bool matchesContext(CodeGenModule &CGM,
+                    const CompleteOMPContextSelectorData &ContextData) {
   for (const OMPContextSelectorData &Data : ContextData) {
     switch (Data.CtxSet) {
     case OMP_CTX_SET_implementation:
       switch (Data.Ctx) {
       case OMP_CTX_vendor:
         if (!checkContext<OMP_CTX_SET_implementation, OMP_CTX_vendor>(Data))
           return false;
         break;
+      case OMP_CTX_kind:
       case OMP_CTX_unknown:
-        llvm_unreachable("Unexpected context selector kind.");
+        llvm_unreachable(
+            "Unexpected context selector kind in implementation set.");
+      }
+      break;
+    case OMP_CTX_SET_device:
+      switch (Data.Ctx) {
+      case OMP_CTX_kind:
+        if (!checkContext<OMP_CTX_SET_device, OMP_CTX_kind, CodeGenModule &>(
+                Data, CGM))
+          return false;
+        break;
+      case OMP_CTX_vendor:
+      case OMP_CTX_unknown:
+        llvm_unreachable("Unexpected context selector kind in device set.");
       }
       break;
     case OMP_CTX_SET_unknown:
       llvm_unreachable("Unexpected context selector set kind.");
     }
   }
   return true;
 }
for (unsigned I = 0, E = A->scores_size(); I < E; ++I) {
     Data.back().Score = Score->EvaluateKnownConstInt(C);
     switch (CtxSet) {
     case OMP_CTX_SET_implementation:
       switch (Ctx) {
       case OMP_CTX_vendor:
         Data.back().Names =
             llvm::makeArrayRef(A->implVendors_begin(), A->implVendors_end());
         break;
+      case OMP_CTX_kind:
       case OMP_CTX_unknown:
-        llvm_unreachable("Unexpected context selector kind.");
+        llvm_unreachable(
+            "Unexpected context selector kind in implementation set.");
+      }
+      break;
+    case OMP_CTX_SET_device:
+      switch (Ctx) {
+      case OMP_CTX_kind:
+        Data.back().Names =
+            llvm::makeArrayRef(A->deviceKinds_begin(), A->deviceKinds_end());
+        break;
+      case OMP_CTX_vendor:
+      case OMP_CTX_unknown:
+        llvm_unreachable("Unexpected context selector kind in device set.");
       }
       break;
     case OMP_CTX_SET_unknown:
       llvm_unreachable("Unexpected context selector set kind.");
     }
   }
   return Data;
 }
 
+static bool isStrictSubset(const CompleteOMPContextSelectorData &LHS,
+                           const CompleteOMPContextSelectorData &RHS) {
+  llvm::SmallDenseMap<std::pair<int, int>, llvm::StringSet<>, 4> RHSData;
+  for (const OMPContextSelectorData &D : RHS) {
+    auto &Pair = RHSData.FindAndConstruct(std::make_pair(D.CtxSet, D.Ctx));
+    Pair.getSecond().insert(D.Names.begin(), D.Names.end());
+  }
+  bool AllSetsAreEqual = true;
+  for (const OMPContextSelectorData &D : LHS) {
+    auto It = RHSData.find(std::make_pair(D.CtxSet, D.Ctx));
+    if (It == RHSData.end())
+      return false;
+    if (D.Names.size() > It->getSecond().size())
+      return false;
+    if (llvm::set_union(It->getSecond(), D.Names))
+      return false;
+    AllSetsAreEqual =
+        AllSetsAreEqual && (D.Names.size() == It->getSecond().size());
+  }
+
+  return LHS.size() != RHS.size() || !AllSetsAreEqual;
+}
+
 static bool greaterCtxScore(const CompleteOMPContextSelectorData &LHS,
                             const CompleteOMPContextSelectorData &RHS) {
   // Score is calculated as sum of all scores + 1.
   llvm::APSInt LHSScore(llvm::APInt(64, 1), /*isUnsigned=*/false);
+  bool RHSIsSubsetOfLHS = isStrictSubset(RHS, LHS);
+  if (RHSIsSubsetOfLHS) {
+    LHSScore = llvm::APSInt::get(0);
+  } else {
     for (const OMPContextSelectorData &Data : LHS) {
       if (Data.Score.getBitWidth() > LHSScore.getBitWidth()) {
         LHSScore = LHSScore.extend(Data.Score.getBitWidth()) + Data.Score;
       } else if (Data.Score.getBitWidth() < LHSScore.getBitWidth()) {
         LHSScore += Data.Score.extend(LHSScore.getBitWidth());
       } else {
         LHSScore += Data.Score;
       }
     }
+  }
   llvm::APSInt RHSScore(llvm::APInt(64, 1), /*isUnsigned=*/false);
+  if (!RHSIsSubsetOfLHS && isStrictSubset(LHS, RHS)) {
+    RHSScore = llvm::APSInt::get(0);
+  } else {
     for (const OMPContextSelectorData &Data : RHS) {
       if (Data.Score.getBitWidth() > RHSScore.getBitWidth()) {
         RHSScore = RHSScore.extend(Data.Score.getBitWidth()) + Data.Score;
       } else if (Data.Score.getBitWidth() < RHSScore.getBitWidth()) {
         RHSScore += Data.Score.extend(RHSScore.getBitWidth());
       } else {
         RHSScore += Data.Score;
       }
     }
+  }
   return llvm::APSInt::compareValues(LHSScore, RHSScore) >= 0;
 }
 
 /// Finds the variant function that matches current context with its context
 /// selector.
-static const FunctionDecl *getDeclareVariantFunction(ASTContext &Ctx,
+static const FunctionDecl *getDeclareVariantFunction(CodeGenModule &CGM,
                                                      const FunctionDecl *FD) {
   if (!FD->hasAttrs() || !FD->hasAttr<OMPDeclareVariantAttr>())
     return FD;
   // Iterate through all DeclareVariant attributes and check context selectors.
   const OMPDeclareVariantAttr *TopMostAttr = nullptr;
   CompleteOMPContextSelectorData TopMostData;
   for (const auto *A : FD->specific_attrs<OMPDeclareVariantAttr>()) {
     CompleteOMPContextSelectorData Data =
-        translateAttrToContextSelectorData(Ctx, A);
-    if (!matchesContext(Data))
+        translateAttrToContextSelectorData(CGM.getContext(), A);
+    if (!matchesContext(CGM, Data))
       continue;
     // If the attribute matches the context, find the attribute with the highest
     // score.
     if (!TopMostAttr || !greaterCtxScore(TopMostData, Data)) {
       TopMostAttr = A;
       TopMostData.swap(Data);
     }
   }
   if (!TopMostAttr)
     return FD;
   return cast<FunctionDecl>(
       cast<DeclRefExpr>(TopMostAttr->getVariantFuncRef()->IgnoreParenImpCasts())
           ->getDecl());
 }
 
 bool CGOpenMPRuntime::emitDeclareVariant(GlobalDecl GD, bool IsForDefinition) {
   const auto *D = cast<FunctionDecl>(GD.getDecl());
   // If the original function is defined already, use its definition.
   StringRef MangledName = CGM.getMangledName(GD);
   llvm::GlobalValue *Orig = CGM.GetGlobalValue(MangledName);
   if (Orig && !Orig->isDeclaration())
     return false;
-  const FunctionDecl *NewFD = getDeclareVariantFunction(CGM.getContext(), D);
+  const FunctionDecl *NewFD = getDeclareVariantFunction(CGM, D);
   // Emit original function if it does not have declare variant attribute or the
   // context does not match.
   if (NewFD == D)
     return false;
   GlobalDecl NewGD = GD.getWithDecl(NewFD);
   if (tryEmitDeclareVariant(NewGD, GD, Orig, IsForDefinition)) {
     DeferredVariantFunction.erase(D);
     return true;

clang/lib/Parse/ParseOpenMP.cpp

do {
       }
     } 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_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> } ')'
+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)
+        << "device";
+    // Skip until either '}', ')', or end of directive.
+    while (!P.SkipUntil(tok::r_brace, tok::r_paren,
+                        tok::annot_pragma_openmp_end, Parser::StopBeforeMatch))
+      ;
+    return;
+  }
+  Sema::OMPCtxStringType Buffer;
+  StringRef CtxSelectorName = P.getPreprocessor().getSpelling(Tok, Buffer);
+  auto Res = UsedCtx.try_emplace(CtxSelectorName, Tok.getLocation());
+  if (!Res.second) {
+    // OpenMP 5.0, 2.3.2 Context Selectors, Restrictions.
+    // Each trait-selector-name can only be specified once.
+    P.Diag(Tok.getLocation(), diag::err_omp_declare_variant_ctx_mutiple_use)
+        << CtxSelectorName << "device";
+    P.Diag(Res.first->getValue(), diag::note_omp_declare_variant_ctx_used_here)
+        << CtxSelectorName;
+  }
+  OpenMPContextSelectorKind CSKind = getOpenMPContextSelector(CtxSelectorName);
+  (void)P.ConsumeToken();
+  switch (CSKind) {
+  case OMP_CTX_kind: {
+    // Parse '('.
+    BalancedDelimiterTracker T(P, tok::l_paren, tok::annot_pragma_openmp_end);
+    (void)T.expectAndConsume(diag::err_expected_lparen_after,
+                             CtxSelectorName.data());
+    llvm::UniqueVector<Sema::OMPCtxStringType> Kinds;
+    do {
+      // Parse <kind>.
+      StringRef KindName;
+      if (Tok.is(tok::identifier)) {
+        Buffer.clear();
+        KindName = P.getPreprocessor().getSpelling(P.getCurToken(), Buffer);
+        SourceLocation SLoc = P.getCurToken().getLocation();
+        (void)P.ConsumeToken();
+        if (llvm::StringSwitch<bool>(KindName)
+                .Case("host", false)
+                .Case("nohost", false)
+                .Case("cpu", false)
+                .Case("gpu", false)
+                .Case("fpga", false)
+                .Default(true)) {
+          P.Diag(SLoc, diag::err_omp_wrong_device_kind_trait) << KindName;
+        } else {
+          Kinds.insert(KindName);
+        }
+      } else {
+        P.Diag(Tok.getLocation(), diag::err_omp_declare_variant_item_expected)
+            << "'host', 'nohost', 'cpu', 'gpu', or 'fpga'"
+            << "kind"
+            << "device";
+      }
+      if (!P.TryConsumeToken(tok::comma) && Tok.isNot(tok::r_paren)) {
+        P.Diag(Tok, diag::err_expected_punc)
+            << (KindName.empty() ? "kind of device" : KindName);
+      }
+    } 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_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))
+      ;
+    return;
+  }
+}
+
 /// Parses clauses for 'declare variant' directive.
 /// clause:
 /// <selector_set_name> '=' '{' <context_selectors> '}'
 /// [ ',' <selector_set_name> '=' '{' <context_selectors> '}' ]
 bool Parser::parseOpenMPContextSelectors(
     SourceLocation Loc, SmallVectorImpl<Sema::OMPCtxSelectorData> &Data) {
   llvm::StringMap<SourceLocation> UsedCtxSets;
   do {
// Unknown selector - just ignore it completely.
       OpenMPContextSelectorSetKind CSSKind =
           getOpenMPContextSelectorSet(CtxSelectorSetName);
       llvm::StringMap<SourceLocation> UsedCtx;
       do {
         switch (CSSKind) {
         case OMP_CTX_SET_implementation:
           parseImplementationSelector(*this, Loc, UsedCtx, Data);
           break;
+        case OMP_CTX_SET_device:
+          parseDeviceSelector(*this, Loc, UsedCtx, Data);
+          break;
         case OMP_CTX_SET_unknown:
           // Skip until either '}', ')', or end of directive.
           while (!SkipUntil(tok::r_brace, tok::r_paren,
                             tok::annot_pragma_openmp_end, StopBeforeMatch))
             ;
           break;
         }
         const Token PrevTok = Tok;

clang/lib/Sema/SemaOpenMP.cpp

 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;
+  SmallVector<StringRef, 4> ImplVendors, DeviceKinds;
   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: 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
         Score = ActOnIntegerConstant(SourceLocation(), 0).get();
     }
     switch (CtxSet) {
     case OMP_CTX_SET_implementation:
       switch (Ctx) {
       case OMP_CTX_vendor:
         ImplVendors.append(D.Names.begin(), D.Names.end());
         break;
+      case OMP_CTX_kind:
+      case OMP_CTX_unknown:
+        llvm_unreachable("Unexpected context selector kind.");
+      }
+      break;
+    case OMP_CTX_SET_device:
+      switch (Ctx) {
+      case OMP_CTX_kind:
+        DeviceKinds.append(D.Names.begin(), D.Names.end());
+        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.");
     }
     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(), SR);
+        ImplVendors.begin(), ImplVendors.size(), DeviceKinds.begin(),
+        DeviceKinds.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

for (unsigned I = 0, E = Attr.scores_size(); I < E; ++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.");
+      }
+      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.");
     }
   }

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)})
+#pragma omp declare variant(foo) match(implementation={vendor(llvm)}, device={kind(fpga)})
 #pragma omp declare variant(foo) match(implementation={vendor(llvm), xxx})
-#pragma omp declare variant(foo) match(implementation={vendor(unknown)})
-#pragma omp declare variant(foo) match(implementation={vendor(score(5): ibm, xxx, ibm)})
+#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(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(implementation={vendor(score(5):ibm, xxx)})
-// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0):unknown)})
-// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0):llvm)})
+// 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(0):llvm)})
+// CHECK-NEXT: #pragma omp declare variant(foo) match(implementation={vendor(score(0):llvm)},device={kind(fpga)})
 // 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)})
+// CHECK:      #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(5):ibm)},device={kind(fpga)})
 // 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)})
+// CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):llvm)},device={kind(cpu)})
 // 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})
+#pragma omp declare variant(foofoo <int>) match(implementation={vendor(llvm), xxx}, device={kind(cpu)})
 #pragma omp declare variant(foofoo <int>) match(implementation={vendor(unknown)})
-#pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(5): ibm)})
+#pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(5): ibm)}, device={kind(fpga)})
 int bar();
 
-// CHECK:      #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(C + 5):ibm, xxx)})
+// CHECK:      #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(C + 5):ibm, xxx)},device={kind(cpu, host)})
 // 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)})
+// CHECK-NEXT: #pragma omp declare variant(foofoo<T>) match(implementation={vendor(score(0):llvm)},device={kind(cpu)})
 // 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)})
+#pragma omp declare variant(foofoo <T>) match(implementation={vendor(llvm)},device={kind(cpu)})
 #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)})
+#pragma omp declare variant(foofoo <T>) match(implementation={vendor(score(C+5): ibm, xxx, ibm)},device={kind(cpu,host)})
 template <typename T, int C>
 T barbar();
 
-// CHECK:      #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(3 + 5):ibm, xxx)})
+// CHECK:      #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(3 + 5):ibm, xxx)},device={kind(cpu, host)})
 // 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)})
+// CHECK-NEXT: #pragma omp declare variant(foofoo<int>) match(implementation={vendor(score(0):llvm)},device={kind(cpu)})
 // 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)})
-// CHECK-NEXT: #pragma omp declare variant(h_ref<C>) match(implementation={vendor(score(0):llvm)})
+// 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-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)})
-#pragma omp declare variant(h_ref <C>) match(implementation={vendor(unknown)})
+#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)})
 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)})
-// CHECK-NEXT: #pragma omp declare variant(h_ref<float>) match(implementation={vendor(score(0):llvm)})
+// 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-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)})
+#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(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)})
+// CHECK:      #pragma omp declare variant(fn) match(implementation={vendor(score(0):unknown)},device={kind(cpu, gpu)})
 // 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)})
+#pragma omp declare variant(fn) match(implementation={vendor(unknown)},device={kind(cpu,gpu)})
 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)})
-// CHECK-NEXT: #pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(score(0):llvm)})
+// 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-NEXT: int fn_deduced();
 #pragma omp declare variant(fn_deduced_variant) match(xxx = {})
-#pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(llvm)})
-#pragma omp declare variant(fn_deduced_variant) match(implementation={vendor(unknown)})
+#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)})
 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)})
-// CHECK-NEXT: #pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(score(0):ibm)})
+// 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-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)})
-#pragma omp declare variant(fn_deduced_variant1) match(implementation={vendor(unknown)})
+#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)})
 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)})
-// CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::bar) match(implementation={vendor(score(0):ibm)})
+// 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: void foo1() {
 // CHECK-NEXT: }
-// CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)})
+// CHECK-NEXT: #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)},device={kind(cpu, host)})
 // 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)})
-#pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(unknown)})
+#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)})
   void foo1() {}
-#pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(unknown)})
+#pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(unknown)},device={kind(cpu, host)})
   void xxx();
 } s;
 
-// CHECK:      #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)})
+// CHECK:      #pragma omp declare variant(SpecialFuncs::baz) match(implementation={vendor(score(0):unknown)},device={kind(cpu, host)})
 // 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)})
+// CHECK-NEXT: #pragma omp declare variant(static_f_variant) match(implementation={vendor(score(0):llvm)},device={kind(fpga)})
 // 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)})
+#pragma omp declare variant(static_f_variant) match(implementation={vendor(llvm)},device={kind(fpga)})
 #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)})
+// CHECK:     #pragma omp declare variant(fn_linkage_variant) match(implementation={vendor(score(0):xxx)},device={kind(cpu, host)})
 // CHECK:     int fn_linkage();
 // CHECK: }
 int fn_linkage_variant();
 extern "C" {
-#pragma omp declare variant(fn_linkage_variant) match(implementation = {vendor(xxx)})
+#pragma omp declare variant(fn_linkage_variant) match(implementation = {vendor(xxx)},device={kind(cpu,host)})
 int fn_linkage();
 }
 
 // CHECK: extern "C" int fn_linkage_variant1()
-// CHECK: #pragma omp declare variant(fn_linkage_variant1) match(implementation={vendor(score(0):xxx)})
+// CHECK: #pragma omp declare variant(fn_linkage_variant1) match(implementation={vendor(score(0):xxx)},device={kind(cpu, host)})
 // CHECK: int fn_linkage1();
 extern "C" int fn_linkage_variant1();
-#pragma omp declare variant(fn_linkage_variant1) match(implementation = {vendor(xxx)})
+#pragma omp declare variant(fn_linkage_variant1) match(implementation = {vendor(xxx)},device={kind(cpu,host)})
 int fn_linkage1();
 

clang/test/OpenMP/declare_variant_device_kind_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 -DHOST | 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 -DHOST
+// 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 -DHOST | 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 -DHOST | 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 -DHOST
+// 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 -DHOST | 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 -DHOST | 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 -DHOST
+// 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 -DHOST | 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 -DCPU | 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 -DCPU
+// 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 -DCPU | 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 -DCPU | 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 -DCPU
+// 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 -DCPU | 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 -DCPU | 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 -DCPU
+// 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 -DCPU | 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 -DCPU
+// 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 - -DCPU | 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 -DCPU
+// 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 - -DCPU | 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 -DCPU
+// 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 - -DCPU | 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 -DCPU
+// 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 - -DCPU | 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 -DNOHOST
+// 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 - -DNOHOST | 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 -DNOHOST
+// 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 - -DNOHOST | 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 -DNOHOST
+// 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 - -DNOHOST | 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 -DNOHOST
+// 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 - -DNOHOST | 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
+#ifdef HOST
+#define CORRECT host
+#define SUBSET host, cpu
+#define WRONG host, nohost
+#endif // HOST
+#ifdef CPU
+#define CORRECT cpu
+#define SUBSET host, cpu
+#define WRONG cpu, gpu
+#endif // CPU
+#ifdef NOHOST
+#define CORRECT nohost
+#define SUBSET nohost, cpu
+#define WRONG nohost, host
+#endif // NOHOST
+
+int foo() { return 2; }
+
+#pragma omp declare variant(foo) match(device = {kind(CORRECT)})
+int bar() { return 1; }
+
+int bazzz();
+#pragma omp declare variant(bazzz) match(device = {kind(CORRECT)})
+int baz() { return 1; }
+
+int test();
+#pragma omp declare variant(test) match(device = {kind(CORRECT)})
+int call() { return 1; }
+
+static int stat_unused_();
+#pragma omp declare variant(stat_unused_) match(device = {kind(CORRECT)})
+static int stat_unused() { return 1; }
+
+static int stat_used_();
+#pragma omp declare variant(stat_used_) match(device = {kind(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 = {kind(CORRECT)})
+  int method() { return 1; }
+#pragma omp declare variant(SpecialFuncs::method_) \
+    match(device = {kind(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 = {kind(CORRECT)})
+  int method() { return 1; }
+#pragma omp declare variant(SpecSpecialFuncs::method_) \
+    match(device = {kind(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 = {kind(SUBSET)})
+#pragma omp declare variant(prio1) match(device = {kind(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 = {kind(SUBSET)})
+#pragma omp declare variant(prio2) match(device = {kind(CORRECT)})
+#pragma omp declare variant(prio3) match(device = {kind(SUBSET)})
+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 = {kind(CORRECT)})
+int fn_linkage() { return 1; }
+}
+
+extern "C" int fn_linkage_variant1() { return 86; }
+#pragma omp declare variant(fn_linkage_variant1) match(device = {kind(CORRECT)})
+int fn_linkage1() { return 1; }
+
+int fn_variant2() { return 1; }
+#pragma omp declare variant(fn_variant2) match(device = {kind(WRONG)})
+int fn2() { return 87; }
+
+#pragma omp end declare target
+#endif // HEADER

clang/test/OpenMP/declare_variant_messages.c

 #pragma omp declare variant(foo) match(implementation={vendor}) // expected-error {{expected '(' after 'vendor'}} expected-error {{expected vendor identifier in 'vendor' context selector of 'implementation' selector set of 'omp declare variant' directive}} expected-error {{expected ')' or ',' after 'vendor name'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
 #pragma omp declare variant(foo) match(implementation={vendor(}) // expected-error {{expected vendor identifier in 'vendor' context selector of 'implementation' selector set of 'omp declare variant' directive}} expected-error {{expected ')' or ',' after 'vendor name'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
 #pragma omp declare variant(foo) match(implementation={vendor()}) // expected-error {{expected vendor identifier in 'vendor' context selector of 'implementation' selector set of 'omp declare variant' directive}}
 #pragma omp declare variant(foo) match(implementation={vendor(score ibm)}) // expected-error {{expected '(' after 'score'}} expected-warning {{missing ':' after context selector score clause - ignoring}}
 #pragma omp declare variant(foo) match(implementation={vendor(score( ibm)}) // expected-error {{expected ')' or ',' after 'vendor name'}} expected-error {{expected ')'}} expected-error {{use of undeclared identifier 'ibm'}} expected-error {{expected vendor identifier in 'vendor' context selector of 'implementation' selector set of 'omp declare variant' directive}} expected-warning {{missing ':' after context selector score clause - ignoring}} expected-note {{to match this '('}}
 #pragma omp declare variant(foo) match(implementation={vendor(score(2 ibm)}) // expected-error {{expected ')' or ',' after 'vendor name'}} expected-error 2 {{expected ')'}} expected-error {{expected vendor identifier in 'vendor' context selector of 'implementation' selector set of 'omp declare variant' directive}} expected-warning {{missing ':' after context selector score clause - ignoring}} expected-note 2 {{to match this '('}}
 #pragma omp declare variant(foo) match(implementation={vendor(score(foo()) ibm)}) // expected-warning {{missing ':' after context selector score clause - ignoring}} expected-error {{expression is not an integer constant expression}}
 #pragma omp declare variant(foo) match(implementation={vendor(score(5): ibm), vendor(llvm)}) // expected-error {{context trait selector 'vendor' is used already in the same 'implementation' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'vendor' used here}}
+#pragma omp declare variant(foo) match(implementation={vendor(score(5): ibm), kind(cpu)}) // expected-warning {{unknown context selector in 'implementation' context selector set of 'omp declare variant' directive, ignored}}
+#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={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'}}}
 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)

 int diff_ret_variant(void);
 // expected-error@+1 {{function with '#pragma omp declare variant' has a different return type}}
 #pragma omp declare variant(diff_ret_variant) match(xxx={})
 void diff_ret(void);
 
 void marked(void);
 void not_marked(void);
 // expected-note@+1 {{marked as 'declare variant' here}}
-#pragma omp declare variant(not_marked) match(implementation={vendor(unknown)})
+#pragma omp declare variant(not_marked) match(implementation={vendor(unknown)}, device={kind(cpu)})
 void marked_variant(void);
 // expected-warning@+1 {{variant function in '#pragma omp declare variant' is itself marked as '#pragma omp declare variant'}}
 #pragma omp declare variant(marked_variant) match(xxx={})
 void marked(void);
 
 // expected-error@+1 {{function declaration is expected after 'declare variant' directive}}
 #pragma omp declare variant
 // expected-error@+1 {{function declaration is expected after 'declare variant' directive}}
 #pragma omp declare variant

clang/test/OpenMP/declare_variant_messages.cpp

 #pragma omp declare variant(foofoo <int>) match(implementation={vendor}) // expected-error {{expected '(' after 'vendor'}} expected-error {{expected vendor identifier in 'vendor' context selector of 'implementation' selector set of 'omp declare variant' directive}} expected-error {{expected ')' or ',' after 'vendor name'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
 #pragma omp declare variant(foofoo <int>) match(implementation={vendor(}) // expected-error {{expected vendor identifier in 'vendor' context selector of 'implementation' selector set of 'omp declare variant' directive}} expected-error {{expected ')' or ',' after 'vendor name'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
 #pragma omp declare variant(foofoo <int>) match(implementation={vendor()}) // expected-error {{expected vendor identifier in 'vendor' context selector of 'implementation' selector set of 'omp declare variant' directive}}
 #pragma omp declare variant(foofoo <int>) match(implementation={vendor(score ibm)}) // expected-error {{expected '(' after 'score'}} expected-warning {{missing ':' after context selector score clause - ignoring}}
 #pragma omp declare variant(foofoo <int>) match(implementation={vendor(score( ibm)}) // expected-error {{expected ')' or ',' after 'vendor name'}} expected-error {{expected ')'}} expected-error {{use of undeclared identifier 'ibm'}} expected-error {{expected vendor identifier in 'vendor' context selector of 'implementation' selector set of 'omp declare variant' directive}} expected-warning {{missing ':' after context selector score clause - ignoring}} expected-note {{to match this '('}}
 #pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(2 ibm)}) // expected-error {{expected ')' or ',' after 'vendor name'}} expected-error 2 {{expected ')'}} expected-error {{expected vendor identifier in 'vendor' context selector of 'implementation' selector set of 'omp declare variant' directive}} expected-warning {{missing ':' after context selector score clause - ignoring}} expected-note 2 {{to match this '('}}
 #pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(foofoo <int>()) ibm)}) // expected-warning {{missing ':' after context selector score clause - ignoring}} expected-error {{expression is not an integral constant expression}} expected-note {{non-constexpr function 'foofoo<int>' cannot be used in a constant expression}}
 #pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(5): ibm), vendor(llvm)}) // expected-error {{context trait selector 'vendor' is used already in the same 'implementation' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'vendor' used here}}
+#pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(5): ibm), kind(cpu)}) // expected-warning {{unknown context selector in 'implementation' context selector set of 'omp declare variant' directive, ignored}}
+#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={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(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(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(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(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(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'}}
 int bar();
 
 #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(foofoo <T>                 // expected-error {{expected ')'}} expected-error {{expected 'match' clause on 'omp declare variant' directive}} expected-note {{to match this '('}}
 #pragma omp declare variant(x)                         // expected-error {{use of undeclared identifier 'x'}}
 #pragma omp declare variant(foo)                       // expected-error {{expected 'match' clause on 'omp declare variant' directive}}
 #pragma omp declare variant(foofoo)                    // expected-error {{expected 'match' clause on 'omp declare variant' directive}}

 #pragma omp declare variant(foofoo <T>) match(user = {condition(<expr>)})
 #pragma omp declare variant(foofoo <T>) match(xxx = {vvv} xxx) // expected-error {{expected ','}} expected-error {{expected '=' after 'xxx' context selector set name on 'omp declare variant' directive}} expected-error {{context selector set 'xxx' is used already in the same 'omp declare variant' directive}} expected-note {{previously context selector set 'xxx' used here}}
 #pragma omp declare variant(foofoo <T>) match(xxx = {vvv}) xxx // expected-warning {{extra tokens at the end of '#pragma omp declare variant' are ignored}}
 #pragma omp declare variant(foofoo <int>) match(implementation={vendor(score ibm)}) // expected-error {{expected '(' after 'score'}} expected-warning {{missing ':' after context selector score clause - ignoring}}
 #pragma omp declare variant(foofoo <int>) match(implementation={vendor(score( ibm)}) // expected-error {{expected ')' or ',' after 'vendor name'}} expected-error {{expected ')'}} expected-error {{use of undeclared identifier 'ibm'}} expected-error {{expected vendor identifier in 'vendor' context selector of 'implementation' selector set of 'omp declare variant' directive}} expected-warning {{missing ':' after context selector score clause - ignoring}} expected-note {{to match this '('}}
 #pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(C ibm)}) // expected-error {{expected ')' or ',' after 'vendor name'}} expected-error 2 {{expected ')'}} expected-error {{expected vendor identifier in 'vendor' context selector of 'implementation' selector set of 'omp declare variant' directive}} expected-warning {{missing ':' after context selector score clause - ignoring}} expected-note 2 {{to match this '('}}
 #pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(foofoo <int>()) ibm)}) // expected-warning {{missing ':' after context selector score clause - ignoring}} expected-error {{expression is not an integral constant expression}} expected-note {{non-constexpr function 'foofoo<int>' cannot be used in a constant expression}}
 #pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(C+5): ibm), vendor(llvm)}) // expected-error {{context trait selector 'vendor' is used already in the same 'implementation' context selector set of 'omp declare variant' directive}} expected-note {{previously context trait selector 'vendor' used here}}
+#pragma omp declare variant(foofoo <int>) match(implementation={vendor(score(5): ibm), kind(cpu)}) // expected-warning {{unknown context selector in 'implementation' context selector set of 'omp declare variant' directive, ignored}}
+#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={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(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'}}
 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
+// 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
+
+int foo() { return 2; }
+
+#pragma omp declare variant(foo) match(implementation = {vendor(llvm)}, device={kind(cpu)})
+int bar() { return 1; }
+
+int bazzz();
+#pragma omp declare variant(bazzz) match(implementation = {vendor(llvm)}, device={kind(host)})
+int baz() { return 1; }
+
+int test();
+#pragma omp declare variant(test) match(implementation = {vendor(llvm)}, device={kind(cpu)})
+int call() { return 1; }
+
+static int stat_unused_();
+#pragma omp declare variant(stat_unused_) match(implementation = {vendor(llvm)}, device={kind(cpu)})
+static int stat_unused() { return 1; }
+
+static int stat_used_();
+#pragma omp declare variant(stat_used_) match(implementation = {vendor(llvm)}, device={kind(host)})
+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)})
+  int method() { return 1; }
+#pragma omp declare variant(SpecialFuncs::method_)                             \
+    match(implementation = {vendor(llvm)}, device={kind(host)})
+  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)})
+  int method() { return 1; }
+#pragma omp declare variant(SpecSpecialFuncs::method_)                         \
+    match(implementation = {vendor(llvm)}, device={kind(host)})
+  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)})
+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)})
+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)})
+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)})
+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)})
+int fn2() { return 87; }
+
+#endif // HEADER

clang/test/OpenMP/nvptx_declare_variant_device_kind_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 -DGPU
+// 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 -DGPU | 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 -DGPU
+// 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 -DGPU | FileCheck %s --implicit-check-not='ret i32 {{1|81|84}}'
+
+// 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 -DNOHOST
+// 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 -DNOHOST | 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 -DNOHOST
+// 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 -DNOHOST | 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
+
+#ifdef GPU
+#define CORRECT gpu
+#define SUBSET nohost, gpu
+#define WRONG cpu, gpu
+#endif // GPU
+#ifdef NOHOST
+#define CORRECT nohost
+#define SUBSET nohost, gpu
+#define WRONG nohost, host
+#endif // NOHOST
+
+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 = {kind(CORRECT)})
+int bar() { return 1; }
+
+#pragma omp declare variant(bazzz) match(device = {kind(CORRECT)})
+int baz() { return 1; }
+
+#pragma omp declare variant(test) match(device = {kind(CORRECT)})
+int call() { return 1; }
+
+#pragma omp declare variant(stat_unused_) match(device = {kind(CORRECT)})
+static int stat_unused() { return 1; }
+
+#pragma omp declare variant(stat_used_) match(device = {kind(CORRECT)})
+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 = {kind(CORRECT)})
+  int method() { return 1; }
+#pragma omp declare variant(SpecialFuncs::method_) \
+    match(device = {kind(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 = {kind(CORRECT)})
+  int method() { return 1; }
+#pragma omp declare variant(SpecSpecialFuncs::method_) \
+    match(device = {kind(CORRECT)})
+  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 = {kind(SUBSET)})
+#pragma omp declare variant(prio1) match(device = {kind(CORRECT)})
+int prio_() { return 1; }
+
+#pragma omp declare variant(prio4) match(device = {kind(SUBSET)})
+#pragma omp declare variant(prio2) match(device = {kind(CORRECT)})
+#pragma omp declare variant(prio3) match(device = {kind(SUBSET)})
+static int prio1_() { return 1; }
+
+int int_fn() { return prio1_(); }
+
+extern "C" {
+#pragma omp declare variant(fn_linkage_variant) match(device = {kind(CORRECT)})
+int fn_linkage() { return 1; }
+}
+
+#pragma omp declare variant(fn_linkage_variant1) match(device = {kind(CORRECT)})
+int fn_linkage1() { return 1; }
+
+#pragma omp declare variant(fn_variant2) match(device = {kind(WRONG)})
+int fn2() { return 87; }
+
+#pragma omp end declare target
+
+#endif // HEADER