[cuda][hip] Fix function overload resolution in the global initiailizer.

[hliao]

Summary

• As global initializers are not under any function body, they need to look into the current variable being initialized. That is not addressed in the current CUDA/HIP overloadable function resolution and ignore target checking. That may result in wrong candidate to be considered as illustrated in the newly added test case.
• In this patch, a non-local variable stack is introduced to keep track the current non-local variable being initialized so that initialization function could be inspected for the target preference.
• Besides newly added tests, existing tests are refined as the current implementation adds extra checks on global initializers to ensure no device functions are used. As the target match checking is enabled in this patch, such check is only necessary for CUDA device global variables. They are not allowed to be non-trivially initialized. As HIP starts to support non-trivial initialization of device initialization, such target matching check is mandatory to be enforced.

[hliao]

refine commit message

[hliao]

refine again

[hliao]

File PR44266 to track that bug.

[tra]

Looks good to me overall. I've pinged rsmith@ to double-check that we're covering all possibilities for non-local variable init.

clang/include/clang/Sema/Sema.h

11632

I'd add a comment describing that it's a wrapper which dispatches the call to one of more specific variants above.

11655–11656

Comment needs updating.

11668

Nit: I'd add an empty line between delarations and the function. Jammed together they are hard to read.

11697–11699

if (const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext))
  if (Kind == SkipImplicitCaller && Caller->isImplicit())
      return true;

11736–11737

Now that we always use getCUDAContextDecl() as the first argument, perhaps we can just always retrieve the context inside the function.

clang/lib/Parse/ParseDecl.cpp

2345

@rsmith -- is this sufficient to catch all attempts to call an initializer for a global?
I wonder if there are other sneaky ways to call an initializer.

clang/test/SemaCUDA/function-overload.cu

458

I'd add more details here.
The problem is that here the overload set has both functions and the one with the integer argument wins, even though it's a device function which we can't execute. We do handle similar cases during overload resolution in other places where we would prefer a callable function over a non-callable function with a better signature match.

[tra]

I wonder if this patch will help with this case:

https://godbolt.org/z/X4KdsV

__device__ float fn(int) { return threadIdx.x; };
__host__ float fn(float);

float gvar1 = []()__device__ { return fn(1);} (); // This ends up calling fn(int) on *host*

We seem to happily let host code call device function from a lambda function used as an initializer.

[hliao]

In D71227#1778136, @tra wrote:

I wonder if this patch will help with this case:

https://godbolt.org/z/X4KdsV

__device__ float fn(int) { return threadIdx.x; };
__host__ float fn(float);

float gvar1 = []()__device__ { return fn(1);} (); // This ends up calling fn(int) on *host*

We seem to happily let host code call device function from a lambda function used as an initializer.

It's turned out that Sema::CheckCUDACall needs to consider global initializer as well. I will revise that part. But, technically, that's irrelevant to overloadable resolution. Should be prepared in another patch enhancing CheckCUDACall to check global initializers.

[hliao]

refinements are made after comments from reviewers.

[hliao]

code refinement after reviewers' comments.

[rsmith]

clang/include/clang/Sema/Sema.h

11656–11658

Please capitalize the first word of each of these parameter descriptions, to match the style used elsewhere in Clang.

11657

"Null" not "nullptr".

11664–11672

Does this really need to be CUDA-specific?

This is (at least) the third time we've needed this. We currently have a ManglingContextDecl on ExpressionEvaluationContextRecord that tracks the non-local variable whose initializer we're parsing. In addition to using this as a lambda context declaration, we also (hackily) use it as the context declaration for DiagRuntimeBehavior. It would seem sensible to use that mechanism here too (and rename it to remove any suggestion that this is specific to lambdas or mangling).

I think we only currently push ExpressionEvaluationContexts for variable initializers in C++. That's presumably fine for CUDA's purposes.

clang/lib/Parse/ParseDecl.cpp

2345

No, this is not sufficient; it's missing (at least) the template instantiation case. (The ExpressionEvaluationContextRecord mechanism does handle that case properly.)

