1 /* Copyright 2018 The TensorFlow Authors. All Rights Reserved.
2
3 Licensed under the Apache License, Version 2.0 (the "License");
4 you may not use this file except in compliance with the License.
5 You may obtain a copy of the License at
6
7 http://www.apache.org/licenses/LICENSE-2.0
8
9 Unless required by applicable law or agreed to in writing, software
10 distributed under the License is distributed on an "AS IS" BASIS,
11 WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12 See the License for the specific language governing permissions and
13 limitations under the License.
14 ==============================================================================*/
15
16 #include "tensorflow/compiler/xla/service/cpu/cpu_compiler.h"
17 #include "tensorflow/compiler/xla/service/cpu/test_target_triple_helper.h"
18 #include "tensorflow/compiler/xla/service/cpu/tests/cpu_codegen_test.h"
19 #include "tensorflow/compiler/xla/service/hlo_parser.h"
20
21 namespace xla {
22 namespace cpu {
23 namespace {
24 class CpuDuplicateConstantsTest : public CpuCodegenTest {};
25
TEST_F(CpuDuplicateConstantsTest,RepeatedArrayConstants)26 TEST_F(CpuDuplicateConstantsTest, RepeatedArrayConstants) {
27 // We use a while loop here to force the two constant HloInstructions to be in
28 // different computations. Otherwise the HLO optimizer itself CSEs them.
29 const std::string hlo_text = R"(
30 HloModule RepeatedConstants
31
32 while_body {
33 arg_body = f32[2,3,2] parameter(0)
34 ROOT const = f32[2,3,2] constant(
35 {{{1, 2}, {1001, 1002}, {2001, 2002}},
36 {{2, 1}, {2001, 3002}, {2001, 2002}}})
37 }
38
39 while_cond {
40 arg_cond = f32[2,3,2] parameter(0)
41 token0 = token[] after-all()
42 infeed = (pred[], token[]) infeed(token0)
43 ROOT unknown = pred[] get-tuple-element((pred[], token[]) infeed), index=0
44 }
45
46 ENTRY main {
47 param = f32[2,3,2] parameter(0)
48 const_a = f32[2,3,2] constant(
49 {{{1, 2}, {1001, 1002}, {2001, 2002}},
50 {{2, 1}, {2001, 3002}, {2001, 2002}}})
51 const_b = f32[2,3,2] while(f32[2,3,2] const_a), condition=while_cond, body=while_body
52
53 token0 = token[] after-all()
54 out0 = token[] outfeed(f32[2,3,2] const_a, token[] token0)
55 ROOT out1 = token[] outfeed(f32[2,3,2] const_b, token[] token0)
56 }
57 )";
58
59 std::string filecheck_pattern = R"(
60 CHECK: private unnamed_addr constant [48 x i8]
61 CHECK-NOT: private unnamed_addr constant [48 x i8]
62 )";
63
64 TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr<HloModule> module,
65 ParseAndReturnVerifiedModule(hlo_text));
66
67 CpuAotCompilationOptions options{
68 /*triple=*/kTargetTripleForHost, /*cpu_name=*/kTargetCpuForHost,
69 /*features=*/"",
70 /*entry_point_name=*/"entry",
71 /*relocation_model=*/CpuAotCompilationOptions::RelocationModel::Static};
72
73 CompileAheadOfTimeAndVerifyIr(std::move(module), options, filecheck_pattern,
74 /*match_optimized_ir=*/false);
75 }
76
TEST_F(CpuDuplicateConstantsTest,RepeatedTupleConstants)77 TEST_F(CpuDuplicateConstantsTest, RepeatedTupleConstants) {
78 // We use a while loop here to force the two constant HloInstructions to be in
79 // different computations. Otherwise the HLO optimizer itself CSEs them.
80 const std::string hlo_text = R"(
81 HloModule RepeatedConstants
82
83 while_body {
84 arg_body = (f32[2,1]{1,0}, f32[1]{0}) parameter(0)
85 ROOT const = (f32[2,1]{1,0}, f32[1]{0}) constant(({ { 1 }, { 2 } }, {2} ))
86 }
87
88 while_cond {
89 arg_cond = (f32[2,1]{1,0}, f32[1]{0}) parameter(0)
90 token0 = token[] after-all()
91 infeed = (pred[], token[]) infeed(token0)
92 ROOT unknown = pred[] get-tuple-element((pred[], token[]) infeed), index=0
93 }
94
95 ENTRY main {
96 param = f32[2,3,2] parameter(0)
97 const_a = (f32[2,1]{1,0}, f32[1]{0}) constant(( { { 1 }, { 2 } }, {2} ))
98 const_b = (f32[2,1]{1,0}, f32[1]{0}) while((f32[2,1]{1,0}, f32[1]{0}) const_a), condition=while_cond, body=while_body
99
100 token0 = token[] after-all()
101 out0 = () outfeed((f32[2,1]{1,0}, f32[1]{0}) const_a, token[] token0)
102 ROOT out1 = () outfeed((f32[2,1]{1,0}, f32[1]{0}) const_b, token[] token0)
103 }
104 )";
105
106 std::string filecheck_pattern = R"(
107 CHECK-DAG: private unnamed_addr constant [4 x i8]
108 CHECK-DAG: private unnamed_addr constant [8 x i8]
109 CHECK-NOT: private unnamed_addr constant [4 x i8]
110 CHECK-NOT: private unnamed_addr constant [8 x i8]
111 )";
112
113 TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr<HloModule> module,
114 ParseAndReturnVerifiedModule(hlo_text));
115
116 CpuAotCompilationOptions options{
117 /*triple=*/kTargetTripleForHost, /*cpu_name=*/kTargetCpuForHost,
118 /*features=*/"",
119 /*entry_point_name=*/"entry",
120 /*relocation_model=*/CpuAotCompilationOptions::RelocationModel::Static};
121
122 CompileAheadOfTimeAndVerifyIr(std::move(module), options, filecheck_pattern,
123 /*match_optimized_ir=*/false);
124 }
125
126 } // namespace
127 } // namespace cpu
128 } // namespace xla
129