diff options
Diffstat (limited to 'tensorflow/compiler/xla/service/liveness_util_test.cc')
| -rw-r--r-- | tensorflow/compiler/xla/service/liveness_util_test.cc | 97 |
1 files changed, 74 insertions, 23 deletions
diff --git a/tensorflow/compiler/xla/service/liveness_util_test.cc b/tensorflow/compiler/xla/service/liveness_util_test.cc index d89dab4a82..b5e15906d3 100644 --- a/tensorflow/compiler/xla/service/liveness_util_test.cc +++ b/tensorflow/compiler/xla/service/liveness_util_test.cc @@ -35,6 +35,8 @@ class PointsToAnalysisTestBase : public HloTestBase { CHECK_NOTNULL(module_.get()); points_to_analysis_ = TuplePointsToAnalysis::Run(module_.get()).ConsumeValueOrDie(); + dataflow_analysis_ = + HloDataflowAnalysis::Run(module_.get()).ConsumeValueOrDie(); } void BuildModuleAndRunAnalysis(std::unique_ptr<HloComputation> computation) { @@ -45,6 +47,7 @@ class PointsToAnalysisTestBase : public HloTestBase { std::unique_ptr<HloModule> module_; HloComputation* computation_ = nullptr; std::unique_ptr<TuplePointsToAnalysis> points_to_analysis_; + std::unique_ptr<HloDataflowAnalysis> dataflow_analysis_; }; class DoesNotUseOperandBufferTest : public PointsToAnalysisTestBase {}; @@ -70,6 +73,11 @@ TEST_F(DoesNotUseOperandBufferTest, GetTupleElement) { EXPECT_TRUE(DoesNotUseOperandBuffer(tuple, {1}, gte1, *points_to_analysis_)); EXPECT_FALSE(DoesNotUseOperandBuffer(tuple, {}, gte0, *points_to_analysis_)); EXPECT_FALSE(DoesNotUseOperandBuffer(tuple, {}, gte1, *points_to_analysis_)); + + EXPECT_TRUE(DoesNotUseOperandBuffer(tuple, {0}, gte0, *dataflow_analysis_)); + EXPECT_TRUE(DoesNotUseOperandBuffer(tuple, {1}, gte1, *dataflow_analysis_)); + EXPECT_FALSE(DoesNotUseOperandBuffer(tuple, {}, gte0, *dataflow_analysis_)); + EXPECT_FALSE(DoesNotUseOperandBuffer(tuple, {}, gte1, *dataflow_analysis_)); } TEST_F(DoesNotUseOperandBufferTest, FusedDynamicUpdateSlice) { @@ -105,6 +113,10 @@ TEST_F(DoesNotUseOperandBufferTest, FusedDynamicUpdateSlice) { DoesNotUseOperandBuffer(tuple, {0}, fusion, *points_to_analysis_)); EXPECT_FALSE( DoesNotUseOperandBuffer(tuple, {1}, fusion, *points_to_analysis_)); + + EXPECT_TRUE(DoesNotUseOperandBuffer(tuple, {0}, fusion, *dataflow_analysis_)); + EXPECT_FALSE( + DoesNotUseOperandBuffer(tuple, {1}, fusion, *dataflow_analysis_)); } class CanShareOperandBufferWithUserTest : public PointsToAnalysisTestBase {}; @@ -122,10 +134,15 @@ TEST_F(CanShareOperandBufferWithUserTest, ElementWiseSameShape) { BuildModuleAndRunAnalysis(builder.Build()); - EXPECT_TRUE(CanShareOperandBufferWithUser(param, {}, exp, {}, - points_to_analysis_.get())); - EXPECT_TRUE(CanShareOperandBufferWithUser(exp, {}, log, {}, - points_to_analysis_.get())); + EXPECT_TRUE( + CanShareOperandBufferWithUser(param, {}, exp, {}, *points_to_analysis_)); + EXPECT_TRUE( + CanShareOperandBufferWithUser(exp, {}, log, {}, *points_to_analysis_)); + + EXPECT_TRUE( + CanShareOperandBufferWithUser(param, {}, exp, {}, *dataflow_analysis_)); + EXPECT_TRUE( + CanShareOperandBufferWithUser(exp, {}, log, {}, *dataflow_analysis_)); } TEST_F(CanShareOperandBufferWithUserTest, ElementWiseDifferentShape) { @@ -143,9 +160,14 @@ TEST_F(CanShareOperandBufferWithUserTest, ElementWiseDifferentShape) { BuildModuleAndRunAnalysis(builder.Build()); EXPECT_FALSE(CanShareOperandBufferWithUser(param0, {}, result, {}, - points_to_analysis_.get())); + *points_to_analysis_)); + EXPECT_FALSE(CanShareOperandBufferWithUser(param1, {}, result, {}, + *points_to_analysis_)); + + EXPECT_FALSE(CanShareOperandBufferWithUser(param0, {}, result, {}, + *dataflow_analysis_)); EXPECT_FALSE(CanShareOperandBufferWithUser(param1, {}, result, {}, - points_to_analysis_.get())); + *dataflow_analysis_)); } TEST_F(CanShareOperandBufferWithUserTest, CopyShares) { @@ -161,10 +183,15 @@ TEST_F(CanShareOperandBufferWithUserTest, CopyShares) { BuildModuleAndRunAnalysis(builder.Build()); - EXPECT_TRUE(CanShareOperandBufferWithUser(param, {}, exp, {}, - points_to_analysis_.get())); - EXPECT_TRUE(CanShareOperandBufferWithUser(exp, {}, copy, {}, - points_to_analysis_.get())); + EXPECT_TRUE( + CanShareOperandBufferWithUser(param, {}, exp, {}, *points_to_analysis_)); + EXPECT_TRUE( + CanShareOperandBufferWithUser(exp, {}, copy, {}, *points_to_analysis_)); + + EXPECT_TRUE( + CanShareOperandBufferWithUser(param, {}, exp, {}, *dataflow_analysis_)); + EXPECT_TRUE( + CanShareOperandBufferWithUser(exp, {}, copy, {}, *dataflow_analysis_)); } TEST_F(CanShareOperandBufferWithUserTest, FusedDynamicUpdateSlice) { @@ -197,9 +224,14 @@ TEST_F(CanShareOperandBufferWithUserTest, FusedDynamicUpdateSlice) { // The fusion instruction can share with tuple element 1. EXPECT_FALSE(CanShareOperandBufferWithUser(tuple, {0}, fusion, {}, - points_to_analysis_.get())); + *points_to_analysis_)); EXPECT_TRUE(CanShareOperandBufferWithUser(tuple, {1}, fusion, {}, - points_to_analysis_.get())); + *points_to_analysis_)); + + EXPECT_FALSE(CanShareOperandBufferWithUser(tuple, {0}, fusion, {}, + *dataflow_analysis_)); + EXPECT_TRUE(CanShareOperandBufferWithUser(tuple, {1}, fusion, {}, + *dataflow_analysis_)); } TEST_F(CanShareOperandBufferWithUserTest, DynamicUpdateSliceCanShare) { @@ -221,12 +253,19 @@ TEST_F(CanShareOperandBufferWithUserTest, DynamicUpdateSliceCanShare) { // The DynamicUpdateSlice instruction can share with the data operand, but not // with update or starts. - EXPECT_TRUE(CanShareOperandBufferWithUser(data, {}, dus, {}, - points_to_analysis_.get())); - EXPECT_FALSE(CanShareOperandBufferWithUser(update, {}, dus, {}, - points_to_analysis_.get())); - EXPECT_FALSE(CanShareOperandBufferWithUser(starts, {}, dus, {}, - points_to_analysis_.get())); + EXPECT_TRUE( + CanShareOperandBufferWithUser(data, {}, dus, {}, *points_to_analysis_)); + EXPECT_FALSE( + CanShareOperandBufferWithUser(update, {}, dus, {}, *points_to_analysis_)); + EXPECT_FALSE( + CanShareOperandBufferWithUser(starts, {}, dus, {}, *points_to_analysis_)); + + EXPECT_TRUE( + CanShareOperandBufferWithUser(data, {}, dus, {}, *dataflow_analysis_)); + EXPECT_FALSE( + CanShareOperandBufferWithUser(update, {}, dus, {}, *dataflow_analysis_)); + EXPECT_FALSE( + CanShareOperandBufferWithUser(starts, {}, dus, {}, *dataflow_analysis_)); } TEST_F(CanShareOperandBufferWithUserTest, FusedDotAdd) { @@ -256,7 +295,10 @@ TEST_F(CanShareOperandBufferWithUserTest, FusedDotAdd) { // Output fused dot add should be able to share buffer with 'add_operand'. EXPECT_TRUE(CanShareOperandBufferWithUser(add_operand, {}, fusion, {}, - points_to_analysis_.get())); + *points_to_analysis_)); + + EXPECT_TRUE(CanShareOperandBufferWithUser(add_operand, {}, fusion, {}, + *dataflow_analysis_)); } TEST_F(CanShareOperandBufferWithUserTest, FusedTransposeDotAdd) { @@ -292,7 +334,10 @@ TEST_F(CanShareOperandBufferWithUserTest, FusedTransposeDotAdd) { // Output fused transpose-dot-add should be share buffer with 'add_operand'. EXPECT_TRUE(CanShareOperandBufferWithUser(add_operand, {}, fusion, {}, - points_to_analysis_.get())); + *points_to_analysis_)); + + EXPECT_TRUE(CanShareOperandBufferWithUser(add_operand, {}, fusion, {}, + *dataflow_analysis_)); } TEST_F(CanShareOperandBufferWithUserTest, OutputFusionCantAliasOperandBuffer) { @@ -320,7 +365,10 @@ TEST_F(CanShareOperandBufferWithUserTest, OutputFusionCantAliasOperandBuffer) { // Output fused operand->reverse->add cannot alias operand buffer 'operand'. EXPECT_FALSE(CanShareOperandBufferWithUser(operand, {}, fusion, {}, - points_to_analysis_.get())); + *points_to_analysis_)); + + EXPECT_FALSE(CanShareOperandBufferWithUser(operand, {}, fusion, {}, + *dataflow_analysis_)); } TEST_F(CanShareOperandBufferWithUserTest, WhileCanShare) { @@ -360,8 +408,11 @@ TEST_F(CanShareOperandBufferWithUserTest, WhileCanShare) { RunAnalysis(); // The While instruction can share with the data operand. - EXPECT_TRUE(CanShareOperandBufferWithUser(data, {}, whil, {}, - points_to_analysis_.get())); + EXPECT_TRUE( + CanShareOperandBufferWithUser(data, {}, whil, {}, *points_to_analysis_)); + + EXPECT_TRUE( + CanShareOperandBufferWithUser(data, {}, whil, {}, *dataflow_analysis_)); } } // namespace |