[CUDA][HIP] Do not externalize implicit constant static variable

[tra]

clang/lib/AST/ASTContext.cpp

11196

I'd include __device__ vars into the !IsImplicit check. While we may not implicitly add the attribute now, for consistency sake they should be treated the same.

[yaxunl]

clang/lib/AST/ASTContext.cpp

11196

will do when committing.

clang/include/clang/AST/ASTContext.h

SectionInfo(DeclaratorDecl *Decl,
           SectionFlags(SectionFlags) {}
   };
 
   llvm::StringMap<SectionInfo> SectionInfos;
 
   /// Return a new OMPTraitInfo object owned by this context.
   OMPTraitInfo &getNewOMPTraitInfo();
 
+  /// Whether a C++ static variable may be externalized.
+  bool mayExternalizeStaticVar(const Decl *D) const;
+
   /// Whether a C++ static variable should be externalized.
   bool shouldExternalizeStaticVar(const Decl *D) const;
 
 private:
   /// All OMPTraitInfo objects live in this collection, one per
   /// `pragma omp [begin] declare variant` directive.
   SmallVector<std::unique_ptr<OMPTraitInfo>, 4> OMPTraitInfoVector;
 };

clang/lib/AST/ASTContext.cpp

 const DiagnosticBuilder &
 clang::operator<<(const DiagnosticBuilder &DB,
                   const ASTContext::SectionInfo &Section) {
   if (Section.Decl)
     return DB << Section.Decl;
   return DB << "a prior #pragma section";
 }
 
-bool ASTContext::shouldExternalizeStaticVar(const Decl *D) const {
+bool ASTContext::mayExternalizeStaticVar(const Decl *D) const {
   return !getLangOpts().GPURelocatableDeviceCode &&
-         (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()) &&
+         ((D->hasAttr<CUDADeviceAttr>() &&
+           !D->getAttr<CUDADeviceAttr>()->isImplicit()) ||
+          (D->hasAttr<CUDAConstantAttr>() &&
+           !D->getAttr<CUDAConstantAttr>()->isImplicit())) &&
          isa<VarDecl>(D) && cast<VarDecl>(D)->isFileVarDecl() &&
-         cast<VarDecl>(D)->getStorageClass() == SC_Static &&
+         cast<VarDecl>(D)->getStorageClass() == SC_Static;
+}
+
+bool ASTContext::shouldExternalizeStaticVar(const Decl *D) const {
+  return mayExternalizeStaticVar(D) &&
          CUDAStaticDeviceVarReferencedByHost.count(cast<VarDecl>(D));
 }

clang/lib/Sema/SemaExpr.cpp

static void DoMarkVarDeclReferenced(Sema &SemaRef, SourceLocation Loc,
   // by host code. This is done conservatively, when the variable is referenced
   // in any of the following contexts:
   //   - a non-function context
   //   - a host function
   //   - a host device function
   // This also requires the reference of the static device/constant variable by
   // host code to be visible in the device compilation for the compiler to be
   // able to externalize the static device/constant variable.
-  if ((Var->hasAttr<CUDADeviceAttr>() || Var->hasAttr<CUDAConstantAttr>()) &&
-      Var->isFileVarDecl() && Var->getStorageClass() == SC_Static) {
+  if (SemaRef.getASTContext().mayExternalizeStaticVar(Var)) {
     auto *CurContext = SemaRef.CurContext;
     if (!CurContext || !isa<FunctionDecl>(CurContext) ||
         cast<FunctionDecl>(CurContext)->hasAttr<CUDAHostAttr>() ||
         (!cast<FunctionDecl>(CurContext)->hasAttr<CUDADeviceAttr>() &&
          !cast<FunctionDecl>(CurContext)->hasAttr<CUDAGlobalAttr>()))
       SemaRef.getASTContext().CUDAStaticDeviceVarReferencedByHost.insert(Var);
   }
 

clang/test/CodeGenCUDA/static-device-var-no-rdc.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: amdgpu-registered-target
 
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
 // RUN:   -emit-llvm -o - -x hip %s | FileCheck \
 // RUN:   -check-prefixes=DEV %s
 
-// RUN: %clang_cc1 -triple x86_64-gnu-linux \
+// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
 // RUN:   -emit-llvm -o - -x hip %s | FileCheck \
 // RUN:   -check-prefixes=HOST %s
 
 #include "Inputs/cuda.h"
 
 // Test function scope static device variable, which should not be externalized.
 // DEV-DAG: @_ZZ6kernelPiPPKiE1w = internal addrspace(4) constant i32 1
 

 
 static __constant__ int y;
 
 // Test static host variable, which should not be externalized nor registered.
 // HOST-DAG: @_ZL1z = internal global i32 0
 // DEV-NOT: @_ZL1z
 static int z;
 
+// Test implicit static constant variable, which should not be externalized.
+// HOST-DAG: @_ZL2z2 = internal constant i32 456
+// DEV-DAG: @_ZL2z2 = internal addrspace(4) constant i32 456
+
+static constexpr int z2 = 456;
+
 // Test static device variable in inline function, which should not be
 // externalized nor registered.
 // DEV-DAG: @_ZZ6devfunPPKiE1p = linkonce_odr addrspace(4) constant i32 2, comdat
 
 inline __device__ void devfun(const int ** b) {
   const static int p = 2;
   b[0] = &p;
   b[1] = &x2;
 }
 
 __global__ void kernel(int *a, const int **b) {
   const static int w = 1;
   a[0] = x;
   a[1] = y;
   a[2] = x2;
   a[3] = x3;
   a[4] = x4;
   a[5] = x5;
   b[0] = &w;
+  b[1] = &z2;
   devfun(b);
 }
 
 __host__ __device__ void hdf(int *a) {
   a[0] = x3;
 }
 
 int* getDeviceSymbol(int *x);
 
-void foo(int *a) {
+void foo(const int **a) {
   getDeviceSymbol(&x);
   getDeviceSymbol(&x5);
   getDeviceSymbol(&y);
   z = 123;
+  a[0] = &z2;
 }
 
 // HOST: __hipRegisterVar({{.*}}@_ZL1x {{.*}}@[[DEVNAMEX]]
 // HOST: __hipRegisterVar({{.*}}@_ZL1y {{.*}}@[[DEVNAMEY]]
 // HOST-NOT: __hipRegisterVar({{.*}}@_ZZ6kernelPiPPKiE1w
 // HOST-NOT: __hipRegisterVar({{.*}}@_ZZ6devfunPPKiE1p