[CUDA][HIP] Fix constexpr variables for C++17

[yaxunl]

Summary

constexpr variables are compile time constants and implicitly const, therefore
they are safe to emit on both device and host side. Besides, in many cases
they are intended for both device and host, therefore it makes sense
to emit them on both device and host sides if necessary.

In most cases constexpr variables are used as rvalue and the variables
themselves do not need to be emitted. However if their address is taken,
then they need to be emitted.

The following example shows constexpr is available on device side without
__device__ or __constant__ attribute

https://godbolt.org/z/Uf7CgK

Which indicates that we need to emit constexpr variables on device side
even without __device__ or __constant__ attribute when necessary.
This should be OK since the initializer is compile time constant and the
variable itself is constant. They can just be emitted in the same way
as the host side.

For C++14, clang is able to handle that since clang emits them with
available_externally linkage together with the initializer.

However for C++17, the constexpr static data member of a class or template class
become inline variables implicitly. Therefore they become definitions with
linkonce_odr or weak_odr linkages. As such, they can not have available_externally
linkage.

This patch fixes that by adding implicit constant attribute to
file scope constexpr variables and constexpr static data members
in device compilation.

[yaxunl]

ping

[tra]

constexpr variables are compile time constants and implicitly const, therefore
they are safe to emit on both device and host side. Besides, in many cases
they are intended for both device and host, therefore it makes sense
to emit them on both device and host sides if necessary.

In most cases constexpr variables are used as rvalue and the variables
themselves do not need to be emitted.

Agreed so far. constexpr values are usable on both sides.

However if their address is taken, then they need to be emitted.

Agreed.

That's where the fun starts. Will this assertion trigger or not?

constexpr int ce = 42;
__global__ void kernel() {
  assert(p == &ce); 
}
void f() {
  kernel<<<1,1>>>(&ce);
}

NVCC does not allow accessing ce, so this does not compile.
Clang allows taking a reference, but the variable is not emitted, so the failure will happen in ptxas.
If compiler does emit ce on device side, the addresses will be different, unless we rely on
CUDA runtime magic to translate addresses of shadows into real device-side address.
If we do that, then we should probably add an implicit constant on the constexpr vars.

The following example shows constexpr is available on device side without
device or constant attribute

https://godbolt.org/z/Uf7CgK

If you clear output filters, you will see that the reference to constexpr is emitted as .extern .global .align 4 .u32 _ZN1B1bE; and the varible's address is not actually available. If you were to actually compile the example, ptxas would complain about ti.
NVCC, predictably, allows using constexpr values, but not the variables themselves:
https://godbolt.org/z/55dEx-

[yaxunl]

In D79237#2031870, @tra wrote:

constexpr variables are compile time constants and implicitly const, therefore
they are safe to emit on both device and host side. Besides, in many cases
they are intended for both device and host, therefore it makes sense
to emit them on both device and host sides if necessary.

In most cases constexpr variables are used as rvalue and the variables
themselves do not need to be emitted.

Agreed so far. constexpr values are usable on both sides.

However if their address is taken, then they need to be emitted.

Agreed.

That's where the fun starts. Will this assertion trigger or not?

constexpr int ce = 42;
__global__ void kernel() {
  assert(p == &ce); 
}
void f() {
  kernel<<<1,1>>>(&ce);
}

NVCC does not allow accessing ce, so this does not compile.
Clang allows taking a reference, but the variable is not emitted, so the failure will happen in ptxas.
If compiler does emit ce on device side, the addresses will be different, unless we rely on
CUDA runtime magic to translate addresses of shadows into real device-side address.
If we do that, then we should probably add an implicit constant on the constexpr vars.

The following example shows constexpr is available on device side without
device or constant attribute

https://godbolt.org/z/Uf7CgK

If you clear output filters, you will see that the reference to constexpr is emitted as .extern .global .align 4 .u32 _ZN1B1bE; and the varible's address is not actually available. If you were to actually compile the example, ptxas would complain about ti.
NVCC, predictably, allows using constexpr values, but not the variables themselves:
https://godbolt.org/z/55dEx-

It seems adding an implicit __constant__ attribute to constexpr variables on device side may be a possible solution. I will give it a try.

[yaxunl]

add implicit constant attribute to constexpr file scope variables and constexpr static data members in device compilation.

[tra]

LGTM in general. Let me check the patch on our tensorflow build.

clang/test/SemaCUDA/constexpr-variables.cu

31–32

Can we verify the diags for bad cases, too?

[yaxunl]

clang/test/SemaCUDA/constexpr-variables.cu

31–32

By bad cases you mean the constexpr var is not compile time constant?

[tra]

In D79237#2039417, @tra wrote:

LGTM in general. Let me check the patch on our tensorflow build.

Bad news -- it breaks the standard C++ library.

Reproducer:

$ bin/clang++ -x cuda /dev/null -fsyntax-only -include algorithm --cuda-path=$HOME/local/cuda-10.1 --cuda-device-only --cuda-gpu-arch=sm_60 -std=c++17 -stdlib=libc++                                              tra@art3:~/work/llvm/build/release+assert+zapcc

In file included from <built-in>:2:
In file included from build/release+assert+zapcc/lib/clang/11.0.0/include/cuda_wrappers/algorithm:55:
In file included from build/release+assert+zapcc/bin/../include/c++/v1/algorithm:642:
build/release+assert+zapcc/bin/../include/c++/v1/utility:937:51: error: dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.
_LIBCPP_INLINE_VAR constexpr in_place_type_t<_Tp> in_place_type{};
                                                  ^            ~~
build/release+assert+zapcc/bin/../include/c++/v1/utility:944:53: error: dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.
_LIBCPP_INLINE_VAR constexpr in_place_index_t<_Idx> in_place_index{};
                                                    ^             ~~
2 errors generated when compiling for sm_60.

[tra]

clang/test/SemaCUDA/constexpr-variables.cu

31–32

It's a general suggestion. It's good to have some negative testing thrown in, too to make sure we didn't do too much by mistake. E.g. allowing all variables to become __constant__ would be bad. :-)

I'm not sure what exactly would be a good addition here. If you can't think of something, I guess we can leave it as is. C++ is largely oblivoius of __constant__, so we may not have much to work with. It's probably observable, but may be an overkill for the test.

[rsmith]

In D79237#2039559, @tra wrote:

In D79237#2039417, @tra wrote:

Bad news -- it breaks the standard C++ library.

[...]

build/release+assert+zapcc/bin/../include/c++/v1/utility:937:51: error: dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.
_LIBCPP_INLINE_VAR constexpr in_place_type_t<_Tp> in_place_type{};
                                                  ^            ~~

This looks like a bug in that diagnostic: the instantiations of this variable template certainly do not have dynamic initialization. The diagnostic is missing an "in instantiation of" note, so I think the bug is that we're performing this check on a dependent variable prior to instantiation. Presumably we should delay the check to instantiation time if either the type of the variable is dependent or the initializer is value-dependent.

[yaxunl]

fix constexpr var in templates

[yaxunl]

clang/test/SemaCUDA/constexpr-variables.cu

31–32

added tests where initializer is not compile time constant

[yaxunl]

In D79237#2039757, @rsmith wrote:

In D79237#2039559, @tra wrote:

In D79237#2039417, @tra wrote:

Bad news -- it breaks the standard C++ library.

[...]

build/release+assert+zapcc/bin/../include/c++/v1/utility:937:51: error: dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.
_LIBCPP_INLINE_VAR constexpr in_place_type_t<_Tp> in_place_type{};
                                                  ^            ~~

This looks like a bug in that diagnostic: the instantiations of this variable template certainly do not have dynamic initialization. The diagnostic is missing an "in instantiation of" note, so I think the bug is that we're performing this check on a dependent variable prior to instantiation. Presumably we should delay the check to instantiation time if either the type of the variable is dependent or the initializer is value-dependent.

right. fixed

[yaxunl]

ping

[tra]

Tested with tensorflow build. The patch- does not seem to break anything now.

clang/include/clang/Sema/Sema.h

bool IsAllowedCUDACall(const FunctionDecl *Caller,
     return IdentifyCUDAPreference(Caller, Callee) != CFP_Never;
   }
 
   /// May add implicit CUDAHostAttr and CUDADeviceAttr attributes to FD,
   /// depending on FD and the current compilation settings.
   void maybeAddCUDAHostDeviceAttrs(FunctionDecl *FD,
                                    const LookupResult &Previous);
 
+  /// May add implicit CUDAConstantAttr attribute to VD, depending on VD
+  /// and current compilation settings.
+  void MaybeAddCUDAConstantAttr(VarDecl *VD);
+
 public:
   /// Check whether we're allowed to call Callee from the current context.
   ///
   /// - If the call is never allowed in a semantically-correct program
   ///   (CFP_Never), emits an error and returns false.
   ///
   /// - If the call is allowed in semantically-correct programs, but only if
   ///   it's never codegen'ed (CFP_WrongSide), creates a deferred diagnostic to

clang/lib/Sema/SemaCUDA.cpp

