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

Artem Belevich via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Thu Apr 21 11:52:51 PDT 2022

tra 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
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.

Comment at: clang/lib/AST/ASTContext.cpp:12300
+         (D->hasAttr<CUDAGlobalAttr>() &&
+          GetGVALinkageForFunction(cast<FunctionDecl>(D),
+                                   /*IgnoreCUDAGlobalAttr=*/true) ==
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))`.



More information about the cfe-commits mailing list