r329127 - [CUDA] Check initializers of instantiated template variables.
Artem Belevich via cfe-commits
cfe-commits at lists.llvm.org
Tue Apr 3 15:41:06 PDT 2018
Author: tra
Date: Tue Apr 3 15:41:06 2018
New Revision: 329127
URL: http://llvm.org/viewvc/llvm-project?rev=329127&view=rev
Log:
[CUDA] Check initializers of instantiated template variables.
We were already performing checks on non-template variables,
but the checks on templated ones were missing.
Differential Revision: https://reviews.llvm.org/D45231
Modified:
cfe/trunk/include/clang/Sema/Sema.h
cfe/trunk/lib/Sema/SemaCUDA.cpp
cfe/trunk/lib/Sema/SemaDecl.cpp
cfe/trunk/lib/Sema/SemaTemplateInstantiateDecl.cpp
cfe/trunk/test/SemaCUDA/device-var-init.cu
Modified: cfe/trunk/include/clang/Sema/Sema.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=329127&r1=329126&r2=329127&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Tue Apr 3 15:41:06 2018
@@ -10150,6 +10150,16 @@ public:
bool isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD);
bool isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *CD);
+ // \brief Checks that initializers of \p Var satisfy CUDA restrictions. In
+ // case of error emits appropriate diagnostic and invalidates \p Var.
+ //
+ // \details CUDA allows only empty constructors as initializers for global
+ // variables (see E.2.3.1, CUDA 7.5). The same restriction also applies to all
+ // __shared__ variables whether they are local or not (they all are implicitly
+ // static in CUDA). One exception is that CUDA allows constant initializers
+ // for __constant__ and __device__ variables.
+ void checkAllowedCUDAInitializer(VarDecl *Var);
+
/// Check whether NewFD is a valid overload for CUDA. Emits
/// diagnostics and invalidates NewFD if not.
void checkCUDATargetOverload(FunctionDecl *NewFD,
Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=329127&r1=329126&r2=329127&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaCUDA.cpp (original)
+++ cfe/trunk/lib/Sema/SemaCUDA.cpp Tue Apr 3 15:41:06 2018
@@ -471,6 +471,59 @@ bool Sema::isEmptyCudaDestructor(SourceL
return true;
}
+void Sema::checkAllowedCUDAInitializer(VarDecl *VD) {
+ if (VD->isInvalidDecl() || !VD->hasInit() || !VD->hasGlobalStorage())
+ return;
+ const Expr *Init = VD->getInit();
+ if (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>() ||
+ VD->hasAttr<CUDASharedAttr>()) {
+ assert(!VD->isStaticLocal() || VD->hasAttr<CUDASharedAttr>());
+ bool AllowedInit = false;
+ if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init))
+ AllowedInit =
+ isEmptyCudaConstructor(VD->getLocation(), CE->getConstructor());
+ // We'll allow constant initializers even if it's a non-empty
+ // 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());
+
+ // Also make sure that destructor, if there is one, is empty.
+ if (AllowedInit)
+ if (CXXRecordDecl *RD = VD->getType()->getAsCXXRecordDecl())
+ AllowedInit =
+ isEmptyCudaDestructor(VD->getLocation(), RD->getDestructor());
+
+ if (!AllowedInit) {
+ Diag(VD->getLocation(), VD->hasAttr<CUDASharedAttr>()
+ ? diag::err_shared_var_init
+ : diag::err_dynamic_var_init)
+ << Init->getSourceRange();
+ VD->setInvalidDecl();
+ }
+ } else {
+ // This is a host-side global variable. Check that the initializer is
+ // callable from the host side.
+ const FunctionDecl *InitFn = nullptr;
+ if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init)) {
+ InitFn = CE->getConstructor();
+ } else if (const CallExpr *CE = dyn_cast<CallExpr>(Init)) {
+ InitFn = CE->getDirectCallee();
+ }
+ if (InitFn) {
+ CUDAFunctionTarget InitFnTarget = IdentifyCUDATarget(InitFn);
+ if (InitFnTarget != CFT_Host && InitFnTarget != CFT_HostDevice) {
+ Diag(VD->getLocation(), diag::err_ref_bad_target_global_initializer)
+ << InitFnTarget << InitFn;
+ Diag(InitFn->getLocation(), diag::note_previous_decl) << InitFn;
+ VD->setInvalidDecl();
+ }
+ }
+ }
+}
+
// With -fcuda-host-device-constexpr, an unattributed constexpr function is
// treated as implicitly __host__ __device__, unless:
// * it is a variadic function (device-side variadic functions are not
Modified: cfe/trunk/lib/Sema/SemaDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=329127&r1=329126&r2=329127&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDecl.cpp Tue Apr 3 15:41:06 2018
@@ -11629,58 +11629,8 @@ void Sema::FinalizeDeclaration(Decl *Thi
// 7.5). We must also apply the same checks to all __shared__
// variables whether they are local or not. CUDA also allows
// constant initializers for __constant__ and __device__ variables.
- if (getLangOpts().CUDA) {
- const Expr *Init = VD->getInit();
- if (Init && VD->hasGlobalStorage()) {
- if (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>() ||
- VD->hasAttr<CUDASharedAttr>()) {
- assert(!VD->isStaticLocal() || VD->hasAttr<CUDASharedAttr>());
- bool AllowedInit = false;
- if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init))
- AllowedInit =
- isEmptyCudaConstructor(VD->getLocation(), CE->getConstructor());
- // We'll allow constant initializers even if it's a non-empty
- // 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());
-
- // Also make sure that destructor, if there is one, is empty.
- if (AllowedInit)
- if (CXXRecordDecl *RD = VD->getType()->getAsCXXRecordDecl())
- AllowedInit =
- isEmptyCudaDestructor(VD->getLocation(), RD->getDestructor());
-
- if (!AllowedInit) {
- Diag(VD->getLocation(), VD->hasAttr<CUDASharedAttr>()
- ? diag::err_shared_var_init
- : diag::err_dynamic_var_init)
- << Init->getSourceRange();
- VD->setInvalidDecl();
- }
- } else {
- // This is a host-side global variable. Check that the initializer is
- // callable from the host side.
- const FunctionDecl *InitFn = nullptr;
- if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init)) {
- InitFn = CE->getConstructor();
- } else if (const CallExpr *CE = dyn_cast<CallExpr>(Init)) {
- InitFn = CE->getDirectCallee();
- }
- if (InitFn) {
- CUDAFunctionTarget InitFnTarget = IdentifyCUDATarget(InitFn);
- if (InitFnTarget != CFT_Host && InitFnTarget != CFT_HostDevice) {
- Diag(VD->getLocation(), diag::err_ref_bad_target_global_initializer)
- << InitFnTarget << InitFn;
- Diag(InitFn->getLocation(), diag::note_previous_decl) << InitFn;
- VD->setInvalidDecl();
- }
- }
- }
- }
- }
+ if (getLangOpts().CUDA)
+ checkAllowedCUDAInitializer(VD);
// Grab the dllimport or dllexport attribute off of the VarDecl.
const InheritableAttr *DLLAttr = getDLLAttr(VD);
Modified: cfe/trunk/lib/Sema/SemaTemplateInstantiateDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaTemplateInstantiateDecl.cpp?rev=329127&r1=329126&r2=329127&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaTemplateInstantiateDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaTemplateInstantiateDecl.cpp Tue Apr 3 15:41:06 2018
@@ -4221,6 +4221,9 @@ void Sema::InstantiateVariableInitialize
ActOnUninitializedDecl(Var);
}
+
+ if (getLangOpts().CUDA)
+ checkAllowedCUDAInitializer(Var);
}
/// \brief Instantiate the definition of the given variable from its
Modified: cfe/trunk/test/SemaCUDA/device-var-init.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/device-var-init.cu?rev=329127&r1=329126&r2=329127&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/device-var-init.cu (original)
+++ cfe/trunk/test/SemaCUDA/device-var-init.cu Tue Apr 3 15:41:06 2018
@@ -225,3 +225,20 @@ inline __host__ __device__ void hd_emitt
static int x = 42; // no error on device because this is never codegen'ed there.
}
void call_hd_emitted_host_only() { hd_emitted_host_only(); }
+
+// Verify that we also check field initializers in instantiated structs.
+struct NontrivialInitializer {
+ __host__ __device__ NontrivialInitializer() : x(43) {}
+ int x;
+};
+
+template <typename T>
+__global__ void bar() {
+ __shared__ T bad;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+}
+
+void instantiate() {
+ bar<NontrivialInitializer><<<1, 1>>>();
+// expected-note at -1 {{in instantiation of function template specialization 'bar<NontrivialInitializer>' requested here}}
+}
More information about the cfe-commits
mailing list