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

[Artem Belevich via Phabricator via cfe-commits]

tra added a subscriber: rsmith.
tra added a comment.

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



================
Comment at: clang/include/clang/Sema/Sema.h:11037
+                                        bool IgnoreImplicitHDAttr = false);
+  CUDAFunctionTarget IdentifyCUDATarget(const Decl *D,
+                                        bool IgnoreImplicitHDAttr = false);
----------------
I'd add a comment describing that it's a wrapper which dispatches the call to one of more specific variants above.


================
Comment at: clang/include/clang/Sema/Sema.h:11061
   /// combination, based on their host/device attributes.
   /// \param Caller function which needs address of \p Callee.
   ///               nullptr in case of global context.
----------------
Comment needs updating.


================
Comment at: clang/include/clang/Sema/Sema.h:11073
+  void popCUDANonLocalVariable(const Decl *D);
+  const Decl *getCUDACurrentNonLocalVariable() const {
+    return CUDANonLocalVariableStack.empty() ? nullptr
----------------
Nit: I'd add an empty line between delarations and the function. Jammed together they are hard to read.


================
Comment at: clang/include/clang/Sema/Sema.h:11096-11098
+    const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext);
+    if (Kind == SkipImplicitCaller && Caller && Caller->isImplicit())
+      return true;
----------------
```
if (const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext))
  if (Kind == SkipImplicitCaller && Caller->isImplicit())
      return true;
```


================
Comment at: clang/include/clang/Sema/Sema.h:11136
   void EraseUnwantedCUDAMatches(
-      const FunctionDecl *Caller,
+      const Decl *ContextDecl,
       SmallVectorImpl<std::pair<DeclAccessPair, FunctionDecl *>> &Matches);
----------------
Now that we always use getCUDAContextDecl() as the first argument, perhaps we can just always retrieve the context inside the function.


================
Comment at: clang/lib/Parse/ParseDecl.cpp:2336
 
+  Actions.pushCUDANonLocalVariable(ThisDecl);
+
----------------
@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. 


================
Comment at: clang/test/SemaCUDA/function-overload.cu:428
+
+// Overload resolution should follow the same rule in the global
+// initialization.
----------------
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.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D71227/new/

https://reviews.llvm.org/D71227