This revision was landed with ongoing or failed builds.
This revision was automatically updated to reflect the committed changes.
yaxunl marked an inline comment as done.
Closed by commit rG73b22935a7a8: [CUDA][HIP] Do not promote constexpr var with
non-constant initializer (authored by yaxunl).
Herald added a project: clang.
Changed prior to commit:
https://reviews.llvm.org/D119615?vs=408145&id=409003#toc
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D119615/new/
https://reviews.llvm.org/D119615
Files:
clang/lib/Sema/SemaCUDA.cpp
clang/test/SemaCUDA/constexpr-var.cu
Index: clang/test/SemaCUDA/constexpr-var.cu
===================================================================
--- /dev/null
+++ clang/test/SemaCUDA/constexpr-var.cu
@@ -0,0 +1,105 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
+// RUN: -fsyntax-only -verify
+// RUN: %clang_cc1 -triple x86_64 -x hip %s \
+// RUN: -fsyntax-only -verify=host
+
+// host-no-diagnostics
+
+#include "Inputs/cuda.h"
+
+// Test constexpr var initialized with address of a const var.
+// Both are promoted to device side.
+
+namespace Test1 {
+const int a = 1;
+
+struct B {
+ static constexpr const int *p = &a;
+ __device__ static constexpr const int *const p2 = &a;
+};
+
+// Const variable 'a' is treated as __constant__ on device side,
+// therefore its address can be used as initializer for another
+// device variable.
+
+__device__ void f() {
+ int y = a;
+ constexpr const int *x = B::p;
+ constexpr const int *z = B::p2;
+}
+}
+
+// Test constexpr var initialized with address of a non-cost var.
+// Neither is promoted to device side.
+
+namespace Test2 {
+int a = 1;
+// expected-note@-1{{host variable declared here}}
+
+struct B {
+ static constexpr int *const p = &a;
+ // expected-note@-1{{const variable cannot be emitted on device side due to dynamic initialization}}
+};
+
+__device__ void f() {
+ int y = a;
+ // expected-error@-1{{reference to __host__ variable 'a' in __device__ function}}
+ const int *const *x = &B::p;
+ // expected-error@-1{{reference to __host__ variable 'p' in __device__ function}}
+ // ToDo: use of non-promotable constexpr variable in device compilation should be treated as
+ // ODR-use and diagnosed.
+ const int *const z = B::p;
+}
+}
+
+// Test constexpr device var initialized with address of a non-const host var, __shared var,
+// __managed__ var, __device__ var, __constant__ var, texture var, surface var.
+
+namespace Test3 {
+struct textureReference {
+ int desc;
+};
+
+enum ReadMode {
+ ElementType = 0,
+ NormalizedFloat = 1
+};
+
+template <typename T, int dim = 1, enum ReadMode mode = ElementType>
+struct __attribute__((device_builtin_texture_type)) texture : public textureReference {
+};
+
+struct surfaceReference {
+ int desc;
+};
+
+template <typename T, int dim = 1>
+struct __attribute__((device_builtin_surface_type)) surface : public surfaceReference {
+};
+
+// Partial specialization over `void`.
+template<int dim>
+struct __attribute__((device_builtin_surface_type)) surface<void, dim> : public surfaceReference {
+};
+
+texture<float, 2, ElementType> tex;
+surface<void, 2> surf;
+
+int a = 1;
+__shared__ int b;
+__managed__ int c = 1;
+__device__ int d = 1;
+__constant__ int e = 1;
+struct B {
+ __device__ static constexpr int *const p1 = &a;
+ // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+ __device__ static constexpr int *const p2 = &b;
+ // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+ __device__ static constexpr int *const p3 = &c;
+ // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+ __device__ static constexpr int *const p4 = &d;
+ __device__ static constexpr int *const p5 = &e;
+ __device__ static constexpr texture<float, 2, ElementType> *const p6 = &tex;
+ __device__ static constexpr surface<void, 2> *const p7 = &surf;
+};
+}
Index: clang/lib/Sema/SemaCUDA.cpp
===================================================================
--- clang/lib/Sema/SemaCUDA.cpp
+++ clang/lib/Sema/SemaCUDA.cpp
@@ -145,9 +145,11 @@
Sema::CUDAVariableTarget Sema::IdentifyCUDATarget(const VarDecl *Var) {
if (Var->hasAttr<HIPManagedAttr>())
return CVT_Unified;
- if (Var->isConstexpr() && !hasExplicitAttr<CUDAConstantAttr>(Var))
- return CVT_Both;
- if (Var->getType().isConstQualified() && Var->hasAttr<CUDAConstantAttr>() &&
+ // Only constexpr and const variabless with implicit constant attribute
+ // are emitted on both sides. Such variables are promoted to device side
+ // only if they have static constant intializers on device side.
+ if ((Var->isConstexpr() || Var->getType().isConstQualified()) &&
+ Var->hasAttr<CUDAConstantAttr>() &&
!hasExplicitAttr<CUDAConstantAttr>(Var))
return CVT_Both;
if (Var->hasAttr<CUDADeviceAttr>() || Var->hasAttr<CUDAConstantAttr>() ||
@@ -718,9 +720,9 @@
!VD->hasAttr<CUDAConstantAttr>() && !VD->hasAttr<CUDASharedAttr>() &&
(VD->isFileVarDecl() || VD->isStaticDataMember()) &&
!IsDependentVar(VD) &&
- (VD->isConstexpr() || (VD->getType().isConstQualified() &&
- HasAllowedCUDADeviceStaticInitializer(
- *this, VD, CICK_DeviceOrConstant)))) {
+ ((VD->isConstexpr() || VD->getType().isConstQualified()) &&
+ HasAllowedCUDADeviceStaticInitializer(*this, VD,
+ CICK_DeviceOrConstant))) {
VD->addAttr(CUDAConstantAttr::CreateImplicit(getASTContext()));
}
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