You should also consider what should happen in default arguments (which are sometimes parsed before we form a FunctionDecl for the function for which they are parameters) and default member initializers (which are parsed after we know whether the enclosing class has a user-declared default constructor, so you could in principle consider the CUDA function kind of the declared constructors, I suppose -- but the constructor bodies are not yet available, so you can't tell which constructors would actually use the initializers). Both of those cases are also tracked by the ExpressionEvaluationContextRecord mechanism, though you may need to track additional information to process default arguments in the same mode as the function for which they are supplied.

[hliao]

• revise comment.
• add tests requiring tempate instantiation.

[hliao]

Sorry for the late reply. Really appreciate your feedback. Thanks!

clang/include/clang/Sema/Sema.h

11664–11672

I tried that before adding the new non-local variable stack. Using ManglingContextDecl on ExpressionEvaluationContextRecord could serve some cases, but it cannot fit the case where the constructor needs resolving as well. When resolving the constructor, ManglingContextDecl scope is already closed and cannot be used to check the target of the global variables. Says the following code

struct EC {
  int ec;
__device__ EC() {}
};
__device__ EC d_ec;

I also tried enlarging the scope of ManglingContextDecl but that triggers even more issues for the generic C++ compilation. I'd appreciate any better solution as I agree that adding CUDA specific facilities should be minimized.

clang/lib/Parse/ParseDecl.cpp

2345

Could you elaborate more? I added new test cases requiring template instantiation. The current code handle them correctly. Do you refer to template variables?

[hliao]

In D71227#1831445, @hliao wrote:

Sorry for the late reply. Really appreciate your feedback. Thanks!

@rsmith Have you chance to review the revised change?

[hliao]

@rsmith do have u the chance to review the revised change again as well as my answers to your comments?

[hliao]

+ @rjmccall

[hliao]

Rebase to the latest trunk code.

[rjmccall]

clang/include/clang/Sema/Sema.h

11678

This is tricky because we could be in a nested context, not just the initializer, and that context just might not be a function. For example, there could be a local class in a lambda or something like that.

[hliao]

Skip non-function or non-TU context so far as more cases need considering.

[hliao]

clang/include/clang/Sema/Sema.h

11678

You are right. Limit that to function and TU context so far. I need more efforts to consider other cases. One case in mind is that default member initialier in a class. But, for local classes in a lambda, they should be in a function (lambda function body) context.

[rjmccall]

clang/include/clang/Sema/Sema.h

11680

You really want this to match whenever we're in a local context, right? How about structuring the function like:

if (CurContext->isFunctionOrMethod())
  return cast<Decl>(CurContext);
if (!CurContext->isFileContext())
  return nullptr;
return getCUDACurrentNonLocalVariable();

As a more general solution, I think Sema funnels all changes to CurContext through a small number of places, and you could make those places save and restore the currently initialized variable as well.

[hliao]

Rebase the code to the latest trunk.

[hliao]

Revise following reviewer comments.

[hliao]

Rebase to the trunk.

[hliao]

Rebase to the trunk.

[hliao]

Rebase to the latest trunk.

[hliao]

@rjmccall @rsmith @tra, could you review on this revision?

[hliao]

Fix pre-merge checks.

[rjmccall]

clang/include/clang/Sema/Sema.h

11680

Richard, I'd like your opinion about this. We have three separate patches right now that would all benefit from being able to track that they're currently within a variable/field initializer in Sema. And it's a general deficiency that it's hard to track declarations in initializers back to their initialized variable.

Swift actually bit the bullet and introduced a kind of DeclContext that's a non-local initializer, and that links back to the variable. That would be hard to bring back to Clang with the current AST because Clang assumes that all DeclContexts are Decls, and I don't think we can reasonably remove that assumption; and of course VarDecl and FieldDecl aren't DeclContexts.

Now, we could try to change that latter point. Making *all* VarDecls and FieldDecls DCs would have prohibitive memory overhead, since the vast majority are local / uninitialized; however, we could introduce a VarDecl subclass for global variables (including static member variables, of course), and similarly we could have a FieldDecl subclass for fields with initializers, which would nicely move some of the other overhead out-of-line and optimize for the C-style/old-style case. (We always know whether a field has an in-class initializer at parse time, right?)

Less invasively, we could forget about trying to track this in the AST and just also track a current initialized variable in Sema. Anything which tried to change the context would have to save and restore that as well. That might be annoying because of PushDeclContext/PopDeclContext, though, which assume that you can restore the old context by just looking at the current context.

clang/lib/Sema/SemaDeclCXX.cpp

16932

The declaration could become invalid while processing its initializer; I think you should drop that condition.

[hliao]

Remove unncessary condition checking.

[hliao]

Rebase to the latest trunk.

[hliao]

Fix warnings from clang-tidy.

[hliao]

Fix more clang-tidy warnings.

[hliao]

Rebase to the latest trunk.

[hliao]

Rebase to trunk.

[yaxunl]

LGTM. Can we get this in? There are other fixes depending on this. Thanks.

[hliao]

Rebase to trunk and resolve the conflict.

[tra]

Is this patch still actual?

[hliao]

Hi @rsmith, @rjmccall and @tra what's your suggestion to make progress on this review?

In D71227#2167596, @tra wrote:

Is this patch still actual?

I need to rebase this to the latest trunk. Interrupt with other heavy loads.

clang/include/clang/Sema/Sema.h

enum CUDAFunctionTarget {
     CFT_InvalidTarget
   };
 
   /// Determines whether the given function is a CUDA device/host/kernel/etc.
   /// function.
   ///
   /// Use this rather than examining the function's attributes yourself -- you
   /// will get it wrong.  Returns CFT_Host if D is null.
+  CUDAFunctionTarget IdentifyCUDATarget(const ParsedAttributesView &Attrs);
   CUDAFunctionTarget IdentifyCUDATarget(const FunctionDecl *D,
                                         bool IgnoreImplicitHDAttr = false);
-  CUDAFunctionTarget IdentifyCUDATarget(const ParsedAttributesView &Attrs);
+  CUDAFunctionTarget IdentifyCUDATarget(const VarDecl *D,
+                                        bool IgnoreImplicitHDAttr = false);
+  // This routine is the top level dispatcher to more specific variants above.
+  CUDAFunctionTarget IdentifyCUDATarget(const Decl *D,
+                                        bool IgnoreImplicitHDAttr = false);
 
   /// Gets the CUDA target for the current context.
   CUDAFunctionTarget CurrentCUDATarget() {
     return IdentifyCUDATarget(dyn_cast<FunctionDecl>(CurContext));
   }
 
   // CUDA function call preference. Must be ordered numerically from
   // worst to best.
   enum CUDAFunctionPreference {
     CFP_Never,      // Invalid caller/callee combination.
     CFP_WrongSide,  // Calls from host-device to host or device
                     // function that do not match current compilation
                     // mode.
     CFP_HostDevice, // Any calls to host/device functions.
     CFP_SameSide,   // Calls from host-device to host or device
                     // function matching current compilation mode.
     CFP_Native,     // host-to-host or device-to-device calls.
   };
 
-  /// Identifies relative preference of a given Caller/Callee
+  /// Identifies relative preference of a given callee and that call context
   /// combination, based on their host/device attributes.
-  /// \param Caller function which needs address of \p Callee.
-  ///               nullptr in case of global context.
-  /// \param Callee target function
+  /// \param CallContextDecl The context decl which needs address of \p Callee.
+  ///                        Null in case of the global context.
+  /// \param Callee          Target function.
   ///
   /// \returns preference value for particular Caller/Callee combination.
-  CUDAFunctionPreference IdentifyCUDAPreference(const FunctionDecl *Caller,
+  CUDAFunctionPreference IdentifyCUDAPreference(const Decl *CallContextDecl,
                                                 const FunctionDecl *Callee);
 
+  SmallVector<const Decl *, 8> CUDANonLocalVariableStack;
+
+  void pushCUDANonLocalVariable(const Decl *D);
+  void popCUDANonLocalVariable(const Decl *D);
+
+  const Decl *getCUDACurrentNonLocalVariable() const {
+    return CUDANonLocalVariableStack.empty() ? nullptr
+                                             : CUDANonLocalVariableStack.back();
+  }
+
+  const Decl *getCUDAContextDecl() const {
+    if (CurContext->isFunctionOrMethod())
+      return cast<Decl>(CurContext);
+    if (!CurContext->isFileContext()) {
+      // TODO: There are cases where proper checking is required, such as the
+      // default member initializer.
+      return nullptr;
+    }
+    // Check the current variable being initialized in the global context.
+    return getCUDACurrentNonLocalVariable();
+  }
+
   /// Determines whether Caller may invoke Callee, based on their CUDA
   /// host/device attributes.  Returns false if the call is not allowed.
   ///
   /// Note: Will return true for CFP_WrongSide calls.  These may appear in
   /// semantically correct CUDA programs, but only if they're never codegen'ed.
-  bool IsAllowedCUDACall(const FunctionDecl *Caller,
-                         const FunctionDecl *Callee) {
-    return IdentifyCUDAPreference(Caller, Callee) != CFP_Never;
+  enum SkipCallerKind_t { SkipNoneCaller, SkipImplicitCaller };
+  bool isCUDACallAllowed(const FunctionDecl *Callee,
+                         SkipCallerKind_t Kind = SkipNoneCaller) {
+    // Skip contexts where no real call could be performed.
+    if (!CurContext->isFileContext() && !CurContext->isFunctionOrMethod())
+      return true;
+    if (const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext))
+      if (Kind == SkipImplicitCaller && Caller->isImplicit())
+        return true;
+    return IdentifyCUDAPreference(getCUDAContextDecl(), Callee) != CFP_Never;
   }
 
   /// May add implicit CUDAHostAttr and CUDADeviceAttr attributes to FD,
   /// depending on FD and the current compilation settings.
   void maybeAddCUDAHostDeviceAttrs(FunctionDecl *FD,
                                    const LookupResult &Previous);
 
 public:
public:
   /// operator() method.
   ///
   /// CUDA lambdas declared inside __device__ or __global__ functions inherit
   /// the __device__ attribute.  Similarly, lambdas inside __host__ __device__
   /// functions become __host__ __device__ themselves.
   void CUDASetLambdaAttrs(CXXMethodDecl *Method);
 
   /// Finds a function in \p Matches with highest calling priority
-  /// from \p Caller context and erases all functions with lower
+  /// from the current context and erases all functions with lower
   /// calling priority.
   void EraseUnwantedCUDAMatches(
-      const FunctionDecl *Caller,
       SmallVectorImpl<std::pair<DeclAccessPair, FunctionDecl *>> &Matches);
 
   /// Given a implicit special member, infer its CUDA target from the
   /// calls it needs to make to underlying base/field special members.
   /// \param ClassDecl the class for which the member is being created.
   /// \param CSM the kind of special member.
   /// \param MemberDecl the special member itself.
   /// \param ConstRHS true if this is a copy operation with a const object on
   ///        its RHS.

clang/lib/Parse/ParseDecl.cpp

if (Tok.is(tok::semi)) {
         ThisDecl =
             Actions.ActOnTemplateDeclarator(getCurScope(), FakedParamLists, D);
       }
     }
     break;
     }
   }
 
