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
@@ -15,14 +15,14 @@
 
 #include "Inputs/cuda.h"
 
-// Check device variables used by neither host nor device functioins are not kept.
-
-// DEV-NEG-NOT: @v1
+// DEV-DAG: @v1
 __device__ int v1;
 
-// DEV-NEG-NOT: @v2
+// DEV-DAG: @v2
 __constant__ int v2;
 
+// Check device variables used by neither host nor device functioins are not kept.
+
 // 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
@@ -15,14 +15,14 @@
 // DCE before internalization. This test makes sure unused global variables
 // are eliminated.
 
-// Check unused device/constant variables are eliminated.
-
-// NEGCHK-NOT: @v1
+// CHECK-DAG: @v1
 __device__ int v1;
 
-// NEGCHK-NOT: @v2
+// CHECK-DAG: @v2
 __constant__ int v2;
 
+// Check unused device/constant variables are eliminated.
+
 // 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