diff options
| author | Michael Liao <michael.hliao@gmail.com> | 2019-05-29 17:23:27 +0000 |
|---|---|---|
| committer | Michael Liao <michael.hliao@gmail.com> | 2019-05-29 17:23:27 +0000 |
| commit | 4b7a713accdce4e46a45f34bd8705e43b495ddec (patch) | |
| tree | 24b70efb61c7c31df61306d94d90ad09883a6947 /clang/lib/CodeGen/CodeGenModule.cpp | |
| parent | c98b288b030cde85ae8d4e188677546efcbe0d29 (diff) | |
| download | bcm5719-llvm-4b7a713accdce4e46a45f34bd8705e43b495ddec.tar.gz bcm5719-llvm-4b7a713accdce4e46a45f34bd8705e43b495ddec.zip | |
[CUDA][HIP] Skip setting `externally_initialized` for static device variables.
Summary:
- By declaring device variables as `static`, we assume they won't be
addressable from the host side. Thus, no `externally_initialized` is
required.
Reviewers: yaxunl
Subscribers: cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D62603
llvm-svn: 361994
Diffstat (limited to 'clang/lib/CodeGen/CodeGenModule.cpp')
| -rw-r--r-- | clang/lib/CodeGen/CodeGenModule.cpp | 3 |
1 files changed, 2 insertions, 1 deletions
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 8c9e240a680..8cfb4e60e0d 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -3869,7 +3869,8 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D, // / cudaMemcpyToSymbol() / cudaMemcpyFromSymbol())." if (GV && LangOpts.CUDA) { if (LangOpts.CUDAIsDevice) { - if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()) + if (Linkage != llvm::GlobalValue::InternalLinkage && + (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>())) GV->setExternallyInitialized(true); } else { // Host-side shadows of external declarations of device-side |
