[OpenMP] `omp begin/end declare variant` - part 2, sema (+"CG")

[jdoerfert]

Summary

This is the second part loosely extracted from D71179 and cleaned up.

This patch provides semantic analysis support for `omp begin/end declare
variant`, mostly as defined in OpenMP technical report 8 (TR8) [0].
The sema handling makes code generation obsolete as we generate "the
right" calls that can just be handled as usual. This handling also
applies to the existing, albeit problematic, `omp declare variant
support`. As a consequence a lot of unneeded code generation and
complexity is removed.

A major purpose of this patch is to provide proper `math.h`/`cmath`
support for OpenMP target offloading. See PR42061, PR42798, PR42799. The
current code was developed with this feature in mind, see [1].

The logic is as follows:

If we have seen a `#pragma omp begin declare variant match(<SELECTOR>)`
but not the corresponding `end declare variant`, and we find a function
definition we will:
  1) Create a function declaration for the definition we were about to generate.
  2) Create a function definition but with a mangled name (according to
     `<SELECTOR>`).
  3) Annotate the declaration with the `OMPDeclareVariantAttr`, the same
     one used already for `omp declare variant`, using and the mangled
     function definition as specialization for the context defined by
     `<SELECTOR>`.

When a call is created we inspect it. If the target has an
`OMPDeclareVariantAttr` attribute we try to specialize the call. To this
end, all variants are checked, the best applicable one is picked and a
new call to the specialization is created. The new call is used instead
of the original one to the base function. To keep the AST printing and
tooling possible we utilize the PseudoObjectExpr. The original call is
the syntactic expression, the specialized call is the semantic
expression.

[0] https://www.openmp.org/wp-content/uploads/openmp-TR8.pdf
[1] https://reviews.llvm.org/D61399#change-496lQkg0mhRN

[jdoerfert]

There will be an update later today or tomorrow, I am making the actual target offloading CUDA interoperability "work" right now.

[jdoerfert]

I'll try to remove the extra namespace and I'll split the function(is_definition) matcher out.

[jdoerfert]

Direct implementation of the standard wording, replace broken declare variant handling

[jdoerfert]

Update comment

[jdoerfert]

Minor corrections

[jdoerfert]

Minor fix

[jdoerfert]

Update tests to match new mangling

[jdoerfert]

unconditionally create explicit declarations

[hfinkel]

clang/include/clang/Basic/DiagnosticParseKinds.td

1239

We're not extremely consistent about how we phrase these kinds of errors, but I grepped for stray and we have only one message about stray tokens, but grepped for matching and we have a bunch more. Phrasing this as:

'#pragma omp end declare variant' with no matching ''#pragma omp begin declare variant'

might be better.

clang/include/clang/Sema/Sema.h

9831

scope

clang/lib/Sema/SemaDecl.cpp

13579

Please add a comment here explaining why the TemplateParameterLists.empty() check is here.

clang/lib/Sema/SemaOpenMP.cpp

5465

Is there any way in which this name might become visible to users (e.g., in error messages)?

5540

Please provide an example or otherwise explain why this failure behavior is correct.

5556

Why is the ordering here useful? Don't you collect all of the variant clauses from all of the declarations? Can't there be duplicates? (and does the relative order need always be the same?) Are we effectively supporting here, as an extension, cases where not all of the declarations have the same set of variants declared (it loos like we are because there's no break in the while (CalleeFnDecl) loop, but this makes me wonder if that would still have an odd behavior.

[jdoerfert]

Rebase and address comments

[jdoerfert]

clang/lib/Sema/SemaOpenMP.cpp

5465

Yes, I think so. One way to trigger it would be to define the same function in the same omp begin declare variant scope (or two with the same context). I haven't verified this though.
TBH, I'm unsure how bad this actually is in the short term. The original name is still a at the beginning. We should obviously make sure the error message is appropriate eventually, e.g., it de-mangles the name.

5540

Tried to improve the comment.

5556

Why is the ordering here useful?

I don't think it is ordered. We have a conceptual set and BestIdx determines the which of the set is the best. All others are equal.

Don't you collect all of the variant clauses from all of the declarations?

Yes.

Can't there be duplicates? (and does the relative order need always be the same?)

Yes (and no). getBestVariantMatchForContext should determine the best regardless of duplicates, we might just try it multiple times if we didn't manage to create a call.

Are we effectively supporting here, as an extension, cases where not all of the declarations have the same set of variants declared (it loos like we are because there's no break in the while (CalleeFnDecl) loop, but this makes me wonder if that would still have an odd behavior.

We are supporting that. All declarations are scanned and all variants are collected. What odd behavior do you refer to?

[hfinkel]

clang/lib/Sema/SemaOpenMP.cpp

5465

I think that we have to consider diagnostic quality as a first-class citizen. There are a few different ways that we might approach this. A minimal diff from what you have might be:

1. Replace the "." above with ".ompvariant."
2. Add logic to FunctionDecl::getNameForDiagnostic (or NamedDecl::getNameForDiagnostic) that recognizes that magic string in the name and alters the printing to do something intelligible for the user with the name.

(I think that (1) is a good idea anyway).

Another way would be to add some first-class property to the function and then leave the mangling to CodeGen. Are we mangling these in Sema in other places too?

Also, these can have external linkage, right?

5556

Makes sense, thanks.

[jdoerfert]

hide the mangling consistently from the user (and fix rebase bug)

clang/lib/Sema/SemaOpenMP.cpp

5465

I added a test for this in begin_declare_variant_messages.c and confirmed my earlier statement. Then I implemented what you describe above (in spirit).

Another way would be to add some first-class property to the function and then leave the mangling to CodeGen. Are we mangling these in Sema in other places too?

That would potentially work but require us to touch a lot of places that deal with redefinitions and overload resolution. Getting the right clashes between specialized versions with the same name and openmp context but not getting the clashes for the same name but different openmp context will be non-trivial. Similarly, we need to ignore specialized versions during overload resolution, etc. etc.

FWIW, this is design 8 or 9, or even more. I tried to keep the names till codegen in the beginning, using an approach similar to what you described. I tried to use namespaces (even in C) to get the right kind of name clashes and interactions, I tried ... This is the only thing that (so far) worked reliably and is (IMHO) very non-intrusive.

Also, these can have external linkage, right?

Yes. My name clash test has 3 different linkages, see begin_declare_variant_messages.c.

[hfinkel]

clang/lib/AST/DeclarationName.cpp

148

Would it make sense to print " (omp variant)" after the name to disambiguate?

[jdoerfert]

pretty print variant names as "variant[device={kind(cpu)}]"

[jdoerfert]

Split off NFC namespace stuff

[hfinkel]

LGTM

clang/include/clang/AST/Decl.h

 /// Check if the given decl is scoped.
 ///
 /// We use this function to break a cycle between the inline definitions in
 /// Type.h and Decl.h.
 inline bool IsEnumDeclScoped(EnumDecl *ED) {
   return ED->isScoped();
 }
 
+/// OpenMP variants are mangled early based on their OpenMP context selector.
+/// The new name looks likes this:
+///  <name> + OpenMPVariantManglingSeparatorStr + <mangled OpenMP context>
+static constexpr StringRef getOpenMPVariantManglingSeparatorStr() {
+  return ".ompvariant";
+}
+
 } // namespace clang
 
 #endif // LLVM_CLANG_AST_DECL_H

clang/include/clang/AST/OpenMPClause.h

 /// collection of selectors. A selector has a kind, an optional score/condition,
 /// and an ordered collection of properties.
 class OMPTraitInfo {
   /// Private constructor accesible only by ASTContext.
   OMPTraitInfo() {}
   friend class ASTContext;
 
 public:
+  /// Reconstruct a (partial) OMPTraitInfo object from a mangled name.
+  OMPTraitInfo(StringRef MangledName);
+
   struct OMPTraitProperty {
     llvm::omp::TraitProperty Kind = llvm::omp::TraitProperty::invalid;
   };
   struct OMPTraitSelector {
     Expr *ScoreOrCondition = nullptr;
     llvm::omp::TraitSelector Kind = llvm::omp::TraitSelector::invalid;
     llvm::SmallVector<OMPTraitProperty, 1> Properties;
   };
public:
   /// latter uses clang::Expr to store the score/condition while the former is
   /// independent of clang. Thus, expressions and conditions are evaluated in
   /// this method. If \p DeviceSetOnly is true, only the device selector set, if
   /// present, is put in \p VMI, otherwise all selector sets are put in \p VMI.
   void getAsVariantMatchInfo(ASTContext &ASTCtx,
                              llvm::omp::VariantMatchInfo &VMI,
                              bool DeviceSetOnly) const;
 
+  /// Return a string representation identifying this context selector.
+  std::string getMangledName() const;
+
   /// Print a human readable representation into \p OS.
   void print(llvm::raw_ostream &OS, const PrintingPolicy &Policy) const;
 };
 llvm::raw_ostream &operator<<(llvm::raw_ostream &OS, const OMPTraitInfo &TI);
 llvm::raw_ostream &operator<<(llvm::raw_ostream &OS, const OMPTraitInfo *TI);
 
 } // namespace clang
 
 #endif // LLVM_CLANG_AST_OPENMPCLAUSE_H

clang/include/clang/Basic/DiagnosticParseKinds.td

 def err_omp_map_type_missing : Error<
   "missing map type">;
 def err_omp_map_type_modifier_missing : Error<
   "missing map type modifier">;
 def err_omp_declare_simd_inbranch_notinbranch : Error<
   "unexpected '%0' clause, '%1' is specified already">;
 def err_expected_end_declare_target_or_variant : Error<
   "expected '#pragma omp end declare %select{target|variant}0'">;
+def err_expected_begin_declare_variant
+    : Error<"'#pragma omp end declare variant' with no matching '#pragma omp "
+            "begin declare variant'">;
 def err_omp_declare_target_unexpected_clause: Error<
   "unexpected '%0' clause, only %select{'to' or 'link'|'to', 'link' or 'device_type'}1 clauses expected">;
 def err_omp_expected_clause: Error<
   "expected at least one clause on '#pragma omp %0' directive">;
 def err_omp_mapper_illegal_identifier : Error<
   "illegal OpenMP user-defined mapper identifier">;
 def err_omp_mapper_expected_declarator : Error<
   "expected declarator on 'omp declare mapper' directive">;

clang/include/clang/Basic/DiagnosticSemaKinds.td

def err_omp_flush_order_clause_and_list : Error<
   "'flush' directive with memory order clause '%0' cannot have the list">;
 def note_omp_flush_order_clause_here : Note<
   "memory order clause '%0' is specified here">;
 def err_omp_non_lvalue_in_map_or_motion_clauses: Error<
   "expected addressable lvalue in '%0' clause"
   >;
 def err_omp_event_var_expected : Error<
   "expected variable of the 'omp_event_handle_t' type%select{|, not %1}0">;
+def warn_nested_declare_variant
+    : Warning<"nesting `omp begin/end declare variant` is not supported yet; "
+              "nested context ignored">,
+      InGroup<SourceUsesOpenMP>;
 } // 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/IdentifierTable.h

class alignas(IdentifierInfoAlignment) IdentifierInfo {
 
   // True if there may be additional information about
   // this identifier stored externally.
   unsigned OutOfDate : 1;
 
   // True if this is the 'import' contextual keyword.
   unsigned IsModulesImport : 1;
 
-  // 29 bits left in a 64-bit word.
+  // True if this is a mangled OpenMP variant name.
+  unsigned IsMangledOpenMPVariantName : 1;
+
+  // 28 bits left in a 64-bit word.
 
   // Managed by the language front-end.
   void *FETokenInfo = nullptr;
 
   llvm::StringMapEntry<IdentifierInfo *> *Entry = nullptr;
 
   IdentifierInfo()
       : TokenID(tok::identifier), ObjCOrBuiltinID(0), HasMacro(false),
         HadMacro(false), IsExtension(false), IsFutureCompatKeyword(false),
         IsPoisoned(false), IsCPPOperatorKeyword(false),
         NeedsHandleIdentifier(false), IsFromAST(false), ChangedAfterLoad(false),
         FEChangedAfterLoad(false), RevertedTokenID(false), OutOfDate(false),
-        IsModulesImport(false) {}
+        IsModulesImport(false), IsMangledOpenMPVariantName(false) {}
 
 public:
   IdentifierInfo(const IdentifierInfo &) = delete;
   IdentifierInfo &operator=(const IdentifierInfo &) = delete;
   IdentifierInfo(IdentifierInfo &&) = delete;
   IdentifierInfo &operator=(IdentifierInfo &&) = delete;
 
   /// Return true if this is the identifier for the specified string.
public:
   void setModulesImport(bool I) {
     IsModulesImport = I;
     if (I)
       NeedsHandleIdentifier = true;
     else
       RecomputeNeedsHandleIdentifier();
   }
 
+  /// Determine whether this is the mangled name of an OpenMP variant.
+  bool isMangledOpenMPVariantName() const { return IsMangledOpenMPVariantName; }
+
+  /// Set whether this is the mangled name of an OpenMP variant.
+  void setMangledOpenMPVariantName(bool I) { IsMangledOpenMPVariantName = I; }
+
   /// Return true if this identifier is an editor placeholder.
   ///
   /// Editor placeholders are produced by the code-completion engine and are
   /// represented as characters between '<#' and '#>' in the source code. An
   /// example of auto-completed call with a placeholder parameter is shown
   /// below:
   /// \code
   ///   function(<#int x#>);

clang/include/clang/Sema/Sema.h

private:
   /// \param Selector selects diagnostic message: 0 for type and 1 for
   ///        declaration.
   /// \return true if the type or declaration is disabled.
   template <typename T, typename DiagLocT, typename DiagInfoT, typename MapT>
   bool checkOpenCLDisabledTypeOrDecl(T D, DiagLocT DiagLoc, DiagInfoT DiagInfo,
                                      MapT &Map, unsigned Selector = 0,
                                      SourceRange SrcRange = SourceRange());
 
-  /// Marks all the functions that might be required for the currently active
-  /// OpenMP context.
-  void markOpenMPDeclareVariantFuncsReferenced(SourceLocation Loc,
-                                               FunctionDecl *Func,
-                                               bool MightBeOdrUse);
+  /// Helper to keep information about the current `omp begin/end declare
+  /// variant` nesting.
+  struct OMPDeclareVariantScope {
+    /// The associated OpenMP context selector.
+    OMPTraitInfo *TI;
+
+    /// The associated OpenMP context selector mangling.
+    std::string NameSuffix;
+
+    OMPDeclareVariantScope(OMPTraitInfo &TI)
+        : TI(&TI), NameSuffix(TI.getMangledName()) {}
+  };
+
+  /// The current `omp begin/end declare variant` scopes.
+  SmallVector<OMPDeclareVariantScope, 4> OMPDeclareVariantScopes;
+
+  /// The declarator \p D defines a function in the scope \p S which is nested
+  /// in an `omp begin/end declare variant` scope. In this method we create a
+  /// declaration for \p D and rename \p D according to the OpenMP context
+  /// selector of the surrounding scope.
+  FunctionDecl *
+  ActOnStartOfFunctionDefinitionInOpenMPDeclareVariantScope(Scope *S,
+                                                            Declarator &D);
+
+  /// Register \p FD as specialization of \p BaseFD in the current `omp
+  /// begin/end declare variant` scope.
+  void ActOnFinishedFunctionDefinitionInOpenMPDeclareVariantScope(
+      FunctionDecl *FD, FunctionDecl *BaseFD);
 
 public:
-  /// Struct to store the context selectors info for declare variant directive.
+
+  /// Can we exit a scope at the moment.
+  bool isInOpenMPDeclareVariantScope() {
+    return !OMPDeclareVariantScopes.empty();
+  }
+
+  /// Given the potential call expression \p Call, determine if there is a
+  /// specialization via the OpenMP declare variant mechanism available. If
+  /// there is, return the specialized call expression, otherwise return the
+  /// original \p Call.
+  ExprResult ActOnOpenMPCall(Sema &S, ExprResult Call, Scope *Scope,
+                             SourceLocation LParenLoc, MultiExprArg ArgExprs,
+                             SourceLocation RParenLoc, Expr *ExecConfig);
+
+  /// Handle a `omp begin declare variant`.
+  void ActOnOpenMPBeginDeclareVariant(SourceLocation Loc, OMPTraitInfo &TI);
+
+  /// Handle a `omp end declare variant`.
+  void ActOnOpenMPEndDeclareVariant();
 
   /// 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/AST/DeclarationName.cpp

 //===----------------------------------------------------------------------===//
 
 #include "clang/AST/DeclarationName.h"
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/Decl.h"
 #include "clang/AST/DeclBase.h"
 #include "clang/AST/DeclCXX.h"
 #include "clang/AST/DeclTemplate.h"
+#include "clang/AST/OpenMPClause.h"
 #include "clang/AST/PrettyPrinter.h"
 #include "clang/AST/Type.h"
 #include "clang/AST/TypeLoc.h"
 #include "clang/AST/TypeOrdering.h"
 #include "clang/Basic/IdentifierTable.h"
 #include "clang/Basic/LLVM.h"
 #include "clang/Basic/LangOptions.h"
 #include "clang/Basic/OperatorKinds.h"
static void printCXXConstructorDestructorName(QualType ClassType,
   }
   ClassType.print(OS, Policy);
 }
 
 void DeclarationName::print(raw_ostream &OS,
                             const PrintingPolicy &Policy) const {
   switch (getNameKind()) {
   case DeclarationName::Identifier:
-    if (const IdentifierInfo *II = getAsIdentifierInfo())
-      OS << II->getName();
+    if (const IdentifierInfo *II = getAsIdentifierInfo()) {
+      StringRef Name = II->getName();
+      // If this is a mangled OpenMP variant name we strip off the mangling for
+      // printing. It should not be visible to the user at all.
+      if (II->isMangledOpenMPVariantName()) {
+        std::pair<StringRef, StringRef> NameContextPair =
+            Name.split(getOpenMPVariantManglingSeparatorStr());
+        OS << NameContextPair.first << "["
+           << OMPTraitInfo(NameContextPair.second) << "]";
+      } else {
+        OS << Name;
+      }
+    }
     return;
 
   case DeclarationName::ObjCZeroArgSelector:
   case DeclarationName::ObjCOneArgSelector:
   case DeclarationName::ObjCMultiArgSelector:
     getObjCSelector().print(OS);
     return;
 

clang/lib/AST/OpenMPClause.cpp

for (const OMPTraitInfo::OMPTraitSelector &Selector : Set.Selectors) {
         }
       }
       OS << ")";
     }
     OS << "}";
   }
 }
 
+std::string OMPTraitInfo::getMangledName() const {
+  std::string MangledName;
+  llvm::raw_string_ostream OS(MangledName);
+  for (const OMPTraitInfo::OMPTraitSet &Set : Sets) {
+    OS << '.' << 'S' << unsigned(Set.Kind);
+    for (const OMPTraitInfo::OMPTraitSelector &Selector : Set.Selectors) {
+
+      bool AllowsTraitScore = false;
+      bool RequiresProperty = false;
+      isValidTraitSelectorForTraitSet(
+          Selector.Kind, Set.Kind, AllowsTraitScore, RequiresProperty);
+      OS << '.' << 's' << unsigned(Selector.Kind);
+
+      if (!RequiresProperty ||
+          Selector.Kind == TraitSelector::user_condition)
+        continue;
+
+      for (const OMPTraitInfo::OMPTraitProperty &Property : Selector.Properties)
+        OS << '.' << 'P'
+           << getOpenMPContextTraitPropertyName(Property.Kind);
+    }
+  }
+  return OS.str();
+}
+
+OMPTraitInfo::OMPTraitInfo(StringRef MangledName) {
+  unsigned long U;
+  do {
+    if (!MangledName.consume_front(".S"))
+      break;
+    if (MangledName.consumeInteger(10, U))
+      break;
+    Sets.push_back(OMPTraitSet());
+    OMPTraitSet &Set = Sets.back();
+    Set.Kind = TraitSet(U);
+    do {
+      if (!MangledName.consume_front(".s"))
+        break;
+      if (MangledName.consumeInteger(10, U))
+        break;
+      Set.Selectors.push_back(OMPTraitSelector());
+      OMPTraitSelector &Selector = Set.Selectors.back();
+      Selector.Kind = TraitSelector(U);
+      do {
+        if (!MangledName.consume_front(".P"))
+          break;
+        Selector.Properties.push_back(OMPTraitProperty());
+        OMPTraitProperty &Property = Selector.Properties.back();
+        std::pair<StringRef, StringRef> PropRestPair = MangledName.split('.');
+        Property.Kind =
+            getOpenMPContextTraitPropertyKind(Set.Kind, PropRestPair.first);
+        MangledName = PropRestPair.second;
+      } while (true);
+    } while (true);
+  } while (true);
+}
+
 llvm::raw_ostream &clang::operator<<(llvm::raw_ostream &OS,
                                      const OMPTraitInfo &TI) {
   LangOptions LO;
   PrintingPolicy Policy(LO);
   TI.print(OS, Policy);
   return OS;
 }
 llvm::raw_ostream &clang::operator<<(llvm::raw_ostream &OS,
                                      const OMPTraitInfo *TI) {
   return TI ? OS << *TI : OS;
 }

clang/lib/CodeGen/CGOpenMPRuntime.h

protected:
   /// Check if the default location must be constant.
   /// Default is false to support OMPT/OMPD.
   virtual bool isDefaultLocationConstant() const { return false; }
 
   /// Returns additional flags that can be stored in reserved_2 field of the
   /// default location.
   virtual unsigned getDefaultLocationReserved2Flags() const { return 0; }
 
-  /// Tries to emit declare variant function for \p OldGD from \p NewGD.
-  /// \param OrigAddr LLVM IR value for \p OldGD.
-  /// \param IsForDefinition true, if requested emission for the definition of
-  /// \p OldGD.
-  /// \returns true, was able to emit a definition function for \p OldGD, which
-  /// points to \p NewGD.
-  virtual bool tryEmitDeclareVariant(const GlobalDecl &NewGD,
-                                     const GlobalDecl &OldGD,
-                                     llvm::GlobalValue *OrigAddr,
-                                     bool IsForDefinition);
-
   /// Returns default flags for the barriers depending on the directive, for
   /// which this barier is going to be emitted.
   static unsigned getDefaultFlagsForBarriers(OpenMPDirectiveKind Kind);
 
   /// Get the LLVM type for the critical name.
   llvm::ArrayType *getKmpCriticalNameTy() const {return KmpCriticalNameTy;}
 
   /// Returns corresponding lock object for the specified critical region
private:
   /// List of the global variables with their addresses that should not be
   /// emitted for the target.
   llvm::StringMap<llvm::WeakTrackingVH> EmittedNonTargetVariables;
 
   /// List of variables that can become declare target implicitly and, thus,
   /// must be emitted.
   llvm::SmallDenseSet<const VarDecl *> DeferredGlobalVariables;
 
-  /// Mapping of the original functions to their variants and original global
-  /// decl.
-  llvm::MapVector<CanonicalDeclPtr<const FunctionDecl>,
-                  std::pair<GlobalDecl, GlobalDecl>>
-      DeferredVariantFunction;
-
   using NontemporalDeclsSet = llvm::SmallDenseSet<CanonicalDeclPtr<const Decl>>;
   /// Stack for list of declarations in current context marked as nontemporal.
   /// The set is the union of all current stack elements.
   llvm::SmallVector<NontemporalDeclsSet, 4> NontemporalDeclsStack;
 
   /// Stack for list of addresses of declarations in current context marked as
   /// lastprivate conditional. The set is the union of all current stack
   /// elements.
public:
   /// Checks if the variable has associated OMPAllocateDeclAttr attribute with
   /// the predefined allocator and translates it into the corresponding address
   /// space.
   virtual bool hasAllocateAttributeForGlobalVar(const VarDecl *VD, LangAS &AS);
 
   /// Return whether the unified_shared_memory has been specified.
   bool hasRequiresUnifiedSharedMemory() const;
 
-  /// Emits the definition of the declare variant function.
-  virtual bool emitDeclareVariant(GlobalDecl GD, bool IsForDefinition);
-
   /// Checks if the \p VD variable is marked as nontemporal declaration in
   /// current context.
   bool isNontemporalDecl(const ValueDecl *VD) const;
 
   /// Create specialized alloca to handle lastprivate conditionals.
   Address emitLastprivateConditionalInit(CodeGenFunction &CGF,
                                          const VarDecl *VD);
 

clang/lib/CodeGen/CGOpenMPRuntime.cpp

CGOpenMPRuntime::CGOpenMPRuntime(CodeGenModule &CGM, StringRef FirstSeparator,
   RD->completeDefinition();
   IdentQTy = C.getRecordType(RD);
   IdentTy = CGM.getTypes().ConvertRecordDeclType(RD);
   KmpCriticalNameTy = llvm::ArrayType::get(CGM.Int32Ty, /*NumElements*/ 8);
 
   loadOffloadInfoMetadata();
 }
 