if (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>() ||
     bool AllowedInit = false;
     if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init))
       AllowedInit =
           isEmptyCudaConstructor(VD->getLocation(), CE->getConstructor());
     // We'll allow constant initializers even if it's a non-empty
     // constructor according to CUDA rules. This deviates from NVCC,
     // but allows us to handle things like constexpr constructors.
     if (!AllowedInit &&
-        (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>()))
-      AllowedInit = VD->getInit()->isConstantInitializer(
-          Context, VD->getType()->isReferenceType());
+        (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>())) {
+      auto *Init = VD->getInit();
+      AllowedInit =
+          ((VD->getType()->isDependentType() || Init->isValueDependent()) &&
+           VD->isConstexpr()) ||
+          Init->isConstantInitializer(Context,
+                                      VD->getType()->isReferenceType());
+    }
 
     // Also make sure that destructor, if there is one, is empty.
     if (AllowedInit)
       if (CXXRecordDecl *RD = VD->getType()->getAsCXXRecordDecl())
         AllowedInit =
             isEmptyCudaDestructor(VD->getLocation(), RD->getDestructor());
 
     if (!AllowedInit) {
if (It != Previous.end()) {
     }
     return;
   }
 
   NewD->addAttr(CUDAHostAttr::CreateImplicit(Context));
   NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
 }
 
