From ced275ae4a9488eef8f65dc85d25fa749394e823 Mon Sep 17 00:00:00 2001
From: Michael Liao <michael.hliao@gmail.com>
Date: Wed, 29 May 2019 17:23:27 +0000
Subject: [PATCH] [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

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@361994 91177308-0d34-0410-b5e6-96231b3b80d8
---
 lib/CodeGen/CodeGenModule.cpp       |  3 ++-
 test/CodeGenCUDA/device-var-init.cu | 10 ++++++++++
 2 files changed, 12 insertions(+), 1 deletion(-)

diff --git a/lib/CodeGen/CodeGenModule.cpp b/lib/CodeGen/CodeGenModule.cpp
index 8c9e240a68..8cfb4e60e0 100644
--- a/lib/CodeGen/CodeGenModule.cpp
+++ b/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
diff --git a/test/CodeGenCUDA/device-var-init.cu b/test/CodeGenCUDA/device-var-init.cu
index af42e698cf..fd236bb842 100644
--- a/test/CodeGenCUDA/device-var-init.cu
+++ b/test/CodeGenCUDA/device-var-init.cu
@@ -33,6 +33,16 @@ __device__ int d_v_i = 1;
 // DEVICE: @d_v_i = addrspace(1) externally_initialized global i32 1,
 // HOST:   @d_v_i = internal global i32 undef,
 
+// For `static` device variables, assume they won't be addressed from the host
+// side.
+static __device__ int d_s_v_i = 1;
+// DEVICE: @_ZL7d_s_v_i = internal addrspace(1) global i32 1,
+
+// Dummy function to keep static variables referenced.
+__device__ int foo() {
+  return d_s_v_i;
+}
+
 // trivial constructor -- allowed
 __device__ T d_t;
 // DEVICE: @d_t = addrspace(1) externally_initialized global %struct.T zeroinitializer
-- 
2.50.1