[PATCH] D105135: [Internalize] Preserve variables externally initialized.

[Michael Liao via Phabricator via cfe-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rG4e5d9c88033f: [Internalize] Preserve variables externally initialized. (authored by hliao).

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D105135/new/

https://reviews.llvm.org/D105135

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


Index: llvm/test/Transforms/Internalize/externally-initialized.ll
===================================================================
--- /dev/null
+++ 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
Index: llvm/lib/Transforms/IPO/Internalize.cpp
===================================================================
--- llvm/lib/Transforms/IPO/Internalize.cpp
+++ 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;
Index: clang/test/CodeGenCUDA/unused-global-var.cu
===================================================================
--- clang/test/CodeGenCUDA/unused-global-var.cu
+++ 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;
 
Index: clang/test/CodeGenCUDA/host-used-device-var.cu
===================================================================
--- clang/test/CodeGenCUDA/host-used-device-var.cu
+++ 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;
 


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D105135.357224.patch
Type: text/x-patch
Size: 2327 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20210708/4423d50b/attachment.bin>