+void Sema::MaybeAddCUDAConstantAttr(VarDecl *VD) {
+  if (getLangOpts().CUDAIsDevice && VD->isConstexpr() &&
+      (VD->isFileVarDecl() || VD->isStaticDataMember())) {
+    VD->addAttr(CUDAConstantAttr::CreateImplicit(getASTContext()));
+  }
+}
+
 Sema::DeviceDiagBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc,
                                                    unsigned DiagID) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
   DeviceDiagBuilder::Kind DiagKind = [this] {
     switch (CurrentCUDATarget()) {
     case CFT_Global:
     case CFT_Device:
       return DeviceDiagBuilder::K_Immediate;

clang/lib/Sema/SemaDecl.cpp

NamedDecl *Sema::ActOnVariableDeclarator(
   case CSK_consteval:
     Diag(D.getDeclSpec().getConstexprSpecLoc(),
         diag::err_constexpr_wrong_decl_kind)
       << D.getDeclSpec().getConstexprSpecifier();
     LLVM_FALLTHROUGH;
 
   case CSK_constexpr:
     NewVD->setConstexpr(true);
+    MaybeAddCUDAConstantAttr(NewVD);
     // C++1z [dcl.spec.constexpr]p1:
     //   A static data member declared with the constexpr specifier is
     //   implicitly an inline variable.
     if (NewVD->isStaticDataMember() &&
         (getLangOpts().CPlusPlus17 ||
          Context.getTargetInfo().getCXXABI().isMicrosoft()))
       NewVD->setImplicitlyInline();
     break;

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

if (OldVar->isLocalExternDecl()) {
     NewVar->setLexicalDeclContext(Owner);
   } else if (OldVar->isOutOfLine())
     NewVar->setLexicalDeclContext(OldVar->getLexicalDeclContext());
   NewVar->setTSCSpec(OldVar->getTSCSpec());
   NewVar->setInitStyle(OldVar->getInitStyle());
   NewVar->setCXXForRangeDecl(OldVar->isCXXForRangeDecl());
   NewVar->setObjCForDecl(OldVar->isObjCForDecl());
   NewVar->setConstexpr(OldVar->isConstexpr());
+  MaybeAddCUDAConstantAttr(NewVar);
   NewVar->setInitCapture(OldVar->isInitCapture());
   NewVar->setPreviousDeclInSameBlockScope(
       OldVar->isPreviousDeclInSameBlockScope());
   NewVar->setAccess(OldVar->getAccess());
 
   if (!OldVar->isStaticDataMember()) {
     if (OldVar->isUsed(false))
       NewVar->setIsUsed();

clang/test/CodeGenCUDA/constexpr-variables.cu

+// RUN: %clang_cc1 -std=c++14 %s -emit-llvm -o - -triple nvptx \
+// RUN:   -fcuda-is-device | FileCheck --check-prefixes=CXX14 %s
+// RUN: %clang_cc1 -std=c++17 %s -emit-llvm -o - -triple nvptx \
+// RUN:   -fcuda-is-device | FileCheck --check-prefixes=CXX17 %s
+
+#include "Inputs/cuda.h"
+
+// COM: @_ZL1a = internal {{.*}}constant i32 7
+constexpr int a = 7;
+__constant__ const int &use_a = a;
+
+namespace B {
+ // COM: @_ZN1BL1bE = internal {{.*}}constant i32 9
+  constexpr int b = 9;
+}
+__constant__ const int &use_B_b = B::b;
+
+struct Q {
+  // CXX14: @_ZN1Q2k2E = {{.*}}externally_initialized constant i32 6
+  // CXX17: @_ZN1Q2k2E = internal {{.*}}constant i32 6
+  // CXX14: @_ZN1Q2k1E = available_externally {{.*}}constant i32 5
+  // CXX17: @_ZN1Q2k1E = linkonce_odr {{.*}}constant i32 5
+  static constexpr int k1 = 5;
+  static constexpr int k2 = 6;
+};
+constexpr int Q::k2;
+
+__constant__ const int &use_Q_k1 = Q::k1;
+__constant__ const int &use_Q_k2 = Q::k2;
+
+template<typename T> struct X {
+  // CXX14: @_ZN1XIiE1aE = available_externally {{.*}}constant i32 123
+  // CXX17: @_ZN1XIiE1aE = linkonce_odr {{.*}}constant i32 123
+  static constexpr int a = 123;
+};
+__constant__ const int &use_X_a = X<int>::a;
+
+template <typename T, T a, T b> struct A {
+  // CXX14: @_ZN1AIiLi1ELi2EE1xE = available_externally {{.*}}constant i32 2
+  // CXX17: @_ZN1AIiLi1ELi2EE1xE = linkonce_odr {{.*}}constant i32 2
+  constexpr static T x = a * b;
+};
+__constant__ const int &y = A<int, 1, 2>::x;

clang/test/SemaCUDA/constexpr-variables.cu

+// RUN: %clang_cc1 -std=c++14 %s -emit-llvm -o - -triple nvptx64-nvidia-cuda \
+// RUN:   -fcuda-is-device -verify -fsyntax-only
+// RUN: %clang_cc1 -std=c++17 %s -emit-llvm -o - -triple nvptx64-nvidia-cuda \
+// RUN:   -fcuda-is-device -verify -fsyntax-only
+// RUN: %clang_cc1 -std=c++14 %s -emit-llvm -o - \
+// RUN:   -triple x86_64-unknown-linux-gnu -verify -fsyntax-only
+// RUN: %clang_cc1 -std=c++17 %s -emit-llvm -o - \
+// RUN:   -triple x86_64-unknown-linux-gnu -verify -fsyntax-only
+#include "Inputs/cuda.h"
+
+template<typename T>
+__host__ __device__ void foo(const T **a) {
+  // expected-note@-1 {{declared here}}
+  static const T b = sizeof(a);
+  static constexpr T c = sizeof(a);
+  const T d = sizeof(a);
+  constexpr T e = sizeof(a);
+  constexpr T f = **a;
+  // expected-error@-1 {{constexpr variable 'f' must be initialized by a constant expression}}
+  // expected-note@-2 {{read of non-constexpr variable 'a' is not allowed in a constant expression}}
+  a[0] = &b;
+  a[1] = &c;
+  a[2] = &d;
+  a[3] = &e;
+}
+
+__device__ void device_fun(const int **a) {
+  // expected-note@-1 {{declared here}}
+  constexpr int b = sizeof(a);
+  static constexpr int c = sizeof(a);
+  constexpr int d = **a;
+  // expected-error@-1 {{constexpr variable 'd' must be initialized by a constant expression}}
+  // expected-note@-2 {{read of non-constexpr variable 'a' is not allowed in a constant expression}}
+  a[0] = &b;
+  a[1] = &c;
+  foo(a);
+  // expected-note@-1 {{in instantiation of function template specialization 'foo<int>' requested here}}
+}
+
+void host_fun(const int **a) {
+  // expected-note@-1 {{declared here}}
+  constexpr int b = sizeof(a);
+  static constexpr int c = sizeof(a);
+  constexpr int d = **a;
+  // expected-error@-1 {{constexpr variable 'd' must be initialized by a constant expression}}
+  // expected-note@-2 {{read of non-constexpr variable 'a' is not allowed in a constant expression}}
+  a[0] = &b;
+  a[1] = &c;
+  foo(a);
+}
+
+__host__ __device__ void host_device_fun(const int **a) {
+  // expected-note@-1 {{declared here}}
+  constexpr int b = sizeof(a);
+  static constexpr int c = sizeof(a);
+  constexpr int d = **a;
+  // expected-error@-1 {{constexpr variable 'd' must be initialized by a constant expression}}
+  // expected-note@-2 {{read of non-constexpr variable 'a' is not allowed in a constant expression}}
+  a[0] = &b;
+  a[1] = &c;
+  foo(a);
+}
+
+template <class T>
+struct A {
+  explicit A() = default;
+};
+template <class T>
+constexpr A<T> a{};
+
+struct B {
+  static constexpr bool value = true;
+};
+
+template<typename T>
+struct C {
+  static constexpr bool value = T::value;
+};
+
+__constant__ const bool &x = C<B>::value;