| 190 | } |
| 191 | |
| 192 | XLA_TEST_F(CpuGpuFusionTest, Test) { |
| 193 | // test expression: |
| 194 | // slice(select({{T, F, T}, {F, T, F}}, |
| 195 | // concat(transpose({{1.0}, {2.0}, {3.0}} + |
| 196 | // {{-1.0}, {-1.0}, {-1.0}}), |
| 197 | // {{1.62, 2.72, 3.14}}) + |
| 198 | // (-{{1.0, 1.0, 1.0}, {0.0, 0.0, 0.0}}), |
| 199 | // {{0.5, 0.5, 0.5}, {0.5, 0.5, 0.5}})) = {{0.5}, {2.72}} |
| 200 | auto builder = HloComputation::Builder(TestName()); |
| 201 | auto hlo_module = CreateNewVerifiedModule(); |
| 202 | auto const0 = builder.AddInstruction(HloInstruction::CreateConstant( |
| 203 | LiteralUtil::CreateR2<float>({{1.0}, {2.0}, {3.0}}))); |
| 204 | auto const1 = builder.AddInstruction(HloInstruction::CreateConstant( |
| 205 | LiteralUtil::CreateR2<float>({{-1.0}, {-1.0}, {-1.0}}))); |
| 206 | auto add2 = builder.AddInstruction(HloInstruction::CreateBinary( |
| 207 | ShapeUtil::MakeShape(F32, {3, 1}), HloOpcode::kAdd, const0, const1)); |
| 208 | auto reshape3 = builder.AddInstruction(HloInstruction::CreateTranspose( |
| 209 | ShapeUtil::MakeShape(F32, {1, 3}), add2, {1, 0})); |
| 210 | auto const4 = builder.AddInstruction(HloInstruction::CreateConstant( |
| 211 | LiteralUtil::CreateR2<float>({{1.62, 2.72, 3.14}}))); |
| 212 | auto concat5 = builder.AddInstruction(HloInstruction::CreateConcatenate( |
| 213 | ShapeUtil::MakeShape(F32, {2, 3}), {reshape3, const4}, 0)); |
| 214 | auto const6 = builder.AddInstruction(HloInstruction::CreateConstant( |
| 215 | LiteralUtil::CreateR2<float>({{1.0, 1.0, 1.0}, {0.0, 0.0, 0.0}}))); |
| 216 | auto negate7 = builder.AddInstruction(HloInstruction::CreateUnary( |
| 217 | ShapeUtil::MakeShape(F32, {2, 3}), HloOpcode::kNegate, const6)); |
| 218 | auto add8 = builder.AddInstruction(HloInstruction::CreateBinary( |
| 219 | ShapeUtil::MakeShape(F32, {2, 3}), HloOpcode::kAdd, concat5, negate7)); |
| 220 | auto const9 = builder.AddInstruction(HloInstruction::CreateConstant( |
| 221 | LiteralUtil::CreateR2<float>({{0.5, 0.5, 0.5}, {0.5, 0.5, 0.5}}))); |
| 222 | auto const10 = builder.AddInstruction( |
| 223 | HloInstruction::CreateConstant(LiteralUtil::CreateR2<bool>( |
| 224 | {{true, false, true}, {false, true, false}}))); |
| 225 | auto select11 = builder.AddInstruction( |
| 226 | HloInstruction::CreateTernary(ShapeUtil::MakeShape(F32, {2, 3}), |
| 227 | HloOpcode::kSelect, const10, add8, const9)); |
| 228 | auto slice12 = builder.AddInstruction(HloInstruction::CreateSlice( |
| 229 | ShapeUtil::MakeShape(F32, {2, 1}), select11, {0, 1}, {2, 2}, {1, 1})); |
| 230 | // CreateFusionInstruction needs the `instructions_to_fuse` argument in |
| 231 | // reverse topological order, so the first element in `instructions_to_fuse` |
| 232 | // must be the root. |
| 233 | hlo_module->AddEntryComputation(builder.Build()) |
| 234 | ->CreateFusionInstruction( |
| 235 | {slice12, select11, const10, const9, add8, negate7, const6, concat5, |
| 236 | const4, reshape3, add2, const1, const0}, |
| 237 | HloInstruction::FusionKind::kLoop); |
| 238 | |
| 239 | EXPECT_TRUE(LiteralTestUtil::Near( |
| 240 | LiteralUtil::CreateR2<float>({{0.5}, {2.72}}), |
| 241 | ExecuteAndTransfer(std::move(hlo_module), {}), ErrorSpec(1e-4))); |
| 242 | } |
| 243 | |
| 244 | // Test whether we emit appropriate code for parameters of fusion instructions. |
| 245 | XLA_TEST_F(CpuGpuFusionTest, Parameter) { |
nothing calls this directly
no test coverage detected