| 58 | } |
| 59 | |
| 60 | TEST_F(HloSchedulingTest, LastUseScheduledFirst) { |
| 61 | // Tests scheduling of the following HLO code: |
| 62 | // |
| 63 | // %ab = abs(%param) |
| 64 | // %exp = exp(%param) |
| 65 | // %add = add(%ab, %exp) |
| 66 | // %negate = negate(%exp) |
| 67 | // %sub = subtract(%add, %negate) |
| 68 | // |
| 69 | // %add should be scheduled before %negate because %add is the last (and only) |
| 70 | // use of %ab. Scheduling %add first then frees up %ab's buffer. |
| 71 | const Shape vec = ShapeUtil::MakeShape(xla::F32, {42}); |
| 72 | auto builder = HloComputation::Builder(TestName()); |
| 73 | auto param = |
| 74 | builder.AddInstruction(HloInstruction::CreateParameter(0, vec, "param")); |
| 75 | auto ab = builder.AddInstruction( |
| 76 | HloInstruction::CreateUnary(vec, HloOpcode::kAbs, param)); |
| 77 | auto exp = builder.AddInstruction( |
| 78 | HloInstruction::CreateUnary(vec, HloOpcode::kExp, param)); |
| 79 | |
| 80 | auto add = builder.AddInstruction( |
| 81 | HloInstruction::CreateBinary(vec, HloOpcode::kAdd, ab, exp)); |
| 82 | auto negate = builder.AddInstruction( |
| 83 | HloInstruction::CreateUnary(vec, HloOpcode::kNegate, exp)); |
| 84 | auto sub = builder.AddInstruction( |
| 85 | HloInstruction::CreateBinary(vec, HloOpcode::kSubtract, add, negate)); |
| 86 | |
| 87 | auto module = CreateNewVerifiedModule(); |
| 88 | module->AddEntryComputation(builder.Build()); |
| 89 | |
| 90 | HloMemoryScheduler scheduler([](const BufferValue& buffer) { |
| 91 | return ShapeUtil::ByteSizeOf(buffer.shape()); |
| 92 | }); |
| 93 | ASSERT_FALSE(module->has_schedule()); |
| 94 | TF_ASSERT_OK_AND_ASSIGN(bool changed, scheduler.Run(module.get())); |
| 95 | EXPECT_TRUE(changed); |
| 96 | ASSERT_TRUE(module->has_schedule()); |
| 97 | TF_ASSERT_OK(module->schedule().Verify()); |
| 98 | |
| 99 | // Verify that all instructions are in the sequence. |
| 100 | const std::vector<HloInstruction*>& sequence = |
| 101 | module->schedule().sequence(module->entry_computation()).instructions(); |
| 102 | EXPECT_EQ(module->entry_computation()->instruction_count(), sequence.size()); |
| 103 | |
| 104 | // The first instruction should be the parameter and the last the root "sub". |
| 105 | EXPECT_EQ(param, sequence.front()); |
| 106 | EXPECT_EQ(sub, sequence.back()); |
| 107 | |
| 108 | SequentialHloOrdering ordering(module->schedule()); |
| 109 | EXPECT_TRUE(ordering.ExecutesBefore(add, negate)); |
| 110 | |
| 111 | // Clear the schedule using the descheduling pass. |
| 112 | HloDescheduler descheduler; |
| 113 | EXPECT_TRUE(module->has_schedule()); |
| 114 | TF_ASSERT_OK_AND_ASSIGN(bool descheduler_changed, |
| 115 | descheduler.Run(module.get())); |
| 116 | EXPECT_TRUE(descheduler_changed); |
| 117 | EXPECT_FALSE(module->has_schedule()); |
nothing calls this directly
no test coverage detected