-bool CGOpenMPRuntime::tryEmitDeclareVariant(const GlobalDecl &NewGD,
-                                            const GlobalDecl &OldGD,
-                                            llvm::GlobalValue *OrigAddr,
-                                            bool IsForDefinition) {
-  // Emit at least a definition for the aliasee if the the address of the
-  // original function is requested.
-  if (IsForDefinition || OrigAddr)
-    (void)CGM.GetAddrOfGlobal(NewGD);
-  StringRef NewMangledName = CGM.getMangledName(NewGD);
-  llvm::GlobalValue *Addr = CGM.GetGlobalValue(NewMangledName);
-  if (Addr && !Addr->isDeclaration()) {
-    const auto *D = cast<FunctionDecl>(OldGD.getDecl());
-    const CGFunctionInfo &FI = CGM.getTypes().arrangeGlobalDeclaration(NewGD);
-    llvm::Type *DeclTy = CGM.getTypes().GetFunctionType(FI);
-
-    // Create a reference to the named value.  This ensures that it is emitted
-    // if a deferred decl.
-    llvm::GlobalValue::LinkageTypes LT = CGM.getFunctionLinkage(OldGD);
-
-    // Create the new alias itself, but don't set a name yet.
-    auto *GA =
-        llvm::GlobalAlias::create(DeclTy, 0, LT, "", Addr, &CGM.getModule());
-
-    if (OrigAddr) {
-      assert(OrigAddr->isDeclaration() && "Expected declaration");
-
-      GA->takeName(OrigAddr);
-      OrigAddr->replaceAllUsesWith(
-          llvm::ConstantExpr::getBitCast(GA, OrigAddr->getType()));
-      OrigAddr->eraseFromParent();
-    } else {
-      GA->setName(CGM.getMangledName(OldGD));
-    }
-
-    // Set attributes which are particular to an alias; this is a
-    // specialization of the attributes which may be set on a global function.
-    if (D->hasAttr<WeakAttr>() || D->hasAttr<WeakRefAttr>() ||
-        D->isWeakImported())
-      GA->setLinkage(llvm::Function::WeakAnyLinkage);
-
-    CGM.SetCommonAttributes(OldGD, GA);
-    return true;
-  }
-  return false;
-}
-
 void CGOpenMPRuntime::clear() {
   InternalVars.clear();
   // Clean non-target variable declarations possibly used only in debug info.
   for (const auto &Data : EmittedNonTargetVariables) {
     if (!Data.getValue().pointsToAliveValue())
       continue;
     auto *GV = dyn_cast<llvm::GlobalVariable>(Data.getValue());
     if (!GV)
       continue;
     if (!GV->isDeclaration() || GV->getNumUses() > 0)
       continue;
     GV->eraseFromParent();
   }
-  // Emit aliases for the deferred aliasees.
-  for (const auto &Pair : DeferredVariantFunction) {
-    StringRef MangledName = CGM.getMangledName(Pair.second.second);
-    llvm::GlobalValue *Addr = CGM.GetGlobalValue(MangledName);
-    // If not able to emit alias, just emit original declaration.
-    (void)tryEmitDeclareVariant(Pair.second.first, Pair.second.second, Addr,
-                                /*IsForDefinition=*/false);
-  }
 }
 
 std::string CGOpenMPRuntime::getName(ArrayRef<StringRef> Parts) const {
   SmallString<128> Buffer;
   llvm::raw_svector_ostream OS(Buffer);
   StringRef Sep = FirstSeparator;
   for (StringRef Part : Parts) {
     OS << Sep << Part;
CGF.EHStack.pushCleanup<OMPAllocateCleanupTy>(NormalAndEHCleanup, FiniRTLFn,
                                                 llvm::makeArrayRef(FiniArgs));
   Addr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
       Addr,
       CGF.ConvertTypeForMem(CGM.getContext().getPointerType(CVD->getType())),
       getName({CVD->getName(), ".addr"}));
   return Address(Addr, Align);
 }
 
-/// Finds the variant function that matches current context with its context
-/// selector.
-static const FunctionDecl *getDeclareVariantFunction(CodeGenModule &CGM,
-                                                     const FunctionDecl *FD) {
-  if (!FD->hasAttrs() || !FD->hasAttr<OMPDeclareVariantAttr>())
-    return FD;
-
-  SmallVector<Expr *, 8> VariantExprs;
-  SmallVector<VariantMatchInfo, 8> VMIs;
-  for (const auto *A : FD->specific_attrs<OMPDeclareVariantAttr>()) {
-    const OMPTraitInfo &TI = *A->getTraitInfos();
-    VMIs.push_back(VariantMatchInfo());
-    TI.getAsVariantMatchInfo(CGM.getContext(), VMIs.back(),
-                             /* DeviceSetOnly */ false);
-    VariantExprs.push_back(A->getVariantFuncRef());
-  }
-
-  OMPContext Ctx(CGM.getLangOpts().OpenMPIsDevice, CGM.getTriple());
-  // FIXME: Keep the context in the OMPIRBuilder so we can add constructs as we
-  //        build them.
-
-  int BestMatchIdx = getBestVariantMatchForContext(VMIs, Ctx);
-  if (BestMatchIdx < 0)
-    return FD;
-
-  return cast<FunctionDecl>(
-      cast<DeclRefExpr>(VariantExprs[BestMatchIdx]->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, 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;
-  }
-  DeferredVariantFunction.insert(std::make_pair(D, std::make_pair(NewGD, GD)));
-  return true;
-}
-
 CGOpenMPRuntime::NontemporalDeclsRAII::NontemporalDeclsRAII(
     CodeGenModule &CGM, const OMPLoopDirective &S)
     : CGM(CGM), NeedToPush(S.hasClausesOfKind<OMPNontemporalClause>()) {
   assert(CGM.getLangOpts().OpenMP && "Not in OpenMP mode.");
   if (!NeedToPush)
     return;
   NontemporalDeclsSet &DS =
       CGM.getOpenMPRuntime().NontemporalDeclsStack.emplace_back();

clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.h

protected:
   bool isDefaultLocationConstant() const override { return true; }
 
   /// Returns additional flags that can be stored in reserved_2 field of the
   /// default location.
   /// For NVPTX target contains data about SPMD/Non-SPMD execution mode +
   /// Full/Lightweight runtime mode. Used for better optimization.
   unsigned getDefaultLocationReserved2Flags() const override;
 
-  /// Tries to emit declare variant function for \p OldGD from \p NewGD.
-  /// \param OrigAddr LLVM IR value for \p OldGD.
-  /// \param IsForDefinition true, if requested emission for the definition of
-  /// \p OldGD.
-  /// \returns true, was able to emit a definition function for \p OldGD, which
-  /// points to \p NewGD.
-  /// NVPTX backend does not support global aliases, so just use the function,
-  /// emitted for \p NewGD instead of \p OldGD.
-  bool tryEmitDeclareVariant(const GlobalDecl &NewGD, const GlobalDecl &OldGD,
-                             llvm::GlobalValue *OrigAddr,
-                             bool IsForDefinition) override;
-
 public:
   explicit CGOpenMPRuntimeNVPTX(CodeGenModule &CGM);
   void clear() override;
 
   /// Emit call to void __kmpc_push_proc_bind(ident_t *loc, kmp_int32
   /// global_tid, int proc_bind) to generate code for 'proc_bind' clause.
   virtual void emitProcBindClause(CodeGenFunction &CGF,
                                   llvm::omp::ProcBindKind ProcBind,

clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

case EM_NonSPMD:
     assert(requiresFullRuntime() && "Expected full runtime.");
     return (~KMP_IDENT_SPMD_MODE) & (~KMP_IDENT_SIMPLE_RT_MODE);
   case EM_Unknown:
     return UndefinedMode;
   }
   llvm_unreachable("Unknown flags are requested.");
 }
 
-bool CGOpenMPRuntimeNVPTX::tryEmitDeclareVariant(const GlobalDecl &NewGD,
-                                                 const GlobalDecl &OldGD,
-                                                 llvm::GlobalValue *OrigAddr,
-                                                 bool IsForDefinition) {
-  // Emit the function in OldGD with the body from NewGD, if NewGD is defined.
-  auto *NewFD = cast<FunctionDecl>(NewGD.getDecl());
-  if (NewFD->isDefined()) {
-    CGM.emitOpenMPDeviceFunctionRedefinition(OldGD, NewGD, OrigAddr);
-    return true;
-  }
-  return false;
-}
-
 CGOpenMPRuntimeNVPTX::CGOpenMPRuntimeNVPTX(CodeGenModule &CGM)
     : CGOpenMPRuntime(CGM, "_", "$") {
   if (!CGM.getLangOpts().OpenMPIsDevice)
     llvm_unreachable("OpenMP NVPTX can only handle device code.");
 }
 
 void CGOpenMPRuntimeNVPTX::emitProcBindClause(CodeGenFunction &CGF,
                                               ProcBindKind ProcBind,

clang/lib/CodeGen/CodeGenModule.h

public:
   /// Emit a code for declare mapper construct.
   void EmitOMPDeclareMapper(const OMPDeclareMapperDecl *D,
                             CodeGenFunction *CGF = nullptr);
 
   /// Emit a code for requires directive.
   /// \param D Requires declaration
   void EmitOMPRequiresDecl(const OMPRequiresDecl *D);
 
-  /// Emits the definition of \p OldGD function with body from \p NewGD.
-  /// Required for proper handling of declare variant directive on the GPU.
-  void emitOpenMPDeviceFunctionRedefinition(GlobalDecl OldGD, GlobalDecl NewGD,
-                                            llvm::GlobalValue *GV);
-
   /// Returns whether the given record has hidden LTO visibility and therefore
   /// may participate in (single-module) CFI and whole-program vtable
   /// optimization.
   bool HasHiddenLTOVisibility(const CXXRecordDecl *RD);
 
   /// Returns whether the given record has public std LTO visibility
   /// and therefore may not participate in (single-module) CFI and whole-program
   /// vtable optimization.

clang/lib/CodeGen/CodeGenModule.cpp

void CodeGenModule::EmitGlobal(GlobalDecl GD) {
   // function. If the global must always be emitted, do it eagerly if possible
   // to benefit from cache locality.
   if (MustBeEmitted(Global) && MayBeEmittedEagerly(Global)) {
     // Emit the definition if it can't be deferred.
     EmitGlobalDefinition(GD);
     return;
   }
 
-    // Check if this must be emitted as declare variant.
-  if (LangOpts.OpenMP && isa<FunctionDecl>(Global) && OpenMPRuntime &&
-      OpenMPRuntime->emitDeclareVariant(GD, /*IsForDefinition=*/false))
-    return;
-
   // If we're deferring emission of a C++ variable with an
   // initializer, remember the order in which it appeared in the file.
   if (getLangOpts().CPlusPlus && isa<VarDecl>(Global) &&
       cast<VarDecl>(Global)->hasInit()) {
     DelayedCXXInitPosition[Global] = CXXGlobalInits.size();
     CXXGlobalInits.push_back(nullptr);
   }
 
if (FD->isCPUSpecificMultiVersion()) {
     auto *Spec = FD->getAttr<CPUSpecificAttr>();
     for (unsigned I = 0; I < Spec->cpus_size(); ++I)
       EmitGlobalFunctionDefinition(GD.getWithMultiVersionIndex(I), nullptr);
     // Requires multiple emits.
   } else
     EmitGlobalFunctionDefinition(GD, GV);
 }
 
-void CodeGenModule::emitOpenMPDeviceFunctionRedefinition(
-    GlobalDecl OldGD, GlobalDecl NewGD, llvm::GlobalValue *GV) {
-  assert(getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice &&
-         OpenMPRuntime && "Expected OpenMP device mode.");
-  const auto *D = cast<FunctionDecl>(OldGD.getDecl());
-
-  // Compute the function info and LLVM type.
-  const CGFunctionInfo &FI = getTypes().arrangeGlobalDeclaration(OldGD);
-  llvm::FunctionType *Ty = getTypes().GetFunctionType(FI);
-
-  // Get or create the prototype for the function.
-  if (!GV || (GV->getType()->getElementType() != Ty)) {
-    GV = cast<llvm::GlobalValue>(GetOrCreateLLVMFunction(
-        getMangledName(OldGD), Ty, GlobalDecl(), /*ForVTable=*/false,
-        /*DontDefer=*/true, /*IsThunk=*/false, llvm::AttributeList(),
-        ForDefinition));
-    SetFunctionAttributes(OldGD, cast<llvm::Function>(GV),
-                          /*IsIncompleteFunction=*/false,
-                          /*IsThunk=*/false);
-  }
-  // We need to set linkage and visibility on the function before
-  // generating code for it because various parts of IR generation
-  // want to propagate this information down (e.g. to local static
-  // declarations).
-  auto *Fn = cast<llvm::Function>(GV);
-  setFunctionLinkage(OldGD, Fn);
-
-  // FIXME: this is redundant with part of
-  // setFunctionDefinitionAttributes
-  setGVProperties(Fn, OldGD);
-
-  MaybeHandleStaticInExternC(D, Fn);
-
-  maybeSetTrivialComdat(*D, *Fn);
-
-  CodeGenFunction(*this).GenerateCode(NewGD, Fn, FI);
-
-  setNonAliasAttributes(OldGD, Fn);
-  SetLLVMFunctionAttributesForDefinition(D, Fn);
-
-  if (D->hasAttr<AnnotateAttr>())
-    AddGlobalAnnotations(D, Fn);
-}
-
 void CodeGenModule::EmitGlobalDefinition(GlobalDecl GD, llvm::GlobalValue *GV) {
   const auto *D = cast<ValueDecl>(GD.getDecl());
 
   PrettyStackTraceDecl CrashInfo(const_cast<ValueDecl *>(D), D->getLocation(),
                                  Context.getSourceManager(),
                                  "Generating code for declaration");
 
   if (const auto *FD = dyn_cast<FunctionDecl>(D)) {
if (getLangOpts().OpenMPIsDevice && OpenMPRuntime &&
           GDDef = GlobalDecl(CD, GD.getCtorType());
         else if (const auto *DD = dyn_cast<CXXDestructorDecl>(FDDef))
           GDDef = GlobalDecl(DD, GD.getDtorType());
         else
           GDDef = GlobalDecl(FDDef);
         EmitGlobal(GDDef);
       }
     }
-    // Check if this must be emitted as declare variant and emit reference to
-    // the the declare variant function.
-    if (LangOpts.OpenMP && OpenMPRuntime)
-      (void)OpenMPRuntime->emitDeclareVariant(GD, /*IsForDefinition=*/true);
 
     if (FD->isMultiVersion()) {
       if (FD->hasAttr<TargetAttr>())
         UpdateMultiVersionNames(GD, FD);
       if (!IsForDefinition)
         return GetOrCreateMultiVersionResolver(GD, Ty, FD);
     }
   }
void CodeGenModule::HandleCXXStaticMemberVarInstantiation(VarDecl *VD) {
   if (VD->getDefinition() && TSK == TSK_ExplicitInstantiationDefinition)
     GetAddrOfGlobalVar(VD);
 
   EmitTopLevelDecl(VD);
 }
 
 void CodeGenModule::EmitGlobalFunctionDefinition(GlobalDecl GD,
                                                  llvm::GlobalValue *GV) {
-  // Check if this must be emitted as declare variant.
-  if (LangOpts.OpenMP && OpenMPRuntime &&
-      OpenMPRuntime->emitDeclareVariant(GD, /*IsForDefinition=*/true))
-    return;
-
   const auto *D = cast<FunctionDecl>(GD.getDecl());
 
   // Compute the function info and LLVM type.
   const CGFunctionInfo &FI = getTypes().arrangeGlobalDeclaration(GD);
   llvm::FunctionType *Ty = getTypes().GetFunctionType(FI);
 
   // Get or create the prototype for the function.
   if (!GV || (GV->getType()->getElementType() != Ty))

clang/lib/Parse/ParseOpenMP.cpp

 //===--- ParseOpenMP.cpp - OpenMP directives parsing ----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 /// \file
 /// This file implements parsing of all OpenMP directives and clauses.
 ///
 //===----------------------------------------------------------------------===//
 
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/StmtOpenMP.h"
 #include "clang/Basic/OpenMPKinds.h"
+#include "clang/Basic/TargetInfo.h"
 #include "clang/Basic/TokenKinds.h"
 #include "clang/Parse/ParseDiagnostic.h"
 #include "clang/Parse/Parser.h"
 #include "clang/Parse/RAIIObjectsForParser.h"
 #include "clang/Sema/Scope.h"
 #include "llvm/ADT/PointerIntPair.h"
 #include "llvm/ADT/UniqueVector.h"
 #include "llvm/Frontend/OpenMP/OMPContext.h"
if (!AssociatedFunction.isUsable()) {
     if (!Tok.is(tok::annot_pragma_openmp_end))
       while (!SkipUntil(tok::annot_pragma_openmp_end, StopBeforeMatch))
         ;
     // Skip the last annot_pragma_openmp_end.
     (void)ConsumeAnnotationToken();
     return;
   }
 
-  OMPTraitInfo TI;
+  OMPTraitInfo &TI = Actions.getASTContext().getNewOMPTraitInfo();
   if (parseOMPDeclareVariantMatchClause(Loc, TI))
     return;
 
   Optional<std::pair<FunctionDecl *, Expr *>> DeclVarData =
       Actions.checkOpenMPDeclareVariantFunction(
           Ptr, AssociatedFunction.get(), TI,
           SourceRange(Loc, Tok.getLocation()));
 
case OMPD_begin_declare_variant: {
     skipUntilPragmaOpenMPEnd(OMPD_begin_declare_variant);
 
     VariantMatchInfo VMI;
     ASTContext &ASTCtx = Actions.getASTContext();
     TI.getAsVariantMatchInfo(ASTCtx, VMI, /* DeviceSetOnly */ true);
     OMPContext OMPCtx(ASTCtx.getLangOpts().OpenMPIsDevice,
                       ASTCtx.getTargetInfo().getTriple());
 
-    if (isVariantApplicableInContext(VMI, OMPCtx))
+    if (isVariantApplicableInContext(VMI, OMPCtx)) {
+      Actions.ActOnOpenMPBeginDeclareVariant(Loc, TI);
       break;
+    }
 
     // Elide all the code till the matching end declare variant was found.
     unsigned Nesting = 1;
     SourceLocation DKLoc;
     OpenMPDirectiveKind DK = OMPD_unknown;
     do {
       DKLoc = Tok.getLocation();
       DK = parseOpenMPDirectiveKind(*this);
       if (DK == OMPD_end_declare_variant)
         --Nesting;
       else if (DK == OMPD_begin_declare_variant)
         ++Nesting;
       if (!Nesting || isEofOrEom())
         break;
       ConsumeAnyToken();
     } while (true);
 
     parseOMPEndDirective(OMPD_begin_declare_variant, OMPD_end_declare_variant,
                          DK, Loc, DKLoc, /* SkipUntilOpenMPEnd */ true);
     if (isEofOrEom())
       return nullptr;
     break;
   }
-  case OMPD_end_declare_variant:
-    // FIXME: With the sema changes we will keep track of nesting and be able to
-    // diagnose unmatchend OMPD_end_declare_variant.
+  case OMPD_end_declare_variant: {
+    if (Actions.isInOpenMPDeclareVariantScope())
+      Actions.ActOnOpenMPEndDeclareVariant();
+    else
+      Diag(Loc, diag::err_expected_begin_declare_variant);
     ConsumeToken();
     break;
+  }
   case OMPD_declare_variant:
   case OMPD_declare_simd: {
     // The syntax is:
     // { #pragma omp declare {simd|variant} }
     // <function-declaration-or-definition>
     //
     CachedTokens Toks;
     Toks.push_back(Tok);

clang/lib/Sema/SemaDecl.cpp

if (isa<C11NoReturnAttr>(NewAttribute)) {
                cast<VarDecl>(New)->isInline() &&
                !cast<VarDecl>(New)->isInlineSpecified()) {
       // Don't warn about applying selectany to implicitly inline variables.
       // Older compilers and language modes would require the use of selectany
       // to make such variables inline, and it would have no effect if we
       // honored it.
       ++I;
       continue;
+    } else if (isa<OMPDeclareVariantAttr>(NewAttribute)) {
+      // We allow to add OMP[Begin]DeclareVariantAttr to be added to
+      // declarations after defintions.
+      ++I;
+      continue;
     }
 
     S.Diag(NewAttribute->getLocation(),
            diag::warn_attribute_precede_definition);
     S.Diag(Def->getLocation(), diag::note_previous_definition);
     NewAttributes.erase(NewAttributes.begin() + I);
     --E;
   }

 Decl *
 Sema::ActOnStartOfFunctionDef(Scope *FnBodyScope, Declarator &D,
                               MultiTemplateParamsArg TemplateParameterLists,
                               SkipBodyInfo *SkipBody) {
   assert(getCurFunctionDecl() == nullptr && "Function parsing confused");
   assert(D.isFunctionDeclarator() && "Not a function declarator!");
   Scope *ParentScope = FnBodyScope->getParent();
 
+  // Check if we are in an `omp begin/end declare variant` scope. If we are, and
+  // we define a non-templated function definition, we will create a declaration
+  // instead (=BaseFD), and emit the definition with a mangled name afterwards.
+  // The base function declaration will have the equivalent of an `omp declare
+  // variant` annotation which specifies the mangled definition as a
+  // specialization function under the OpenMP context defined as part of the
+  // `omp begin declare variant`.
+  FunctionDecl *BaseFD = nullptr;
+  if (LangOpts.OpenMP && isInOpenMPDeclareVariantScope() &&
+      TemplateParameterLists.empty())
+    BaseFD = ActOnStartOfFunctionDefinitionInOpenMPDeclareVariantScope(
+        ParentScope, D);
+
   D.setFunctionDefinitionKind(FDK_Definition);
   Decl *DP = HandleDeclarator(ParentScope, D, TemplateParameterLists);
-  return ActOnStartOfFunctionDef(FnBodyScope, DP, SkipBody);
+  Decl *Dcl = ActOnStartOfFunctionDef(FnBodyScope, DP, SkipBody);
+
+  if (BaseFD)
+    ActOnFinishedFunctionDefinitionInOpenMPDeclareVariantScope(
+        cast<FunctionDecl>(Dcl), BaseFD);
+
+  return Dcl;
 }
 
 void Sema::ActOnFinishInlineFunctionDef(FunctionDecl *D) {
   Consumer.HandleInlineFunctionDefinition(D);
 }
 
 static bool
 ShouldWarnAboutMissingPrototype(const FunctionDecl *FD,

clang/lib/Sema/SemaExpr.cpp

if (ULE->hasExplicitTemplateArgs() &&
         ULE->decls_begin() == ULE->decls_end()) {
       Diag(Fn->getExprLoc(), getLangOpts().CPlusPlus2a
                                  ? diag::warn_cxx17_compat_adl_only_template_id
                                  : diag::ext_adl_only_template_id)
           << ULE->getName();
     }
   }
 
+  if (LangOpts.OpenMP)
+    Call = ActOnOpenMPCall(*this, Call, Scope, LParenLoc, ArgExprs, RParenLoc,
+                           ExecConfig);
+
   return Call;
 }
 
 /// BuildCallExpr - Handle a call to Fn with the specified array of arguments.
 /// This provides the location of the left/right parens and a list of comma
 /// locations.
 ExprResult Sema::BuildCallExpr(Scope *Scope, Expr *Fn, SourceLocation LParenLoc,
                                MultiExprArg ArgExprs, SourceLocation RParenLoc,
if (OdrUse == OdrUseContext::Used && !Func->isUsed(/*CheckUsedAttr=*/false)) {
     // Some x86 Windows calling conventions mangle the size of the parameter
     // pack into the name. Computing the size of the parameters requires the
     // parameter types to be complete. Check that now.
     if (funcHasParameterSizeMangling(*this, Func))
       CheckCompleteParameterTypesForMangler(*this, Func, Loc);
 
     Func->markUsed(Context);
   }
-
-  if (LangOpts.OpenMP)
-    markOpenMPDeclareVariantFuncsReferenced(Loc, Func, MightBeOdrUse);
 }
 
 /// Directly mark a variable odr-used. Given a choice, prefer to use
 /// MarkVariableReferenced since it does additional checks and then
 /// calls MarkVarDeclODRUsed.
 /// If the variable must be captured:
 ///  - if FunctionScopeIndexToStopAt is null, capture it in the CurContext
 ///  - else capture it in the DeclContext that maps to the

clang/lib/Sema/SemaOpenMP.cpp

if (!VD->hasLocalStorage()) {
                  getOpenMPCaptureLevels(DSAStack->getDirective(Level - 1)) - 1);
     }
   }
   return true;
 }
 
 void Sema::DestroyDataSharingAttributesStack() { delete DSAStack; }
 
