[PATCH] D20140: [CUDA] Do not allow non-empty destructors for global device-side variables.

Artem Belevich via cfe-commits cfe-commits at lists.llvm.org
Tue May 10 15:38:49 PDT 2016


tra created this revision.
tra added reviewers: jlebar, rsmith, jingyue.
tra added a subscriber: cfe-commits.

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.
> 
> 
> A destructor for a class is considered empty at a point in the translation unit, if it is either a trivial destructor or it satisfies all of the following conditions:
> 
> * The destructor function has been defined.
> * The destructor function body is an empty compound statement.
> * Its class has no virtual functions and no virtual base classes.
> * The destructors of all base classes of its class can be considered empty.
>  * For all the nonstatic data members of its class that are of class type (or array thereof), the destructor can be considered empty.
> 

Clang already deals with device-side constructors (see D15305).
This patch enforces similar rules for destructors.

http://reviews.llvm.org/D20140

Files:
  include/clang/Sema/Sema.h
  lib/Sema/SemaCUDA.cpp
  lib/Sema/SemaDecl.cpp
  test/CodeGenCUDA/device-var-init.cu
  test/SemaCUDA/device-var-init.cu

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D20140.56829.patch
Type: text/x-patch
Size: 13073 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20160510/52589ea5/attachment.bin>


More information about the cfe-commits mailing list