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

Yaxun Liu via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Thu Apr 21 12:17:56 PDT 2022

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.



More information about the cfe-commits mailing list