+void Sema::ActOnOpenMPBeginDeclareVariant(SourceLocation Loc,
+                                          OMPTraitInfo &TI) {
+  if (!OMPDeclareVariantScopes.empty()) {
+    Diag(Loc, diag::warn_nested_declare_variant);
+    return;
+  }
+  OMPDeclareVariantScopes.push_back(OMPDeclareVariantScope(TI));
+}
+
+void Sema::ActOnOpenMPEndDeclareVariant() {
+  assert(isInOpenMPDeclareVariantScope() &&
+         "Not in OpenMP declare variant scope!");
+
+  OMPDeclareVariantScopes.pop_back();
+}
+
 void Sema::finalizeOpenMPDelayedAnalysis(const FunctionDecl *Caller,
                                          const FunctionDecl *Callee,
                                          SourceLocation Loc) {
   assert(LangOpts.OpenMP && "Expected OpenMP compilation mode.");
   Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
       OMPDeclareTargetDeclAttr::getDeviceType(Caller->getMostRecentDecl());
   // Ignore host functions during device analyzis.
   if (LangOpts.OpenMPIsDevice && DevTy &&
for (const ParmVarDecl *P : FDWithProto->parameters()) {
     Param->setScopeInfo(0, Params.size());
     Param->setImplicit();
     Params.push_back(Param);
   }
 
   FD->setParams(Params);
 }
 
+FunctionDecl *
+Sema::ActOnStartOfFunctionDefinitionInOpenMPDeclareVariantScope(Scope *S,
+                                                                Declarator &D) {
+  auto *BaseFD = cast<FunctionDecl>(ActOnDeclarator(S, D));
+  OMPDeclareVariantScope &DVScope = OMPDeclareVariantScopes.back();
+  std::string MangledName;
+  MangledName += D.getIdentifier()->getName();
+  MangledName += getOpenMPVariantManglingSeparatorStr();
+  MangledName += DVScope.NameSuffix;
+  IdentifierInfo &VariantII = Context.Idents.get(MangledName);
+
+  VariantII.setMangledOpenMPVariantName(true);
+  D.SetIdentifier(&VariantII, D.getBeginLoc());
+  return BaseFD;
+}
+
+void Sema::ActOnFinishedFunctionDefinitionInOpenMPDeclareVariantScope(
+    FunctionDecl *FD, FunctionDecl *BaseFD) {
+  // Do not mark function as is used to prevent its emission if this is the
+  // only place where it is used.
+  EnterExpressionEvaluationContext Unevaluated(
+      *this, Sema::ExpressionEvaluationContext::Unevaluated);
+
+  Expr *VariantFuncRef = DeclRefExpr::Create(
+      Context, NestedNameSpecifierLoc(), SourceLocation(), FD,
+      /* RefersToEnclosingVariableOrCapture */ false,
+      /* NameLoc */ FD->getLocation(), FD->getType(), ExprValueKind::VK_RValue);
+
+  OMPDeclareVariantScope &DVScope = OMPDeclareVariantScopes.back();
+  auto *OMPDeclareVariantA = OMPDeclareVariantAttr::CreateImplicit(
+      Context, VariantFuncRef, DVScope.TI);
+  BaseFD->addAttr(OMPDeclareVariantA);
+
+  BaseFD->setImplicit(true);
+}
+
+ExprResult Sema::ActOnOpenMPCall(Sema &S, ExprResult Call, Scope *Scope,
+                                 SourceLocation LParenLoc,
+                                 MultiExprArg ArgExprs,
+                                 SourceLocation RParenLoc, Expr *ExecConfig) {
+  // The common case is a regular call we do not want to specialize at all. Try
+  // to make that case fast by bailing early.
+  CallExpr *CE = dyn_cast<CallExpr>(Call.get());
+  if (!CE)
+    return Call;
+
+  FunctionDecl *CalleeFnDecl = CE->getDirectCallee();
+  if (!CalleeFnDecl)
+    return Call;
+
+  if (!CalleeFnDecl->hasAttr<OMPDeclareVariantAttr>())
+    return Call;
+
+  ASTContext &Context = S.getASTContext();
+  OMPContext OMPCtx(S.getLangOpts().OpenMPIsDevice,
+                    Context.getTargetInfo().getTriple());
+
+  SmallVector<Expr *, 4> Exprs;
+  SmallVector<VariantMatchInfo, 4> VMIs;
+  while (CalleeFnDecl) {
+    for (OMPDeclareVariantAttr *A :
+         CalleeFnDecl->specific_attrs<OMPDeclareVariantAttr>()) {
+      Expr *VariantRef = A->getVariantFuncRef();
+
+      VariantMatchInfo VMI;
+      OMPTraitInfo &TI = A->getTraitInfo();
+      TI.getAsVariantMatchInfo(Context, VMI, /* DeviceSetOnly */ false);
+      if (!isVariantApplicableInContext(VMI, OMPCtx))
+        continue;
+
+      VMIs.push_back(VMI);
+      Exprs.push_back(VariantRef);
+    }
+
+    CalleeFnDecl = CalleeFnDecl->getPreviousDecl();
+  }
+
+  ExprResult NewCall;
+  do {
+    int BestIdx = getBestVariantMatchForContext(VMIs, OMPCtx);
+    if (BestIdx < 0)
+      return Call;
+    Expr *BestExpr = cast<DeclRefExpr>(Exprs[BestIdx]);
+    Decl *BestDecl = cast<DeclRefExpr>(BestExpr)->getDecl();
+
+    {
+      // Try to build a (member) call expression for the current best applicable
+      // variant expression. We allow this to fail in which case we continue
+      // with the next best variant expression. The fail case is part of the
+      // implementation defined behavior in the OpenMP standard when it talks
+      // about what differences in the function prototypes: "Any differences
+      // that the specific OpenMP context requires in the prototype of the
+      // variant from the base function prototype are implementation defined."
+      // This wording is there to allow the specialized variant to have a
+      // different type than the base function. This is intended and OK but if
+      // we cannot create a call the difference is not in the "implementation
+      // defined range" we allow.
+      Sema::TentativeAnalysisScope Trap(*this);
+
+      if (auto *SpecializedMethod = dyn_cast<CXXMethodDecl>(BestDecl)) {
+        auto *MemberCall = dyn_cast<CXXMemberCallExpr>(CE);
+        BestExpr = MemberExpr::CreateImplicit(
+            S.Context, MemberCall->getImplicitObjectArgument(),
+            /* IsArrow */ false, SpecializedMethod, S.Context.BoundMemberTy,
+            MemberCall->getValueKind(), MemberCall->getObjectKind());
+      }
+      NewCall = S.BuildCallExpr(Scope, BestExpr, LParenLoc, ArgExprs, RParenLoc,
+                                ExecConfig);
+      if (NewCall.isUsable())
+        break;
+    }
+
+    VMIs.erase(VMIs.begin() + BestIdx);
+    Exprs.erase(Exprs.begin() + BestIdx);
+  } while (!VMIs.empty());
+
+  if (!NewCall.isUsable())
+    return Call;
+
+  return PseudoObjectExpr::Create(Context, CE, {NewCall.get()}, 0);
+}
+
 Optional<std::pair<FunctionDecl *, Expr *>>
 Sema::checkOpenMPDeclareVariantFunction(Sema::DeclGroupPtrTy DG,
                                         Expr *VariantRef, OMPTraitInfo &TI,
                                         SourceRange SR) {
   if (!DG || DG.get().isNull())
     return None;
 
   const int VariantId = 1;
void Sema::ActOnOpenMPDeclareVariantDirective(FunctionDecl *FD,
                                               Expr *VariantRef,
                                               OMPTraitInfo &TI,
                                               SourceRange SR) {
   auto *NewAttr =
       OMPDeclareVariantAttr::CreateImplicit(Context, VariantRef, &TI, SR);
   FD->addAttr(NewAttr);
 }
 
-void Sema::markOpenMPDeclareVariantFuncsReferenced(SourceLocation Loc,
-                                                   FunctionDecl *Func,
-                                                   bool MightBeOdrUse) {
-  assert(LangOpts.OpenMP && "Expected OpenMP mode.");
-
-  if (!Func->isDependentContext() && Func->hasAttrs()) {
-    for (OMPDeclareVariantAttr *A :
-         Func->specific_attrs<OMPDeclareVariantAttr>()) {
-      // TODO: add checks for active OpenMP context where possible.
-      Expr *VariantRef = A->getVariantFuncRef();
-      auto *DRE = cast<DeclRefExpr>(VariantRef->IgnoreParenImpCasts());
-      auto *F = cast<FunctionDecl>(DRE->getDecl());
-      if (!F->isDefined() && F->isTemplateInstantiation())
-        InstantiateFunctionDefinition(Loc, F->getFirstDecl());
-      MarkFunctionReferenced(Loc, F, MightBeOdrUse);
-    }
-  }
-}
-
 StmtResult Sema::ActOnOpenMPParallelDirective(ArrayRef<OMPClause *> Clauses,
                                               Stmt *AStmt,
                                               SourceLocation StartLoc,
                                               SourceLocation EndLoc) {
   if (!AStmt)
     return StmtError();
 
   auto *CS = cast<CapturedStmt>(AStmt);

clang/lib/Sema/SemaPseudoObject.cpp

return new (Context) CompoundAssignOperator(lhs, rhs, cop->getOpcode(),
                                                 FPOptions());
   } else if (BinaryOperator *bop = dyn_cast<BinaryOperator>(syntax)) {
     Expr *lhs = stripOpaqueValuesFromPseudoObjectRef(*this, bop->getLHS());
     Expr *rhs = cast<OpaqueValueExpr>(bop->getRHS())->getSourceExpr();
     return new (Context) BinaryOperator(lhs, rhs, bop->getOpcode(),
                                         bop->getType(), bop->getValueKind(),
                                         bop->getObjectKind(),
                                         bop->getOperatorLoc(), FPOptions());
+  } else if (isa<CallExpr>(syntax)) {
+    return syntax;
   } else {
     assert(syntax->hasPlaceholderType(BuiltinType::PseudoObject));
     return stripOpaqueValuesFromPseudoObjectRef(*this, syntax);
   }
 }

clang/lib/Sema/SemaStmt.cpp

if (const ObjCMessageExpr *ME = dyn_cast<ObjCMessageExpr>(E)) {
     const ObjCMethodDecl *MD = ME->getMethodDecl();
     if (MD) {
       if (DiagnoseNoDiscard(*this, MD->getAttr<WarnUnusedResultAttr>(), Loc, R1,
                             R2, /*isCtor=*/false))
         return;
     }
   } else if (const PseudoObjectExpr *POE = dyn_cast<PseudoObjectExpr>(E)) {
     const Expr *Source = POE->getSyntacticForm();
+    // Handle the actually selected call of an OpenMP specialized call.
+    if (LangOpts.OpenMP && isa<CallExpr>(Source) &&
+        POE->getNumSemanticExprs() == 1 &&
+        isa<CallExpr>(POE->getSemanticExpr(0)))
+      return DiagnoseUnusedExprResult(POE->getSemanticExpr(0));
     if (isa<ObjCSubscriptRefExpr>(Source))
       DiagID = diag::warn_unused_container_subscript_expr;
     else
       DiagID = diag::warn_unused_property_expr;
   } else if (const CXXFunctionalCastExpr *FC
                                        = dyn_cast<CXXFunctionalCastExpr>(E)) {
     const Expr *E = FC->getSubExpr();
     if (const CXXBindTemporaryExpr *TE = dyn_cast<CXXBindTemporaryExpr>(E))

clang/test/AST/ast-dump-openmp-begin-declare-variant.c

-// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -ast-dump %s | FileCheck %s
-
-#pragma omp begin declare variant match(device={kind(cpu)})
-int also_before(void) {
-  return 0;
-}
-#pragma omp end declare variant
-
-#pragma omp begin declare variant match(device={kind(gpu)})
-int also_after(void) {
-  return 2;
-}
-int also_before(void) {
-  return 2;
-}
-#pragma omp end declare variant
-
-
-#pragma omp begin declare variant match(device={kind(fpga)})
-
-This text is never parsed!
-
-#pragma omp end declare variant
-
-int also_after(void) {
-  return 0;
-}
-
-int test() {
-  return also_after() + also_before();
-}
-
-// Make sure:
-//  - we do see the ast nodes for the cpu kind
-//  - we do not see the ast nodes for the gpu kind
-//  - we do not choke on the text in the kind(fpga) guarded scope.
-
-// CHECK:       -FunctionDecl {{.*}} <{{.*}}4:1, line:{{.*}}:1> line:{{.*}}:5 used also_before 'int (void)'
-// CHECK-NEXT:   -CompoundStmt {{.*}} <col:23, line:{{.*}}:1>
-// CHECK-NEXT:    `-ReturnStmt {{.*}} <line:{{.*}}:3, col:10>
-// CHECK-NEXT:      `-IntegerLiteral {{.*}} <col:10> 'int' 0
-// CHECK-NEXT:  -FunctionDecl {{.*}} <line:{{.*}}:1, line:{{.*}}:1> line:{{.*}}:5 used also_after 'int (void)'
-// CHECK-NEXT:   -CompoundStmt {{.*}} <col:22, line:{{.*}}:1>
-// CHECK-NEXT:    `-ReturnStmt {{.*}} <line:{{.*}}:3, col:10>
-// CHECK-NEXT:      `-IntegerLiteral {{.*}} <col:10> 'int' 0
-// CHECK-NEXT:  `-FunctionDecl {{.*}} <line:{{.*}}:1, line:{{.*}}:1> line:{{.*}}:5 test 'int ()'
-// CHECK-NEXT:    `-CompoundStmt {{.*}} <col:12, line:{{.*}}:1>
-// CHECK-NEXT:      `-ReturnStmt {{.*}} <line:{{.*}}:3, col:37>
-// CHECK-NEXT:        `-BinaryOperator {{.*}} <col:10, col:37> 'int' '+'
-// CHECK-NEXT:          |-CallExpr {{.*}} <col:10, col:21> 'int'
-// CHECK-NEXT:          | `-ImplicitCastExpr {{.*}} <col:10> 'int (*)(void)' <FunctionToPointerDecay>
-// CHECK-NEXT:          |   `-DeclRefExpr {{.*}} <col:10> 'int (void)' Function {{.*}} 'also_after' 'int (void)'
-// CHECK-NEXT:          `-CallExpr {{.*}} <col:25, col:37> 'int'
-// CHECK-NEXT:            `-ImplicitCastExpr {{.*}} <col:25> 'int (*)(void)' <FunctionToPointerDecay>
-// CHECK-NEXT:              `-DeclRefExpr {{.*}} <col:25> 'int (void)' Function {{.*}} 'also_before' 'int (void)'

clang/test/AST/ast-dump-openmp-begin-declare-variant_1.c

+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s       | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s -x c++| FileCheck %s
+// expected-no-diagnostics
+
+int also_before(void) {
+  return 0;
+}
+
+#pragma omp begin declare variant match(device={kind(gpu)})
+int also_after(void) {
+  return 1;
+}
+int also_before(void) {
+  return 2;
+}
+#pragma omp end declare variant
+
+
+#pragma omp begin declare variant match(device={kind(fpga)})
+
+This text is never parsed!
+
+#pragma omp end declare variant
+
+int also_after(void) {
+  return 0;
+}
+
+int test() {
+  // Should return 0.
+  return also_after() + also_before();
+}
+
+// Make sure:
+//  - we do not see the ast nodes for the gpu kind
+//  - we do not choke on the text in the kind(fpga) guarded scopes
+
+// CHECK:      |-FunctionDecl [[ADDR_0:0x[a-z0-9]*]] <{{.*}}, line:7:1> line:5:5 used also_before 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_1:0x[a-z0-9]*]] <col:23, line:7:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_2:0x[a-z0-9]*]] <line:6:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_3:0x[a-z0-9]*]] <col:10> 'int' 0
+// CHECK-NEXT: |-FunctionDecl [[ADDR_4:0x[a-z0-9]*]] <line:25:1, line:27:1> line:25:5 used also_after 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_5:0x[a-z0-9]*]] <col:22, line:27:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_6:0x[a-z0-9]*]] <line:26:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_7:0x[a-z0-9]*]] <col:10> 'int' 0
+// CHECK-NEXT: `-FunctionDecl [[ADDR_8:0x[a-z0-9]*]] <line:29:1, line:32:1> line:29:5 test 'int ({{.*}})'
+// CHECK-NEXT:   `-CompoundStmt [[ADDR_9:0x[a-z0-9]*]] <col:12, line:32:1>
+// CHECK-NEXT:     `-ReturnStmt [[ADDR_10:0x[a-z0-9]*]] <line:31:3, col:37>
+// CHECK-NEXT:       `-BinaryOperator [[ADDR_11:0x[a-z0-9]*]] <col:10, col:37> 'int' '+'
+// CHECK-NEXT:         |-CallExpr [[ADDR_12:0x[a-z0-9]*]] <col:10, col:21> 'int'
+// CHECK-NEXT:         | `-ImplicitCastExpr [[ADDR_13:0x[a-z0-9]*]] <col:10> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:         |   `-DeclRefExpr [[ADDR_14:0x[a-z0-9]*]] <col:10> 'int ({{.*}})' {{.*}}Function [[ADDR_4]] 'also_after' 'int ({{.*}})'
+// CHECK-NEXT:         `-CallExpr [[ADDR_15:0x[a-z0-9]*]] <col:25, col:37> 'int'
+// CHECK-NEXT:           `-ImplicitCastExpr [[ADDR_16:0x[a-z0-9]*]] <col:25> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:             `-DeclRefExpr [[ADDR_17:0x[a-z0-9]*]] <col:25> 'int ({{.*}})' {{.*}}Function [[ADDR_0]] 'also_before' 'int ({{.*}})'

clang/test/AST/ast-dump-openmp-begin-declare-variant_2.c

+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s       | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s -x c++| FileCheck %s
+// expected-no-diagnostics
+
+#pragma omp begin declare variant match(device={kind(cpu)})
+int also_before(void) {
+  return 0;
+}
+#pragma omp end declare variant
+
+#pragma omp begin declare variant match(implementation={vendor(score(100):llvm)})
+int also_after(void) {
+  return 0;
+}
+#pragma omp end declare variant
+#pragma omp begin declare variant match(implementation={vendor(score(0):llvm)})
+int also_before(void) {
+  return 1;
+}
+#pragma omp end declare variant
+
+int also_after(void) {
+  return 2;
+}
+
+int test() {
+  // Should return 0.
+  return also_after() + also_before();
+}
+
+// Make sure:
+//  - we do see the ast nodes for the cpu kind
+//  - we do see the ast nodes for the llvm vendor
+//  - we pick the right callees
+
+// CHECK:      |-FunctionDecl [[ADDR_0:0x[a-z0-9]*]] <{{.*}}, col:21> col:5 implicit used also_before 'int ({{.*}})'
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_1:0x[a-z0-9]*]] <<invalid sloc>> Implicit device={kind(cpu)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_2:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_3:0x[a-z0-9]*]] 'also_before[device={kind(cpu)}]' 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_3]] <col:1, line:8:1> line:6:1 also_before[device={kind(cpu)}] 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_4:0x[a-z0-9]*]] <col:23, line:8:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_5:0x[a-z0-9]*]] <line:7:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_6:0x[a-z0-9]*]] <col:10> 'int' 0
+// CHECK-NEXT: |-FunctionDecl [[ADDR_7:0x[a-z0-9]*]] <line:12:1, col:20> col:5 implicit used also_after 'int ({{.*}})'
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_8:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(score(100): llvm)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_9:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_10:0x[a-z0-9]*]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_10]] <col:1, line:14:1> line:12:1 also_after[implementation={vendor(llvm)}] 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_11:0x[a-z0-9]*]] <col:22, line:14:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_12:0x[a-z0-9]*]] <line:13:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_13:0x[a-z0-9]*]] <col:10> 'int' 0
+// CHECK-NEXT: |-FunctionDecl [[ADDR_14:0x[a-z0-9]*]] prev [[ADDR_0]] <line:17:1, col:21> col:5 implicit used also_before 'int ({{.*}})'
+// CHECK-NEXT: | |-OMPDeclareVariantAttr [[ADDR_15:0x[a-z0-9]*]] <<invalid sloc>> Inherited Implicit device={kind(cpu)}
+// CHECK-NEXT: | | `-DeclRefExpr [[ADDR_2]] <line:6:1> 'int ({{.*}})' Function [[ADDR_3]] 'also_before[device={kind(cpu)}]' 'int ({{.*}})'
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_16:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(score(0): llvm)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_17:0x[a-z0-9]*]] <line:17:1> 'int ({{.*}})' Function [[ADDR_18:0x[a-z0-9]*]] 'also_before[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_18]] <col:1, line:19:1> line:17:1 also_before[implementation={vendor(llvm)}] 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_19:0x[a-z0-9]*]] <col:23, line:19:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_20:0x[a-z0-9]*]] <line:18:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_21:0x[a-z0-9]*]] <col:10> 'int' 1
+// CHECK-NEXT: |-FunctionDecl [[ADDR_22:0x[a-z0-9]*]] prev [[ADDR_7]] <line:22:1, line:24:1> line:22:5 used also_after 'int ({{.*}})'
+// CHECK-NEXT: | |-CompoundStmt [[ADDR_23:0x[a-z0-9]*]] <col:22, line:24:1>
+// CHECK-NEXT: | | `-ReturnStmt [[ADDR_24:0x[a-z0-9]*]] <line:23:3, col:10>
+// CHECK-NEXT: | |   `-IntegerLiteral [[ADDR_25:0x[a-z0-9]*]] <col:10> 'int' 2
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_26:0x[a-z0-9]*]] <<invalid sloc>> Inherited Implicit implementation={vendor(score(100): llvm)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_9]] <line:12:1> 'int ({{.*}})' Function [[ADDR_10]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CHECK-NEXT: `-FunctionDecl [[ADDR_27:0x[a-z0-9]*]] <line:26:1, line:29:1> line:26:5 test 'int ({{.*}})'
+// CHECK-NEXT:   `-CompoundStmt [[ADDR_28:0x[a-z0-9]*]] <col:12, line:29:1>
+// CHECK-NEXT:     `-ReturnStmt [[ADDR_29:0x[a-z0-9]*]] <line:28:3, col:37>
+// CHECK-NEXT:       `-BinaryOperator [[ADDR_30:0x[a-z0-9]*]] <col:10, col:37> 'int' '+'
+// CHECK-NEXT:         |-PseudoObjectExpr [[ADDR_31:0x[a-z0-9]*]] <col:10, col:21> 'int'
+// CHECK-NEXT:         | |-CallExpr [[ADDR_32:0x[a-z0-9]*]] <col:10, col:21> 'int'
+// CHECK-NEXT:         | | `-ImplicitCastExpr [[ADDR_33:0x[a-z0-9]*]] <col:10> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:         | |   `-DeclRefExpr [[ADDR_34:0x[a-z0-9]*]] <col:10> 'int ({{.*}})' {{.*}}Function [[ADDR_22]] 'also_after' 'int ({{.*}})'
+// CHECK-NEXT:         | `-CallExpr [[ADDR_35:0x[a-z0-9]*]] <line:12:1, line:28:21> 'int'
+// CHECK-NEXT:         |   `-ImplicitCastExpr [[ADDR_36:0x[a-z0-9]*]] <line:12:1> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:         |     `-DeclRefExpr [[ADDR_9]] <col:1> 'int ({{.*}})' Function [[ADDR_10]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CHECK-NEXT:         `-PseudoObjectExpr [[ADDR_37:0x[a-z0-9]*]] <line:28:25, col:37> 'int'
+// CHECK-NEXT:           |-CallExpr [[ADDR_38:0x[a-z0-9]*]] <col:25, col:37> 'int'
+// CHECK-NEXT:           | `-ImplicitCastExpr [[ADDR_39:0x[a-z0-9]*]] <col:25> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:           |   `-DeclRefExpr [[ADDR_40:0x[a-z0-9]*]] <col:25> 'int ({{.*}})' {{.*}}Function [[ADDR_14]] 'also_before' 'int ({{.*}})'
+// CHECK-NEXT:           `-CallExpr [[ADDR_41:0x[a-z0-9]*]] <line:6:1, line:28:37> 'int'
+// CHECK-NEXT:             `-ImplicitCastExpr [[ADDR_42:0x[a-z0-9]*]] <line:6:1> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:               `-DeclRefExpr [[ADDR_2]] <col:1> 'int ({{.*}})' Function [[ADDR_3]] 'also_before[device={kind(cpu)}]' 'int ({{.*}})'

clang/test/AST/ast-dump-openmp-begin-declare-variant_3.c

+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s       | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s -x c++| FileCheck %s
+// expected-no-diagnostics
+
+#pragma omp begin declare variant match(device={kind(cpu)})
+int also_before(void) {
+  return 1;
+}
+#pragma omp end declare variant
+
+#pragma omp begin declare variant match(implementation={vendor(score(0):llvm)})
+int also_after(void) {
+  return 0;
+}
+#pragma omp end declare variant
+#pragma omp begin declare variant match(implementation={vendor(score(100):llvm)})
+int also_before(void) {
+  return 0;
+}
+#pragma omp end declare variant
+
+int also_after(void) {
+  return 2;
+}
+
+int test() {
+  // Should return 0.
+  return also_after() + also_before();
+}
+
+// Make sure:
+//  - we do see the ast nodes for the cpu kind
+//  - we do see the ast nodes for the llvm vendor
+//  - we pick the right callees
+
+// CHECK:      |-FunctionDecl [[ADDR_0:0x[a-z0-9]*]] <{{.*}}, col:21> col:5 implicit used also_before 'int ({{.*}})'
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_1:0x[a-z0-9]*]] <<invalid sloc>> Implicit device={kind(cpu)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_2:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_3:0x[a-z0-9]*]] 'also_before[device={kind(cpu)}]' 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_3]] <col:1, line:8:1> line:6:1 also_before[device={kind(cpu)}] 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_4:0x[a-z0-9]*]] <col:23, line:8:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_5:0x[a-z0-9]*]] <line:7:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_6:0x[a-z0-9]*]] <col:10> 'int' 1
+// CHECK-NEXT: |-FunctionDecl [[ADDR_7:0x[a-z0-9]*]] <line:12:1, col:20> col:5 implicit used also_after 'int ({{.*}})'
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_8:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(score(0): llvm)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_9:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_10:0x[a-z0-9]*]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_10]] <col:1, line:14:1> line:12:1 also_after[implementation={vendor(llvm)}] 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_11:0x[a-z0-9]*]] <col:22, line:14:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_12:0x[a-z0-9]*]] <line:13:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_13:0x[a-z0-9]*]] <col:10> 'int' 0
+// CHECK-NEXT: |-FunctionDecl [[ADDR_14:0x[a-z0-9]*]] prev [[ADDR_0]] <line:17:1, col:21> col:5 implicit used also_before 'int ({{.*}})'
+// CHECK-NEXT: | |-OMPDeclareVariantAttr [[ADDR_15:0x[a-z0-9]*]] <<invalid sloc>> Inherited Implicit device={kind(cpu)}
+// CHECK-NEXT: | | `-DeclRefExpr [[ADDR_2]] <line:6:1> 'int ({{.*}})' Function [[ADDR_3]] 'also_before[device={kind(cpu)}]' 'int ({{.*}})'
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_16:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(score(100): llvm)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_17:0x[a-z0-9]*]] <line:17:1> 'int ({{.*}})' Function [[ADDR_18:0x[a-z0-9]*]] 'also_before[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_18]] <col:1, line:19:1> line:17:1 also_before[implementation={vendor(llvm)}] 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_19:0x[a-z0-9]*]] <col:23, line:19:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_20:0x[a-z0-9]*]] <line:18:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_21:0x[a-z0-9]*]] <col:10> 'int' 0
+// CHECK-NEXT: |-FunctionDecl [[ADDR_22:0x[a-z0-9]*]] prev [[ADDR_7]] <line:22:1, line:24:1> line:22:5 used also_after 'int ({{.*}})'
+// CHECK-NEXT: | |-CompoundStmt [[ADDR_23:0x[a-z0-9]*]] <col:22, line:24:1>
+// CHECK-NEXT: | | `-ReturnStmt [[ADDR_24:0x[a-z0-9]*]] <line:23:3, col:10>
+// CHECK-NEXT: | |   `-IntegerLiteral [[ADDR_25:0x[a-z0-9]*]] <col:10> 'int' 2
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_26:0x[a-z0-9]*]] <<invalid sloc>> Inherited Implicit implementation={vendor(score(0): llvm)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_9]] <line:12:1> 'int ({{.*}})' Function [[ADDR_10]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CHECK-NEXT: `-FunctionDecl [[ADDR_27:0x[a-z0-9]*]] <line:26:1, line:29:1> line:26:5 test 'int ({{.*}})'
+// CHECK-NEXT:   `-CompoundStmt [[ADDR_28:0x[a-z0-9]*]] <col:12, line:29:1>
+// CHECK-NEXT:     `-ReturnStmt [[ADDR_29:0x[a-z0-9]*]] <line:28:3, col:37>
+// CHECK-NEXT:       `-BinaryOperator [[ADDR_30:0x[a-z0-9]*]] <col:10, col:37> 'int' '+'
+// CHECK-NEXT:         |-PseudoObjectExpr [[ADDR_31:0x[a-z0-9]*]] <col:10, col:21> 'int'
+// CHECK-NEXT:         | |-CallExpr [[ADDR_32:0x[a-z0-9]*]] <col:10, col:21> 'int'
+// CHECK-NEXT:         | | `-ImplicitCastExpr [[ADDR_33:0x[a-z0-9]*]] <col:10> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:         | |   `-DeclRefExpr [[ADDR_34:0x[a-z0-9]*]] <col:10> 'int ({{.*}})' {{.*}}Function [[ADDR_22]] 'also_after' 'int ({{.*}})'
+// CHECK-NEXT:         | `-CallExpr [[ADDR_35:0x[a-z0-9]*]] <line:12:1, line:28:21> 'int'
+// CHECK-NEXT:         |   `-ImplicitCastExpr [[ADDR_36:0x[a-z0-9]*]] <line:12:1> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:         |     `-DeclRefExpr [[ADDR_9]] <col:1> 'int ({{.*}})' Function [[ADDR_10]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CHECK-NEXT:         `-PseudoObjectExpr [[ADDR_37:0x[a-z0-9]*]] <line:28:25, col:37> 'int'
+// CHECK-NEXT:           |-CallExpr [[ADDR_38:0x[a-z0-9]*]] <col:25, col:37> 'int'
+// CHECK-NEXT:           | `-ImplicitCastExpr [[ADDR_39:0x[a-z0-9]*]] <col:25> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:           |   `-DeclRefExpr [[ADDR_40:0x[a-z0-9]*]] <col:25> 'int ({{.*}})' {{.*}}Function [[ADDR_14]] 'also_before' 'int ({{.*}})'
+// CHECK-NEXT:           `-CallExpr [[ADDR_41:0x[a-z0-9]*]] <line:17:1, line:28:37> 'int'
+// CHECK-NEXT:             `-ImplicitCastExpr [[ADDR_42:0x[a-z0-9]*]] <line:17:1> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:               `-DeclRefExpr [[ADDR_17]] <col:1> 'int ({{.*}})' Function [[ADDR_18]] 'also_before[implementation={vendor(llvm)}]' 'int ({{.*}})'

clang/test/AST/ast-dump-openmp-begin-declare-variant_4.c

+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s       | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s -x c++| FileCheck %s
+// expected-no-diagnostics
+
+#pragma omp begin declare variant match(device={kind(cpu)})
+int also_before(void) {
+  return 0;
+}
+#pragma omp end declare variant
+
+int also_after(void) {
+  return 0;
+}
+
+int test() {
+  // Should return 0.
+  return also_after() + also_before();
+}
+
+// Make sure:
+//  - we do see the ast nodes for the cpu kind
+//  - we pick the right callees
+
+// CHECK:      |-FunctionDecl [[ADDR_0:0x[a-z0-9]*]] <{{.*}}, col:21> col:5 implicit used also_before 'int ({{.*}})'
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_1:0x[a-z0-9]*]] <<invalid sloc>> Implicit device={kind(cpu)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_2:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_3:0x[a-z0-9]*]] 'also_before[device={kind(cpu)}]' 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_3]] <col:1, line:8:1> line:6:1 also_before[device={kind(cpu)}] 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_4:0x[a-z0-9]*]] <col:23, line:8:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_5:0x[a-z0-9]*]] <line:7:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_6:0x[a-z0-9]*]] <col:10> 'int' 0
+// CHECK-NEXT: |-FunctionDecl [[ADDR_7:0x[a-z0-9]*]] <line:11:1, line:13:1> line:11:5 used also_after 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_8:0x[a-z0-9]*]] <col:22, line:13:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_9:0x[a-z0-9]*]] <line:12:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_10:0x[a-z0-9]*]] <col:10> 'int' 0
+// CHECK-NEXT: `-FunctionDecl [[ADDR_11:0x[a-z0-9]*]] <line:15:1, line:18:1> line:15:5 test 'int ({{.*}})'
+// CHECK-NEXT:   `-CompoundStmt [[ADDR_12:0x[a-z0-9]*]] <col:12, line:18:1>
+// CHECK-NEXT:     `-ReturnStmt [[ADDR_13:0x[a-z0-9]*]] <line:17:3, col:37>
+// CHECK-NEXT:       `-BinaryOperator [[ADDR_14:0x[a-z0-9]*]] <col:10, col:37> 'int' '+'
+// CHECK-NEXT:         |-CallExpr [[ADDR_15:0x[a-z0-9]*]] <col:10, col:21> 'int'
+// CHECK-NEXT:         | `-ImplicitCastExpr [[ADDR_16:0x[a-z0-9]*]] <col:10> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:         |   `-DeclRefExpr [[ADDR_17:0x[a-z0-9]*]] <col:10> 'int ({{.*}})' {{.*}}Function [[ADDR_7]] 'also_after' 'int ({{.*}})'
+// CHECK-NEXT:         `-PseudoObjectExpr [[ADDR_18:0x[a-z0-9]*]] <col:25, col:37> 'int'
+// CHECK-NEXT:           |-CallExpr [[ADDR_19:0x[a-z0-9]*]] <col:25, col:37> 'int'
+// CHECK-NEXT:           | `-ImplicitCastExpr [[ADDR_20:0x[a-z0-9]*]] <col:25> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:           |   `-DeclRefExpr [[ADDR_21:0x[a-z0-9]*]] <col:25> 'int ({{.*}})' {{.*}}Function [[ADDR_0]] 'also_before' 'int ({{.*}})'
+// CHECK-NEXT:           `-CallExpr [[ADDR_22:0x[a-z0-9]*]] <line:6:1, line:17:37> 'int'
+// CHECK-NEXT:             `-ImplicitCastExpr [[ADDR_23:0x[a-z0-9]*]] <line:6:1> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:               `-DeclRefExpr [[ADDR_2]] <col:1> 'int ({{.*}})' Function [[ADDR_3]] 'also_before[device={kind(cpu)}]' 'int ({{.*}})'

