diff options
| author | Artem Belevich <tra@google.com> | 2016-05-19 20:13:53 +0000 |
|---|---|---|
| committer | Artem Belevich <tra@google.com> | 2016-05-19 20:13:53 +0000 |
| commit | 3650bbeebc07f3d4cfd026bbb5088432bb1766e6 (patch) | |
| tree | a8eb02e2565dec791c0ef8e33dfe5a35ab2c874e /clang/lib/Sema/SemaDecl.cpp | |
| parent | 85b6f63f425ba7da6f644625f2c01f207cb93279 (diff) | |
| download | bcm5719-llvm-3650bbeebc07f3d4cfd026bbb5088432bb1766e6.tar.gz bcm5719-llvm-3650bbeebc07f3d4cfd026bbb5088432bb1766e6.zip | |
[CUDA] Do not allow non-empty destructors for global device-side variables.
According to Cuda Programming guide (v7.5, E2.3.1):
> __device__, __constant__ and __shared__ variables defined in namespace
> scope, that are of class type, cannot have a non-empty constructor or a
> non-empty destructor.
Clang already deals with device-side constructors (see D15305).
This patch enforces similar rules for destructors.
Differential Revision: http://reviews.llvm.org/D20140
llvm-svn: 270108
Diffstat (limited to 'clang/lib/Sema/SemaDecl.cpp')
| -rw-r--r-- | clang/lib/Sema/SemaDecl.cpp | 6 |
1 files changed, 6 insertions, 0 deletions
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index d34d9b5688c..9c442a61ce2 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -10442,6 +10442,12 @@ Sema::FinalizeDeclaration(Decl *ThisDecl) { 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 |
