[PATCH] D105135: [Internalize] Preserve variables externally initialized.
Michael Liao via Phabricator via cfe-commits
cfe-commits at lists.llvm.org
Wed Jun 30 20:52:24 PDT 2021
hliao updated this revision to Diff 355764.
hliao added a comment.
Fix typo. Kindly PING for review.
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
@@ -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.355764.patch
Type: text/x-patch
Size: 1981 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20210701/98c39c75/attachment.bin>
More information about the cfe-commits
mailing list