clang/test/AST/ast-dump-openmp-begin-declare-variant_5.c

+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s       | FileCheck %s --check-prefix=C
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s -x c++| FileCheck %s --check-prefix=CXX
+// expected-no-diagnostics
+
+int also_before(void) {
+  return 1;
+}
+
+#pragma omp begin declare variant match(implementation={vendor(llvm)})
+int also_after(void) {
+  return 0;
+}
+int also_before(void) {
+  return 0;
+}
+#pragma omp end declare variant
+
+int also_after(void) {
+  return 2;
+}
+
+int main() {
+  // Should return 0.
+  return (also_after)() +
+         (also_before)() +
+         (&also_after)() +
+         (&also_before)();
+}
+
+// Make sure:
+//  - we see the specialization in the AST
+//  - we pick the right callees
+
+// C:      |-FunctionDecl [[ADDR_0:0x[a-z0-9]*]] <{{.*}}, line:7:1> line:5:5 used also_before 'int ({{.*}})'
+// C-NEXT: | `-CompoundStmt [[ADDR_1:0x[a-z0-9]*]] <col:23, line:7:1>
+// C-NEXT: |   `-ReturnStmt [[ADDR_2:0x[a-z0-9]*]] <line:6:3, col:10>
+// C-NEXT: |     `-IntegerLiteral [[ADDR_3:0x[a-z0-9]*]] <col:10> 'int' 1
+// C-NEXT: |-FunctionDecl [[ADDR_4:0x[a-z0-9]*]] <line:10:1, col:20> col:5 implicit used also_after 'int ({{.*}})'
+// C-NEXT: | `-OMPDeclareVariantAttr [[ADDR_5:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(llvm)}
+// C-NEXT: |   `-DeclRefExpr [[ADDR_6:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_7:0x[a-z0-9]*]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// C-NEXT: |-FunctionDecl [[ADDR_7]] <col:1, line:12:1> line:10:1 also_after[implementation={vendor(llvm)}] 'int ({{.*}})'
+// C-NEXT: | `-CompoundStmt [[ADDR_8:0x[a-z0-9]*]] <col:22, line:12:1>
+// C-NEXT: |   `-ReturnStmt [[ADDR_9:0x[a-z0-9]*]] <line:11:3, col:10>
+// C-NEXT: |     `-IntegerLiteral [[ADDR_10:0x[a-z0-9]*]] <col:10> 'int' 0
+// C-NEXT: |-FunctionDecl [[ADDR_11:0x[a-z0-9]*]] prev [[ADDR_0]] <line:13:1, col:21> col:5 implicit used also_before 'int ({{.*}})'
+// C-NEXT: | `-OMPDeclareVariantAttr [[ADDR_12:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(llvm)}
+// C-NEXT: |   `-DeclRefExpr [[ADDR_13:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_14:0x[a-z0-9]*]] 'also_before[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// C-NEXT: |-FunctionDecl [[ADDR_14]] <col:1, line:15:1> line:13:1 also_before[implementation={vendor(llvm)}] 'int ({{.*}})'
+// C-NEXT: | `-CompoundStmt [[ADDR_15:0x[a-z0-9]*]] <col:23, line:15:1>
+// C-NEXT: |   `-ReturnStmt [[ADDR_16:0x[a-z0-9]*]] <line:14:3, col:10>
+// C-NEXT: |     `-IntegerLiteral [[ADDR_17:0x[a-z0-9]*]] <col:10> 'int' 0
+// C-NEXT: |-FunctionDecl [[ADDR_18:0x[a-z0-9]*]] prev [[ADDR_4]] <line:18:1, line:20:1> line:18:5 used also_after 'int ({{.*}})'
+// C-NEXT: | |-CompoundStmt [[ADDR_19:0x[a-z0-9]*]] <col:22, line:20:1>
+// C-NEXT: | | `-ReturnStmt [[ADDR_20:0x[a-z0-9]*]] <line:19:3, col:10>
+// C-NEXT: | |   `-IntegerLiteral [[ADDR_21:0x[a-z0-9]*]] <col:10> 'int' 2
+// C-NEXT: | `-OMPDeclareVariantAttr [[ADDR_22:0x[a-z0-9]*]] <<invalid sloc>> Inherited Implicit implementation={vendor(llvm)}
+// C-NEXT: |   `-DeclRefExpr [[ADDR_6]] <line:10:1> 'int ({{.*}})' Function [[ADDR_7]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// C-NEXT: `-FunctionDecl [[ADDR_23:0x[a-z0-9]*]] <line:22:1, line:28:1> line:22:5 main 'int ({{.*}})'
+// C-NEXT:   `-CompoundStmt [[ADDR_24:0x[a-z0-9]*]] <col:12, line:28:1>
+// C-NEXT:     `-ReturnStmt [[ADDR_25:0x[a-z0-9]*]] <line:24:3, line:27:25>
+// C-NEXT:       `-BinaryOperator [[ADDR_26:0x[a-z0-9]*]] <line:24:10, line:27:25> 'int' '+'
+// C-NEXT:         |-BinaryOperator [[ADDR_27:0x[a-z0-9]*]] <line:24:10, line:26:24> 'int' '+'
+// C-NEXT:         | |-BinaryOperator [[ADDR_28:0x[a-z0-9]*]] <line:24:10, line:25:24> 'int' '+'
+// C-NEXT:         | | |-PseudoObjectExpr [[ADDR_29:0x[a-z0-9]*]] <line:24:10, col:23> 'int'
+// C-NEXT:         | | | |-CallExpr [[ADDR_30:0x[a-z0-9]*]] <col:10, col:23> 'int'
+// C-NEXT:         | | | | `-ImplicitCastExpr [[ADDR_31:0x[a-z0-9]*]] <col:10, col:21> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// C-NEXT:         | | | |   `-ParenExpr [[ADDR_32:0x[a-z0-9]*]] <col:10, col:21> 'int ({{.*}})'
+// C-NEXT:         | | | |     `-DeclRefExpr [[ADDR_33:0x[a-z0-9]*]] <col:11> 'int ({{.*}})' Function [[ADDR_18]] 'also_after' 'int ({{.*}})'
+// C-NEXT:         | | | `-CallExpr [[ADDR_34:0x[a-z0-9]*]] <line:10:1, line:24:23> 'int'
+// C-NEXT:         | | |   `-ImplicitCastExpr [[ADDR_35:0x[a-z0-9]*]] <line:10:1> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// C-NEXT:         | | |     `-DeclRefExpr [[ADDR_6]] <col:1> 'int ({{.*}})' Function [[ADDR_7]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// C-NEXT:         | | `-PseudoObjectExpr [[ADDR_36:0x[a-z0-9]*]] <line:25:10, col:24> 'int'
+// C-NEXT:         | |   |-CallExpr [[ADDR_37:0x[a-z0-9]*]] <col:10, col:24> 'int'
+// C-NEXT:         | |   | `-ImplicitCastExpr [[ADDR_38:0x[a-z0-9]*]] <col:10, col:22> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// C-NEXT:         | |   |   `-ParenExpr [[ADDR_39:0x[a-z0-9]*]] <col:10, col:22> 'int ({{.*}})'
+// C-NEXT:         | |   |     `-DeclRefExpr [[ADDR_40:0x[a-z0-9]*]] <col:11> 'int ({{.*}})' Function [[ADDR_11]] 'also_before' 'int ({{.*}})'
+// C-NEXT:         | |   `-CallExpr [[ADDR_41:0x[a-z0-9]*]] <line:13:1, line:25:24> 'int'
+// C-NEXT:         | |     `-ImplicitCastExpr [[ADDR_42:0x[a-z0-9]*]] <line:13:1> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// C-NEXT:         | |       `-DeclRefExpr [[ADDR_13]] <col:1> 'int ({{.*}})' Function [[ADDR_14]] 'also_before[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// C-NEXT:         | `-PseudoObjectExpr [[ADDR_43:0x[a-z0-9]*]] <line:26:10, col:24> 'int'
+// C-NEXT:         |   |-CallExpr [[ADDR_44:0x[a-z0-9]*]] <col:10, col:24> 'int'
+// C-NEXT:         |   | `-ParenExpr [[ADDR_45:0x[a-z0-9]*]] <col:10, col:22> 'int (*)({{.*}})'
+// C-NEXT:         |   |   `-UnaryOperator [[ADDR_46:0x[a-z0-9]*]] <col:11, col:12> 'int (*)({{.*}})' prefix '&' cannot overflow
+// C-NEXT:         |   |     `-DeclRefExpr [[ADDR_47:0x[a-z0-9]*]] <col:12> 'int ({{.*}})' Function [[ADDR_18]] 'also_after' 'int ({{.*}})'
+// C-NEXT:         |   `-CallExpr [[ADDR_48:0x[a-z0-9]*]] <line:10:1, line:26:24> 'int'
+// C-NEXT:         |     `-ImplicitCastExpr [[ADDR_49:0x[a-z0-9]*]] <line:10:1> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// C-NEXT:         |       `-DeclRefExpr [[ADDR_6]] <col:1> 'int ({{.*}})' Function [[ADDR_7]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// C-NEXT:         `-PseudoObjectExpr [[ADDR_50:0x[a-z0-9]*]] <line:27:10, col:25> 'int'
+// C-NEXT:           |-CallExpr [[ADDR_51:0x[a-z0-9]*]] <col:10, col:25> 'int'
+// C-NEXT:           | `-ParenExpr [[ADDR_52:0x[a-z0-9]*]] <col:10, col:23> 'int (*)({{.*}})'
+// C-NEXT:           |   `-UnaryOperator [[ADDR_53:0x[a-z0-9]*]] <col:11, col:12> 'int (*)({{.*}})' prefix '&' cannot overflow
+// C-NEXT:           |     `-DeclRefExpr [[ADDR_54:0x[a-z0-9]*]] <col:12> 'int ({{.*}})' Function [[ADDR_11]] 'also_before' 'int ({{.*}})'
+// C-NEXT:           `-CallExpr [[ADDR_55:0x[a-z0-9]*]] <line:13:1, line:27:25> 'int'
+// C-NEXT:             `-ImplicitCastExpr [[ADDR_56:0x[a-z0-9]*]] <line:13:1> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// C-NEXT:               `-DeclRefExpr [[ADDR_13]] <col:1> 'int ({{.*}})' Function [[ADDR_14]] 'also_before[implementation={vendor(llvm)}]' 'int ({{.*}})'
+
+// CXX:      |-FunctionDecl [[ADDR_0:0x[a-z0-9]*]] <{{.*}}, line:7:1> line:5:5 used also_before 'int ({{.*}})'
+// CXX-NEXT: | `-CompoundStmt [[ADDR_1:0x[a-z0-9]*]] <col:23, line:7:1>
+// CXX-NEXT: |   `-ReturnStmt [[ADDR_2:0x[a-z0-9]*]] <line:6:3, col:10>
+// CXX-NEXT: |     `-IntegerLiteral [[ADDR_3:0x[a-z0-9]*]] <col:10> 'int' 1
+// CXX-NEXT: |-FunctionDecl [[ADDR_4:0x[a-z0-9]*]] <line:10:1, col:20> col:5 implicit used also_after 'int ({{.*}})'
+// CXX-NEXT: | `-OMPDeclareVariantAttr [[ADDR_5:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(llvm)}
+// CXX-NEXT: |   `-DeclRefExpr [[ADDR_6:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_7:0x[a-z0-9]*]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CXX-NEXT: |-FunctionDecl [[ADDR_7]] <col:1, line:12:1> line:10:1 also_after[implementation={vendor(llvm)}] 'int ({{.*}})'
+// CXX-NEXT: | `-CompoundStmt [[ADDR_8:0x[a-z0-9]*]] <col:22, line:12:1>
+// CXX-NEXT: |   `-ReturnStmt [[ADDR_9:0x[a-z0-9]*]] <line:11:3, col:10>
+// CXX-NEXT: |     `-IntegerLiteral [[ADDR_10:0x[a-z0-9]*]] <col:10> 'int' 0
+// CXX-NEXT: |-FunctionDecl [[ADDR_11:0x[a-z0-9]*]] prev [[ADDR_0]] <line:13:1, col:21> col:5 implicit used also_before 'int ({{.*}})'
+// CXX-NEXT: | `-OMPDeclareVariantAttr [[ADDR_12:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(llvm)}
+// CXX-NEXT: |   `-DeclRefExpr [[ADDR_13:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_14:0x[a-z0-9]*]] 'also_before[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CXX-NEXT: |-FunctionDecl [[ADDR_14]] <col:1, line:15:1> line:13:1 also_before[implementation={vendor(llvm)}] 'int ({{.*}})'
+// CXX-NEXT: | `-CompoundStmt [[ADDR_15:0x[a-z0-9]*]] <col:23, line:15:1>
+// CXX-NEXT: |   `-ReturnStmt [[ADDR_16:0x[a-z0-9]*]] <line:14:3, col:10>
+// CXX-NEXT: |     `-IntegerLiteral [[ADDR_17:0x[a-z0-9]*]] <col:10> 'int' 0
+// CXX-NEXT: |-FunctionDecl [[ADDR_18:0x[a-z0-9]*]] prev [[ADDR_4]] <line:18:1, line:20:1> line:18:5 used also_after 'int ({{.*}})'
+// CXX-NEXT: | |-CompoundStmt [[ADDR_19:0x[a-z0-9]*]] <col:22, line:20:1>
+// CXX-NEXT: | | `-ReturnStmt [[ADDR_20:0x[a-z0-9]*]] <line:19:3, col:10>
+// CXX-NEXT: | |   `-IntegerLiteral [[ADDR_21:0x[a-z0-9]*]] <col:10> 'int' 2
+// CXX-NEXT: | `-OMPDeclareVariantAttr [[ADDR_22:0x[a-z0-9]*]] <<invalid sloc>> Inherited Implicit implementation={vendor(llvm)}
+// CXX-NEXT: |   `-DeclRefExpr [[ADDR_6]] <line:10:1> 'int ({{.*}})' Function [[ADDR_7]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CXX-NEXT: `-FunctionDecl [[ADDR_23:0x[a-z0-9]*]] <line:22:1, line:28:1> line:22:5 main 'int ({{.*}})'
+// CXX-NEXT:   `-CompoundStmt [[ADDR_24:0x[a-z0-9]*]] <col:12, line:28:1>
+// CXX-NEXT:     `-ReturnStmt [[ADDR_25:0x[a-z0-9]*]] <line:24:3, line:27:25>
+// CXX-NEXT:       `-BinaryOperator [[ADDR_26:0x[a-z0-9]*]] <line:24:10, line:27:25> 'int' '+'
+// CXX-NEXT:         |-BinaryOperator [[ADDR_27:0x[a-z0-9]*]] <line:24:10, line:26:24> 'int' '+'
+// CXX-NEXT:         | |-BinaryOperator [[ADDR_28:0x[a-z0-9]*]] <line:24:10, line:25:24> 'int' '+'
+// CXX-NEXT:         | | |-PseudoObjectExpr [[ADDR_29:0x[a-z0-9]*]] <line:24:10, col:23> 'int'
+// CXX-NEXT:         | | | |-CallExpr [[ADDR_30:0x[a-z0-9]*]] <col:10, col:23> 'int'
+// CXX-NEXT:         | | | | `-ImplicitCastExpr [[ADDR_31:0x[a-z0-9]*]] <col:10, col:21> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CXX-NEXT:         | | | |   `-ParenExpr [[ADDR_32:0x[a-z0-9]*]] <col:10, col:21> 'int ({{.*}})' lvalue
+// CXX-NEXT:         | | | |     `-DeclRefExpr [[ADDR_33:0x[a-z0-9]*]] <col:11> 'int ({{.*}})' {{.*}}Function [[ADDR_18]] 'also_after' 'int ({{.*}})'
+// CXX-NEXT:         | | | `-CallExpr [[ADDR_34:0x[a-z0-9]*]] <line:10:1, line:24:23> 'int'
+// CXX-NEXT:         | | |   `-ImplicitCastExpr [[ADDR_35:0x[a-z0-9]*]] <line:10:1> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CXX-NEXT:         | | |     `-DeclRefExpr [[ADDR_6]] <col:1> 'int ({{.*}})' Function [[ADDR_7]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CXX-NEXT:         | | `-PseudoObjectExpr [[ADDR_36:0x[a-z0-9]*]] <line:25:10, col:24> 'int'
+// CXX-NEXT:         | |   |-CallExpr [[ADDR_37:0x[a-z0-9]*]] <col:10, col:24> 'int'
+// CXX-NEXT:         | |   | `-ImplicitCastExpr [[ADDR_38:0x[a-z0-9]*]] <col:10, col:22> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CXX-NEXT:         | |   |   `-ParenExpr [[ADDR_39:0x[a-z0-9]*]] <col:10, col:22> 'int ({{.*}})' lvalue
+// CXX-NEXT:         | |   |     `-DeclRefExpr [[ADDR_40:0x[a-z0-9]*]] <col:11> 'int ({{.*}})' {{.*}}Function [[ADDR_11]] 'also_before' 'int ({{.*}})'
+// CXX-NEXT:         | |   `-CallExpr [[ADDR_41:0x[a-z0-9]*]] <line:13:1, line:25:24> 'int'
+// CXX-NEXT:         | |     `-ImplicitCastExpr [[ADDR_42:0x[a-z0-9]*]] <line:13:1> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CXX-NEXT:         | |       `-DeclRefExpr [[ADDR_13]] <col:1> 'int ({{.*}})' Function [[ADDR_14]] 'also_before[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CXX-NEXT:         | `-PseudoObjectExpr [[ADDR_43:0x[a-z0-9]*]] <line:26:10, col:24> 'int'
+// CXX-NEXT:         |   |-CallExpr [[ADDR_44:0x[a-z0-9]*]] <col:10, col:24> 'int'
+// CXX-NEXT:         |   | `-ParenExpr [[ADDR_45:0x[a-z0-9]*]] <col:10, col:22> 'int (*)({{.*}})'
+// CXX-NEXT:         |   |   `-UnaryOperator [[ADDR_46:0x[a-z0-9]*]] <col:11, col:12> 'int (*)({{.*}})' prefix '&' cannot overflow
+// CXX-NEXT:         |   |     `-DeclRefExpr [[ADDR_47:0x[a-z0-9]*]] <col:12> 'int ({{.*}})' {{.*}}Function [[ADDR_18]] 'also_after' 'int ({{.*}})'
+// CXX-NEXT:         |   `-CallExpr [[ADDR_48:0x[a-z0-9]*]] <line:10:1, line:26:24> 'int'
+// CXX-NEXT:         |     `-ImplicitCastExpr [[ADDR_49:0x[a-z0-9]*]] <line:10:1> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CXX-NEXT:         |       `-DeclRefExpr [[ADDR_6]] <col:1> 'int ({{.*}})' Function [[ADDR_7]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CXX-NEXT:         `-PseudoObjectExpr [[ADDR_50:0x[a-z0-9]*]] <line:27:10, col:25> 'int'
+// CXX-NEXT:           |-CallExpr [[ADDR_51:0x[a-z0-9]*]] <col:10, col:25> 'int'
+// CXX-NEXT:           | `-ParenExpr [[ADDR_52:0x[a-z0-9]*]] <col:10, col:23> 'int (*)({{.*}})'
+// CXX-NEXT:           |   `-UnaryOperator [[ADDR_53:0x[a-z0-9]*]] <col:11, col:12> 'int (*)({{.*}})' prefix '&' cannot overflow
+// CXX-NEXT:           |     `-DeclRefExpr [[ADDR_54:0x[a-z0-9]*]] <col:12> 'int ({{.*}})' {{.*}}Function [[ADDR_11]] 'also_before' 'int ({{.*}})'
+// CXX-NEXT:           `-CallExpr [[ADDR_55:0x[a-z0-9]*]] <line:13:1, line:27:25> 'int'
+// CXX-NEXT:             `-ImplicitCastExpr [[ADDR_56:0x[a-z0-9]*]] <line:13:1> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CXX-NEXT:               `-DeclRefExpr [[ADDR_13]] <col:1> 'int ({{.*}})' Function [[ADDR_14]] 'also_before[implementation={vendor(llvm)}]' 'int ({{.*}})'