+  Actions.pushCUDANonLocalVariable(ThisDecl);
+
   // Parse declarator '=' initializer.
   // If a '==' or '+=' is found, suggest a fixit to '='.
   if (isTokenEqualOrEqualTypo()) {
     SourceLocation EqualLoc = ConsumeToken();
 
     if (Tok.is(tok::kw_delete)) {
       if (D.isFunctionDeclarator())
         Diag(ConsumeToken(), diag::err_default_delete_in_multiple_declaration)
if (Init.isInvalid()) {
       Actions.ActOnInitializerError(ThisDecl);
     } else
       Actions.AddInitializerToDecl(ThisDecl, Init.get(), /*DirectInit=*/true);
 
   } else {
     Actions.ActOnUninitializedDecl(ThisDecl);
   }
 
+  Actions.popCUDANonLocalVariable(ThisDecl);
+
   Actions.FinalizeDeclaration(ThisDecl);
 
   return ThisDecl;
 }
 
 /// ParseSpecifierQualifierList
 ///        specifier-qualifier-list:
 ///          type-specifier specifier-qualifier-list[opt]

clang/lib/Sema/SemaCUDA.cpp

Sema::IdentifyCUDATarget(const ParsedAttributesView &Attrs) {
 
   if (HasDeviceAttr)
     return CFT_Device;
 
   return CFT_Host;
 }
 
 template <typename A>
