6f81dbf07a
Also enable some tests that are now passing on the GPU backend. Finally, remove some unused hlo_parser.h includes and the corresponding dependency. PiperOrigin-RevId: 275224726 Change-Id: Icc206d85b40c439abe8232aaa748ed4a07b50b09
105 lines
3.3 KiB
C++
105 lines
3.3 KiB
C++
/* Copyright 2018 The TensorFlow Authors. All Rights Reserved.
|
|
|
|
Licensed under the Apache License, Version 2.0 (the "License");
|
|
you may not use this file except in compliance with the License.
|
|
You may obtain a copy of the License at
|
|
|
|
http://www.apache.org/licenses/LICENSE-2.0
|
|
|
|
Unless required by applicable law or agreed to in writing, software
|
|
distributed under the License is distributed on an "AS IS" BASIS,
|
|
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
|
See the License for the specific language governing permissions and
|
|
limitations under the License.
|
|
==============================================================================*/
|
|
|
|
#include "tensorflow/compiler/xla/literal.h"
|
|
#include "tensorflow/compiler/xla/shape_util.h"
|
|
#include "tensorflow/compiler/xla/test.h"
|
|
#include "tensorflow/compiler/xla/test_helpers.h"
|
|
#include "tensorflow/compiler/xla/tests/hlo_test_base.h"
|
|
#include "tensorflow/compiler/xla/tests/test_macros.h"
|
|
|
|
namespace xla {
|
|
namespace {
|
|
|
|
class TrivialAllReduceTest : public HloTestBase {};
|
|
|
|
// Currently the CPU and GPU backends only support AllReduce with one
|
|
// replica. But we can at least check this.
|
|
|
|
XLA_TEST_F(TrivialAllReduceTest, OneOperand) {
|
|
const char* module_str = R"(
|
|
HloModule test
|
|
|
|
add {
|
|
x = f32[] parameter(0)
|
|
y = f32[] parameter(1)
|
|
add = f32[] add(x, y)
|
|
}
|
|
|
|
ENTRY test_computation {
|
|
p = f32[3] parameter(0)
|
|
ROOT crs = f32[3] all-reduce(p), to_apply=add
|
|
})";
|
|
auto module =
|
|
ParseAndReturnVerifiedModule(module_str, GetModuleConfigForTest())
|
|
.ValueOrDie();
|
|
auto literal = LiteralUtil::CreateR1<float>({1, 2, 3});
|
|
EXPECT_EQ(literal, ExecuteAndTransfer(std::move(module), {&literal}));
|
|
}
|
|
|
|
XLA_TEST_F(TrivialAllReduceTest, MultipleOperands) {
|
|
const char* module_str = R"(
|
|
HloModule test
|
|
|
|
add {
|
|
x = f32[] parameter(0)
|
|
y = f32[] parameter(1)
|
|
add = f32[] add(x, y)
|
|
}
|
|
|
|
ENTRY test_computation {
|
|
p0 = f32[3] parameter(0)
|
|
p1 = f32[2] parameter(1)
|
|
ROOT crs = (f32[3], f32[2]) all-reduce(p0, p1), to_apply=add
|
|
})";
|
|
auto module =
|
|
ParseAndReturnVerifiedModule(module_str, GetModuleConfigForTest())
|
|
.ValueOrDie();
|
|
auto literal0 = LiteralUtil::CreateR1<float>({1, 2, 3});
|
|
auto literal1 = LiteralUtil::CreateR1<float>({10, 20});
|
|
EXPECT_EQ(LiteralUtil::MakeTuple({&literal0, &literal1}),
|
|
ExecuteAndTransfer(std::move(module), {&literal0, &literal1}));
|
|
}
|
|
|
|
// On the GPU backend, constants get special handling. Someone might pass a
|
|
// constant to CRS to e.g. count the number of replicas -- we need to make sure
|
|
// it works.
|
|
XLA_TEST_F(TrivialAllReduceTest, ConstantOperand) {
|
|
const char* module_str = R"(
|
|
HloModule test
|
|
|
|
add {
|
|
x = f32[] parameter(0)
|
|
y = f32[] parameter(1)
|
|
add = f32[] add(x, y)
|
|
}
|
|
|
|
ENTRY test_computation {
|
|
p0 = f32[3] parameter(0)
|
|
p1 = f32[2] constant({10, 20})
|
|
ROOT crs = (f32[3], f32[2]) all-reduce(p0, p1), to_apply=add
|
|
})";
|
|
auto module =
|
|
ParseAndReturnVerifiedModule(module_str, GetModuleConfigForTest())
|
|
.ValueOrDie();
|
|
auto literal0 = LiteralUtil::CreateR1<float>({1, 2, 3});
|
|
auto literal1 = LiteralUtil::CreateR1<float>({10, 20});
|
|
EXPECT_EQ(LiteralUtil::MakeTuple({&literal0, &literal1}),
|
|
ExecuteAndTransfer(std::move(module), {&literal0}));
|
|
}
|
|
|
|
} // namespace
|
|
} // namespace xla
|