clang/test/AST/ast-dump-openmp-begin-declare-variant_6.c

+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s       | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s -x c++| FileCheck %s
+// expected-no-diagnostics
+
+int also_before(void) {
+  return 0;
+}
+
+#pragma omp begin declare variant match(implementation={vendor(ibm)})
+int also_after(void) {
+  return 1;
+}
+int also_before(void) {
+  return 2;
+}
+#pragma omp end declare variant
+
+int also_after(void) {
+  return 0;
+}
+
+int main() {
+  // Should return 0.
+  return also_after() + also_before();
+}
+
+// Make sure:
+//  - we see the specialization in the AST
+//  - we do use the original pointers for the calls as the variants are not applicable (this is not the ibm compiler).
+
+// CHECK:      |-FunctionDecl [[ADDR_0:0x[a-z0-9]*]] <{{.*}}, line:7:1> line:5:5 used also_before 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_1:0x[a-z0-9]*]] <col:23, line:7:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_2:0x[a-z0-9]*]] <line:6:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_3:0x[a-z0-9]*]] <col:10> 'int' 0
+// CHECK-NEXT: |-FunctionDecl [[ADDR_4:0x[a-z0-9]*]] <line:10:1, col:20> col:5 implicit used also_after 'int ({{.*}})'
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_5:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(ibm)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_6:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_7:0x[a-z0-9]*]] 'also_after[implementation={vendor(ibm)}]' 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_7]] <col:1, line:12:1> line:10:1 also_after[implementation={vendor(ibm)}] 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_8:0x[a-z0-9]*]] <col:22, line:12:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_9:0x[a-z0-9]*]] <line:11:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_10:0x[a-z0-9]*]] <col:10> 'int' 1
+// CHECK-NEXT: |-FunctionDecl [[ADDR_11:0x[a-z0-9]*]] prev [[ADDR_0]] <line:13:1, col:21> col:5 implicit used also_before 'int ({{.*}})'
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_12:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(ibm)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_13:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_14:0x[a-z0-9]*]] 'also_before[implementation={vendor(ibm)}]' 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_14]] <col:1, line:15:1> line:13:1 also_before[implementation={vendor(ibm)}] 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_15:0x[a-z0-9]*]] <col:23, line:15:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_16:0x[a-z0-9]*]] <line:14:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_17:0x[a-z0-9]*]] <col:10> 'int' 2
+// CHECK-NEXT: |-FunctionDecl [[ADDR_18:0x[a-z0-9]*]] prev [[ADDR_4]] <line:18:1, line:20:1> line:18:5 used also_after 'int ({{.*}})'
+// CHECK-NEXT: | |-CompoundStmt [[ADDR_19:0x[a-z0-9]*]] <col:22, line:20:1>
+// CHECK-NEXT: | | `-ReturnStmt [[ADDR_20:0x[a-z0-9]*]] <line:19:3, col:10>
+// CHECK-NEXT: | |   `-IntegerLiteral [[ADDR_21:0x[a-z0-9]*]] <col:10> 'int' 0
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_22:0x[a-z0-9]*]] <<invalid sloc>> Inherited Implicit implementation={vendor(ibm)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_6]] <line:10:1> 'int ({{.*}})' Function [[ADDR_7]] 'also_after[implementation={vendor(ibm)}]' 'int ({{.*}})'
+// CHECK-NEXT: `-FunctionDecl [[ADDR_23:0x[a-z0-9]*]] <line:22:1, line:25:1> line:22:5 main 'int ({{.*}})'
+// CHECK-NEXT:   `-CompoundStmt [[ADDR_24:0x[a-z0-9]*]] <col:12, line:25:1>
+// CHECK-NEXT:     `-ReturnStmt [[ADDR_25:0x[a-z0-9]*]] <line:24:3, col:37>
+// CHECK-NEXT:       `-BinaryOperator [[ADDR_26:0x[a-z0-9]*]] <col:10, col:37> 'int' '+'
+// CHECK-NEXT:         |-CallExpr [[ADDR_27:0x[a-z0-9]*]] <col:10, col:21> 'int'
+// CHECK-NEXT:         | `-ImplicitCastExpr [[ADDR_28:0x[a-z0-9]*]] <col:10> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:         |   `-DeclRefExpr [[ADDR_29:0x[a-z0-9]*]] <col:10> 'int ({{.*}})' {{.*}}Function [[ADDR_18]] 'also_after' 'int ({{.*}})'
+// CHECK-NEXT:         `-CallExpr [[ADDR_30:0x[a-z0-9]*]] <col:25, col:37> 'int'
+// CHECK-NEXT:           `-ImplicitCastExpr [[ADDR_31:0x[a-z0-9]*]] <col:25> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:             `-DeclRefExpr [[ADDR_32:0x[a-z0-9]*]] <col:25> 'int ({{.*}})' {{.*}}Function [[ADDR_11]] 'also_before' 'int ({{.*}})'

clang/test/AST/ast-dump-openmp-begin-declare-variant_7.c

+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s       | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s -x c++| FileCheck %s
+// expected-no-diagnostics
+
+int OK_1(void);
+
+#pragma omp begin declare variant match(implementation={vendor(intel)})
+int OK_1(void) {
+  return 1;
+}
+int OK_2(void) {
+  return 1;
+}
+int not_OK(void) {
+  return 1;
+}
+int OK_3(void) {
+  return 1;
+}
+#pragma omp end declare variant
+
+int OK_3(void);
+
+int test() {
+  // Should cause an error due to not_OK()
+  return OK_1() + not_OK() + OK_3();
+}
+
+// Make sure:
+//  - we see a single error for `not_OK`
+//  - we do not see errors for OK_{1,2,3}
+//  FIXME: We actually do not see there error here.
+//         This case is unlikely to happen in practise and hard to diagnose during SEMA.
+//         We will issue an error during code generation instead. This is similar to the
+//         diagnosis in other multi-versioning schemes.
+
+// CHECK:      |-FunctionDecl [[ADDR_0:0x[a-z0-9]*]] <{{.*}}, col:14> col:5 used OK_1 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_1:0x[a-z0-9]*]] prev [[ADDR_0]] <line:8:1, col:14> col:5 implicit used OK_1 'int ({{.*}})'
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_2:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(intel)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_3:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_4:0x[a-z0-9]*]] 'OK_1[implementation={vendor(intel)}]' 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_4]] <col:1, line:10:1> line:8:1 OK_1[implementation={vendor(intel)}] 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_5:0x[a-z0-9]*]] <col:16, line:10:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_6:0x[a-z0-9]*]] <line:9:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_7:0x[a-z0-9]*]] <col:10> 'int' 1
+// CHECK-NEXT: |-FunctionDecl [[ADDR_8:0x[a-z0-9]*]] <line:11:1, col:14> col:5 implicit OK_2 'int ({{.*}})'
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_9:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(intel)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_10:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_11:0x[a-z0-9]*]] 'OK_2[implementation={vendor(intel)}]' 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_11]] <col:1, line:13:1> line:11:1 OK_2[implementation={vendor(intel)}] 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_12:0x[a-z0-9]*]] <col:16, line:13:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_13:0x[a-z0-9]*]] <line:12:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_14:0x[a-z0-9]*]] <col:10> 'int' 1
+// CHECK-NEXT: |-FunctionDecl [[ADDR_15:0x[a-z0-9]*]] <line:14:1, col:16> col:5 implicit used not_OK 'int ({{.*}})'
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_16:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(intel)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_17:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_18:0x[a-z0-9]*]] 'not_OK[implementation={vendor(intel)}]' 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_18]] <col:1, line:16:1> line:14:1 not_OK[implementation={vendor(intel)}] 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_19:0x[a-z0-9]*]] <col:18, line:16:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_20:0x[a-z0-9]*]] <line:15:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_21:0x[a-z0-9]*]] <col:10> 'int' 1
+// CHECK-NEXT: |-FunctionDecl [[ADDR_22:0x[a-z0-9]*]] <line:17:1, col:14> col:5 implicit used OK_3 'int ({{.*}})'
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_23:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(intel)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_24:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_25:0x[a-z0-9]*]] 'OK_3[implementation={vendor(intel)}]' 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_25]] <col:1, line:19:1> line:17:1 OK_3[implementation={vendor(intel)}] 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_26:0x[a-z0-9]*]] <col:16, line:19:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_27:0x[a-z0-9]*]] <line:18:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_28:0x[a-z0-9]*]] <col:10> 'int' 1
+// CHECK-NEXT: |-FunctionDecl [[ADDR_29:0x[a-z0-9]*]] prev [[ADDR_22]] <line:22:1, col:14> col:5 used OK_3 'int ({{.*}})'
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_30:0x[a-z0-9]*]] <<invalid sloc>> Inherited Implicit implementation={vendor(intel)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_24]] <line:17:1> 'int ({{.*}})' Function [[ADDR_25]] 'OK_3[implementation={vendor(intel)}]' 'int ({{.*}})'
+// CHECK-NEXT: `-FunctionDecl [[ADDR_31:0x[a-z0-9]*]] <line:24:1, line:27:1> line:24:5 test 'int ({{.*}})'
+// CHECK-NEXT:   `-CompoundStmt [[ADDR_32:0x[a-z0-9]*]] <col:12, line:27:1>
+// CHECK-NEXT:     `-ReturnStmt [[ADDR_33:0x[a-z0-9]*]] <line:26:3, col:35>
+// CHECK-NEXT:       `-BinaryOperator [[ADDR_34:0x[a-z0-9]*]] <col:10, col:35> 'int' '+'
+// CHECK-NEXT:         |-BinaryOperator [[ADDR_35:0x[a-z0-9]*]] <col:10, col:26> 'int' '+'
+// CHECK-NEXT:         | |-CallExpr [[ADDR_36:0x[a-z0-9]*]] <col:10, col:15> 'int'
+// CHECK-NEXT:         | | `-ImplicitCastExpr [[ADDR_37:0x[a-z0-9]*]] <col:10> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:         | |   `-DeclRefExpr [[ADDR_38:0x[a-z0-9]*]] <col:10> 'int ({{.*}})' {{.*}}Function [[ADDR_1]] 'OK_1' 'int ({{.*}})'
+// CHECK-NEXT:         | `-CallExpr [[ADDR_39:0x[a-z0-9]*]] <col:19, col:26> 'int'
+// CHECK-NEXT:         |   `-ImplicitCastExpr [[ADDR_40:0x[a-z0-9]*]] <col:19> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:         |     `-DeclRefExpr [[ADDR_41:0x[a-z0-9]*]] <col:19> 'int ({{.*}})' {{.*}}Function [[ADDR_15]] 'not_OK' 'int ({{.*}})'
+// CHECK-NEXT:         `-CallExpr [[ADDR_42:0x[a-z0-9]*]] <col:30, col:35> 'int'
+// CHECK-NEXT:           `-ImplicitCastExpr [[ADDR_43:0x[a-z0-9]*]] <col:30> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:             `-DeclRefExpr [[ADDR_44:0x[a-z0-9]*]] <col:30> 'int ({{.*}})' {{.*}}Function [[ADDR_29]] 'OK_3' 'int ({{.*}})'

clang/test/AST/ast-dump-openmp-begin-declare-variant_8.c

+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s       | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s -x c++| FileCheck %s
+// expected-no-diagnostics
+
+#pragma omp begin declare variant match(device={kind(cpu)})
+int also_before(void) {
+  return 1;
+}
+#pragma omp end declare variant
+
+#pragma omp begin declare variant match(implementation={vendor(score(0):llvm)})
+int also_after(void) {
+  return 0;
+}
+#pragma omp end declare variant
+#pragma omp begin declare variant match(implementation={vendor(score(100):llvm)})
+int also_before(void) {
+  return 0;
+}
+#pragma omp end declare variant
+
+int also_after(void) {
+  return 2;
+}
+
+int test() {
+  // Should return 0.
+  return also_after() + also_before();
+}
+
+// Make sure:
+//  - we do see the ast nodes for the cpu kind
+//  - we do see the ast nodes for the llvm vendor
+//  - we pick the right callees
+
+// CHECK:      |-FunctionDecl [[ADDR_0:0x[a-z0-9]*]] <{{.*}}, col:21> col:5 implicit used also_before 'int ({{.*}})'
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_1:0x[a-z0-9]*]] <<invalid sloc>> Implicit device={kind(cpu)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_2:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_3:0x[a-z0-9]*]] 'also_before[device={kind(cpu)}]' 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_3]] <col:1, line:8:1> line:6:1 also_before[device={kind(cpu)}] 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_4:0x[a-z0-9]*]] <col:23, line:8:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_5:0x[a-z0-9]*]] <line:7:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_6:0x[a-z0-9]*]] <col:10> 'int' 1
+// CHECK-NEXT: |-FunctionDecl [[ADDR_7:0x[a-z0-9]*]] <line:12:1, col:20> col:5 implicit used also_after 'int ({{.*}})'
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_8:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(score(0): llvm)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_9:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_10:0x[a-z0-9]*]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_10]] <col:1, line:14:1> line:12:1 also_after[implementation={vendor(llvm)}] 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_11:0x[a-z0-9]*]] <col:22, line:14:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_12:0x[a-z0-9]*]] <line:13:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_13:0x[a-z0-9]*]] <col:10> 'int' 0
+// CHECK-NEXT: |-FunctionDecl [[ADDR_14:0x[a-z0-9]*]] prev [[ADDR_0]] <line:17:1, col:21> col:5 implicit used also_before 'int ({{.*}})'
+// CHECK-NEXT: | |-OMPDeclareVariantAttr [[ADDR_15:0x[a-z0-9]*]] <<invalid sloc>> Inherited Implicit device={kind(cpu)}
+// CHECK-NEXT: | | `-DeclRefExpr [[ADDR_2]] <line:6:1> 'int ({{.*}})' Function [[ADDR_3]] 'also_before[device={kind(cpu)}]' 'int ({{.*}})'
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_16:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(score(100): llvm)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_17:0x[a-z0-9]*]] <line:17:1> 'int ({{.*}})' Function [[ADDR_18:0x[a-z0-9]*]] 'also_before[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_18]] <col:1, line:19:1> line:17:1 also_before[implementation={vendor(llvm)}] 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_19:0x[a-z0-9]*]] <col:23, line:19:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_20:0x[a-z0-9]*]] <line:18:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_21:0x[a-z0-9]*]] <col:10> 'int' 0
+// CHECK-NEXT: |-FunctionDecl [[ADDR_22:0x[a-z0-9]*]] prev [[ADDR_7]] <line:22:1, line:24:1> line:22:5 used also_after 'int ({{.*}})'
+// CHECK-NEXT: | |-CompoundStmt [[ADDR_23:0x[a-z0-9]*]] <col:22, line:24:1>
+// CHECK-NEXT: | | `-ReturnStmt [[ADDR_24:0x[a-z0-9]*]] <line:23:3, col:10>
+// CHECK-NEXT: | |   `-IntegerLiteral [[ADDR_25:0x[a-z0-9]*]] <col:10> 'int' 2
+// CHECK-NEXT: | `-OMPDeclareVariantAttr [[ADDR_26:0x[a-z0-9]*]] <<invalid sloc>> Inherited Implicit implementation={vendor(score(0): llvm)}
+// CHECK-NEXT: |   `-DeclRefExpr [[ADDR_9]] <line:12:1> 'int ({{.*}})' Function [[ADDR_10]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CHECK-NEXT: `-FunctionDecl [[ADDR_27:0x[a-z0-9]*]] <line:26:1, line:29:1> line:26:5 test 'int ({{.*}})'
+// CHECK-NEXT:   `-CompoundStmt [[ADDR_28:0x[a-z0-9]*]] <col:12, line:29:1>
+// CHECK-NEXT:     `-ReturnStmt [[ADDR_29:0x[a-z0-9]*]] <line:28:3, col:37>
+// CHECK-NEXT:       `-BinaryOperator [[ADDR_30:0x[a-z0-9]*]] <col:10, col:37> 'int' '+'
+// CHECK-NEXT:         |-PseudoObjectExpr [[ADDR_31:0x[a-z0-9]*]] <col:10, col:21> 'int'
+// CHECK-NEXT:         | |-CallExpr [[ADDR_32:0x[a-z0-9]*]] <col:10, col:21> 'int'
+// CHECK-NEXT:         | | `-ImplicitCastExpr [[ADDR_33:0x[a-z0-9]*]] <col:10> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:         | |   `-DeclRefExpr [[ADDR_34:0x[a-z0-9]*]] <col:10> 'int ({{.*}})' {{.*}}Function [[ADDR_22]] 'also_after' 'int ({{.*}})'
+// CHECK-NEXT:         | `-CallExpr [[ADDR_35:0x[a-z0-9]*]] <line:12:1, line:28:21> 'int'
+// CHECK-NEXT:         |   `-ImplicitCastExpr [[ADDR_36:0x[a-z0-9]*]] <line:12:1> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:         |     `-DeclRefExpr [[ADDR_9]] <col:1> 'int ({{.*}})' Function [[ADDR_10]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CHECK-NEXT:         `-PseudoObjectExpr [[ADDR_37:0x[a-z0-9]*]] <line:28:25, col:37> 'int'
+// CHECK-NEXT:           |-CallExpr [[ADDR_38:0x[a-z0-9]*]] <col:25, col:37> 'int'
+// CHECK-NEXT:           | `-ImplicitCastExpr [[ADDR_39:0x[a-z0-9]*]] <col:25> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:           |   `-DeclRefExpr [[ADDR_40:0x[a-z0-9]*]] <col:25> 'int ({{.*}})' {{.*}}Function [[ADDR_14]] 'also_before' 'int ({{.*}})'
+// CHECK-NEXT:           `-CallExpr [[ADDR_41:0x[a-z0-9]*]] <line:17:1, line:28:37> 'int'
+// CHECK-NEXT:             `-ImplicitCastExpr [[ADDR_42:0x[a-z0-9]*]] <line:17:1> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:               `-DeclRefExpr [[ADDR_17]] <col:1> 'int ({{.*}})' Function [[ADDR_18]] 'also_before[implementation={vendor(llvm)}]' 'int ({{.*}})'

clang/test/AST/ast-dump-openmp-begin-declare-variant_9.c

