diff options
Diffstat (limited to 'clang/lib/CodeGen/CodeGenModule.cpp')
-rw-r--r-- | clang/lib/CodeGen/CodeGenModule.cpp | 17 |
1 files changed, 6 insertions, 11 deletions
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<CUDASharedAttr>()) { - if (InitExpr) { - const auto *C = dyn_cast<CXXConstructExpr>(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<CUDASharedAttr>()) 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 }. // |