[PATCH] D85223: [CUDA][HIP] Support accessing static device variable in host code for -fgpu-rdc

[Yaxun Liu via Phabricator via cfe-commits]

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