Automated g4 rollback of changelist 179506342

PiperOrigin-RevId: 179552496

diff --git a/tensorflow/compiler/xla/service/buffer_liveness.cc b/tensorflow/compiler/xla/service/buffer_liveness.cc
index e7749252ce44f0daf7016f72d80401695eaaacb9..513bfa3b7f7b45696093d03c1dd8250c548d260a 100644 (file)
--- a/tensorflow/compiler/xla/service/buffer_liveness.cc
+++ b/tensorflow/compiler/xla/service/buffer_liveness.cc
@@ -102,8 +102,8 @@ bool BufferLiveness::live_range_strictly_before(const LogicalBuffer& a,
     return false;
   }
 
+  // Every user of 'a' must be a predecessor of 'b' or 'b' itself.
   for (const BufferAlias& alias : points_to_analysis_->GetBufferAliases(a)) {
-    // Every user of 'a' must be a predecessor of 'b' or 'b' itself.
     for (auto user : alias.instruction()->users()) {
       if (DoesNotUseOperandBuffer(alias.instruction(), alias.index(), user,
                                   points_to_analysis())) {
@@ -114,16 +114,6 @@ bool BufferLiveness::live_range_strictly_before(const LogicalBuffer& a,
         return false;
       }
     }
-
-    // If the root instruction aliases the buffer 'a', the live range of 'a' is
-    // until the end of the computation and can never be strictly before another
-    // buffer. This is needed to prevent the root instruction's buffers from
-    // being reused by later instructions even when the root is not the last
-    // instruction in the schedule.
-    if (alias.instruction()->parent()->root_instruction() ==
-        alias.instruction()) {
-      return false;
-    }
   }
 
   // If 'b' is a user of 'a' then the buffers interfere unless 'a.instruction'

diff --git a/tensorflow/compiler/xla/service/buffer_liveness_test.cc b/tensorflow/compiler/xla/service/buffer_liveness_test.cc
index f623aef67a4f98b447a9a15634a78deb60cfe6f1..13825fe05bb1b98045f1a3dac3d7272a2d1151fb 100644 (file)
--- a/tensorflow/compiler/xla/service/buffer_liveness_test.cc
+++ b/tensorflow/compiler/xla/service/buffer_liveness_test.cc
@@ -311,48 +311,6 @@ TEST_F(BufferLivenessTest, OverlappedBuffersSequentialOrder) {
   EXPECT_FALSE(InstructionsMayInterfere(*liveness, add, exp));
 }
 
-TEST_F(BufferLivenessTest, RootInstructionIsNotLastInSequentialOrder) {
-  // Tests that when the root instruction is not the last instruction in the
-  // schedule, the live range of its buffers interfere with the buffers of the
-  // later instructions.
-  //
-  // Two sets of independent instructions are executed in the computation.
-  // param --> add (root)
-  // recv --> recv-done --> send --> send-done
-  //
-  // Sequential order:
-  //  param, add (root), recv, recv-done, send, send-done
-  auto builder = HloComputation::Builder(TestName());
-  auto param =
-      builder.AddInstruction(HloInstruction::CreateParameter(0, vec_, "param"));
-  auto add = builder.AddInstruction(
-      HloInstruction::CreateBinary(vec_, HloOpcode::kAdd, param, param));
-  auto recv = builder.AddInstruction(
-      HloInstruction::CreateRecv(vec_, /*channel_id=*/0));
-  auto recv_done = builder.AddInstruction(HloInstruction::CreateRecvDone(recv));
-  auto send = builder.AddInstruction(
-      HloInstruction::CreateSend(recv_done, /*channel_id=*/1));
-  auto send_done = builder.AddInstruction(HloInstruction::CreateSendDone(send));
-
-  auto module = CreateNewModule();
-  auto computation = module->AddEntryComputation(builder.Build(add));
-
-  SequentialHloOrdering::HloModuleSequence module_sequence;
-  std::vector<const HloInstruction*> order = {param,     add,  recv,
-                                              recv_done, send, send_done};
-  module_sequence.emplace(computation, order);
-  auto liveness =
-      BufferLiveness::Run(module.get(), xla::MakeUnique<SequentialHloOrdering>(
-                                            module.get(), module_sequence))
-          .ConsumeValueOrDie();
-
-  EXPECT_FALSE(InstructionsMayInterfere(*liveness, param, add));
-  // Check the root instruction (add) buffer interferes with the recv buffer.
-  EXPECT_TRUE(
-      liveness->MayInterfere(GetBuffer(*liveness, add, /*index=*/{}),
-                             GetBuffer(*liveness, recv, /*index=*/{0})));
-}
-
 TEST_F(BufferLivenessTest, TupleLiveOut) {
   // Verify MaybeLiveOut with nested tuples. Result of computation looks like:
   //