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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