-static bool hasAttr(const FunctionDecl *D, bool IgnoreImplicitAttr) {
+static bool hasAttr(const Decl *D, bool IgnoreImplicitAttr) {
   return D->hasAttrs() && llvm::any_of(D->getAttrs(), [&](Attr *Attribute) {
            return isa<A>(Attribute) &&
                   !(IgnoreImplicitAttr && Attribute->isImplicit());
          });
 }
 
 /// IdentifyCUDATarget - Determine the CUDA compilation target for this function
 Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D,
if (hasAttr<CUDADeviceAttr>(D, IgnoreImplicitHDAttr)) {
     // Some implicit declarations (like intrinsic functions) are not marked.
     // Set the most lenient target on them for maximal flexibility.
     return CFT_HostDevice;
   }
 
   return CFT_Host;
 }
 
+Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const VarDecl *D,
+                                                  bool IgnoreImplicitHDAttr) {
+  if (D == nullptr)
+    return CFT_Host;
+
+  assert(D->hasGlobalStorage() && "Only non-local variable needs identifying.");
+
+  if (D->hasAttr<CUDAInvalidTargetAttr>())
+    return CFT_InvalidTarget;
+
+  if (hasAttr<CUDAConstantAttr>(D, IgnoreImplicitHDAttr) ||
+      hasAttr<CUDADeviceAttr>(D, IgnoreImplicitHDAttr) ||
+      hasAttr<CUDASharedAttr>(D, IgnoreImplicitHDAttr))
+    return CFT_Device;
+
+  return CFT_Host;
+}
+
+Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const Decl *D,
+                                                  bool IgnoreImplicitHDAttr) {
+  if (D == nullptr)
+    return CFT_Host;
+
+  if (const auto *FD = dyn_cast<FunctionDecl>(D))
+    return IdentifyCUDATarget(FD, IgnoreImplicitHDAttr);
+
+  if (const auto *VD = dyn_cast<VarDecl>(D))
+    return IdentifyCUDATarget(VD, IgnoreImplicitHDAttr);
+
+  llvm_unreachable("Unexpected decl for CUDA target identification.");
+}
+
 // * CUDA Call preference table
 //
 // F - from,
 // T - to
 // Ph - preference in host mode
 // Pd - preference in device mode
 // H  - handled in (x)
 // Preferences: N:native, SS:same side, HD:host-device, WS:wrong side, --:never.

 // | h  | h  | N   | N   | (c) |
 // | h  | hd | HD  | HD  | (b) |
 // | hd | d  | WS  | SS  | (d) |
 // | hd | g  | SS  | --  |(d/a)|
 // | hd | h  | SS  | WS  | (d) |
 // | hd | hd | HD  | HD  | (b) |
 
 Sema::CUDAFunctionPreference
