diff options
author | Artem Belevich <tra@google.com> | 2016-02-02 22:29:48 +0000 |
---|---|---|
committer | Artem Belevich <tra@google.com> | 2016-02-02 22:29:48 +0000 |
commit | 44f73317fc0d168da26f72b16bd75106df113515 (patch) | |
tree | b36208985914e3c551575977af8799f54fafba6b /lib/CodeGen/CGDeclCXX.cpp | |
parent | 2fcd11968f1040e9d89dc32fed1148f667f75890 (diff) |
[CUDA] Do not allow dynamic initialization of global device side variables.
In general CUDA does not allow dynamic initialization of
global device-side variables. One exception is that CUDA allows
records with empty constructors as described in section E2.2.1 of
CUDA 7.5 Programming guide.
This patch applies initializer checks for all device-side variables.
Empty constructors are accepted, but no code is generated for them.
Differential Revision: http://reviews.llvm.org/D15305
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@259592 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'lib/CodeGen/CGDeclCXX.cpp')
-rw-r--r-- | lib/CodeGen/CGDeclCXX.cpp | 11 |
1 files changed, 11 insertions, 0 deletions
diff --git a/lib/CodeGen/CGDeclCXX.cpp b/lib/CodeGen/CGDeclCXX.cpp index adba731687..98d904b156 100644 --- a/lib/CodeGen/CGDeclCXX.cpp +++ b/lib/CodeGen/CGDeclCXX.cpp @@ -304,6 +304,17 @@ void CodeGenModule::EmitCXXGlobalVarDeclInitFunc(const VarDecl *D, llvm::GlobalVariable *Addr, bool PerformInit) { + + // According to E.2.3.1 in 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. All + // the checks have been done in Sema by now. Whatever initializers + // are allowed are empty and we just need to ignore them here. + if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice && + (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() || + D->hasAttr<CUDASharedAttr>())) + return; + // Check if we've already initialized this decl. auto I = DelayedCXXInitPosition.find(D); if (I != DelayedCXXInitPosition.end() && I->second == ~0U) |