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 <algorithm>
17 #include <string>
18
19 #include "absl/strings/ascii.h"
20 #include "absl/strings/str_cat.h"
21 #include "llvm-c/Target.h"
22 #include "tensorflow/compiler/xla/service/cpu/cpu_compiler.h"
23 #include "tensorflow/compiler/xla/service/cpu/tests/cpu_codegen_test.h"
24 #include "tensorflow/compiler/xla/service/hlo_computation.h"
25 #include "tensorflow/core/platform/test.h"
26
27 namespace xla {
28 namespace cpu {
29 namespace {
30
31 const char* const kTriple_x86_64 = "x86_64-pc-linux";
32 const char* const kTriple_android_arm = "armv7-none-android";
33
34 struct IntrinsicTestSpec {
35 HloOpcode opcode;
36 absl::string_view triple;
37 absl::string_view features;
38 absl::string_view check_lines;
39 };
40
41 // Tests that unary functions get lowered using intrinsic calls.
42 class CpuUnaryIntrinsicTest
43 : public CpuCodegenTest,
44 public ::testing::WithParamInterface<IntrinsicTestSpec> {
45 public:
Name(const::testing::TestParamInfo<IntrinsicTestSpec> & info)46 static std::string Name(
47 const ::testing::TestParamInfo<IntrinsicTestSpec>& info) {
48 auto spec = info.param;
49
50 std::string opcode = HloOpcodeString(spec.opcode);
51 opcode[0] = toupper(opcode[0]);
52
53 std::string triple{spec.triple.data(), spec.triple.size()};
54 if (triple == kTriple_x86_64) {
55 triple = "x86_64";
56 } else if (triple == kTriple_android_arm) {
57 triple = "android_arm";
58 } else {
59 triple = "Unknown";
60 }
61
62 std::string features{spec.features.data(), spec.features.size()};
63 if (!features.empty()) {
64 std::replace_if(
65 features.begin(), features.end(),
66 [](char c) { return c != '_' && !absl::ascii_isalnum(c); }, '_');
67 } else {
68 features = "";
69 }
70
71 return absl::StrCat(opcode, "_On_", triple,
72 (features.empty() ? "" : "_With"), features);
73 }
74
75 private:
GetDebugOptionsForTest()76 DebugOptions GetDebugOptionsForTest() override {
77 DebugOptions debug_options = HloTestBase::GetDebugOptionsForTest();
78 HloTestBase::SetAotFastMathDebugOptions(&debug_options);
79 return debug_options;
80 }
81 };
82
83 // Creates a module with a call to the unary op, and tests if the
84 // compiler replaced it with a call to the intrinsic.
TEST_P(CpuUnaryIntrinsicTest,DoIt)85 TEST_P(CpuUnaryIntrinsicTest, DoIt) {
86 HloComputation::Builder builder(TestName());
87 IntrinsicTestSpec spec = GetParam();
88
89 LLVMInitializeX86Target();
90 LLVMInitializeX86TargetInfo();
91 LLVMInitializeX86TargetMC();
92 LLVMInitializeARMTarget();
93 LLVMInitializeARMTargetInfo();
94 LLVMInitializeARMTargetMC();
95
96 auto param_shape = ShapeUtil::MakeShape(F32, {1024});
97 HloInstruction* param = builder.AddInstruction(
98 HloInstruction::CreateParameter(0, param_shape, "input"));
99 builder.AddInstruction(
100 HloInstruction::CreateUnary(param_shape, spec.opcode, param));
101 std::unique_ptr<HloComputation> computation = builder.Build();
102
103 std::string triple{spec.triple.data(), spec.triple.size()};
104 std::string features{spec.features.data(), spec.features.size()};
105
106 CpuAotCompilationOptions options{
107 /*triple=*/triple, /*cpu_name=*/"", /*features=*/features,
108 /*entry_point_name=*/"entry",
109 /*relocation_model=*/CpuAotCompilationOptions::RelocationModel::Static};
110
111 auto hlo_module = CreateNewVerifiedModule();
112 hlo_module->AddEntryComputation(std::move(computation));
113
114 std::string check_lines{spec.check_lines.data(), spec.check_lines.size()};
115
116 CompileAheadOfTimeAndVerifyIr(std::move(hlo_module), options, check_lines,
117 /*match_optimized_ir=*/true);
118 }
119
120 IntrinsicTestSpec CpuUnaryIntrinsicTestCases[] = {
121 // The intrinsics are always inlined, so we match a line from it instead of
122 // a function call.
123
124 IntrinsicTestSpec{
125 HloOpcode::kExp, kTriple_x86_64, "",
126 R"(CHECK: fmul fast <4 x float> <float 0xBF2BD01060000000, float 0xBF2BD01060000000, float 0xBF2BD01060000000, float 0xBF2BD01060000000>)"},
127
128 IntrinsicTestSpec{
129 HloOpcode::kExp, kTriple_x86_64, "+avx",
130 R"(CHECK: fmul fast <8 x float> <float 0xBF2BD01060000000, float 0xBF2BD01060000000, float 0xBF2BD01060000000, float 0xBF2BD01060000000, float 0xBF2BD01060000000, float 0xBF2BD01060000000, float 0xBF2BD01060000000, float 0xBF2BD01060000000>)"},
131
132 IntrinsicTestSpec{
133 HloOpcode::kExp, kTriple_android_arm, "+neon",
134 R"(CHECK: fmul fast <4 x float> <float 0xBF2BD01060000000, float 0xBF2BD01060000000, float 0xBF2BD01060000000, float 0xBF2BD01060000000>)"},
135
136 IntrinsicTestSpec{
137 HloOpcode::kTanh, kTriple_x86_64, "",
138 R"(CHECK: fcmp fast uge <4 x float> %wide.load, <float -9.000000e+00, float -9.000000e+00, float -9.000000e+00, float -9.000000e+00>)"},
139
140 IntrinsicTestSpec{
141 HloOpcode::kTanh, kTriple_x86_64, "+avx",
142 R"(CHECK: fcmp fast uge <8 x float> %wide.load, <float -9.000000e+00, float -9.000000e+00, float -9.000000e+00, float -9.000000e+00, float -9.000000e+00, float -9.000000e+00, float -9.000000e+00, float -9.000000e+00>)"},
143
144 IntrinsicTestSpec{
145 HloOpcode::kTanh, kTriple_android_arm, "",
146 R"(CHECK: fcmp fast uge <4 x float> %wide.load, <float -9.000000e+00, float -9.000000e+00, float -9.000000e+00, float -9.000000e+00>)"},
147
148 IntrinsicTestSpec{
149 HloOpcode::kLog, kTriple_x86_64, "",
150 R"(CHECK: fadd fast <4 x float> <float 0x3FBDE4A340000000, float 0x3FBDE4A340000000, float 0x3FBDE4A340000000, float 0x3FBDE4A340000000>)"},
151
152 IntrinsicTestSpec{
153 HloOpcode::kLog, kTriple_x86_64, "+avx",
154 R"(CHECK: fadd fast <8 x float> <float 0x3FBDE4A340000000, float 0x3FBDE4A340000000, float 0x3FBDE4A340000000, float 0x3FBDE4A340000000, float 0x3FBDE4A340000000, float 0x3FBDE4A340000000, float 0x3FBDE4A340000000, float 0x3FBDE4A340000000>)"},
155
156 IntrinsicTestSpec{
157 HloOpcode::kLog, kTriple_android_arm, "",
158 R"(CHECK: fadd fast <4 x float> <float 0x3FBDE4A340000000, float 0x3FBDE4A340000000, float 0x3FBDE4A340000000, float 0x3FBDE4A340000000>)"}};
159
160 INSTANTIATE_TEST_SUITE_P(CpuUnaryIntrinsicTestInstantiation,
161 CpuUnaryIntrinsicTest,
162 ::testing::ValuesIn(CpuUnaryIntrinsicTestCases),
163 CpuUnaryIntrinsicTest::Name);
164
165 } // namespace
166 } // namespace cpu
167 } // namespace xla
168