[clang] 4e5d9c8 - [Internalize] Preserve variables externally initialized.

[Michael Liao via cfe-commits]

Author: Michael Liao
Date: 2021-07-08T10:48:19-04:00
New Revision: 4e5d9c88033f1fc5d5206a02d8303bc6de43cf2b

URL: https://github.com/llvm/llvm-project/commit/4e5d9c88033f1fc5d5206a02d8303bc6de43cf2b
DIFF: https://github.com/llvm/llvm-project/commit/4e5d9c88033f1fc5d5206a02d8303bc6de43cf2b.diff

LOG: [Internalize] Preserve variables externally initialized.

- ``externally_initialized`` variables would be initialized or modified
  elsewhere. Particularly, CUDA or HIP may have host code to initialize
  or modify ``externally_initialized`` device variables, which may not
  be explicitly referenced on the device side but may still be used
  through the host side interfaces. Not preserving them triggers the
  elimination of them in the GlobalDCE and breaks the user code.

Reviewed By: yaxunl

Differential Revision: https://reviews.llvm.org/D105135

Added: 
    llvm/test/Transforms/Internalize/externally-initialized.ll

Modified: 
    clang/test/CodeGenCUDA/host-used-device-var.cu
    clang/test/CodeGenCUDA/unused-global-var.cu
    llvm/lib/Transforms/IPO/Internalize.cpp

Removed: 
    


################################################################################
diff  --git a/clang/test/CodeGenCUDA/host-used-device-var.cu b/clang/test/CodeGenCUDA/host-used-device-var.cu
index b94ef689b3162..6bb5757052946 100644
--- 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
index 1dbb3a22563c8..c091e83eda70a 100644
--- 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
index 008712c87988b..cf8da0baebe41 100644
--- a/llvm/lib/Transforms/IPO/Internalize.cpp
+++ b/llvm/lib/Transforms/IPO/Internalize.cpp
@@ -101,6 +101,12 @@ bool InternalizePass::shouldPreserveGV(const GlobalValue &GV) {
   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
index 0000000000000..4c24e53543db9
--- /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
+ at G0 = protected externally_initialized global i32 0, align 4