[clang] 049d860 - [CUDA][HIP] Fix constexpr variables for C++17
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Wed Jun 3 18:57:28 PDT 2020
Author: Yaxun (Sam) Liu
Date: 2020-06-03T21:56:52-04:00
New Revision: 049d860707ef22978b9379fee6dce38c66a22671
URL: https://github.com/llvm/llvm-project/commit/049d860707ef22978b9379fee6dce38c66a22671
DIFF: https://github.com/llvm/llvm-project/commit/049d860707ef22978b9379fee6dce38c66a22671.diff
LOG: [CUDA][HIP] Fix constexpr variables for C++17
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.
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.
Differential Revision: https://reviews.llvm.org/D79237
Added:
clang/test/CodeGenCUDA/constexpr-variables.cu
clang/test/SemaCUDA/constexpr-variables.cu
Modified:
clang/include/clang/Sema/Sema.h
clang/lib/Sema/SemaCUDA.cpp
clang/lib/Sema/SemaDecl.cpp
clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 8409abc4caab..c87777c0a6a6 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 73d190891b0f..5d6c15196750 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<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)
@@ -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 76754adbf20b..aec3d551701b 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 327022218e01..519d9128037d 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 000000000000..b8b0782b4f62
--- /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<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;
diff --git a/clang/test/SemaCUDA/constexpr-variables.cu b/clang/test/SemaCUDA/constexpr-variables.cu
new file mode 100644
index 000000000000..6e17a0856838
--- /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<typename T>
+__host__ __device__ void foo(const T **a) {
+ // expected-note at -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 at -1 {{constexpr variable 'f' must be initialized by a constant expression}}
+ // expected-note at -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 at -1 {{declared here}}
+ constexpr int b = sizeof(a);
+ static constexpr int c = sizeof(a);
+ constexpr int d = **a;
+ // expected-error at -1 {{constexpr variable 'd' must be initialized by a constant expression}}
+ // expected-note at -2 {{read of non-constexpr variable 'a' is not allowed in a constant expression}}
+ a[0] = &b;
+ a[1] = &c;
+ foo(a);
+ // expected-note at -1 {{in instantiation of function template specialization 'foo<int>' requested here}}
+}
+
+void host_fun(const int **a) {
+ // expected-note at -1 {{declared here}}
+ constexpr int b = sizeof(a);
+ static constexpr int c = sizeof(a);
+ constexpr int d = **a;
+ // expected-error at -1 {{constexpr variable 'd' must be initialized by a constant expression}}
+ // expected-note at -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 at -1 {{declared here}}
+ constexpr int b = sizeof(a);
+ static constexpr int c = sizeof(a);
+ constexpr int d = **a;
+ // expected-error at -1 {{constexpr variable 'd' must be initialized by a constant expression}}
+ // expected-note at -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;
More information about the cfe-commits
mailing list