From 4b7a713accdce4e46a45f34bd8705e43b495ddec Mon Sep 17 00:00:00 2001
From: Michael Liao <michael.hliao@gmail.com>
Date: Wed, 29 May 2019 17:23:27 +0000
Subject: [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
---
 clang/lib/CodeGen/CodeGenModule.cpp | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

(limited to 'clang/lib/CodeGen/CodeGenModule.cpp')

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
-- 
cgit v1.2.3