r329229 - Revert "[CUDA] Check initializers of instantiated template variables."
Artem Belevich via cfe-commits
cfe-commits at lists.llvm.org
Wed Apr 4 13:48:42 PDT 2018
Author: tra
Date: Wed Apr 4 13:48:42 2018
New Revision: 329229
URL: http://llvm.org/viewvc/llvm-project?rev=329229&view=rev
Log:
Revert "[CUDA] Check initializers of instantiated template variables."
This (temporarily) reverts commit r329127 due to the problems
it exposed in TensorFlow.
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=329229&r1=329228&r2=329229&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Wed Apr 4 13:48:42 2018
@@ -10150,16 +10150,6 @@ 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=329229&r1=329228&r2=329229&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaCUDA.cpp (original)
+++ cfe/trunk/lib/Sema/SemaCUDA.cpp Wed Apr 4 13:48:42 2018
@@ -471,59 +471,6 @@ 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=329229&r1=329228&r2=329229&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDecl.cpp Wed Apr 4 13:48:42 2018
@@ -11629,8 +11629,58 @@ 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)
- checkAllowedCUDAInitializer(VD);
+ 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();
+ }
+ }
+ }
+ }
+ }
// 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=329229&r1=329228&r2=329229&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaTemplateInstantiateDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaTemplateInstantiateDecl.cpp Wed Apr 4 13:48:42 2018
@@ -4221,9 +4221,6 @@ 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=329229&r1=329228&r2=329229&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/device-var-init.cu (original)
+++ cfe/trunk/test/SemaCUDA/device-var-init.cu Wed Apr 4 13:48:42 2018
@@ -225,20 +225,3 @@ 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