[PATCH] D13382: GlobalOpt should maintain externally_initialized when splitting aggregates

[Oliver Stannard via llvm-commits]

olista01 created this revision.
olista01 added a subscriber: llvm-commits.
olista01 set the repository for this revision to rL LLVM.

When GlobalOpt splits an internal, global variable with an aggregate type, it should propagate the externally_initialized flag to the newly created globals.

This makes the pass safe for our downstream use of this flag, while still allowing some useful optimisations (such as removing dead parts of the split aggregate) to be performed. However, I'm not familiar enough with the other uses of this flag (Objective-C selectors and the CUDA __device__ qualifier) to know if splitting an internal, externally_initialized global is valid for them. I'm guessing that it is safe for internal globals, but could anyone confirm this?

I've got a related patch up for review at D13343, the test in this patch depends on that change.

Repository:
  rL LLVM

http://reviews.llvm.org/D13382

Files:
  lib/Transforms/IPO/GlobalOpt.cpp
  test/Transforms/GlobalOpt/externally-initialized-aggregate.ll

Index: test/Transforms/GlobalOpt/externally-initialized-aggregate.ll
===================================================================
--- /dev/null
+++ test/Transforms/GlobalOpt/externally-initialized-aggregate.ll
@@ -0,0 +1,50 @@
+; RUN: opt < %s -S -globalopt | FileCheck %s
+
+; This global is externally_initialized, so if we split it into scalars we
+; should keep that flag set on all of the new globals. This will prevent the
+; store to @a[0] from being constant propagated to the load in @foo, but will not
+; prevent @a[1] from being removed since it is dead.
+; CHECK: @a.0 = internal unnamed_addr externally_initialized global i32 undef
+; CHECK-NOT @a.1
+ at a = internal externally_initialized global [2 x i32] undef, align 4
+; This is the same, but a struct rather than an array.
+; CHECK: @b.0 = internal unnamed_addr externally_initialized global i32 undef
+; CHECK-NOT @b.1
+ at b = internal externally_initialized global {i32, i32} undef, align 4
+
+define i32 @foo() {
+; CHECK-LABEL: define i32 @foo
+entry:
+; This load uses the split global, but cannot be constant-propagated away.
+; CHECK: %0 = load i32, i32* @a.0
+  %0 = load i32, i32* getelementptr inbounds ([2 x i32], [2 x i32]* @a, i32 0, i32 0), align 4
+  ret i32 %0
+}
+
+define i32 @bar() {
+; CHECK-LABEL: define i32 @bar
+entry:
+; This load uses the split global, but cannot be constant-propagated away.
+; CHECK: %0 = load i32, i32* @b.0
+  %0 = load i32, i32* getelementptr inbounds ({i32, i32}, {i32, i32}* @b, i32 0, i32 0), align 4
+  ret i32 %0
+}
+
+define void @init() {
+; CHECK-LABEL: define void @init
+entry:
+; This store uses the split global, but cannot be constant-propagated away.
+; CHECK: store i32 1, i32* @a.0
+  store i32 1, i32* getelementptr inbounds ([2 x i32], [2 x i32]* @a, i32 0, i32 0), align 4
+; This store can be removed, because the second element of @a is never read.
+; CHECK-NOT: store i32 2, i32* @a.1
+  store i32 2, i32* getelementptr inbounds ([2 x i32], [2 x i32]* @a, i32 0, i32 1), align 4
+
+; This store uses the split global, but cannot be constant-propagated away.
+; CHECK: store i32 3, i32* @b.0
+  store i32 3, i32* getelementptr inbounds ({i32, i32}, {i32, i32}* @b, i32 0, i32 0), align 4
+; This store can be removed, because the second element of @b is never read.
+; CHECK-NOT: store i32 4, i32* @b.1
+  store i32 4, i32* getelementptr inbounds ({i32, i32}, {i32, i32}* @b, i32 0, i32 1), align 4
+  ret void
+}
Index: lib/Transforms/IPO/GlobalOpt.cpp
===================================================================
--- lib/Transforms/IPO/GlobalOpt.cpp
+++ lib/Transforms/IPO/GlobalOpt.cpp
@@ -497,6 +497,7 @@
                                                In, GV->getName()+"."+Twine(i),
                                                GV->getThreadLocalMode(),
                                               GV->getType()->getAddressSpace());
+      NGV->setExternallyInitialized(GV->isExternallyInitialized());
       Globals.insert(GV, NGV);
       NewGlobals.push_back(NGV);
 
@@ -530,6 +531,7 @@
                                                In, GV->getName()+"."+Twine(i),
                                                GV->getThreadLocalMode(),
                                               GV->getType()->getAddressSpace());
+      NGV->setExternallyInitialized(GV->isExternallyInitialized());
       Globals.insert(GV, NGV);
       NewGlobals.push_back(NGV);
 


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D13382.36345.patch
Type: text/x-patch
Size: 3422 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20151002/0468a13a/attachment.bin>