[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))`.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D124189/new/
https://reviews.llvm.org/D124189
More information about the cfe-commits
mailing list