[PATCH] D15305: [CUDA] Do not allow dynamic initialization of global device side variables.

[Artem Belevich via cfe-commits]

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

In general CUDA does not allow dynamic initialization of
global device-side variables except for records with empty constructors as described in section [[ http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#device-memory-qualifiers | E.2.3.1 of
CUDA 7.5 Programming guide ]]:

> __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 constructor for a class type is considered empty at a point in the translation unit, 
> if it is either a trivial constructor or it satisfies all of the following conditions:

> * The constructor function has been defined.
> * The constructor function has no parameters, the initializer list is empty and the function body is an empty compound statement.
> * Its class has no virtual functions and no virtual base classes.
> * The default constructors 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 default constructors can be considered empty.

Clang is already enforcing no-initializers for __shared__ variables, but currently allows dynamic initialization for __device__ and __constant__ variables. 

This patch applies initializer checks for all device-side variables.
Empty constructors are accepted, but no code is generated for them.

http://reviews.llvm.org/D15305

Files:
  lib/CodeGen/CGDeclCXX.cpp
  lib/CodeGen/CodeGenModule.cpp
  lib/CodeGen/CodeGenModule.h
  test/CodeGenCUDA/device-var-init.cu

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D15305.42101.patch
Type: text/x-patch
Size: 20020 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20151207/99aa2a3c/attachment-0001.bin>