diff --git a/clang/test/CodeGenCUDA/host-used-device-var.cu b/clang/test/CodeGenCUDA/host-used-device-var.cu
--- a/clang/test/CodeGenCUDA/host-used-device-var.cu
+++ b/clang/test/CodeGenCUDA/host-used-device-var.cu
@@ -17,12 +17,6 @@
 
 // Check device variables used by neither host nor device functioins are not kept.
 
-// DEV-NEG-NOT: @v1
-__device__ int v1;
-
-// DEV-NEG-NOT: @v2
-__constant__ int v2;
-
 // DEV-NEG-NOT: @_ZL2v3
 static __device__ int v3;
 
diff --git a/clang/test/CodeGenCUDA/unused-global-var.cu b/clang/test/CodeGenCUDA/unused-global-var.cu
--- a/clang/test/CodeGenCUDA/unused-global-var.cu
+++ b/clang/test/CodeGenCUDA/unused-global-var.cu
@@ -17,12 +17,6 @@
 
 // Check unused device/constant variables are eliminated.
 
-// NEGCHK-NOT: @v1
-__device__ int v1;
-
-// NEGCHK-NOT: @v2
-__constant__ int v2;
-
 // NEGCHK-NOT: @_ZL2v3
 constexpr int v3 = 1;
 
diff --git a/llvm/lib/Transforms/IPO/Internalize.cpp b/llvm/lib/Transforms/IPO/Internalize.cpp
--- a/llvm/lib/Transforms/IPO/Internalize.cpp
+++ b/llvm/lib/Transforms/IPO/Internalize.cpp
@@ -101,6 +101,12 @@
   if (GV.hasDLLExportStorageClass())
     return true;
 
+  // As the name suggests, externally initialized variables need preserving as
+  // they would be initialized elsewhere externally.
+  if (const auto *G = dyn_cast<GlobalVariable>(&GV))
+    if (G->isExternallyInitialized())
+      return true;
+
   // Already local, has nothing to do.
   if (GV.hasLocalLinkage())
     return false;
diff --git a/llvm/test/Transforms/Internalize/externally-initialized.ll b/llvm/test/Transforms/Internalize/externally-initialized.ll
new file mode 100644
--- /dev/null
+++ b/llvm/test/Transforms/Internalize/externally-initialized.ll
@@ -0,0 +1,7 @@
+; RUN: opt < %s -internalize -S | FileCheck %s
+; RUN: opt < %s -passes=internalize -S | FileCheck %s
+
+; CHECK: @G0
+; CHECK-NOT: internal
+; CHECK-SAME: global i32
+@G0 = protected externally_initialized global i32 0, align 4