diff options
| author | Artem Belevich <tra@google.com> | 2018-12-13 21:43:04 +0000 |
|---|---|---|
| committer | Artem Belevich <tra@google.com> | 2018-12-13 21:43:04 +0000 |
| commit | 7b05666a19c1310d2d3b7e05d08313695b4eaa92 (patch) | |
| tree | 07a6e0635f2578e826eb792e86d079f6dcc58e7c /clang/lib/CodeGen/CodeGenModule.cpp | |
| parent | 82b6f98f9a1240ca8a4907a738f0ec6d0c5d9fbb (diff) | |
| download | bcm5719-llvm-7b05666a19c1310d2d3b7e05d08313695b4eaa92.tar.gz bcm5719-llvm-7b05666a19c1310d2d3b7e05d08313695b4eaa92.zip | |
[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
Diffstat (limited to 'clang/lib/CodeGen/CodeGenModule.cpp')
| -rw-r--r-- | clang/lib/CodeGen/CodeGenModule.cpp | 11 |
1 files changed, 9 insertions, 2 deletions
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<CUDASharedAttr>()) + bool IsCUDASharedVar = + getLangOpts().CUDAIsDevice && D->hasAttr<CUDASharedAttr>(); + // Shadows of initialized device-side global variables are also left + // undefined. + bool IsCUDAShadowVar = + !getLangOpts().CUDAIsDevice && + (D->hasAttr<CUDAConstantAttr>() || D->hasAttr<CUDADeviceAttr>() || + D->hasAttr<CUDASharedAttr>()); + if (getLangOpts().CUDA && (IsCUDASharedVar || IsCUDAShadowVar)) Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy)); else if (!InitExpr) { // This is a tentative definition; tentative definitions are |