+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s       | FileCheck %s --check-prefix=C
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s -x c++| FileCheck %s --check-prefix=CXX
+// expected-no-diagnostics
+
+int also_before(void) {
+  return 0;
+}
+
+#pragma omp begin declare variant match(implementation={vendor(llvm)})
+int also_after(void) {
+  return 1;
+}
+int also_before(void) {
+  return 2;
+}
+#pragma omp end declare variant
+
+int also_after(void) {
+  return 0;
+}
+
+void foo();
+typedef int(*fd)(void);
+int main() {
+  // Should return 0.
+  fd fns[2];
+  fns[0] = &also_before;
+  fns[1] = also_after;
+  return (foo(), also_after)() +
+         (fns[0])() +
+         (1[fns])();
+}
+
+// Make sure:
+//  - we see the specialization in the AST
+//  - we pick the right callees
+
+// C:      |-FunctionDecl [[ADDR_0:0x[a-z0-9]*]] <{{.*}}, line:7:1> line:5:5 used also_before 'int ({{.*}})'
+// C-NEXT: | `-CompoundStmt [[ADDR_1:0x[a-z0-9]*]] <col:23, line:7:1>
+// C-NEXT: |   `-ReturnStmt [[ADDR_2:0x[a-z0-9]*]] <line:6:3, col:10>
+// C-NEXT: |     `-IntegerLiteral [[ADDR_3:0x[a-z0-9]*]] <col:10> 'int' 0
+// C-NEXT: |-FunctionDecl [[ADDR_4:0x[a-z0-9]*]] <line:10:1, col:20> col:5 implicit used also_after 'int ({{.*}})'
+// C-NEXT: | `-OMPDeclareVariantAttr [[ADDR_5:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(llvm)}
+// C-NEXT: |   `-DeclRefExpr [[ADDR_6:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_7:0x[a-z0-9]*]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// C-NEXT: |-FunctionDecl [[ADDR_7]] <col:1, line:12:1> line:10:1 also_after[implementation={vendor(llvm)}] 'int ({{.*}})'
+// C-NEXT: | `-CompoundStmt [[ADDR_8:0x[a-z0-9]*]] <col:22, line:12:1>
+// C-NEXT: |   `-ReturnStmt [[ADDR_9:0x[a-z0-9]*]] <line:11:3, col:10>
+// C-NEXT: |     `-IntegerLiteral [[ADDR_10:0x[a-z0-9]*]] <col:10> 'int' 1
+// C-NEXT: |-FunctionDecl [[ADDR_11:0x[a-z0-9]*]] prev [[ADDR_0]] <line:13:1, col:21> col:5 implicit used also_before 'int ({{.*}})'
+// C-NEXT: | `-OMPDeclareVariantAttr [[ADDR_12:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(llvm)}
+// C-NEXT: |   `-DeclRefExpr [[ADDR_13:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_14:0x[a-z0-9]*]] 'also_before[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// C-NEXT: |-FunctionDecl [[ADDR_14]] <col:1, line:15:1> line:13:1 also_before[implementation={vendor(llvm)}] 'int ({{.*}})'
+// C-NEXT: | `-CompoundStmt [[ADDR_15:0x[a-z0-9]*]] <col:23, line:15:1>
+// C-NEXT: |   `-ReturnStmt [[ADDR_16:0x[a-z0-9]*]] <line:14:3, col:10>
+// C-NEXT: |     `-IntegerLiteral [[ADDR_17:0x[a-z0-9]*]] <col:10> 'int' 2
+// C-NEXT: |-FunctionDecl [[ADDR_18:0x[a-z0-9]*]] prev [[ADDR_4]] <line:18:1, line:20:1> line:18:5 used also_after 'int ({{.*}})'
+// C-NEXT: | |-CompoundStmt [[ADDR_19:0x[a-z0-9]*]] <col:22, line:20:1>
+// C-NEXT: | | `-ReturnStmt [[ADDR_20:0x[a-z0-9]*]] <line:19:3, col:10>
+// C-NEXT: | |   `-IntegerLiteral [[ADDR_21:0x[a-z0-9]*]] <col:10> 'int' 0
+// C-NEXT: | `-OMPDeclareVariantAttr [[ADDR_22:0x[a-z0-9]*]] <<invalid sloc>> Inherited Implicit implementation={vendor(llvm)}
+// C-NEXT: |   `-DeclRefExpr [[ADDR_6]] <line:10:1> 'int ({{.*}})' Function [[ADDR_7]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// C-NEXT: |-FunctionDecl [[ADDR_23:0x[a-z0-9]*]] <line:22:1, col:10> col:6 used foo 'void ({{.*}})'
+// C-NEXT: |-TypedefDecl [[ADDR_24:0x[a-z0-9]*]] <line:23:1, col:22> col:14 referenced fd 'int (*)({{.*}})'
+// C-NEXT: | `-PointerType [[ADDR_25:0x[a-z0-9]*]] 'int (*)({{.*}})'
+// C-NEXT: |   `-ParenType [[ADDR_26:0x[a-z0-9]*]] 'int ({{.*}})' sugar
+// C-NEXT: |     `-FunctionProtoType [[ADDR_27:0x[a-z0-9]*]] 'int ({{.*}})' cdecl
+// C-NEXT: |       `-BuiltinType [[ADDR_28:0x[a-z0-9]*]] 'int'
+// C-NEXT: `-FunctionDecl [[ADDR_29:0x[a-z0-9]*]] <line:24:1, line:32:1> line:24:5 main 'int ({{.*}})'
+// C-NEXT:   `-CompoundStmt [[ADDR_30:0x[a-z0-9]*]] <col:12, line:32:1>
+// C-NEXT:     |-DeclStmt [[ADDR_31:0x[a-z0-9]*]] <line:26:3, col:12>
+// C-NEXT:     | `-VarDecl [[ADDR_32:0x[a-z0-9]*]] <col:3, col:11> col:6 used fns 'fd [2]'
+// C-NEXT:     |-BinaryOperator [[ADDR_33:0x[a-z0-9]*]] <line:27:3, col:13> 'fd':'int (*)({{.*}})' '='
+// C-NEXT:     | |-ArraySubscriptExpr [[ADDR_34:0x[a-z0-9]*]] <col:3, col:8> 'fd':'int (*)({{.*}})' lvalue
+// C-NEXT:     | | |-ImplicitCastExpr [[ADDR_35:0x[a-z0-9]*]] <col:3> 'fd *' <ArrayToPointerDecay>
+// C-NEXT:     | | | `-DeclRefExpr [[ADDR_36:0x[a-z0-9]*]] <col:3> 'fd [2]' {{.*}}Var [[ADDR_32]] 'fns' 'fd [2]'
+// C-NEXT:     | | `-IntegerLiteral [[ADDR_37:0x[a-z0-9]*]] <col:7> 'int' 0
+// C-NEXT:     | `-UnaryOperator [[ADDR_38:0x[a-z0-9]*]] <col:12, col:13> 'int (*)({{.*}})' prefix '&' cannot overflow
+// C-NEXT:     |   `-DeclRefExpr [[ADDR_39:0x[a-z0-9]*]] <col:13> 'int ({{.*}})' Function [[ADDR_11]] 'also_before' 'int ({{.*}})'
+// C-NEXT:     |-BinaryOperator [[ADDR_40:0x[a-z0-9]*]] <line:28:3, col:12> 'fd':'int (*)({{.*}})' '='
+// C-NEXT:     | |-ArraySubscriptExpr [[ADDR_41:0x[a-z0-9]*]] <col:3, col:8> 'fd':'int (*)({{.*}})' lvalue
+// C-NEXT:     | | |-ImplicitCastExpr [[ADDR_42:0x[a-z0-9]*]] <col:3> 'fd *' <ArrayToPointerDecay>
+// C-NEXT:     | | | `-DeclRefExpr [[ADDR_43:0x[a-z0-9]*]] <col:3> 'fd [2]' {{.*}}Var [[ADDR_32]] 'fns' 'fd [2]'
+// C-NEXT:     | | `-IntegerLiteral [[ADDR_44:0x[a-z0-9]*]] <col:7> 'int' 1
+// C-NEXT:     | `-ImplicitCastExpr [[ADDR_45:0x[a-z0-9]*]] <col:12> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// C-NEXT:     |   `-DeclRefExpr [[ADDR_46:0x[a-z0-9]*]] <col:12> 'int ({{.*}})' Function [[ADDR_18]] 'also_after' 'int ({{.*}})'
+// C-NEXT:     `-ReturnStmt [[ADDR_47:0x[a-z0-9]*]] <line:29:3, line:31:19>
+// C-NEXT:       `-BinaryOperator [[ADDR_48:0x[a-z0-9]*]] <line:29:10, line:31:19> 'int' '+'
+// C-NEXT:         |-BinaryOperator [[ADDR_49:0x[a-z0-9]*]] <line:29:10, line:30:19> 'int' '+'
+// C-NEXT:         | |-CallExpr [[ADDR_50:0x[a-z0-9]*]] <line:29:10, col:30> 'int'
+// C-NEXT:         | | `-ParenExpr [[ADDR_51:0x[a-z0-9]*]] <col:10, col:28> 'int (*)({{.*}})'
+// C-NEXT:         | |   `-BinaryOperator [[ADDR_52:0x[a-z0-9]*]] <col:11, col:18> 'int (*)({{.*}})' ','
+// C-NEXT:         | |     |-CallExpr [[ADDR_53:0x[a-z0-9]*]] <col:11, col:15> 'void'
+// C-NEXT:         | |     | `-ImplicitCastExpr [[ADDR_54:0x[a-z0-9]*]] <col:11> 'void (*)({{.*}})' <FunctionToPointerDecay>
+// C-NEXT:         | |     |   `-DeclRefExpr [[ADDR_55:0x[a-z0-9]*]] <col:11> 'void ({{.*}})' Function [[ADDR_23]] 'foo' 'void ({{.*}})'
+// C-NEXT:         | |     `-ImplicitCastExpr [[ADDR_56:0x[a-z0-9]*]] <col:18> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// C-NEXT:         | |       `-DeclRefExpr [[ADDR_57:0x[a-z0-9]*]] <col:18> 'int ({{.*}})' Function [[ADDR_18]] 'also_after' 'int ({{.*}})'
+// C-NEXT:         | `-CallExpr [[ADDR_58:0x[a-z0-9]*]] <line:30:10, col:19> 'int'
+// C-NEXT:         |   `-ImplicitCastExpr [[ADDR_59:0x[a-z0-9]*]] <col:10, col:17> 'fd':'int (*)({{.*}})' <LValueToRValue>
+// C-NEXT:         |     `-ParenExpr [[ADDR_60:0x[a-z0-9]*]] <col:10, col:17> 'fd':'int (*)({{.*}})' lvalue
+// C-NEXT:         |       `-ArraySubscriptExpr [[ADDR_61:0x[a-z0-9]*]] <col:11, col:16> 'fd':'int (*)({{.*}})' lvalue
+// C-NEXT:         |         |-ImplicitCastExpr [[ADDR_62:0x[a-z0-9]*]] <col:11> 'fd *' <ArrayToPointerDecay>
+// C-NEXT:         |         | `-DeclRefExpr [[ADDR_63:0x[a-z0-9]*]] <col:11> 'fd [2]' {{.*}}Var [[ADDR_32]] 'fns' 'fd [2]'
+// C-NEXT:         |         `-IntegerLiteral [[ADDR_64:0x[a-z0-9]*]] <col:15> 'int' 0
+// C-NEXT:         `-CallExpr [[ADDR_65:0x[a-z0-9]*]] <line:31:10, col:19> 'int'
+// C-NEXT:           `-ImplicitCastExpr [[ADDR_66:0x[a-z0-9]*]] <col:10, col:17> 'fd':'int (*)({{.*}})' <LValueToRValue>
+// C-NEXT:             `-ParenExpr [[ADDR_67:0x[a-z0-9]*]] <col:10, col:17> 'fd':'int (*)({{.*}})' lvalue
+// C-NEXT:               `-ArraySubscriptExpr [[ADDR_68:0x[a-z0-9]*]] <col:11, col:16> 'fd':'int (*)({{.*}})' lvalue
+// C-NEXT:                 |-IntegerLiteral [[ADDR_69:0x[a-z0-9]*]] <col:11> 'int' 1
+// C-NEXT:                 `-ImplicitCastExpr [[ADDR_70:0x[a-z0-9]*]] <col:13> 'fd *' <ArrayToPointerDecay>
+// C-NEXT:                   `-DeclRefExpr [[ADDR_71:0x[a-z0-9]*]] <col:13> 'fd [2]' {{.*}}Var [[ADDR_32]] 'fns' 'fd [2]'
+
+// CXX:      |-FunctionDecl [[ADDR_0:0x[a-z0-9]*]] <{{.*}}, line:7:1> line:5:5 used also_before 'int ({{.*}})'
+// CXX-NEXT: | `-CompoundStmt [[ADDR_1:0x[a-z0-9]*]] <col:23, line:7:1>
+// CXX-NEXT: |   `-ReturnStmt [[ADDR_2:0x[a-z0-9]*]] <line:6:3, col:10>
+// CXX-NEXT: |     `-IntegerLiteral [[ADDR_3:0x[a-z0-9]*]] <col:10> 'int' 0
+// CXX-NEXT: |-FunctionDecl [[ADDR_4:0x[a-z0-9]*]] <line:10:1, col:20> col:5 implicit used also_after 'int ({{.*}})'
+// CXX-NEXT: | `-OMPDeclareVariantAttr [[ADDR_5:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(llvm)}
+// CXX-NEXT: |   `-DeclRefExpr [[ADDR_6:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_7:0x[a-z0-9]*]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CXX-NEXT: |-FunctionDecl [[ADDR_7]] <col:1, line:12:1> line:10:1 also_after[implementation={vendor(llvm)}] 'int ({{.*}})'
+// CXX-NEXT: | `-CompoundStmt [[ADDR_8:0x[a-z0-9]*]] <col:22, line:12:1>
+// CXX-NEXT: |   `-ReturnStmt [[ADDR_9:0x[a-z0-9]*]] <line:11:3, col:10>
+// CXX-NEXT: |     `-IntegerLiteral [[ADDR_10:0x[a-z0-9]*]] <col:10> 'int' 1
+// CXX-NEXT: |-FunctionDecl [[ADDR_11:0x[a-z0-9]*]] prev [[ADDR_0]] <line:13:1, col:21> col:5 implicit used also_before 'int ({{.*}})'
+// CXX-NEXT: | `-OMPDeclareVariantAttr [[ADDR_12:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(llvm)}
+// CXX-NEXT: |   `-DeclRefExpr [[ADDR_13:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_14:0x[a-z0-9]*]] 'also_before[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CXX-NEXT: |-FunctionDecl [[ADDR_14]] <col:1, line:15:1> line:13:1 also_before[implementation={vendor(llvm)}] 'int ({{.*}})'
+// CXX-NEXT: | `-CompoundStmt [[ADDR_15:0x[a-z0-9]*]] <col:23, line:15:1>
+// CXX-NEXT: |   `-ReturnStmt [[ADDR_16:0x[a-z0-9]*]] <line:14:3, col:10>
+// CXX-NEXT: |     `-IntegerLiteral [[ADDR_17:0x[a-z0-9]*]] <col:10> 'int' 2
+// CXX-NEXT: |-FunctionDecl [[ADDR_18:0x[a-z0-9]*]] prev [[ADDR_4]] <line:18:1, line:20:1> line:18:5 used also_after 'int ({{.*}})'
+// CXX-NEXT: | |-CompoundStmt [[ADDR_19:0x[a-z0-9]*]] <col:22, line:20:1>
+// CXX-NEXT: | | `-ReturnStmt [[ADDR_20:0x[a-z0-9]*]] <line:19:3, col:10>
+// CXX-NEXT: | |   `-IntegerLiteral [[ADDR_21:0x[a-z0-9]*]] <col:10> 'int' 0
+// CXX-NEXT: | `-OMPDeclareVariantAttr [[ADDR_22:0x[a-z0-9]*]] <<invalid sloc>> Inherited Implicit implementation={vendor(llvm)}
+// CXX-NEXT: |   `-DeclRefExpr [[ADDR_6]] <line:10:1> 'int ({{.*}})' Function [[ADDR_7]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CXX-NEXT: |-FunctionDecl [[ADDR_23:0x[a-z0-9]*]] <line:22:1, col:10> col:6 used foo 'void ({{.*}})'
+// CXX-NEXT: |-TypedefDecl [[ADDR_24:0x[a-z0-9]*]] <line:23:1, col:22> col:14 referenced fd 'int (*)({{.*}})'
+// CXX-NEXT: | `-PointerType [[ADDR_25:0x[a-z0-9]*]] 'int (*)({{.*}})'
+// CXX-NEXT: |   `-ParenType [[ADDR_26:0x[a-z0-9]*]] 'int ({{.*}})' sugar
+// CXX-NEXT: |     `-FunctionProtoType [[ADDR_27:0x[a-z0-9]*]] 'int ({{.*}})' cdecl
+// CXX-NEXT: |       `-BuiltinType [[ADDR_28:0x[a-z0-9]*]] 'int'
+// CXX-NEXT: `-FunctionDecl [[ADDR_29:0x[a-z0-9]*]] <line:24:1, line:32:1> line:24:5 main 'int ({{.*}})'
+// CXX-NEXT:   `-CompoundStmt [[ADDR_30:0x[a-z0-9]*]] <col:12, line:32:1>
+// CXX-NEXT:     |-DeclStmt [[ADDR_31:0x[a-z0-9]*]] <line:26:3, col:12>
+// CXX-NEXT:     | `-VarDecl [[ADDR_32:0x[a-z0-9]*]] <col:3, col:11> col:6 used fns 'fd [2]'
+// CXX-NEXT:     |-BinaryOperator [[ADDR_33:0x[a-z0-9]*]] <line:27:3, col:13> 'fd':'int (*)({{.*}})' {{.*}}'='
+// CXX-NEXT:     | |-ArraySubscriptExpr [[ADDR_34:0x[a-z0-9]*]] <col:3, col:8> 'fd':'int (*)({{.*}})' lvalue
+// CXX-NEXT:     | | |-ImplicitCastExpr [[ADDR_35:0x[a-z0-9]*]] <col:3> 'fd *' <ArrayToPointerDecay>
+// CXX-NEXT:     | | | `-DeclRefExpr [[ADDR_36:0x[a-z0-9]*]] <col:3> 'fd [2]' {{.*}}Var [[ADDR_32]] 'fns' 'fd [2]'
+// CXX-NEXT:     | | `-IntegerLiteral [[ADDR_37:0x[a-z0-9]*]] <col:7> 'int' 0
+// CXX-NEXT:     | `-UnaryOperator [[ADDR_38:0x[a-z0-9]*]] <col:12, col:13> 'int (*)({{.*}})' prefix '&' cannot overflow
+// CXX-NEXT:     |   `-DeclRefExpr [[ADDR_39:0x[a-z0-9]*]] <col:13> 'int ({{.*}})' {{.*}}Function [[ADDR_11]] 'also_before' 'int ({{.*}})'
+// CXX-NEXT:     |-BinaryOperator [[ADDR_40:0x[a-z0-9]*]] <line:28:3, col:12> 'fd':'int (*)({{.*}})' {{.*}}'='
+// CXX-NEXT:     | |-ArraySubscriptExpr [[ADDR_41:0x[a-z0-9]*]] <col:3, col:8> 'fd':'int (*)({{.*}})' lvalue
+// CXX-NEXT:     | | |-ImplicitCastExpr [[ADDR_42:0x[a-z0-9]*]] <col:3> 'fd *' <ArrayToPointerDecay>
+// CXX-NEXT:     | | | `-DeclRefExpr [[ADDR_43:0x[a-z0-9]*]] <col:3> 'fd [2]' {{.*}}Var [[ADDR_32]] 'fns' 'fd [2]'
+// CXX-NEXT:     | | `-IntegerLiteral [[ADDR_44:0x[a-z0-9]*]] <col:7> 'int' 1
+// CXX-NEXT:     | `-ImplicitCastExpr [[ADDR_45:0x[a-z0-9]*]] <col:12> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CXX-NEXT:     |   `-DeclRefExpr [[ADDR_46:0x[a-z0-9]*]] <col:12> 'int ({{.*}})' {{.*}}Function [[ADDR_18]] 'also_after' 'int ({{.*}})'
+// CXX-NEXT:     `-ReturnStmt [[ADDR_47:0x[a-z0-9]*]] <line:29:3, line:31:19>
+// CXX-NEXT:       `-BinaryOperator [[ADDR_48:0x[a-z0-9]*]] <line:29:10, line:31:19> 'int' '+'
+// CXX-NEXT:         |-BinaryOperator [[ADDR_49:0x[a-z0-9]*]] <line:29:10, line:30:19> 'int' '+'
+// CXX-NEXT:         | |-CallExpr [[ADDR_50:0x[a-z0-9]*]] <line:29:10, col:30> 'int'
+// CXX-NEXT:         | | `-ImplicitCastExpr [[ADDR_51:0x[a-z0-9]*]] <col:10, col:28> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CXX-NEXT:         | |   `-ParenExpr [[ADDR_52:0x[a-z0-9]*]] <col:10, col:28> 'int ({{.*}})' lvalue
+// CXX-NEXT:         | |     `-BinaryOperator [[ADDR_53:0x[a-z0-9]*]] <col:11, col:18> 'int ({{.*}})' {{.*}}','
+// CXX-NEXT:         | |       |-CallExpr [[ADDR_54:0x[a-z0-9]*]] <col:11, col:15> 'void'
+// CXX-NEXT:         | |       | `-ImplicitCastExpr [[ADDR_55:0x[a-z0-9]*]] <col:11> 'void (*)({{.*}})' <FunctionToPointerDecay>
+// CXX-NEXT:         | |       |   `-DeclRefExpr [[ADDR_56:0x[a-z0-9]*]] <col:11> 'void ({{.*}})' {{.*}}Function [[ADDR_23]] 'foo' 'void ({{.*}})'
+// CXX-NEXT:         | |       `-DeclRefExpr [[ADDR_57:0x[a-z0-9]*]] <col:18> 'int ({{.*}})' {{.*}}Function [[ADDR_18]] 'also_after' 'int ({{.*}})'
+// CXX-NEXT:         | `-CallExpr [[ADDR_58:0x[a-z0-9]*]] <line:30:10, col:19> 'int'
+// CXX-NEXT:         |   `-ImplicitCastExpr [[ADDR_59:0x[a-z0-9]*]] <col:10, col:17> 'fd':'int (*)({{.*}})' <LValueToRValue>
+// CXX-NEXT:         |     `-ParenExpr [[ADDR_60:0x[a-z0-9]*]] <col:10, col:17> 'fd':'int (*)({{.*}})' lvalue
+// CXX-NEXT:         |       `-ArraySubscriptExpr [[ADDR_61:0x[a-z0-9]*]] <col:11, col:16> 'fd':'int (*)({{.*}})' lvalue
+// CXX-NEXT:         |         |-ImplicitCastExpr [[ADDR_62:0x[a-z0-9]*]] <col:11> 'fd *' <ArrayToPointerDecay>
+// CXX-NEXT:         |         | `-DeclRefExpr [[ADDR_63:0x[a-z0-9]*]] <col:11> 'fd [2]' {{.*}}Var [[ADDR_32]] 'fns' 'fd [2]'
+// CXX-NEXT:         |         `-IntegerLiteral [[ADDR_64:0x[a-z0-9]*]] <col:15> 'int' 0
+// CXX-NEXT:         `-CallExpr [[ADDR_65:0x[a-z0-9]*]] <line:31:10, col:19> 'int'
+// CXX-NEXT:           `-ImplicitCastExpr [[ADDR_66:0x[a-z0-9]*]] <col:10, col:17> 'fd':'int (*)({{.*}})' <LValueToRValue>
+// CXX-NEXT:             `-ParenExpr [[ADDR_67:0x[a-z0-9]*]] <col:10, col:17> 'fd':'int (*)({{.*}})' lvalue
+// CXX-NEXT:               `-ArraySubscriptExpr [[ADDR_68:0x[a-z0-9]*]] <col:11, col:16> 'fd':'int (*)({{.*}})' lvalue
+// CXX-NEXT:                 |-IntegerLiteral [[ADDR_69:0x[a-z0-9]*]] <col:11> 'int' 1
+// CXX-NEXT:                 `-ImplicitCastExpr [[ADDR_70:0x[a-z0-9]*]] <col:13> 'fd *' <ArrayToPointerDecay>
+// CXX-NEXT:                   `-DeclRefExpr [[ADDR_71:0x[a-z0-9]*]] <col:13> 'fd [2]' {{.*}}Var [[ADDR_32]] 'fns' 'fd [2]'

clang/test/AST/ast-dump-openmp-begin-declare-variant_addr_1.c

+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s       | FileCheck %s --check-prefix=C
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s -x c++| FileCheck %s --check-prefix=CXX
+// expected-no-diagnostics
+
+int also_before(void) {
+  return 0;
+}
+
+#pragma omp begin declare variant match(implementation={vendor(llvm)})
+int also_after(void) {
+  return 1;
+}
+int also_before(void) {
+  return 2;
+}
+#pragma omp end declare variant
+
+int also_after(void) {
+  return 0;
+}
+
+int test(int (*fd)(void)) {
+  return fd();
+}
+int main() {
+  // Should return 0.
+  return test(also_after) +
+         test(also_before) +
+         test(&also_after) +
+         test(&also_before);
+}
+
+// Make sure:
+//  - we see the specialization in the AST
+//  - we pick the right callees
+
+// CXX:      |-FunctionDecl [[ADDR_0:0x[a-z0-9]*]] <{{.*}}, line:7:1> line:5:5 used also_before 'int ({{.*}})'
+// CXX-NEXT: | `-CompoundStmt [[ADDR_1:0x[a-z0-9]*]] <col:23, line:7:1>
+// CXX-NEXT: |   `-ReturnStmt [[ADDR_2:0x[a-z0-9]*]] <line:6:3, col:10>
+// CXX-NEXT: |     `-IntegerLiteral [[ADDR_3:0x[a-z0-9]*]] <col:10> 'int' 0
+// CXX-NEXT: |-FunctionDecl [[ADDR_4:0x[a-z0-9]*]] <line:10:1, col:20> col:5 implicit used also_after 'int ({{.*}})'
+// CXX-NEXT: | `-OMPDeclareVariantAttr [[ADDR_5:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(llvm)}
+// CXX-NEXT: |   `-DeclRefExpr [[ADDR_6:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_7:0x[a-z0-9]*]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CXX-NEXT: |-FunctionDecl [[ADDR_7]] <col:1, line:12:1> line:10:1 also_after[implementation={vendor(llvm)}] 'int ({{.*}})'
+// CXX-NEXT: | `-CompoundStmt [[ADDR_8:0x[a-z0-9]*]] <col:22, line:12:1>
+// CXX-NEXT: |   `-ReturnStmt [[ADDR_9:0x[a-z0-9]*]] <line:11:3, col:10>
+// CXX-NEXT: |     `-IntegerLiteral [[ADDR_10:0x[a-z0-9]*]] <col:10> 'int' 1
+// CXX-NEXT: |-FunctionDecl [[ADDR_11:0x[a-z0-9]*]] prev [[ADDR_0]] <line:13:1, col:21> col:5 implicit used also_before 'int ({{.*}})'
+// CXX-NEXT: | `-OMPDeclareVariantAttr [[ADDR_12:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(llvm)}
+// CXX-NEXT: |   `-DeclRefExpr [[ADDR_13:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_14:0x[a-z0-9]*]] 'also_before[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CXX-NEXT: |-FunctionDecl [[ADDR_14]] <col:1, line:15:1> line:13:1 also_before[implementation={vendor(llvm)}] 'int ({{.*}})'
+// CXX-NEXT: | `-CompoundStmt [[ADDR_15:0x[a-z0-9]*]] <col:23, line:15:1>
+// CXX-NEXT: |   `-ReturnStmt [[ADDR_16:0x[a-z0-9]*]] <line:14:3, col:10>
+// CXX-NEXT: |     `-IntegerLiteral [[ADDR_17:0x[a-z0-9]*]] <col:10> 'int' 2
+// CXX-NEXT: |-FunctionDecl [[ADDR_18:0x[a-z0-9]*]] prev [[ADDR_4]] <line:18:1, line:20:1> line:18:5 used also_after 'int ({{.*}})'
+// CXX-NEXT: | |-CompoundStmt [[ADDR_19:0x[a-z0-9]*]] <col:22, line:20:1>
+// CXX-NEXT: | | `-ReturnStmt [[ADDR_20:0x[a-z0-9]*]] <line:19:3, col:10>
+// CXX-NEXT: | |   `-IntegerLiteral [[ADDR_21:0x[a-z0-9]*]] <col:10> 'int' 0
+// CXX-NEXT: | `-OMPDeclareVariantAttr [[ADDR_22:0x[a-z0-9]*]] <<invalid sloc>> Inherited Implicit implementation={vendor(llvm)}
+// CXX-NEXT: |   `-DeclRefExpr [[ADDR_6]] <line:10:1> 'int ({{.*}})' Function [[ADDR_7]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// CXX-NEXT: |-FunctionDecl [[ADDR_23:0x[a-z0-9]*]] <line:22:1, line:24:1> line:22:5 used test 'int (int (*)({{.*}}))'
+// CXX-NEXT: | |-ParmVarDecl [[ADDR_24:0x[a-z0-9]*]] <col:10, col:24> col:16 used fd 'int (*)({{.*}})'
+// CXX-NEXT: | `-CompoundStmt [[ADDR_25:0x[a-z0-9]*]] <col:27, line:24:1>
+// CXX-NEXT: |   `-ReturnStmt [[ADDR_26:0x[a-z0-9]*]] <line:23:3, col:13>
+// CXX-NEXT: |     `-CallExpr [[ADDR_27:0x[a-z0-9]*]] <col:10, col:13> 'int'
+// CXX-NEXT: |       `-ImplicitCastExpr [[ADDR_28:0x[a-z0-9]*]] <col:10> 'int (*)({{.*}})' <LValueToRValue>
+// CXX-NEXT: |         `-DeclRefExpr [[ADDR_29:0x[a-z0-9]*]] <col:10> 'int (*)({{.*}})' {{.*}}ParmVar [[ADDR_24]] 'fd' 'int (*)({{.*}})'
+// CXX-NEXT: `-FunctionDecl [[ADDR_30:0x[a-z0-9]*]] <line:25:1, line:31:1> line:25:5 main 'int ({{.*}})'
+// CXX-NEXT:   `-CompoundStmt [[ADDR_31:0x[a-z0-9]*]] <col:12, line:31:1>
+// CXX-NEXT:     `-ReturnStmt [[ADDR_32:0x[a-z0-9]*]] <line:27:3, line:30:27>
+// CXX-NEXT:       `-BinaryOperator [[ADDR_33:0x[a-z0-9]*]] <line:27:10, line:30:27> 'int' '+'
+// CXX-NEXT:         |-BinaryOperator [[ADDR_34:0x[a-z0-9]*]] <line:27:10, line:29:26> 'int' '+'
+// CXX-NEXT:         | |-BinaryOperator [[ADDR_35:0x[a-z0-9]*]] <line:27:10, line:28:26> 'int' '+'
+// CXX-NEXT:         | | |-CallExpr [[ADDR_36:0x[a-z0-9]*]] <line:27:10, col:25> 'int'
+// CXX-NEXT:         | | | |-ImplicitCastExpr [[ADDR_37:0x[a-z0-9]*]] <col:10> 'int (*)(int (*)({{.*}}))' <FunctionToPointerDecay>
+// CXX-NEXT:         | | | | `-DeclRefExpr [[ADDR_38:0x[a-z0-9]*]] <col:10> 'int (int (*)({{.*}}))' {{.*}}Function [[ADDR_23]] 'test' 'int (int (*)({{.*}}))'
+// CXX-NEXT:         | | | `-ImplicitCastExpr [[ADDR_39:0x[a-z0-9]*]] <col:15> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CXX-NEXT:         | | |   `-DeclRefExpr [[ADDR_40:0x[a-z0-9]*]] <col:15> 'int ({{.*}})' {{.*}}Function [[ADDR_18]] 'also_after' 'int ({{.*}})'
+// CXX-NEXT:         | | `-CallExpr [[ADDR_41:0x[a-z0-9]*]] <line:28:10, col:26> 'int'
+// CXX-NEXT:         | |   |-ImplicitCastExpr [[ADDR_42:0x[a-z0-9]*]] <col:10> 'int (*)(int (*)({{.*}}))' <FunctionToPointerDecay>
+// CXX-NEXT:         | |   | `-DeclRefExpr [[ADDR_43:0x[a-z0-9]*]] <col:10> 'int (int (*)({{.*}}))' {{.*}}Function [[ADDR_23]] 'test' 'int (int (*)({{.*}}))'
+// CXX-NEXT:         | |   `-ImplicitCastExpr [[ADDR_44:0x[a-z0-9]*]] <col:15> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CXX-NEXT:         | |     `-DeclRefExpr [[ADDR_45:0x[a-z0-9]*]] <col:15> 'int ({{.*}})' {{.*}}Function [[ADDR_11]] 'also_before' 'int ({{.*}})'
+// CXX-NEXT:         | `-CallExpr [[ADDR_46:0x[a-z0-9]*]] <line:29:10, col:26> 'int'
+// CXX-NEXT:         |   |-ImplicitCastExpr [[ADDR_47:0x[a-z0-9]*]] <col:10> 'int (*)(int (*)({{.*}}))' <FunctionToPointerDecay>
+// CXX-NEXT:         |   | `-DeclRefExpr [[ADDR_48:0x[a-z0-9]*]] <col:10> 'int (int (*)({{.*}}))' {{.*}}Function [[ADDR_23]] 'test' 'int (int (*)({{.*}}))'
+// CXX-NEXT:         |   `-UnaryOperator [[ADDR_49:0x[a-z0-9]*]] <col:15, col:16> 'int (*)({{.*}})' prefix '&' cannot overflow
+// CXX-NEXT:         |     `-DeclRefExpr [[ADDR_50:0x[a-z0-9]*]] <col:16> 'int ({{.*}})' {{.*}}Function [[ADDR_18]] 'also_after' 'int ({{.*}})'
+// CXX-NEXT:         `-CallExpr [[ADDR_51:0x[a-z0-9]*]] <line:30:10, col:27> 'int'
+// CXX-NEXT:           |-ImplicitCastExpr [[ADDR_52:0x[a-z0-9]*]] <col:10> 'int (*)(int (*)({{.*}}))' <FunctionToPointerDecay>
+// CXX-NEXT:           | `-DeclRefExpr [[ADDR_53:0x[a-z0-9]*]] <col:10> 'int (int (*)({{.*}}))' {{.*}}Function [[ADDR_23]] 'test' 'int (int (*)({{.*}}))'
+// CXX-NEXT:           `-UnaryOperator [[ADDR_54:0x[a-z0-9]*]] <col:15, col:16> 'int (*)({{.*}})' prefix '&' cannot overflow
+// CXX-NEXT:             `-DeclRefExpr [[ADDR_55:0x[a-z0-9]*]] <col:16> 'int ({{.*}})' {{.*}}Function [[ADDR_11]] 'also_before' 'int ({{.*}})'
+
+// C:      |-FunctionDecl [[ADDR_0:0x[a-z0-9]*]] <{{.*}}, line:7:1> line:5:5 used also_before 'int ({{.*}})'
+// C-NEXT: | `-CompoundStmt [[ADDR_1:0x[a-z0-9]*]] <col:23, line:7:1>
+// C-NEXT: |   `-ReturnStmt [[ADDR_2:0x[a-z0-9]*]] <line:6:3, col:10>
+// C-NEXT: |     `-IntegerLiteral [[ADDR_3:0x[a-z0-9]*]] <col:10> 'int' 0
+// C-NEXT: |-FunctionDecl [[ADDR_4:0x[a-z0-9]*]] <line:10:1, col:20> col:5 implicit used also_after 'int ({{.*}})'
+// C-NEXT: | `-OMPDeclareVariantAttr [[ADDR_5:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(llvm)}
+// C-NEXT: |   `-DeclRefExpr [[ADDR_6:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_7:0x[a-z0-9]*]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// C-NEXT: |-FunctionDecl [[ADDR_7]] <col:1, line:12:1> line:10:1 also_after[implementation={vendor(llvm)}] 'int ({{.*}})'
+// C-NEXT: | `-CompoundStmt [[ADDR_8:0x[a-z0-9]*]] <col:22, line:12:1>
+// C-NEXT: |   `-ReturnStmt [[ADDR_9:0x[a-z0-9]*]] <line:11:3, col:10>
+// C-NEXT: |     `-IntegerLiteral [[ADDR_10:0x[a-z0-9]*]] <col:10> 'int' 1
+// C-NEXT: |-FunctionDecl [[ADDR_11:0x[a-z0-9]*]] prev [[ADDR_0]] <line:13:1, col:21> col:5 implicit used also_before 'int ({{.*}})'
+// C-NEXT: | `-OMPDeclareVariantAttr [[ADDR_12:0x[a-z0-9]*]] <<invalid sloc>> Implicit implementation={vendor(llvm)}
+// C-NEXT: |   `-DeclRefExpr [[ADDR_13:0x[a-z0-9]*]] <col:1> 'int ({{.*}})' Function [[ADDR_14:0x[a-z0-9]*]] 'also_before[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// C-NEXT: |-FunctionDecl [[ADDR_14]] <col:1, line:15:1> line:13:1 also_before[implementation={vendor(llvm)}] 'int ({{.*}})'
+// C-NEXT: | `-CompoundStmt [[ADDR_15:0x[a-z0-9]*]] <col:23, line:15:1>
+// C-NEXT: |   `-ReturnStmt [[ADDR_16:0x[a-z0-9]*]] <line:14:3, col:10>
+// C-NEXT: |     `-IntegerLiteral [[ADDR_17:0x[a-z0-9]*]] <col:10> 'int' 2
+// C-NEXT: |-FunctionDecl [[ADDR_18:0x[a-z0-9]*]] prev [[ADDR_4]] <line:18:1, line:20:1> line:18:5 used also_after 'int ({{.*}})'
+// C-NEXT: | |-CompoundStmt [[ADDR_19:0x[a-z0-9]*]] <col:22, line:20:1>
+// C-NEXT: | | `-ReturnStmt [[ADDR_20:0x[a-z0-9]*]] <line:19:3, col:10>
+// C-NEXT: | |   `-IntegerLiteral [[ADDR_21:0x[a-z0-9]*]] <col:10> 'int' 0
+// C-NEXT: | `-OMPDeclareVariantAttr [[ADDR_22:0x[a-z0-9]*]] <<invalid sloc>> Inherited Implicit implementation={vendor(llvm)}
+// C-NEXT: |   `-DeclRefExpr [[ADDR_6]] <line:10:1> 'int ({{.*}})' Function [[ADDR_7]] 'also_after[implementation={vendor(llvm)}]' 'int ({{.*}})'
+// C-NEXT: |-FunctionDecl [[ADDR_23:0x[a-z0-9]*]] <line:22:1, line:24:1> line:22:5 used test 'int (int (*)({{.*}}))'
+// C-NEXT: | |-ParmVarDecl [[ADDR_24:0x[a-z0-9]*]] <col:10, col:24> col:16 used fd 'int (*)({{.*}})'
+// C-NEXT: | `-CompoundStmt [[ADDR_25:0x[a-z0-9]*]] <col:27, line:24:1>
+// C-NEXT: |   `-ReturnStmt [[ADDR_26:0x[a-z0-9]*]] <line:23:3, col:13>
+// C-NEXT: |     `-CallExpr [[ADDR_27:0x[a-z0-9]*]] <col:10, col:13> 'int'
+// C-NEXT: |       `-ImplicitCastExpr [[ADDR_28:0x[a-z0-9]*]] <col:10> 'int (*)({{.*}})' <LValueToRValue>
+// C-NEXT: |         `-DeclRefExpr [[ADDR_29:0x[a-z0-9]*]] <col:10> 'int (*)({{.*}})' {{.*}}ParmVar [[ADDR_24]] 'fd' 'int (*)({{.*}})'
+// C-NEXT: `-FunctionDecl [[ADDR_30:0x[a-z0-9]*]] <line:25:1, line:31:1> line:25:5 main 'int ({{.*}})'
+// C-NEXT:   `-CompoundStmt [[ADDR_31:0x[a-z0-9]*]] <col:12, line:31:1>
+// C-NEXT:     `-ReturnStmt [[ADDR_32:0x[a-z0-9]*]] <line:27:3, line:30:27>
+// C-NEXT:       `-BinaryOperator [[ADDR_33:0x[a-z0-9]*]] <line:27:10, line:30:27> 'int' '+'
+// C-NEXT:         |-BinaryOperator [[ADDR_34:0x[a-z0-9]*]] <line:27:10, line:29:26> 'int' '+'
+// C-NEXT:         | |-BinaryOperator [[ADDR_35:0x[a-z0-9]*]] <line:27:10, line:28:26> 'int' '+'
+// C-NEXT:         | | |-CallExpr [[ADDR_36:0x[a-z0-9]*]] <line:27:10, col:25> 'int'
+// C-NEXT:         | | | |-ImplicitCastExpr [[ADDR_37:0x[a-z0-9]*]] <col:10> 'int (*)(int (*)({{.*}}))' <FunctionToPointerDecay>
+// C-NEXT:         | | | | `-DeclRefExpr [[ADDR_38:0x[a-z0-9]*]] <col:10> 'int (int (*)({{.*}}))' Function [[ADDR_23]] 'test' 'int (int (*)({{.*}}))'
+// C-NEXT:         | | | `-ImplicitCastExpr [[ADDR_39:0x[a-z0-9]*]] <col:15> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// C-NEXT:         | | |   `-DeclRefExpr [[ADDR_40:0x[a-z0-9]*]] <col:15> 'int ({{.*}})' Function [[ADDR_18]] 'also_after' 'int ({{.*}})'
+// C-NEXT:         | | `-CallExpr [[ADDR_41:0x[a-z0-9]*]] <line:28:10, col:26> 'int'
+// C-NEXT:         | |   |-ImplicitCastExpr [[ADDR_42:0x[a-z0-9]*]] <col:10> 'int (*)(int (*)({{.*}}))' <FunctionToPointerDecay>
+// C-NEXT:         | |   | `-DeclRefExpr [[ADDR_43:0x[a-z0-9]*]] <col:10> 'int (int (*)({{.*}}))' Function [[ADDR_23]] 'test' 'int (int (*)({{.*}}))'
+// C-NEXT:         | |   `-ImplicitCastExpr [[ADDR_44:0x[a-z0-9]*]] <col:15> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// C-NEXT:         | |     `-DeclRefExpr [[ADDR_45:0x[a-z0-9]*]] <col:15> 'int ({{.*}})' Function [[ADDR_11]] 'also_before' 'int ({{.*}})'
+// C-NEXT:         | `-CallExpr [[ADDR_46:0x[a-z0-9]*]] <line:29:10, col:26> 'int'
+// C-NEXT:         |   |-ImplicitCastExpr [[ADDR_47:0x[a-z0-9]*]] <col:10> 'int (*)(int (*)({{.*}}))' <FunctionToPointerDecay>
+// C-NEXT:         |   | `-DeclRefExpr [[ADDR_48:0x[a-z0-9]*]] <col:10> 'int (int (*)({{.*}}))' Function [[ADDR_23]] 'test' 'int (int (*)({{.*}}))'
+// C-NEXT:         |   `-UnaryOperator [[ADDR_49:0x[a-z0-9]*]] <col:15, col:16> 'int (*)({{.*}})' prefix '&' cannot overflow
+// C-NEXT:         |     `-DeclRefExpr [[ADDR_50:0x[a-z0-9]*]] <col:16> 'int ({{.*}})' Function [[ADDR_18]] 'also_after' 'int ({{.*}})'
+// C-NEXT:         `-CallExpr [[ADDR_51:0x[a-z0-9]*]] <line:30:10, col:27> 'int'
+// C-NEXT:           |-ImplicitCastExpr [[ADDR_52:0x[a-z0-9]*]] <col:10> 'int (*)(int (*)({{.*}}))' <FunctionToPointerDecay>
+// C-NEXT:           | `-DeclRefExpr [[ADDR_53:0x[a-z0-9]*]] <col:10> 'int (int (*)({{.*}}))' Function [[ADDR_23]] 'test' 'int (int (*)({{.*}}))'
+// C-NEXT:           `-UnaryOperator [[ADDR_54:0x[a-z0-9]*]] <col:15, col:16> 'int (*)({{.*}})' prefix '&' cannot overflow
+// C-NEXT:             `-DeclRefExpr [[ADDR_55:0x[a-z0-9]*]] <col:16> 'int ({{.*}})' Function [[ADDR_11]] 'also_before' 'int ({{.*}})'

clang/test/AST/ast-dump-openmp-begin-declare-variant_decl_1.c

+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s       | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify -ast-dump %s -x c++| FileCheck %s
+// expected-no-diagnostics
+// FIXME: We have to improve the warnings here as nothing is impacted by the declare variant.
+int also_before(void) {
+  return 0;
+}
+
+#pragma omp begin declare variant match(device={kind(cpu)})
+int also_before(void);
+#pragma omp end declare variant
+#pragma omp begin declare variant match(implementation={vendor(score(100):llvm)})
+int also_after(void);
+#pragma omp end declare variant
+#pragma omp begin declare variant match(implementation={vendor(score(0):llvm)})
+int also_before(void);
+#pragma omp end declare variant
+
+int also_after(void) {
+  return 0;
+}
+
+int test() {
+  // Should return 0.
+  return also_after() + also_before();
+}
+
+// Make sure:
+//  - we do see the ast nodes for the cpu kind
+//  - we do see the ast nodes for the llvm vendor
+//  - we pick the right callees
+
+// CHECK:      |-FunctionDecl [[ADDR_0:0x[a-z0-9]*]] <{{.*}}, line:7:1> line:5:5 used also_before 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_1:0x[a-z0-9]*]] <col:23, line:7:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_2:0x[a-z0-9]*]] <line:6:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_3:0x[a-z0-9]*]] <col:10> 'int' 0
+// CHECK-NEXT: |-FunctionDecl [[ADDR_4:0x[a-z0-9]*]] prev [[ADDR_0]] <line:10:1, col:21> col:5 used also_before 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_5:0x[a-z0-9]*]] <line:13:1, col:20> col:5 used also_after 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_6:0x[a-z0-9]*]] prev [[ADDR_4]] <line:16:1, col:21> col:5 used also_before 'int ({{.*}})'
+// CHECK-NEXT: |-FunctionDecl [[ADDR_7:0x[a-z0-9]*]] prev [[ADDR_5]] <line:19:1, line:21:1> line:19:5 used also_after 'int ({{.*}})'
+// CHECK-NEXT: | `-CompoundStmt [[ADDR_8:0x[a-z0-9]*]] <col:22, line:21:1>
+// CHECK-NEXT: |   `-ReturnStmt [[ADDR_9:0x[a-z0-9]*]] <line:20:3, col:10>
+// CHECK-NEXT: |     `-IntegerLiteral [[ADDR_10:0x[a-z0-9]*]] <col:10> 'int' 0
+// CHECK-NEXT: `-FunctionDecl [[ADDR_11:0x[a-z0-9]*]] <line:23:1, line:26:1> line:23:5 test 'int ({{.*}})'
+// CHECK-NEXT:   `-CompoundStmt [[ADDR_12:0x[a-z0-9]*]] <col:12, line:26:1>
+// CHECK-NEXT:     `-ReturnStmt [[ADDR_13:0x[a-z0-9]*]] <line:25:3, col:37>
+// CHECK-NEXT:       `-BinaryOperator [[ADDR_14:0x[a-z0-9]*]] <col:10, col:37> 'int' '+'
+// CHECK-NEXT:         |-CallExpr [[ADDR_15:0x[a-z0-9]*]] <col:10, col:21> 'int'
+// CHECK-NEXT:         | `-ImplicitCastExpr [[ADDR_16:0x[a-z0-9]*]] <col:10> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:         |   `-DeclRefExpr [[ADDR_17:0x[a-z0-9]*]] <col:10> 'int ({{.*}})' {{.*}}Function [[ADDR_7]] 'also_after' 'int ({{.*}})'
+// CHECK-NEXT:         `-CallExpr [[ADDR_18:0x[a-z0-9]*]] <col:25, col:37> 'int'
+// CHECK-NEXT:           `-ImplicitCastExpr [[ADDR_19:0x[a-z0-9]*]] <col:25> 'int (*)({{.*}})' <FunctionToPointerDecay>
+// CHECK-NEXT:             `-DeclRefExpr [[ADDR_20:0x[a-z0-9]*]] <col:25> 'int ({{.*}})' {{.*}}Function [[ADDR_6]] 'also_before' 'int ({{.*}})'

clang/test/OpenMP/begin_declare_variant_elided_range_withouth_end.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-simd -x c -std=c99 -fms-extensions -Wno-pragma-pack %s
 
-#pragma omp begin declare variant match(device={kind(cpu)})
+#pragma omp begin declare variant match(device={kind(gpu)}) // expected-note {{to match this '#pragma omp begin declare variant'}}
 int also_before(void) {
   return 0;
 }
-#pragma omp end declare variant
 
-#pragma omp begin declare variant match(device={kind(gpu)}) // expected-note {{to match this '#pragma omp begin declare variant'}}
-// The matching end is missing. Since the device clause is not matching we will
-// cause us to elide the rest of the file and emit and error.
+#pragma omp begin declare variant match(device={kind(gpu)})
 int also_after(void) {
   return 2;
 }
 int also_before(void) {
   return 2;
 }
+#pragma omp end declare variant
 
 
 #pragma omp begin declare variant match(device={kind(fpga)})
 
 This text is never parsed!
 
 #pragma omp end declare variant
 
 int also_after(void) {
   return 0;
 }
 
 int test() {
   return also_after() + also_before();
 } // expected-error {{expected '#pragma omp end declare variant'}}

clang/test/OpenMP/begin_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-simd -x c -std=c99 -fms-extensions -Wno-pragma-pack %s
 
 
 #pragma omp begin // expected-error {{expected an OpenMP directive}}
-#pragma omp end declare variant
+#pragma omp end declare variant // expected-error {{'#pragma omp end declare variant' with no matching '#pragma omp begin declare variant'}}
 #pragma omp begin declare // expected-error {{expected an OpenMP directive}}
-#pragma omp end declare variant
+#pragma omp end declare variant // expected-error {{'#pragma omp end declare variant' with no matching '#pragma omp begin declare variant'}}
 #pragma omp begin variant // expected-error {{expected an OpenMP directive}}
-#pragma omp end declare variant
+#pragma omp end declare variant // expected-error {{'#pragma omp end declare variant' with no matching '#pragma omp begin declare variant'}}
 #pragma omp variant begin // expected-error {{expected an OpenMP directive}}
 #pragma omp declare variant end // expected-error {{function declaration is expected after 'declare variant' directive}}
 #pragma omp begin declare variant // expected-error {{expected 'match' clause on 'omp declare variant' directive}}
 #pragma omp end declare variant
 // TODO: Issue an error message
-#pragma omp end declare variant
-#pragma omp end declare variant
-#pragma omp end declare variant
-#pragma omp end declare variant
+#pragma omp end declare variant // expected-error {{'#pragma omp end declare variant' with no matching '#pragma omp begin declare variant'}}
+#pragma omp end declare variant // expected-error {{'#pragma omp end declare variant' with no matching '#pragma omp begin declare variant'}}
+#pragma omp end declare variant // expected-error {{'#pragma omp end declare variant' with no matching '#pragma omp begin declare variant'}}
+#pragma omp end declare variant // expected-error {{'#pragma omp end declare variant' with no matching '#pragma omp begin declare variant'}}
 
 int foo(void);
 const int var;
 
 #pragma omp begin declare variant // expected-error {{expected 'match' clause on 'omp declare variant' directive}}
 #pragma omp end declare variant
 #pragma omp begin declare variant xxx // expected-error {{expected 'match' clause on 'omp declare variant' directive}}
 #pragma omp end declare variant

 #pragma omp end declare variant
 #pragma omp begin declare variant match(device = {kind(score(foo): cpu)}) // expected-warning {{the context selector 'kind' in the context set 'device' cannot have a score ('foo'); score ignored}}
 #pragma omp end declare variant
 #pragma omp begin declare variant match(device = {kind(score(foo()): cpu)}) // expected-warning {{the context selector 'kind' in the context set 'device' cannot have a score ('foo()'); score ignored}}
 #pragma omp end declare variant
 #pragma omp begin declare variant match(device = {kind(score(<expr>): cpu)}) // expected-error {{expected expression}} expected-error {{use of undeclared identifier 'expr'}} expected-error {{expected expression}} expected-warning {{the context selector 'kind' in the context set 'device' cannot have a score ('<invalid>'); score ignored}}
 #pragma omp end declare variant
 
+#pragma omp begin declare variant match(device={kind(cpu)})
+static int defined_twice_a(void) { // expected-note {{previous definition is here}}
+  return 0;
+}
+int defined_twice_b(void) { // expected-note {{previous definition is here}}
+  return 0;
+}
+inline int defined_twice_c(void) { // expected-note {{previous definition is here}}
+  return 0;
+}
+#pragma omp end declare variant
+#pragma omp begin declare variant match(device={kind(cpu)})
+static int defined_twice_a(void) { // expected-error {{redefinition of 'defined_twice_a[device={kind(cpu)}]'}}
+  return 1;
+}
+int defined_twice_b(void) { // expected-error {{redefinition of 'defined_twice_b[device={kind(cpu)}]'}}
+  return 1;
+}
+inline int defined_twice_c(void) { // expected-error {{redefinition of 'defined_twice_c[device={kind(cpu)}]'}}
+  return 1;
+}
+#pragma omp end declare variant
 
 
 // TODO: Issue an error message
 #pragma omp begin declare variant match(device={kind(cpu)})
 // The matching end is missing. Since the device clause is matching we will
 // emit and error.
 int also_before(void) {
   return 0;

clang/test/OpenMP/begin_declare_variant_range_withouth_end.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-simd -x c -std=c99 -fms-extensions -Wno-pragma-pack %s
 
+// TODO: Issue an eror message as the end is missing
+// expected-no-diagnostics
 #pragma omp begin declare variant match(device={kind(cpu)})
 int also_before(void) {
   return 0;
 }
-#pragma omp end declare variant
 
-#pragma omp begin declare variant match(device={kind(gpu)}) // expected-note {{to match this '#pragma omp begin declare variant'}}
-// The matching end is missing. Since the device clause is not matching we will
-// cause us to elide the rest of the file and emit and error.
+#pragma omp begin declare variant match(device={kind(gpu)})
 int also_after(void) {
   return 2;
 }
 int also_before(void) {
   return 2;
 }
+#pragma omp end declare variant
 
 
 #pragma omp begin declare variant match(device={kind(fpga)})
 
 This text is never parsed!
 
 #pragma omp end declare variant
 
 int also_after(void) {
   return 0;
 }
 
 int test() {
   return also_after() + also_before();
-} // expected-error {{expected '#pragma omp end declare variant'}}
+}

clang/test/OpenMP/begin_declare_variant_using_messages.cpp

+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -verify %s
+
+namespace BEFORE_AND_1 {
+void before_and_1();
+}
+namespace AFTER_AND_2 {
+void after_and_2(); // expected-note {{'AFTER_AND_2::after_and_2' declared here}} expected-note {{'AFTER_AND_2::after_and_2' declared here}}
+}
+namespace ONLY_1 {
+void only_1(); // expected-note {{'ONLY_1::only_1' declared here}}
+}
+namespace BEFORE_1_AND_2 {
+void before_1_and_2();
+}
+
+using BEFORE_1_AND_2::before_1_and_2;
+using BEFORE_AND_1::before_and_1;
+
+void test_before() {
+  before_and_1();
+  after_and_2(); // expected-error {{use of undeclared identifier 'after_and_2'; did you mean 'AFTER_AND_2::after_and_2'?}}
+  only_1(); // expected-error {{use of undeclared identifier 'only_1'; did you mean 'ONLY_1::only_1'?}}
+  before_1_and_2();
+}
+
+#pragma omp begin declare variant match(implementation = {vendor(llvm)})
+using BEFORE_1_AND_2::before_1_and_2;
+using BEFORE_AND_1::before_and_1;
+using ONLY_1::only_1;
+void test_1() {
+  before_and_1();
+  after_and_2(); // expected-error {{use of undeclared identifier 'after_and_2'; did you mean 'AFTER_AND_2::after_and_2'?}}
+  only_1();
+  before_1_and_2();
+}
+#pragma omp end declare variant
+
+#pragma omp begin declare variant match(implementation = {vendor(llvm)})
+using AFTER_AND_2::after_and_2;
+using BEFORE_1_AND_2::before_1_and_2;
+void test_2() {
+  before_and_1();
+  after_and_2();
+  only_1();
+  before_1_and_2();
+}
+#pragma omp end declare variant
+
+void test_after() {
+  before_and_1();
+  after_and_2();
+  only_1();
+  before_1_and_2();
+}
+
+using AFTER_AND_2::after_and_2;
+
+// Make sure:
+//  - we do not see the ast nodes for the gpu kind
+//  - we do not choke on the text in the kind(fpga) guarded scopes
+//  - we pick the right cbefore_1_and_2ees

clang/test/OpenMP/declare_variant_ast_print.cpp

 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: }
+// CHECK-NEXT: template<> void h_ref<float>(float *hp, float *hp2, float *hq, float *lin)
 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(unknown)}, device={kind(nohost)})
 // CHECK-NEXT: #pragma omp declare variant(h_ref<C>) match(implementation={vendor(llvm)}, device={kind(gpu)})
 // CHECK-NEXT: template <class C> void h(C *hp, C *hp2, C *hq, C *lin) {
 // CHECK-NEXT: }

clang/test/OpenMP/declare_variant_device_kind_codegen.cpp

 
 // 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: alias
+
 // 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: define {{.*}}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}}
+// CHECK-NOT: ret i32 {{4|81|84}}
 
 #ifndef HEADER
 #define HEADER
 
 #pragma omp declare target
 #ifdef HOST
 #define SUBSET host
 #define CORRECT host, cpu

clang/test/OpenMP/declare_variant_implementation_vendor_codegen.cpp

-// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple %itanium_abi_triple -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple %itanium_abi_triple -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope | FileCheck %s --implicit-check-not='ret i32 {{6|7|8|9|10|12|13|14|15|19|21|22|23|24}}'
 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -emit-pch -o %t -fopenmp-version=50 %s
-// RUN: %clang_cc1 -fopenmp -x c++ -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 | FileCheck %s --implicit-check-not='ret i32 {{6|7|8|9|10|12|13|14|15|19|21|22|23|24}}'
 // 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}}
+// CHECK-DAG:  ret i32 2
+// CHECK-DAG:  ret i32 3
+// CHECK-DAG:  ret i32 4
+// CHECK-DAG:  ret i32 5
+// CHECK-DAG:  ret i32 11
+// CHECK-DAG:  ret i32 16
+// CHECK-DAG:  ret i32 17
+// CHECK-DAG:  ret i32 18
+// CHECK-DAG:  ret i32 19
+// CHECK-DAG:  ret i32 20
+// CHECK-DAG:  ret i32 25
+// CHECK-DAG:  ret i32 26
+// CHECK-DAG:  ret i32 27
+// CHECK-DAG:  ret i32 28
+// CHECK-DAG:  ret i32 29
 
 #ifndef HEADER
 #define HEADER
 
 int foo() { return 2; }
 
 #pragma omp declare variant(foo) match(implementation = {vendor(llvm)})
-int bar() { return 1; }
+int bar() { return 3; }
 
 int bazzz();
 #pragma omp declare variant(bazzz) match(implementation = {vendor(llvm)})
-int baz() { return 1; }
+int baz() { return 4; }
 
 int test();
 #pragma omp declare variant(test) match(implementation = {vendor(llvm)})
-int call() { return 1; }
+int call() { return 5; }
 
 static int stat_unused_();
 #pragma omp declare variant(stat_unused_) match(implementation = {vendor(llvm)})
-static int stat_unused() { return 1; }
+static int stat_unused() { return 6; }
 
 static int stat_used_();
 #pragma omp declare variant(stat_used_) match(implementation = {vendor(llvm)})
-static int stat_used() { return 1; }
+static int stat_used() { return 7; }
 
 int main() { return bar() + baz() + call() + stat_used(); }
 
-int test() { return 3; }
-static int stat_unused_() { return 4; }
-static int stat_used_() { return 5; }
+int test() { return 8; }
+static int stat_unused_() { return 9; }
+static int stat_used_() { return 10; }
 
 struct SpecialFuncs {
   void vd() {}
   SpecialFuncs();
   ~SpecialFuncs();
 
-  int method_() { return 6; }
+  int method_() { return 11; }
 #pragma omp declare variant(SpecialFuncs::method_)                             \
     match(implementation = {vendor(llvm)})
-  int method() { return 1; }
+  int method() { return 12; }
 #pragma omp declare variant(SpecialFuncs::method_)                             \
     match(implementation = {vendor(llvm)})
   int Method();
 } s;
 
-int SpecialFuncs::Method() { return 1; }
+int SpecialFuncs::Method() { return 13; }
 
 struct SpecSpecialFuncs {
   void vd() {}
   SpecSpecialFuncs();
   ~SpecSpecialFuncs();
 
   int method_();
 #pragma omp declare variant(SpecSpecialFuncs::method_)                         \
     match(implementation = {vendor(llvm)})
-  int method() { return 1; }
+  int method() { return 14; }
 #pragma omp declare variant(SpecSpecialFuncs::method_)                         \
     match(implementation = {vendor(llvm)})
   int Method();
 } s1;
 
-int SpecSpecialFuncs::method_() { return 7; }
-int SpecSpecialFuncs::Method() { return 1; }
+int SpecSpecialFuncs::method_() { return 15; }
+int SpecSpecialFuncs::Method() { return 16; }
 
 void xxx() {
   (void)s.method();
   (void)s1.method();
 }
 
-int prio() { return 81; }
-int prio1() { return 82; }
+int prio() { return 17; }
+int prio1() { return 18; }
 
 #pragma omp declare variant(prio) match(implementation = {vendor(llvm)})
 #pragma omp declare variant(prio1) match(implementation = {vendor(score(1): llvm)})
-int prio_() { return 1; }
+int prio_() { return 19; }
 
-static int prio2() { return 83; }
-static int prio3() { return 84; }
-static int prio4() { return 84; }
+static int prio2() { return 20; }
+static int prio3() { return 21; }
+static int prio4() { return 22; }
 
 #pragma omp declare variant(prio4) match(implementation = {vendor(score(3): llvm)})
 #pragma omp declare variant(prio2) match(implementation = {vendor(score(5): llvm)})
 #pragma omp declare variant(prio3) match(implementation = {vendor(score(1): llvm)})
-static int prio1_() { return 1; }
+static int prio1_() { return 23; }
 
 int int_fn() { return prio1_(); }
 
-int fn_linkage_variant() { return 85; }
+int fn_linkage_variant() { return 24; }
 extern "C" {
 #pragma omp declare variant(fn_linkage_variant) match(implementation = {vendor(llvm)})
-int fn_linkage() { return 1; }
+int fn_linkage() { return 25; }
 }
 
-extern "C" int fn_linkage_variant1() { return 86; }
+extern "C" int fn_linkage_variant1() { return 26; }
 #pragma omp declare variant(fn_linkage_variant1) match(implementation = {vendor(llvm)})
-int fn_linkage1() { return 1; }
+int fn_linkage1() { return 27; }
 
-int fn_variant2() { return 1; }
+int fn_variant2() { return 28; }
 #pragma omp declare variant(fn_variant2) match(implementation = {vendor(llvm, ibm)})
-int fn2() { return 87; }
+int fn2() { return 29; }
 
 #endif // HEADER

clang/test/OpenMP/declare_variant_messages.c

 
 #pragma omp declare variant(foo) match(xxx={}) // expected-error {{function declaration is expected after 'declare variant' directive}}
 #pragma omp declare variant(foo) match(xxx={}) // expected-error {{function declaration is expected after 'declare variant' directive}}
 #pragma options align=packed
 int main();
 
 
 
-#pragma omp declare variant(foo) match(xxx={}) // expected-error {{function declaration is expected after 'declare variant' directive}}
-#pragma omp declare variant(foo) match(xxx={}) // expected-error {{function declaration is expected after 'declare variant' directive}}
+#pragma omp declare variant(foo) match(implementation={vendor(llvm)}) // expected-error {{function declaration is expected after 'declare variant' directive}}
+#pragma omp declare variant(foo) match(implementation={vendor(llvm)}) // expected-error {{function declaration is expected after 'declare variant' directive}}
 #pragma init_seg(compiler)
 int main();
 
 
 #pragma omp declare variant(foo) match(xxx={}) // expected-error {{single declaration is expected after 'declare variant' directive}} expected-warning {{'xxx' is not a valid context set in a `declare variant`; set ignored}} expected-note {{context set options are: 'construct' 'device' 'implementation' 'user'}} expected-note {{the ignored set spans until here}}
 int b, c;
 
 int no_proto();

clang/test/OpenMP/declare_variant_messages.cpp

 
 #pragma omp declare variant(foo) match(implementation = {vendor(score(constexpr_fn(3)) : llvm)})
 #pragma omp declare variant(foo) match(implementation = {vendor(score(constexpr_fn(C)) : llvm)})
 #pragma omp declare variant(foo) match(user = {condition(constexpr_fn(1))})
 #pragma omp declare variant(foo) match(user = {condition(constexpr_fn(C))})
 template <int C>
 int score_and_cond_const_inst();
 
+__attribute__((pure)) int pure() { return 0; }
+
+#pragma omp declare variant(pure) match(user = {condition(1)}) // expected-warning {{ignoring return value of function declared with pure attribute}}
+int unused_warning_after_specialization() { return foo(); }
+
 void score_and_cond_inst() {
   score_and_cond_non_const();
   score_and_cond_non_const_inst<8>(); // expected-note {{in instantiation of function template specialization 'score_and_cond_non_const_inst<8>' requested here}}
   score_and_cond_const_inst<9>();
+  unused_warning_after_specialization();
 }
 
 #pragma omp declare variant(barbar <int>) match(implementation = {}) // expected-warning {{expected identifier or string literal describing a context selector; selector skipped}} expected-note {{context selector options are: 'vendor' 'extension' 'unified_address' 'unified_shared_memory' 'reverse_offload' 'dynamic_allocators' 'atomic_default_mem_order'}} expected-note {{the ignored selector spans until here}}
 int a; // expected-error {{'#pragma omp declare variant' can only be applied to functions}}
 
 #pragma omp declare variant(barbar <int>) match(implementation = {}) // expected-warning {{expected identifier or string literal describing a context selector; selector skipped}} expected-note {{context selector options are: 'vendor' 'extension' 'unified_address' 'unified_shared_memory' 'reverse_offload' 'dynamic_allocators' 'atomic_default_mem_order'}} expected-note {{the ignored selector spans until here}}
 #pragma omp threadprivate(a) // expected-error {{'#pragma omp declare variant' can only be applied to functions}}
 int var;

clang/test/OpenMP/declare_variant_mixed_codegen.c

 // RUN: %clang_cc1 -verify -fopenmp -x c -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --check-prefix HOST
 // RUN: %clang_cc1 -fopenmp -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=50 %s
 // RUN: %clang_cc1 -fopenmp -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=50 | FileCheck %s --check-prefix HOST
 // 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 --check-prefix GPU
 // 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 --check-prefix GPU
 // expected-no-diagnostics
 
-// HOST: @base = alias i32 (double), i32 (double)* @hst
 #ifndef HEADER
 #define HEADER
 
 int dev(double i) { return 0; }
 
 int hst(double i) { return 1; }
 
 #pragma omp declare variant(hst) match(device = {kind(host)})
 #pragma omp declare variant(dev) match(device = {kind(gpu)})
 int base();
 
 // HOST-LABEL: define void @foo()
-// HOST: call i32 (double, ...) bitcast (i32 (double)* @base to i32 (double, ...)*)(double -1.000000e+00)
+// HOST: call i32 @hst(double -1.000000e+00)
 // HOST: call i32 @hst(double -2.000000e+00)
-// HOST: call void [[OFFL:@.+_foo_l29]]()
+// HOST: call void [[OFFL:@.+_foo_l28]]()
 void foo() {
   base(-1);
   hst(-2);
 #pragma omp target
   {
     base(-3);
     dev(-4);
   }
 }
 
 // HOST: define {{.*}}void [[OFFL]]()
-// HOST: call i32 (double, ...) bitcast (i32 (double)* @base to i32 (double, ...)*)(double -3.000000e+00)
+// HOST: call i32 @hst(double -3.000000e+00)
 // HOST: call i32 @dev(double -4.000000e+00)
 
-// GPU: define {{.*}}void @__omp_offloading_{{.+}}_foo_l29()
-// GPU: call i32 @base(double -3.000000e+00)
+// GPU: define {{.*}}void @__omp_offloading_{{.+}}_foo_l28()
+// GPU: call i32 @dev(double -3.000000e+00)
 // GPU: call i32 @dev(double -4.000000e+00)
 
-// GPU: define {{.*}}i32 @base(double
-// GPU: ret i32 0
+// GPU-NOT: @base
 // GPU: define {{.*}}i32 @dev(double
 // GPU: ret i32 0
 
 #endif // HEADER

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 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope | FileCheck %s --implicit-check-not='ret i32 {{6|7|8|10|13|15|19|22|23|24}}'
 // 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 -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 --implicit-check-not='ret i32 {{6|7|8|10|13|15|19|22|23|24}}'
+
 // 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-DAG: ret i32 88
-// CHECK-NOT: ret i32 {{1|4|81|84}}
+// CHECK-DAG:  ret i32 2
+// CHECK-DAG:  ret i32 3
+// CHECK-DAG:  ret i32 4
+// CHECK-DAG:  ret i32 5
+// CHECK-DAG:  ret i32 9
+// CHECK-DAG:  ret i32 11
+// CHECK-DAG:  ret i32 12
+// CHECK-DAG:  ret i32 14
+// CHECK-DAG:  ret i32 16
+// CHECK-DAG:  ret i32 17
+// CHECK-DAG:  ret i32 18
+// CHECK-DAG:  ret i32 19
+// CHECK-DAG:  ret i32 20
+// CHECK-DAG:  ret i32 21
+// CHECK-DAG:  ret i32 25
+// CHECK-DAG:  ret i32 26
+// CHECK-DAG:  ret i32 27
+// CHECK-DAG:  ret i32 28
+// CHECK-DAG:  ret i32 29
+// CHECK-DAG:  ret i32 30
+// CHECK-DAG:  ret i32 31
 
 #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 bar() { return 3; }
 
 int bazzz();
 #pragma omp declare variant(bazzz) match(implementation = {vendor(llvm)}, device={kind(host)})
-int baz() { return 1; }
+int baz() { return 4; }
 
 int test();
 #pragma omp declare variant(test) match(implementation = {vendor(llvm)}, device={kind(cpu)})
-int call() { return 1; }
+int call() { return 5; }
 
-static int stat_unused_no_emit() { return 1; }
+static int stat_unused_no_emit() { return 6; }
 static int stat_unused_();
 #pragma omp declare variant(stat_unused_) match(implementation = {vendor(llvm)}, device={kind(cpu)})
 #pragma omp declare variant(stat_unused_no_emit) match(implementation = {vendor(unknown)}, device = {kind(gpu)})
-static int stat_unused() { return 1; }
+static int stat_unused() { return 7; }
 
 static int stat_used_();
 #pragma omp declare variant(stat_used_) match(implementation = {vendor(llvm)}, device={kind(host)})
-static int stat_used() { return 1; }
+static int stat_used() { return 8; }
 
 int main() { return bar() + baz() + call() + stat_used(); }
 
-int test() { return 3; }
-static int stat_unused_() { return 4; }
-static int stat_used_() { return 5; }
+int test() { return 9; }
+static int stat_unused_() { return 10; }
+static int stat_used_() { return 11; }
 
 struct SpecialFuncs {
   void vd() {}
   SpecialFuncs();
   ~SpecialFuncs();
 
-  int method_() { return 6; }
+  int method_() { return 12; }
 #pragma omp declare variant(SpecialFuncs::method_)                             \
     match(implementation = {vendor(llvm)}, device={kind(cpu)})
-  int method() { return 1; }
+  int method() { return 13; }
 #pragma omp declare variant(SpecialFuncs::method_)                             \
     match(implementation = {vendor(llvm)}, device={kind(host)})
   int Method();
 } s;
 
-int SpecialFuncs::Method() { return 1; }
+int SpecialFuncs::Method() { return 14; }
 
 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; }
+  int method() { return 15; }
 #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; }
+int SpecSpecialFuncs::method_() { return 16; }
+int SpecSpecialFuncs::Method() { return 17; }
 
 void xxx() {
   (void)s.method();
   (void)s1.method();
 }
 
-int prio() { return 81; }
-int prio1() { return 82; }
+int prio() { return 18; }
+int prio1() { return 19; }
 
 #pragma omp declare variant(prio1) match(implementation = {vendor(score(2): llvm)}, device={kind(cpu,host)})
 #pragma omp declare variant(prio) match(implementation = {vendor(score(1): llvm)}, device={kind(cpu)})
-int prio_() { return 1; }
+int prio_() { return 20; }
 
-static int prio2() { return 83; }
-static int prio3() { return 84; }
-static int prio4() { return 84; }
+static int prio2() { return 21; }
+static int prio3() { return 22; }
+static int prio4() { return 23; }
 
 #pragma omp declare variant(prio4) match(implementation = {vendor(score(5): llvm)})
 #pragma omp declare variant(prio2) match(implementation = {vendor(score(8): llvm)}, device={kind(cpu,host)})
 #pragma omp declare variant(prio3) match(implementation = {vendor(score(7): llvm)}, device={kind(cpu)})
-static int prio1_() { return 1; }
+static int prio1_() { return 24; }
 
 int int_fn() { return prio1_(); }
 
-int fn_linkage_variant() { return 85; }
+int fn_linkage_variant() { return 25; }
 extern "C" {
 #pragma omp declare variant(fn_linkage_variant) match(implementation = {vendor(llvm)}, device={kind(cpu)})
-int fn_linkage() { return 1; }
+int fn_linkage() { return 26; }
 }
 
-extern "C" int fn_linkage_variant1() { return 86; }
+extern "C" int fn_linkage_variant1() { return 27; }
 #pragma omp declare variant(fn_linkage_variant1) match(implementation = {vendor(llvm)}, device={kind(host)})
-int fn_linkage1() { return 1; }
+int fn_linkage1() { return 28; }
 
-int fn_variant2() { return 1; }
+int fn_variant2() { return 29; }
 #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; }
+int fn2() { return 30; }
 
 #pragma omp declare variant(stat_unused_no_emit) match(implementation = {vendor(unknown)}, device = {kind(gpu)})
 template <typename T>
-static T stat_unused_T() { return 88; }
+static T stat_unused_T() { return 31; }
 
 int bazzzzzzzz() {
   return stat_unused_T<int>();
 }
 
 #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 -o - -fopenmp-version=50 -DGPU | FileCheck %s --implicit-check-not='ret i32 {{6|7|9|10|12|14|17|18|20|21|22|23|24|26}}'
 // 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 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 {{6|7|9|10|12|14|17|18|20|21|22|23|24|26}}'
 
 // 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 -o - -fopenmp-version=50 -DNOHOST | FileCheck %s --implicit-check-not='ret i32 {{6|7|9|10|12|14|17|18|20|21|22|23|24|26}}'
 // 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}}'
+// 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 {{6|7|9|10|12|14|17|18|20|21|22|23|24|26}}'
 // 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}}
+// CHECK-DAG: ret i32 8
+// CHECK-DAG: ret i32 11
+// CHECK-DAG: ret i32 13
+// CHECK-DAG: ret i32 15
+// CHECK-DAG: ret i32 16
+// CHECK-DAG: ret i32 19
+// CHECK-DAG: ret i32 25
+
+// Outputs for function members checked via implicit filecheck flag
+
 
 #ifndef HEADER
 #define HEADER
 
 #ifdef GPU
 #define SUBSET gpu
 #define CORRECT nohost, gpu
 #define WRONG cpu, gpu
 #endif // GPU
 #ifdef NOHOST
 #define SUBSET nohost
 #define CORRECT 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; }
+int bar() { return 3; }
 
 #pragma omp declare variant(bazzz) match(device = {kind(CORRECT)})
-int baz() { return 1; }
+int baz() { return 4; }
 
 #pragma omp declare variant(test) match(device = {kind(CORRECT)})
-int call() { return 1; }
+int call() { return 5; }
 
 #pragma omp declare variant(stat_unused_) match(device = {kind(CORRECT)})
-static int stat_unused() { return 1; }
+static int stat_unused() { return 6; }
 
 #pragma omp declare variant(stat_used_) match(device = {kind(CORRECT)})
-static int stat_used() { return 1; }
+static int stat_used() { return 7; }
 
 #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; }
+int test() { return 8; }
+static int stat_unused_() { return 9; }
+static int stat_used_() { return 10; }
 
 #pragma omp declare target
 
 struct SpecialFuncs {
   void vd() {}
   SpecialFuncs();
   ~SpecialFuncs();
-
-  int method_() { return 6; }
+  int method_() { return 11; }
 #pragma omp declare variant(SpecialFuncs::method_) \
     match(device = {kind(CORRECT)})
-  int method() { return 1; }
+  int method() { return 12; }
 #pragma omp declare variant(SpecialFuncs::method_) \
     match(device = {kind(CORRECT)})
   int Method();
 } s;
 
-int SpecialFuncs::Method() { return 1; }
+int SpecialFuncs::Method() { return 13; }
 
 struct SpecSpecialFuncs {
   void vd() {}
   SpecSpecialFuncs();
   ~SpecSpecialFuncs();
 
   int method_();
 #pragma omp declare variant(SpecSpecialFuncs::method_) \
     match(device = {kind(CORRECT)})
-  int method() { return 1; }
+  int method() { return 14; }
 #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 SpecSpecialFuncs::method_() { return 15; }
+int SpecSpecialFuncs::Method() { return 16; }
 
-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; }
+int prio() { return 17; }
+int prio1() { return 18; }
+static int prio2() { return 19; }
+static int prio3() { return 20; }
+static int prio4() { return 21; }
+int fn_linkage_variant() { return 22; }
+extern "C" int fn_linkage_variant1() { return 23; }
+int fn_variant2() { return 24; }
 
 #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; }
+int prio_() { return 25; }
 
 #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; }
+static int prio1_() { return 26; }
 
 int int_fn() { return prio1_(); }
 
 extern "C" {
 #pragma omp declare variant(fn_linkage_variant) match(device = {kind(CORRECT)})
-int fn_linkage() { return 1; }
+int fn_linkage() { return 27; }
 }
 
 #pragma omp declare variant(fn_linkage_variant1) match(device = {kind(CORRECT)})
-int fn_linkage1() { return 1; }
+int fn_linkage1() { return 28; }
 
 #pragma omp declare variant(fn_variant2) match(device = {kind(WRONG)})
-int fn2() { return 87; }
+int fn2() { return 29; }
 
 #pragma omp end declare target
 
 #endif // HEADER

clang/test/OpenMP/nvptx_declare_variant_implementation_vendor_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 -o - -fopenmp-version=50 | FileCheck %s --implicit-check-not='ret i32 {{6|7|8|9|10|12|13|14|15|17|18|19|20|21|22|23|24|26}}'
 // 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}}'
+// 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 {{6|7|8|9|10|12|13|14|15|17|18|19|20|21|22|23|24|26}}'
 // 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}}
+// CHECK-DAG:  ret i32 2
+// CHECK-DAG:  ret i32 3
+// CHECK-DAG:  ret i32 4
+// CHECK-DAG:  ret i32 5
+// CHECK-DAG:  ret i32 11
+// CHECK-DAG:  ret i32 16
+// CHECK-DAG:  ret i32 19
+// CHECK-DAG:  ret i32 25
+// CHECK-DAG:  ret i32 27
+// CHECK-DAG:  ret i32 28
+// CHECK-DAG:  ret i32 29
 
 #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(implementation = {vendor(llvm)})
-int bar() { return 1; }
+int bar() { return 3; }
 
 #pragma omp declare variant(bazzz) match(implementation = {vendor(llvm)})
-int baz() { return 1; }
+int baz() { return 4; }
 
 #pragma omp declare variant(test) match(implementation = {vendor(llvm)})
-int call() { return 1; }
+int call() { return 5; }
 
 #pragma omp declare variant(stat_unused_) match(implementation = {vendor(llvm)})
-static int stat_unused() { return 1; }
+static int stat_unused() { return 6; }
 
 #pragma omp declare variant(stat_used_) match(implementation = {vendor(llvm)})
-static int stat_used() { return 1; }
+static int stat_used() { return 7; }
 
 #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; }
+int test() { return 8; }
+static int stat_unused_() { return 9; }
+static int stat_used_() { return 10; }
 
 #pragma omp declare target
 
 struct SpecialFuncs {
   void vd() {}
   SpecialFuncs();
   ~SpecialFuncs();
 
-  int method_() { return 6; }
+  int method_() { return 11; }
 #pragma omp declare variant(SpecialFuncs::method_) \
     match(implementation = {vendor(llvm)})
-  int method() { return 1; }
+  int method() { return 12; }
 #pragma omp declare variant(SpecialFuncs::method_) \
     match(implementation = {vendor(llvm)})
   int Method();
 } s;
 
-int SpecialFuncs::Method() { return 1; }
+int SpecialFuncs::Method() { return 13; }
 
 struct SpecSpecialFuncs {
   void vd() {}
   SpecSpecialFuncs();
   ~SpecSpecialFuncs();
 
   int method_();
 #pragma omp declare variant(SpecSpecialFuncs::method_) \
     match(implementation = {vendor(llvm)})
-  int method() { return 1; }
+  int method() { return 14; }
 #pragma omp declare variant(SpecSpecialFuncs::method_) \
     match(implementation = {vendor(llvm)})
   int Method();
 } s1;
 
 #pragma omp end declare target
 
-int SpecSpecialFuncs::method_() { return 7; }
-int SpecSpecialFuncs::Method() { return 1; }
+int SpecSpecialFuncs::method_() { return 15; }
+int SpecSpecialFuncs::Method() { return 16; }
 
-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; }
+int prio() { return 17; }
+int prio1() { return 18; }
+static int prio2() { return 19; }
+static int prio3() { return 20; }
+static int prio4() { return 21; }
+int fn_linkage_variant() { return 22; }
+extern "C" int fn_linkage_variant1() { return 23; }
+int fn_variant2() { return 24; }
 
 #pragma omp declare target
 
 void xxx() {
   (void)s.method();
   (void)s1.method();
 }
 
 #pragma omp declare variant(prio) match(implementation = {vendor(llvm)})
 #pragma omp declare variant(prio1) match(implementation = {vendor(score(1) \
                                                                   : llvm)})
-int prio_() { return 1; }
+int prio_() { return 25; }
 
 #pragma omp declare variant(prio4) match(implementation = {vendor(score(3) \
                                                                   : llvm)})
 #pragma omp declare variant(prio2) match(implementation = {vendor(score(5) \
                                                                   : llvm)})
 #pragma omp declare variant(prio3) match(implementation = {vendor(score(1) \
                                                                   : llvm)})
-static int prio1_() { return 1; }
+static int prio1_() { return 26; }
 
 int int_fn() { return prio1_(); }
 
 extern "C" {
 #pragma omp declare variant(fn_linkage_variant) match(implementation = {vendor(llvm)})
-int fn_linkage() { return 1; }
+int fn_linkage() { return 27; }
 }
 
 #pragma omp declare variant(fn_linkage_variant1) match(implementation = {vendor(llvm)})
-int fn_linkage1() { return 1; }
+int fn_linkage1() { return 28; }
 
 #pragma omp declare variant(fn_variant2) match(implementation = {vendor(llvm, ibm)})
-int fn2() { return 87; }
+int fn2() { return 29; }
 
 #pragma omp end declare target
 
 #endif // HEADER