diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 8409abc4caabc..c87777c0a6a62 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -11724,6 +11724,10 @@ class Sema final { 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. /// diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index 73d190891b0fe..5d6c15196750f 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -513,9 +513,14 @@ void Sema::checkAllowedCUDAInitializer(VarDecl *VD) { // constructor according to CUDA rules. This deviates from NVCC, // but allows us to handle things like constexpr constructors. if (!AllowedInit && - (VD->hasAttr() || VD->hasAttr())) - AllowedInit = VD->getInit()->isConstantInitializer( - Context, VD->getType()->isReferenceType()); + (VD->hasAttr() || VD->hasAttr())) { + 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) @@ -612,6 +617,13 @@ void Sema::maybeAddCUDAHostDeviceAttrs(FunctionDecl *NewD, 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"); diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 76754adbf20bd..aec3d551701be 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -7100,6 +7100,7 @@ NamedDecl *Sema::ActOnVariableDeclarator( 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. diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index 327022218e016..519d9128037db 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -4841,6 +4841,7 @@ void Sema::BuildVariableInstantiation( NewVar->setCXXForRangeDecl(OldVar->isCXXForRangeDecl()); NewVar->setObjCForDecl(OldVar->isObjCForDecl()); NewVar->setConstexpr(OldVar->isConstexpr()); + MaybeAddCUDAConstantAttr(NewVar); NewVar->setInitCapture(OldVar->isInitCapture()); NewVar->setPreviousDeclInSameBlockScope( OldVar->isPreviousDeclInSameBlockScope()); diff --git a/clang/test/CodeGenCUDA/constexpr-variables.cu b/clang/test/CodeGenCUDA/constexpr-variables.cu new file mode 100644 index 0000000000000..b8b0782b4f62f --- /dev/null +++ b/clang/test/CodeGenCUDA/constexpr-variables.cu @@ -0,0 +1,43 @@ +// 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 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::a; + +template 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::x; diff --git a/clang/test/SemaCUDA/constexpr-variables.cu b/clang/test/SemaCUDA/constexpr-variables.cu new file mode 100644 index 0000000000000..6e17a08568381 --- /dev/null +++ b/clang/test/SemaCUDA/constexpr-variables.cu @@ -0,0 +1,80 @@ +// 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 +__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' 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 +struct A { + explicit A() = default; +}; +template +constexpr A a{}; + +struct B { + static constexpr bool value = true; +}; + +template +struct C { + static constexpr bool value = T::value; +}; + +__constant__ const bool &x = C::value;