yaxunl created this revision.
yaxunl added a reviewer: tra.
Herald added a reviewer: aaron.ballman.
yaxunl requested review of this revision.
constexpr variables are implicit constant variables in device compilation.
Not all constexpr variables are valid to be emitted on device side, therefore
we should not force emit implicit constant variables even if they are
ODR-used by host code.
This fixes the regression caused by https://reviews.llvm.org/D102237
https://reviews.llvm.org/D102801
Files:
clang/lib/CodeGen/CodeGenModule.cpp
clang/lib/Sema/SemaCUDA.cpp
clang/lib/Sema/SemaDeclAttr.cpp
clang/test/AST/ast-dump-constant-var.cu
clang/test/CodeGenCUDA/host-used-device-var.cu
Index: clang/test/CodeGenCUDA/host-used-device-var.cu
===================================================================
--- clang/test/CodeGenCUDA/host-used-device-var.cu
+++ clang/test/CodeGenCUDA/host-used-device-var.cu
@@ -66,8 +66,22 @@
template <typename T>
__device__ func_t<T> p_add_func = add_func<T>;
+// Check non-constant constexpr variables ODR-used by host code only is not emitted.
+// DEV-NEG-NOT: constexpr_var1a
+// DEV-NEG-NOT: constexpr_var1b
+constexpr int constexpr_var1a = 1;
+inline constexpr int constexpr_var1b = 1;
+
+// Check constant constexpr variables ODR-used by host code only.
+// Non-inline constexpr variable has internal linkage, therefore it is not accessible by host and not kept.
+// Inline constexpr variable has linkonce_ord linkage, therefore it can be accessed by host and kept.
+// DEV-NEG-NOT: constexpr_var2a
+// DEV-DAG: @constexpr_var2b = linkonce_odr addrspace(4) externally_initialized constant i32 2
+__constant__ constexpr int constexpr_var2a = 2;
+inline __constant__ constexpr int constexpr_var2b = 2;
+
void use(func_t<int> p);
-void use(int *p);
+void use(const int *p);
void fun1() {
use(&u1);
@@ -76,20 +90,58 @@
use(&ext_var);
use(&inline_var);
use(p_add_func<int>);
+ use(&constexpr_var1a);
+ use(&constexpr_var1b);
+ use(&constexpr_var2a);
+ use(&constexpr_var2b);
}
__global__ void kern1(int **x) {
*x = &u4;
}
+// Check implicit constant variable ODR-used by host code is not emitted.
+// DEV-NEG-NOT: _ZN16TestConstexprVar1oE
+namespace TestConstexprVar {
+char o;
+class ou {
+public:
+ ou(char) { __builtin_strlen(&o); }
+};
+template < typename ao > struct aw { static constexpr ao c; };
+class x {
+protected:
+ typedef ou (*y)(const x *);
+ constexpr x(y ag) : ah(ag) {}
+ template < bool * > struct ak;
+ template < typename > struct al {
+ static bool am;
+ static ak< &am > an;
+ };
+ template < typename ao > static x ap() { (void)aw< ao >::c; return x(nullptr); }
+ y ah;
+};
+template < typename ao > bool x::al< ao >::am(&ap< ao >);
+class ar : x {
+public:
+ constexpr ar() : x(as) {}
+ static ou as(const x *) { return 0; }
+ al< ar > av;
+};
+}
+
// Check the exact list of variables to ensure @_ZL2u4 is not among them.
-// DEV: @llvm.compiler.used = {{[^@]*}} @_Z10p_add_funcIiE {{[^@]*}} @_ZL2u3 {{[^@]*}} @inline_var {{[^@]*}} @u1 {{[^@]*}} @u2 {{[^@]*}} @u5
+// DEV: @llvm.compiler.used = {{[^@]*}} @_Z10p_add_funcIiE {{[^@]*}} @_ZL2u3 {{[^@]*}} @constexpr_var2b {{[^@]*}} @inline_var {{[^@]*}} @u1 {{[^@]*}} @u2 {{[^@]*}} @u5
// HOST-DAG: hipRegisterVar{{.*}}@u1
// HOST-DAG: hipRegisterVar{{.*}}@u2
// HOST-DAG: hipRegisterVar{{.*}}@_ZL2u3
+// HOST-DAG: hipRegisterVar{{.*}}@constexpr_var2b
// HOST-DAG: hipRegisterVar{{.*}}@u5
// HOST-DAG: hipRegisterVar{{.*}}@inline_var
// HOST-DAG: hipRegisterVar{{.*}}@_Z10p_add_funcIiE
// HOST-NEG-NOT: hipRegisterVar{{.*}}@ext_var
// HOST-NEG-NOT: hipRegisterVar{{.*}}@_ZL2u4
+// HOST-NEG-NOT: hipRegisterVar{{.*}}@constexpr_var1a
+// HOST-NEG-NOT: hipRegisterVar{{.*}}@constexpr_var1b
+// HOST-NEG-NOT: hipRegisterVar{{.*}}@constexpr_var2a
Index: clang/test/AST/ast-dump-constant-var.cu
===================================================================
--- /dev/null
+++ clang/test/AST/ast-dump-constant-var.cu
@@ -0,0 +1,32 @@
+// RUN: %clang_cc1 -std=c++14 -ast-dump -x hip %s | FileCheck -check-prefixes=CHECK,HOST %s
+// RUN: %clang_cc1 -std=c++14 -ast-dump -fcuda-is-device -x hip %s | FileCheck -check-prefixes=CHECK,DEV %s
+
+#include "Inputs/cuda.h"
+
+// CHECK-LABEL: VarDecl {{.*}} m1 'int'
+// CHECK-NEXT: CUDAConstantAttr {{.*}}cuda.h
+__constant__ int m1;
+
+// CHECK-LABEL: VarDecl {{.*}} m2 'int'
+// CHECK-NEXT: CUDAConstantAttr {{.*}}cuda.h
+// CHECK-NOT: CUDAConstantAttr
+__constant__ __constant__ int m2;
+
+// CHECK-LABEL: VarDecl {{.*}} m3 'const int'
+// HOST-NOT: CUDAConstantAttr
+// DEV-NOT: CUDAConstantAttr {{.*}}cuda.h
+// DEV: CUDAConstantAttr {{.*}}Implicit
+// DEV-NOT: CUDAConstantAttr {{.*}}cuda.h
+constexpr int m3 = 1;
+
+// CHECK-LABEL: VarDecl {{.*}} m3a 'const int'
+// CHECK-NOT: CUDAConstantAttr {{.*}}Implicit
+// CHECK: CUDAConstantAttr {{.*}}cuda.h
+// CHECK-NOT: CUDAConstantAttr {{.*}}Implicit
+constexpr __constant__ int m3a = 2;
+
+// CHECK-LABEL: VarDecl {{.*}} m3b 'const int'
+// CHECK-NOT: CUDAConstantAttr {{.*}}Implicit
+// CHECK: CUDAConstantAttr {{.*}}cuda.h
+// CHECK-NOT: CUDAConstantAttr {{.*}}Implicit
+__constant__ constexpr int m3b = 3;
Index: clang/lib/Sema/SemaDeclAttr.cpp
===================================================================
--- clang/lib/Sema/SemaDeclAttr.cpp
+++ clang/lib/Sema/SemaDeclAttr.cpp
@@ -4410,6 +4410,13 @@
S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev);
return;
}
+ // constexpr variable may already get an implicit constant attr, which should
+ // be replaced by the explicit constant attr.
+ if (auto *A = D->getAttr<CUDAConstantAttr>()) {
+ if (!A->isImplicit())
+ return;
+ D->dropAttr<CUDAConstantAttr>();
+ }
D->addAttr(::new (S.Context) CUDAConstantAttr(S.Context, AL));
}
Index: clang/lib/Sema/SemaCUDA.cpp
===================================================================
--- clang/lib/Sema/SemaCUDA.cpp
+++ clang/lib/Sema/SemaCUDA.cpp
@@ -637,7 +637,8 @@
void Sema::MaybeAddCUDAConstantAttr(VarDecl *VD) {
if (getLangOpts().CUDAIsDevice && VD->isConstexpr() &&
- (VD->isFileVarDecl() || VD->isStaticDataMember())) {
+ (VD->isFileVarDecl() || VD->isStaticDataMember()) &&
+ !VD->hasAttr<CUDAConstantAttr>()) {
VD->addAttr(CUDAConstantAttr::CreateImplicit(getASTContext()));
}
}
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -2364,9 +2364,28 @@
}
// Emit CUDA/HIP static device variables referenced by host code only.
+ // Note we should not clear CUDADeviceVarODRUsedByHost since it is still
+ // needed for further handling.
if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice)
- for (const auto *V : getContext().CUDADeviceVarODRUsedByHost)
- DeferredDeclsToEmit.push_back(V);
+ for (const auto *V : getContext().CUDADeviceVarODRUsedByHost) {
+ // ToDo: The user of ODR-used variables may not be emitted on host side.
+ // There needs a more accurate way to determine whether a device side
+ // variable ODR-used by host only should be emitted on device side.
+ //
+ // Currently we do it conservatively. However this does not work well
+ // with implicit constant variables since use of implicit constant
+ // variable in host function cannot be easily differentiated from
+ // use of the host variable with the same name. Therefore we do not
+ // force emit implicit constant variable.
+ auto HasImplicitConstantAttr = [](const VarDecl *Var) {
+ auto *A = Var->getAttr<CUDAConstantAttr>();
+ if (!A)
+ return false;
+ return A->isImplicit();
+ };
+ if (!HasImplicitConstantAttr(V))
+ DeferredDeclsToEmit.push_back(V);
+ }
// Stop if we're out of both deferred vtables and deferred declarations.
if (DeferredDeclsToEmit.empty())
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
