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

[Michael Liao via Phabricator via cfe-commits]

hliao created this revision.
hliao added reviewers: yaxunl, bogner.
Herald added subscribers: ormris, hiraditya.
hliao requested review of this revision.
Herald added projects: clang, LLVM.
Herald added subscribers: llvm-commits, cfe-commits.

- ``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.


Repository:
  rG LLVM Github Monorepo

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 suggestes, 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
@@ -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;
 
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
@@ -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;
 


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D105135.355291.patch
Type: text/x-patch
Size: 1982 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20210629/d5110b54/attachment.bin>