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