-Sema::IdentifyCUDAPreference(const FunctionDecl *Caller,
+Sema::IdentifyCUDAPreference(const Decl *ContextDecl,
                              const FunctionDecl *Callee) {
   assert(Callee && "Callee must be valid.");
-  CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller);
+  CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(ContextDecl);
   CUDAFunctionTarget CalleeTarget = IdentifyCUDATarget(Callee);
 
   // If one of the targets is invalid, the check always fails, no matter what
   // the other target is.
   if (CallerTarget == CFT_InvalidTarget || CalleeTarget == CFT_InvalidTarget)
     return CFP_Never;
 
   // (a) Can't call global from some contexts until we support CUDA's
if ((CallerTarget == CFT_Host && CalleeTarget == CFT_Device) ||
       (CallerTarget == CFT_Device && CalleeTarget == CFT_Host) ||
       (CallerTarget == CFT_Global && CalleeTarget == CFT_Host))
     return CFP_Never;
 
   llvm_unreachable("All cases should've been handled by now.");
 }
 
 void Sema::EraseUnwantedCUDAMatches(
-    const FunctionDecl *Caller,
     SmallVectorImpl<std::pair<DeclAccessPair, FunctionDecl *>> &Matches) {
   if (Matches.size() <= 1)
     return;
 
   using Pair = std::pair<DeclAccessPair, FunctionDecl*>;
 
-  // Gets the CUDA function preference for a call from Caller to Match.
+  const Decl *ContextDecl = getCUDAContextDecl();
+
+  // Gets the CUDA function preference for a call from call context to Match.
   auto GetCFP = [&](const Pair &Match) {
-    return IdentifyCUDAPreference(Caller, Match.second);
+    return IdentifyCUDAPreference(ContextDecl, Match.second);
   };
 
   // Find the best call preference among the functions in Matches.
   CUDAFunctionPreference BestCFP = GetCFP(*std::max_element(
       Matches.begin(), Matches.end(),
       [&](const Pair &M1, const Pair &M2) { return GetCFP(M1) < GetCFP(M2); }));
 
   // Erase all functions with lower priority.

clang/lib/Sema/SemaDeclCXX.cpp

 /// static data member.
 static bool isNonlocalVariable(const Decl *D) {
   if (const VarDecl *Var = dyn_cast_or_null<VarDecl>(D))
     return Var->hasGlobalStorage();
 
   return false;
 }
 
+void Sema::pushCUDANonLocalVariable(const Decl *D) {
+  if (!D || !isNonlocalVariable(D))
+    return;
+  CUDANonLocalVariableStack.push_back(D);
+}
+
+void Sema::popCUDANonLocalVariable(const Decl *D) {
+  if (!D || !isNonlocalVariable(D))
+    return;
+  assert(!CUDANonLocalVariableStack.empty() &&
+         CUDANonLocalVariableStack.back() == D);
+  CUDANonLocalVariableStack.pop_back();
+}
+
 /// Invoked when we are about to parse an initializer for the declaration
 /// 'Dcl'.
 ///
 /// After this method is called, according to [C++ 3.4.1p13], if 'Dcl' is a
 /// static data member of class X, names should be looked up in the scope of
 /// class X. If the declaration had a scope specifier, a scope will have
 /// been created and passed in for this purpose. Otherwise, S will be null.
 void Sema::ActOnCXXEnterDeclInitializer(Scope *S, Decl *D) {

clang/lib/Sema/SemaExprCXX.cpp

Result = CXXFunctionalCastExpr::Create(
         Result.get(), /*Path=*/nullptr, Locs.getBegin(), Locs.getEnd());
   }
 
   return Result;
 }
 
 bool Sema::isUsualDeallocationFunction(const CXXMethodDecl *Method) {
   // [CUDA] Ignore this function, if we can't call it.
-  const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext);
+  const Decl *ContextDecl = getCUDAContextDecl();
   if (getLangOpts().CUDA &&
-      IdentifyCUDAPreference(Caller, Method) <= CFP_WrongSide)
+      IdentifyCUDAPreference(ContextDecl, Method) <= CFP_WrongSide)
     return false;
 
   SmallVector<const FunctionDecl*, 4> PreventedBy;
   bool Result = Method->isUsualDeallocationFunction(PreventedBy);
 
   if (Result || !getLangOpts().CUDA || PreventedBy.empty())
     return Result;
 
   // In case of CUDA, return true if none of the 1-argument deallocator
   // functions are actually callable.
   return llvm::none_of(PreventedBy, [&](const FunctionDecl *FD) {
     assert(FD->getNumParams() == 1 &&
            "Only single-operand functions should be in PreventedBy");
-    return IdentifyCUDAPreference(Caller, FD) >= CFP_HostDevice;
+    return IdentifyCUDAPreference(ContextDecl, FD) >= CFP_HostDevice;
   });
 }
 
 /// Determine whether the given function is a non-placement
 /// deallocation function.
 static bool isNonPlacementDeallocationFunction(Sema &S, FunctionDecl *FD) {
   if (CXXMethodDecl *Method = dyn_cast<CXXMethodDecl>(FD))
     return S.isUsualDeallocationFunction(Method);
UsualDeallocFnInfo(Sema &S, DeclAccessPair Found)
       if (NumBaseParams < FD->getNumParams() &&
           FD->getParamDecl(NumBaseParams)->getType()->isAlignValT()) {
         ++NumBaseParams;
         HasAlignValT = true;
       }
 
       // In CUDA, determine how much we'd like / dislike to call this.
       if (S.getLangOpts().CUDA)
-        if (auto *Caller = dyn_cast<FunctionDecl>(S.CurContext))
-          CUDAPref = S.IdentifyCUDAPreference(Caller, FD);
+        CUDAPref = S.IdentifyCUDAPreference(S.getCUDAContextDecl(), FD);
     }
 
     explicit operator bool() const { return FD; }
 
     bool isBetterThan(const UsualDeallocFnInfo &Other, bool WantSize,
                       bool WantAlign) const {
       // C++ P0722:
       //   A destroying operator delete is preferred over a non-destroying
for (LookupResult::iterator D = FoundDelete.begin(),
       if (Context.hasSameType(adjustCCAndNoReturn(Fn->getType(),
                                                   ExpectedFunctionType,
                                                   /*AdjustExcpetionSpec*/true),
                               ExpectedFunctionType))
         Matches.push_back(std::make_pair(D.getPair(), Fn));
     }
 
     if (getLangOpts().CUDA)
-      EraseUnwantedCUDAMatches(dyn_cast<FunctionDecl>(CurContext), Matches);
+      EraseUnwantedCUDAMatches(Matches);
   } else {
     // C++1y [expr.new]p22:
     //   For a non-placement allocation function, the normal deallocation
     //   function lookup is used
     //
     // Per [expr.delete]p10, this lookup prefers a member operator delete
     // without a size_t argument, but prefers a non-member operator delete
     // with a size_t where possible (which it always is in this case).

clang/lib/Sema/SemaOverload.cpp

void Sema::AddOverloadCandidate(
   if (Args.size() < MinRequiredArgs && !PartialOverloading) {
     // Not enough arguments.
     Candidate.Viable = false;
     Candidate.FailureKind = ovl_fail_too_few_arguments;
     return;
   }
 
   // (CUDA B.1): Check for invalid calls between targets.
-  if (getLangOpts().CUDA)
-    if (const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext))
-      // Skip the check for callers that are implicit members, because in this
-      // case we may not yet know what the member's target is; the target is
-      // inferred for the member automatically, based on the bases and fields of
-      // the class.
-      if (!Caller->isImplicit() && !IsAllowedCUDACall(Caller, Function)) {
-        Candidate.Viable = false;
-        Candidate.FailureKind = ovl_fail_bad_target;
-        return;
-      }
+  if (getLangOpts().CUDA &&
+      !isCUDACallAllowed(Function, Sema::SkipImplicitCaller)) {
 
   if (Function->getTrailingRequiresClause()) {
     ConstraintSatisfaction Satisfaction;
     if (CheckFunctionConstraints(Function, Satisfaction) ||
         !Satisfaction.IsSatisfied) {
       Candidate.Viable = false;
       Candidate.FailureKind = ovl_fail_constraints_not_satisfied;
       return;
else {
     if (Candidate.Conversions[ConvIdx].isBad()) {
       Candidate.Viable = false;
       Candidate.FailureKind = ovl_fail_bad_conversion;
       return;
     }
   }
 
   // (CUDA B.1): Check for invalid calls between targets.
-  if (getLangOpts().CUDA)
-    if (const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext))
-      if (!IsAllowedCUDACall(Caller, Method)) {
-        Candidate.Viable = false;
-        Candidate.FailureKind = ovl_fail_bad_target;
-        return;
-      }
+  if (getLangOpts().CUDA && !isCUDACallAllowed(Method)) {
 
   if (Method->getTrailingRequiresClause()) {
     ConstraintSatisfaction Satisfaction;
     if (CheckFunctionConstraints(Method, Satisfaction) ||
         !Satisfaction.IsSatisfied) {
       Candidate.Viable = false;
       Candidate.FailureKind = ovl_fail_constraints_not_satisfied;
       return;
bool clang::isBetterOverloadCandidate(
   // viability of a function. If two functions are both viable, other factors
   // should take precedence in preference, e.g. the standard-defined preferences
   // like argument conversion ranks or enable_if partial-ordering. The
   // preference for pass-object-size parameters is probably most similar to a
   // type-based-overloading decision and so should take priority.
   //
   // If other rules cannot determine which is better, CUDA preference will be
   // used again to determine which is better.
-  //
-  // TODO: Currently IdentifyCUDAPreference does not return correct values
-  // for functions called in global variable initializers due to missing
-  // correct context about device/host. Therefore we can only enforce this
-  // rule when there is a caller. We should enforce this rule for functions
-  // in global variable initializers once proper context is added.
   if (S.getLangOpts().CUDA && Cand1.Function && Cand2.Function) {
-    if (FunctionDecl *Caller = dyn_cast<FunctionDecl>(S.CurContext)) {
-      auto P1 = S.IdentifyCUDAPreference(Caller, Cand1.Function);
-      auto P2 = S.IdentifyCUDAPreference(Caller, Cand2.Function);
-      assert(P1 != Sema::CFP_Never && P2 != Sema::CFP_Never);
-      auto Cand1Emittable = P1 > Sema::CFP_WrongSide;
-      auto Cand2Emittable = P2 > Sema::CFP_WrongSide;
-      if (Cand1Emittable && !Cand2Emittable)
-        return true;
-      if (!Cand1Emittable && Cand2Emittable)
-        return false;
-    }
-  }
+    const Decl *ContextDecl = S.getCUDAContextDecl();
+    auto P1 = S.IdentifyCUDAPreference(ContextDecl, Cand1.Function);
+    auto P2 = S.IdentifyCUDAPreference(ContextDecl, Cand2.Function);
 
   // C++ [over.match.best]p1:
   //
   //   -- if F is a static member function, ICS1(F) is defined such
   //      that ICS1(F) is neither better nor worse than ICS1(G) for
   //      any function G, and, symmetrically, ICS1(G) is neither
   //      better nor worse than ICS1(F).
   unsigned StartArg = 0;
bool clang::isBetterOverloadCandidate(
   if (MV == Comparison::Better)
     return true;
   if (MV == Comparison::Worse)
     return false;
 
   // If other rules cannot determine which is better, CUDA preference is used
   // to determine which is better.
   if (S.getLangOpts().CUDA && Cand1.Function && Cand2.Function) {
-    FunctionDecl *Caller = dyn_cast<FunctionDecl>(S.CurContext);
-    return S.IdentifyCUDAPreference(Caller, Cand1.Function) >
-           S.IdentifyCUDAPreference(Caller, Cand2.Function);
+    const Decl *ContextDecl = S.getCUDAContextDecl();
+    return S.IdentifyCUDAPreference(ContextDecl, Cand1.Function) >
+           S.IdentifyCUDAPreference(ContextDecl, Cand2.Function);
   }
 
   return false;
 }
 
 /// Determine whether two declarations are "equivalent" for the purposes of
 /// name lookup and overload resolution. This applies when the same internal/no
 /// linkage entity is defined by two modules (probably by textually including
if (CheckArityMismatch(S, Cand, NumArgs))
       return;
   }
   DiagnoseBadDeduction(S, Cand->FoundDecl, Cand->Function, // pattern
                        Cand->DeductionFailure, NumArgs, TakingCandidateAddress);
 }
 
 /// CUDA: diagnose an invalid call across targets.
 static void DiagnoseBadTarget(Sema &S, OverloadCandidate *Cand) {
-  FunctionDecl *Caller = cast<FunctionDecl>(S.CurContext);
+  const Decl *ContextDecl = S.getCUDAContextDecl();
   FunctionDecl *Callee = Cand->Function;
 
-  Sema::CUDAFunctionTarget CallerTarget = S.IdentifyCUDATarget(Caller),
+  Sema::CUDAFunctionTarget CallerTarget = S.IdentifyCUDATarget(ContextDecl),
                            CalleeTarget = S.IdentifyCUDATarget(Callee);
 
   std::string FnDesc;
   std::pair<OverloadCandidateKind, OverloadCandidateSelect> FnKindPair =
       ClassifyOverloadCandidate(S, Cand->FoundDecl, Callee,
                                 Cand->getRewriteKind(), FnDesc);
 
   S.Diag(Callee->getLocation(), diag::note_ovl_candidate_bad_target)
if (CXXMethodDecl *Method = dyn_cast<CXXMethodDecl>(Fn)) {
       // when converting to member pointer.
       if (Method->isStatic() == TargetTypeIsNonStaticMemberFunction)
         return false;
     }
     else if (TargetTypeIsNonStaticMemberFunction)
       return false;
 
     if (FunctionDecl *FunDecl = dyn_cast<FunctionDecl>(Fn)) {
-      if (S.getLangOpts().CUDA)
-        if (FunctionDecl *Caller = dyn_cast<FunctionDecl>(S.CurContext))
-          if (!Caller->isImplicit() && !S.IsAllowedCUDACall(Caller, FunDecl))
-            return false;
+      if (S.getLangOpts().CUDA &&
+          !S.isCUDACallAllowed(FunDecl, Sema::SkipImplicitCaller))
       if (FunDecl->isMultiVersion()) {
         const auto *TA = FunDecl->getAttr<TargetAttr>();
         if (TA && !TA->isDefaultVersion())
           return false;
       }
 
       // If any candidate has a placeholder return type, trigger its deduction
       // now.
for (unsigned I = 0, N = Matches.size(); I != N; ) {
         ++I;
       else {
         Matches[I] = Matches[--N];
         Matches.resize(N);
       }
     }
   }
 
-  void EliminateSuboptimalCudaMatches() {
-    S.EraseUnwantedCUDAMatches(dyn_cast<FunctionDecl>(S.CurContext), Matches);
-  }
+  void EliminateSuboptimalCudaMatches() { S.EraseUnwantedCUDAMatches(Matches); }
 
 public:
   void ComplainNoMatchesFound() const {
     assert(Matches.empty());
     S.Diag(OvlExpr->getBeginLoc(), diag::err_addr_ovl_no_viable)
         << OvlExpr->getName() << TargetFunctionType
         << OvlExpr->getSourceRange();
     if (FailedCandidates.empty())

clang/test/SemaCUDA/function-overload.cu

 #if defined(__CUDA_ARCH__)
   // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
 #endif
 }
 
 // Test for address of overloaded function resolution in the global context.
 HostFnPtr fp_h = h;
 HostFnPtr fp_ch = ch;
+#if !defined(__CUDA_ARCH__)
 CurrentFnPtr fp_dh = dh;
 CurrentFnPtr fp_cdh = cdh;
+#endif
 GlobalFnPtr fp_g = g;
 
 
 // Test overloading of destructors
 // Can't mix H and unattributed destructors
 struct d_h {
   ~d_h() {} // expected-note {{previous definition is here}}
   __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}

 // Verify that function overloading doesn't prune candidate wrongly.
 int test_constexpr_overload(C2 &x, C2 &y) {
   return constexpr_overload(x, y);
 }
 
 // Verify no ambiguity for new operator.
 void *a = new int;
 __device__ void *b = new int;
 // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 
 // Verify no ambiguity for new operator.
 template<typename _Tp> _Tp&& f();
 template<typename _Tp, typename = decltype(new _Tp(f<_Tp>()))>
 void __test();
 
 void foo() {
   __test<int>();
 }
+
+// Overload resolution in the global initialization should follow the same rule
+// as the one in other places. That is, we prefer a callable function over a
+// non-callable function with a better signature match. In this test case, even
+// though the device function has exactly matching with the integer argument,
+// it can't be executed.
+
+__device__ float fn(int);
+__host__ float fn(float);
+
+float gvar1 = fn(1);
+
+__device__ float dev_only_fn(int);
+// expected-note@-1 {{candidate function not viable: call to __device__ function from __host__ function}}
+
+float gvar2 = dev_only_fn(1); // expected-error {{no matching function for call to 'dev_only_fn'}}
+
+#ifdef __CUDA_ARCH__
+__device__ DeviceReturnTy gvar3 = template_vs_function(1.f);
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__device__ int gvar4 = template_overload(1);
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+#else
+TemplateReturnTy gvar3 = template_vs_function(2.f);
+int gvar4 = template_overload(1);
+#endif

clang/test/SemaCUDA/global-initializers-host.cu

 // RUN: %clang_cc1 %s --std=c++11 -triple x86_64-linux-unknown -fsyntax-only -o - -verify
 
 #include "Inputs/cuda.h"
 
 // Check that we get an error if we try to call a __device__ function from a
 // module initializer.
 
 struct S {
+  // expected-note@-1 {{candidate constructor (the implicit copy constructor) not viable: requires 1 argument, but 0 were provided}}
+  // expected-note@-2 {{candidate constructor (the implicit move constructor) not viable: requires 1 argument, but 0 were provided}}
   __device__ S() {}
-  // expected-note@-1 {{'S' declared here}}
+  // expected-note@-1 {{candidate constructor not viable: call to __device__ function from __host__ function}}
 };
 
 S s;
-// expected-error@-1 {{reference to __device__ function 'S' in global initializer}}
+// expected-error@-1 {{no matching constructor for initialization of 'S'}}
 
 struct T {
   __host__ __device__ T() {}
 };
 T t;  // No error, this is OK.
 
 struct U {
+  // expected-note@-1 {{candidate constructor (the implicit copy constructor) not viable: no known conversion from 'int' to 'const U' for 1st argument}}
+  // expected-note@-2 {{candidate constructor (the implicit move constructor) not viable: no known conversion from 'int' to 'U' for 1st argument}}
   __host__ U() {}
+  // expected-note@-1 {{candidate constructor not viable: requires 0 arguments, but 1 was provided}}
   __device__ U(int) {}
-  // expected-note@-1 {{'U' declared here}}
+  // expected-note@-1 {{candidate constructor not viable: call to __device__ function from __host__ function}}
 };
 U u(42);
-// expected-error@-1 {{reference to __device__ function 'U' in global initializer}}
+// expected-error@-1 {{no matching constructor for initialization of 'U'}}
 
 __device__ int device_fn() { return 42; }
-// expected-note@-1 {{'device_fn' declared here}}
+// expected-note@-1 {{candidate function not viable: call to __device__ function from __host__ function}}
 int n = device_fn();
-// expected-error@-1 {{reference to __device__ function 'device_fn' in global initializer}}
+// expected-error@-1 {{no matching function for call to 'device_fn'}}