[CUDA][HIP] Do not promote constexpr var with non-constant initializer

[yaxunl]

Summary

constexpr var may be initialized with address of non-const variable.
In this case the initializer is not constant in device compilation.
This has been handled for const vars but not for constexpr vars.

This patch makes handling of const var and constexpr var
consistent.

Fixes: https://github.com/llvm/llvm-project/issues/53780

[tra]

clang/lib/Sema/SemaCUDA.cpp

148–153

So the idea here is that the constexpr vars that have an initializer that may not be static enough for CUDA would not get the implicit const attribute. I.e. it's tied to the change below that checks for allowed initializers.
We could use a comment about that here.

[yaxunl]

clang/lib/Sema/SemaCUDA.cpp

148–153

will add a comment when committing

clang/lib/Sema/SemaCUDA.cpp

Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D,
 
   return CFT_Host;
 }
 
 /// IdentifyTarget - Determine the CUDA compilation target for this variable.
 Sema::CUDAVariableTarget Sema::IdentifyCUDATarget(const VarDecl *Var) {
   if (Var->hasAttr<HIPManagedAttr>())
     return CVT_Unified;
-  if (Var->isConstexpr() && !hasExplicitAttr<CUDAConstantAttr>(Var))
-    return CVT_Both;
-  if (Var->getType().isConstQualified() && Var->hasAttr<CUDAConstantAttr>() &&
+  // Only constexpr and const variabless with implicit constant attribute
+  // are emitted on both sides. Such variables are promoted to device side
+  // only if they have static constant intializers on device side.
+  if ((Var->isConstexpr() || Var->getType().isConstQualified()) &&
+      Var->hasAttr<CUDAConstantAttr>() &&
       !hasExplicitAttr<CUDAConstantAttr>(Var))
     return CVT_Both;
   if (Var->hasAttr<CUDADeviceAttr>() || Var->hasAttr<CUDAConstantAttr>() ||
       Var->hasAttr<CUDASharedAttr>() ||
       Var->getType()->isCUDADeviceBuiltinSurfaceType() ||
       Var->getType()->isCUDADeviceBuiltinTextureType())
     return CVT_Device;
   // Function-scope static variable without explicit device or constant
   // attribute are emitted

 // `__constant__` memory usage goes beyond limit.
 void Sema::MaybeAddCUDAConstantAttr(VarDecl *VD) {
   // Do not promote dependent variables since the cotr/dtor/initializer are
   // not determined. Do it after instantiation.
   if (getLangOpts().CUDAIsDevice && !VD->hasAttr<CUDAConstantAttr>() &&
       !VD->hasAttr<CUDAConstantAttr>() && !VD->hasAttr<CUDASharedAttr>() &&
       (VD->isFileVarDecl() || VD->isStaticDataMember()) &&
       !IsDependentVar(VD) &&
-      (VD->isConstexpr() || (VD->getType().isConstQualified() &&
-                             HasAllowedCUDADeviceStaticInitializer(
-                                 *this, VD, CICK_DeviceOrConstant)))) {
+      ((VD->isConstexpr() || VD->getType().isConstQualified()) &&
+       HasAllowedCUDADeviceStaticInitializer(*this, VD,
+                                             CICK_DeviceOrConstant))) {
     VD->addAttr(CUDAConstantAttr::CreateImplicit(getASTContext()));
   }
 }
 
 Sema::SemaDiagnosticBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc,
                                                        unsigned DiagID) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
   SemaDiagnosticBuilder::Kind DiagKind = [&] {

clang/test/SemaCUDA/constexpr-var.cu

+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
+// RUN:   -fsyntax-only -verify
+// RUN: %clang_cc1 -triple x86_64 -x hip %s \
+// RUN:   -fsyntax-only -verify=host
+
+// host-no-diagnostics
+
+#include "Inputs/cuda.h"
+
+// Test constexpr var initialized with address of a const var.
+// Both are promoted to device side.
+
+namespace Test1 {
+const int a = 1;
+
+struct B {
+    static constexpr const int *p = &a;
+    __device__ static constexpr const int *const p2 = &a;
+};
+
+// Const variable 'a' is treated as __constant__ on device side,
+// therefore its address can be used as initializer for another
+// device variable.
+
+__device__ void f() {
+  int y = a;
+  constexpr const int *x = B::p;
+  constexpr const int *z = B::p2;
+}
+}
+
+// Test constexpr var initialized with address of a non-cost var.
+// Neither is promoted to device side.
+
+namespace Test2 {
+int a = 1;
+// expected-note@-1{{host variable declared here}}
+
+struct B {
+    static constexpr int *const p = &a;
+    // expected-note@-1{{const variable cannot be emitted on device side due to dynamic initialization}}
+};
+
+__device__ void f() {
+  int y = a;
+  // expected-error@-1{{reference to __host__ variable 'a' in __device__ function}}
+  const int *const *x = &B::p;
+  // expected-error@-1{{reference to __host__ variable 'p' in __device__ function}}
+  // ToDo: use of non-promotable constexpr variable in device compilation should be treated as
+  // ODR-use and diagnosed.
+  const int *const z = B::p;
+}
+}
+
+// Test constexpr device var initialized with address of a non-const host var, __shared var,
+// __managed__ var, __device__ var, __constant__ var, texture var, surface var.
+
+namespace Test3 {
+struct textureReference {
+  int desc;
+};
+
+enum ReadMode {
+  ElementType = 0,
+  NormalizedFloat = 1
+};
+
+template <typename T, int dim = 1, enum ReadMode mode = ElementType>
+struct __attribute__((device_builtin_texture_type)) texture : public textureReference {
+};
+
+struct surfaceReference {
+  int desc;
+};
+
+template <typename T, int dim = 1>
+struct __attribute__((device_builtin_surface_type)) surface : public surfaceReference {
+};
+
+// Partial specialization over `void`.
+template<int dim>
+struct __attribute__((device_builtin_surface_type)) surface<void, dim> : public surfaceReference {
+};
+
+texture<float, 2, ElementType> tex;
+surface<void, 2> surf;
+
+int a = 1;
+__shared__ int b;
+__managed__ int c = 1;
+__device__ int d = 1;
+__constant__ int e = 1;
+struct B {
+    __device__ static constexpr int *const p1 = &a;
+    // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+    __device__ static constexpr int *const p2 = &b;
+    // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+    __device__ static constexpr int *const p3 = &c;
+    // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+    __device__ static constexpr int *const p4 = &d;
+    __device__ static constexpr int *const p5 = &e;
+    __device__ static constexpr texture<float, 2, ElementType> *const p6 = &tex;
+    __device__ static constexpr surface<void, 2> *const p7 = &surf;
+};
+}