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

[Yaxun Liu via cfe-commits]

Author: Yaxun (Sam) Liu
Date: 2022-02-15T15:15:55-05:00
New Revision: 73b22935a7a863679021598db6a45fcfb62cd321

URL: https://github.com/llvm/llvm-project/commit/73b22935a7a863679021598db6a45fcfb62cd321
DIFF: https://github.com/llvm/llvm-project/commit/73b22935a7a863679021598db6a45fcfb62cd321.diff

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

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.

Reviewed by: Artem Belevich

Differential Revision: https://reviews.llvm.org/D119615

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

Added: 
    clang/test/SemaCUDA/constexpr-var.cu

Modified: 
    clang/lib/Sema/SemaCUDA.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index efa38554bc83f..e4e34d687dd2b 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -145,9 +145,11 @@ Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D,
 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>() ||
@@ -718,9 +720,9 @@ void Sema::MaybeAddCUDAConstantAttr(VarDecl *VD) {
       !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()));
   }
 }

diff  --git a/clang/test/SemaCUDA/constexpr-var.cu b/clang/test/SemaCUDA/constexpr-var.cu
new file mode 100644
index 0000000000000..a028ba8f6c1a1
--- /dev/null
+++ b/clang/test/SemaCUDA/constexpr-var.cu
@@ -0,0 +1,105 @@
+// 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 at -1{{host variable declared here}}
+
+struct B {
+    static constexpr int *const p = &a;
+    // expected-note at -1{{const variable cannot be emitted on device side due to dynamic initialization}}
+};
+
+__device__ void f() {
+  int y = a;
+  // expected-error at -1{{reference to __host__ variable 'a' in __device__ function}}
+  const int *const *x = &B::p;
+  // expected-error at -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 at -1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+    __device__ static constexpr int *const p2 = &b;
+    // expected-error at -1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+    __device__ static constexpr int *const p3 = &c;
+    // expected-error at -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;
+};
+}