yaxunl updated this revision to Diff 264394. yaxunl marked an inline comment as done. yaxunl added a comment.
fix constexpr var in templates CHANGES SINCE LAST ACTION https://reviews.llvm.org/D79237/new/ https://reviews.llvm.org/D79237 Files: clang/include/clang/Sema/Sema.h clang/lib/Sema/SemaCUDA.cpp clang/lib/Sema/SemaDecl.cpp clang/lib/Sema/SemaTemplateInstantiateDecl.cpp clang/test/CodeGenCUDA/constexpr-variables.cu clang/test/SemaCUDA/constexpr-variables.cu
Index: clang/test/SemaCUDA/constexpr-variables.cu =================================================================== --- /dev/null +++ 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<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; Index: clang/test/CodeGenCUDA/constexpr-variables.cu =================================================================== --- /dev/null +++ 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<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; Index: clang/lib/Sema/SemaTemplateInstantiateDecl.cpp =================================================================== --- clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -4836,6 +4836,7 @@ NewVar->setCXXForRangeDecl(OldVar->isCXXForRangeDecl()); NewVar->setObjCForDecl(OldVar->isObjCForDecl()); NewVar->setConstexpr(OldVar->isConstexpr()); + MaybeAddCUDAConstantAttr(NewVar); NewVar->setInitCapture(OldVar->isInitCapture()); NewVar->setPreviousDeclInSameBlockScope( OldVar->isPreviousDeclInSameBlockScope()); Index: clang/lib/Sema/SemaDecl.cpp =================================================================== --- clang/lib/Sema/SemaDecl.cpp +++ clang/lib/Sema/SemaDecl.cpp @@ -7081,6 +7081,7 @@ 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. Index: clang/lib/Sema/SemaCUDA.cpp =================================================================== --- clang/lib/Sema/SemaCUDA.cpp +++ clang/lib/Sema/SemaCUDA.cpp @@ -528,9 +528,14 @@ // 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) @@ -627,6 +632,13 @@ 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"); Index: clang/include/clang/Sema/Sema.h =================================================================== --- clang/include/clang/Sema/Sema.h +++ clang/include/clang/Sema/Sema.h @@ -11656,6 +11656,10 @@ 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. ///
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits