tra added inline comments.
================
Comment at: clang/lib/AST/ASTContext.cpp:11446-11447
bool ASTContext::mayExternalizeStaticVar(const Decl *D) const {
- return !getLangOpts().GPURelocatableDeviceCode &&
+ return (!getLangOpts().CUID.empty() ||
+ !getLangOpts().GPURelocatableDeviceCode) &&
((D->hasAttr<CUDADeviceAttr>() &&
----------------
`!(getLangOpts().GPURelocatableDeviceCode && getLangOpts().CUID.empty())`.
Maybe this should be broken down into something easier to read.
```
// Applies only to -fgpu-rdc or when we were given a CUID
if (!getLangOpts().GPURelocatableDeviceCode || !getLangOpts().CUID.empty()))
return false;
// .. only file-scope static vars...
auto *VD = dyn_cast<VarDecl>(D);
if (!(VD && VD->isFileVarDecl() && VD->getStorageClass() == SC_Static))
return false;
// .. with explicit __device__ or __constant__ attributes.
return ((D->hasAttr<CUDADeviceAttr>() &&
!D->getAttr<CUDADeviceAttr>()->isImplicit()) ||
(D->hasAttr<CUDAConstantAttr>()
&&!D->getAttr<CUDAConstantAttr>()->isImplicit()));
```
================
Comment at: clang/lib/AST/ASTContext.cpp:11446-11447
bool ASTContext::mayExternalizeStaticVar(const Decl *D) const {
- return !getLangOpts().GPURelocatableDeviceCode &&
+ return (!getLangOpts().CUID.empty() ||
+ !getLangOpts().GPURelocatableDeviceCode) &&
((D->hasAttr<CUDADeviceAttr>() &&
----------------
tra wrote:
> `!(getLangOpts().GPURelocatableDeviceCode && getLangOpts().CUID.empty())`.
>
> Maybe this should be broken down into something easier to read.
> ```
> // Applies only to -fgpu-rdc or when we were given a CUID
> if (!getLangOpts().GPURelocatableDeviceCode || !getLangOpts().CUID.empty()))
> return false;
> // .. only file-scope static vars...
> auto *VD = dyn_cast<VarDecl>(D);
> if (!(VD && VD->isFileVarDecl() && VD->getStorageClass() == SC_Static))
> return false;
> // .. with explicit __device__ or __constant__ attributes.
> return ((D->hasAttr<CUDADeviceAttr>() &&
> !D->getAttr<CUDADeviceAttr>()->isImplicit()) ||
> (D->hasAttr<CUDAConstantAttr>()
> &&!D->getAttr<CUDAConstantAttr>()->isImplicit()));
>
> ```
BTW, does this mean that we'll externalize & uniquify the vars even w/o
`-fgpu-rdc` if CUID is given?
IMO `-fgpu-rdc` should remain the flag to control whether externalization is
needed.
CUID controls the value of a unique suffix, if we need it, but should not
automatically enable externalization.
================
Comment at: clang/lib/CodeGen/CodeGenModule.cpp:2864-2865
+#if 0
+ // We need to decide whether to externalize a static variable after checking
+ // whether it is referenced in host code.
+ if (isa<VarDecl>(Global) && getContext().mayExternalizeStaticVar(
----------------
Is this code needed?
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D85223/new/
https://reviews.llvm.org/D85223
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
