yaxunl marked 3 inline comments as done.
yaxunl 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>() &&
----------------
tra wrote:
> 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.
>
>
done
================
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>() &&
----------------
yaxunl wrote:
> tra wrote:
> > 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.
> >
> >
> done
mayExternalizeStaticVar returns true does not mean the static var must be
externalized. mayExternalizeStaticVar only indicates the static var may be
externalized. It is used to enable checking whether this var is used by host
code.
For -fno-gpu-rdc, we only externalize a static variable if it is referenced by
host code. If a static var is referenced by host code, -fno-gpu-rdc will change
its linkage to external, but does not need to make the symbol unique because
each TU ends up as a different device binary.
================
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(
----------------
tra wrote:
> Is this code needed?
this code is not needed. removed.
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
