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

[Artem Belevich via Phabricator via cfe-commits]

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