From 97c01c35f8da48b47c397fd915a82bd1d881d4ab Mon Sep 17 00:00:00 2001 From: Artem Belevich Date: Tue, 2 Feb 2016 22:29:48 +0000 Subject: [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 llvm-svn: 259592 --- clang/lib/CodeGen/CodeGenModule.cpp | 17 ++++++----------- 1 file changed, 6 insertions(+), 11 deletions(-) (limited to 'clang/lib/CodeGen/CodeGenModule.cpp') diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index de580f2..4d7f627 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -2334,18 +2334,13 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D, const VarDecl *InitDecl; const Expr *InitExpr = D->getAnyInitializer(InitDecl); - // CUDA E.2.4.1 "__shared__ variables cannot have an initialization as part - // of their declaration." - if (getLangOpts().CPlusPlus && getLangOpts().CUDAIsDevice - && D->hasAttr()) { - if (InitExpr) { - const auto *C = dyn_cast(InitExpr); - if (C == nullptr || !C->getConstructor()->hasTrivialBody()) - Error(D->getLocation(), - "__shared__ variable cannot have an initialization."); - } + // CUDA E.2.4.1 "__shared__ variables cannot have an initialization + // as part of their declaration." Sema has already checked for + // error cases, so we just need to set Init to UndefValue. + if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice && + D->hasAttr()) Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy)); - } else if (!InitExpr) { + else if (!InitExpr) { // This is a tentative definition; tentative definitions are // implicitly initialized with { 0 }. // -- cgit v1.1