From 7b05666a19c1310d2d3b7e05d08313695b4eaa92 Mon Sep 17 00:00:00 2001 From: Artem Belevich Date: Thu, 13 Dec 2018 21:43:04 +0000 Subject: [CUDA] Make all host-side shadows of device-side variables undef. The host-side code can't (and should not) access the values that may only exist on the device side. E.g. address of a __device__ function does not exist on the host side as we don't generate the code for it there. Differential Revision: https://reviews.llvm.org/D55663 llvm-svn: 349087 --- clang/lib/CodeGen/CodeGenModule.cpp | 11 +++++++++-- 1 file changed, 9 insertions(+), 2 deletions(-) (limited to 'clang/lib/CodeGen/CodeGenModule.cpp') diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index ce4558fac4b..1b7bb2ce37e 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -3485,8 +3485,15 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D, // 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()) + bool IsCUDASharedVar = + getLangOpts().CUDAIsDevice && D->hasAttr(); + // Shadows of initialized device-side global variables are also left + // undefined. + bool IsCUDAShadowVar = + !getLangOpts().CUDAIsDevice && + (D->hasAttr() || D->hasAttr() || + D->hasAttr()); + if (getLangOpts().CUDA && (IsCUDASharedVar || IsCUDAShadowVar)) Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy)); else if (!InitExpr) { // This is a tentative definition; tentative definitions are -- cgit v1.2.3