[PATCH] D124189: [CUDA][HIP] Externalize kernels with internal linkage

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl marked an inline comment as done.
yaxunl added inline comments.


================
Comment at: clang/lib/AST/ASTContext.cpp:11322
+  } else if (Context.getLangOpts().CUDA && Context.getLangOpts().CUDAIsDevice &&
+             !IgnoreCUDAGlobalAttr) {
     // Device-side functions with __global__ attribute must always be
----------------
tra wrote:
> Nit: I'd phrase it as a positive assertion `ConsiderCudaGlobalAttr` and default it to true.
> 
> `DontDoX` always trips me and gets me to question it -- "what *are* we doing then? what else is there besides X?".
> With  a `DoX` things are usually simpler and limited to `X` -- we're either doing X or not.
> 
will do


================
Comment at: clang/lib/AST/ASTContext.cpp:12300
+         (D->hasAttr<CUDAGlobalAttr>() &&
+          GetGVALinkageForFunction(cast<FunctionDecl>(D),
+                                   /*IgnoreCUDAGlobalAttr=*/true) ==
----------------
tra wrote:
> Perhaps we don't need to change the public AST API and plumb `IgnoreCUDAGlobalAttr` through.
> We cold create CUDA-aware static version of `GetGVALinkageForCudaKernel` instead, which would call `adjustGVALinkageForExternalDefinitionKind(..., adjustGVALinkageForAttributes(IgnoreCUDAGlobalAttr=true))`.
We could have a static function but it would be GetGVALinkageForCUDAKernelWithoutGlobalAttr since we need to know the linkage of the kernel assuming it has no `__global__` attribute.

If you think it is OK I can make the change.


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

https://reviews.llvm.org/D124189