From 85a9058eff03cda6ec9b58849897774c7d56fc4a Mon Sep 17 00:00:00 2001 From: Bixia Zheng Date: Mon, 5 Aug 2019 18:35:36 -0700 Subject: [PATCH] [XLA] Add exhaustive binary tests for F32 and F64. PiperOrigin-RevId: 261816030 --- tensorflow/compiler/xla/tests/BUILD | 40 ++++ .../xla/tests/exhaustive_binary_test.cc | 214 ++++++++++++++++++ .../xla/tests/exhaustive_op_test_utils.cc | 14 ++ .../xla/tests/exhaustive_op_test_utils.h | 1 + 4 files changed, 269 insertions(+) diff --git a/tensorflow/compiler/xla/tests/BUILD b/tensorflow/compiler/xla/tests/BUILD index 8b77c1b864a..5d86433f1bf 100644 --- a/tensorflow/compiler/xla/tests/BUILD +++ b/tensorflow/compiler/xla/tests/BUILD @@ -829,6 +829,46 @@ xla_test( ], ) +xla_test( + name = "exhaustive_binary_test_f32", + srcs = ["exhaustive_binary_test.cc"], + backends = [ + "gpu", + "cpu", + ], + copts = ["-DBINARY_TEST_TARGET_F32"], + real_hardware_only = True, # Very slow on the interpreter. + shard_count = 48, + tags = [ + "optonly", + # This is a big test that we skip for capacity reasons in OSS testing. + "no_oss", + ], + deps = [ + ":exhaustive_op_test_utils", + ], +) + +xla_test( + name = "exhaustive_binary_test_f64", + srcs = ["exhaustive_binary_test.cc"], + backends = [ + "gpu", + "cpu", + ], + copts = ["-DBINARY_TEST_TARGET_F64"], + real_hardware_only = True, # Very slow on the interpreter. + shard_count = 48, + tags = [ + "optonly", + # This is a big test that we skip for capacity reasons in OSS testing. + "no_oss", + ], + deps = [ + ":exhaustive_op_test_utils", + ], +) + xla_test( name = "reduce_precision_test", srcs = ["reduce_precision_test.cc"], diff --git a/tensorflow/compiler/xla/tests/exhaustive_binary_test.cc b/tensorflow/compiler/xla/tests/exhaustive_binary_test.cc index 05738414429..c0f8a0dc626 100644 --- a/tensorflow/compiler/xla/tests/exhaustive_binary_test.cc +++ b/tensorflow/compiler/xla/tests/exhaustive_binary_test.cc @@ -174,5 +174,219 @@ INSTANTIATE_TEST_SUITE_P(BF16, ExhaustiveBF16BinaryTest, #endif #endif +// Exhaustive test for binary operations for float and double. +// +// Test parameter is a tuple of (FpValues, FpValues) describing the possible +// values for each operand. The inputs for the test are the Cartesian product +// of the possible values for the two operands. +template +class Exhaustive32BitOrMoreBinaryTest + : public ExhaustiveBinaryTest, + public ::testing::WithParamInterface> { + protected: + using typename ExhaustiveBinaryTest::NativeT; + using ExhaustiveBinaryTest::ConvertAndReplaceKnownIncorrectValueWith; + + private: + int64 GetInputSize() override { + FpValues values_0; + FpValues values_1; + std::tie(values_0, values_1) = GetParam(); + return values_0.GetTotalNumValues() * values_1.GetTotalNumValues(); + } + + void FillInput(std::array* input_literals) override { + int64 input_size = GetInputSize(); + FpValues values_0; + FpValues values_1; + std::tie(values_0, values_1) = GetParam(); + + VLOG(2) << " testing " << values_0.ToString() << " " << values_1.ToString() + << "total values " << input_size; + CHECK(input_size == (*input_literals)[0].element_count() && + input_size == (*input_literals)[1].element_count()); + + absl::Span input_arr_0 = (*input_literals)[0].data(); + absl::Span input_arr_1 = (*input_literals)[1].data(); + + uint64 i = 0; + for (auto src0 : values_0) { + for (auto src1 : values_1) { + input_arr_0[i] = ConvertAndReplaceKnownIncorrectValueWith(src0, 1); + input_arr_1[i] = ConvertAndReplaceKnownIncorrectValueWith(src1, 1); + ++i; + } + } + CHECK_EQ(i, input_size); + } +}; + +using ExhaustiveF32BinaryTest = Exhaustive32BitOrMoreBinaryTest; +using ExhaustiveF64BinaryTest = Exhaustive32BitOrMoreBinaryTest; + +XLA_TEST_P(ExhaustiveF32BinaryTest, Add) { + auto host_add = [](float x, float y) { return x + y; }; + Run(AddEmptyBroadcastDimension(Add), host_add); +} + +XLA_TEST_P(ExhaustiveF32BinaryTest, Sub) { + auto host_sub = [](float x, float y) { return x - y; }; + Run(AddEmptyBroadcastDimension(Sub), host_sub); +} + +// TODO(bixia): Need to investigate the failure on CPU and file bugs. +XLA_TEST_P(ExhaustiveF32BinaryTest, DISABLED_ON_CPU(Mul)) { + auto host_mul = [](float x, float y) { return x * y; }; + Run(AddEmptyBroadcastDimension(Mul), host_mul); +} + +// TODO(bixia): Need to investigate the failure on CPU and file bugs. +XLA_TEST_P(ExhaustiveF32BinaryTest, DISABLED_ON_CPU(Div)) { + auto host_div = [](float x, float y) { return x / y; }; + Run(AddEmptyBroadcastDimension(Div), host_div); +} + +XLA_TEST_P(ExhaustiveF32BinaryTest, Max) { + Run(AddEmptyBroadcastDimension(Max), ReferenceMax); +} + +XLA_TEST_P(ExhaustiveF32BinaryTest, Min) { + Run(AddEmptyBroadcastDimension(Min), ReferenceMin); +} + +// It is more convenient to implement Abs(complex) as a binary op than a unary +// op, as the operations we currently support all have the same data type for +// the source operands and the results. +// TODO(bixia): May want to move this test to unary test if we will be able to +// implement Abs(complex) as unary conveniently. +// +// TODO(bixia): Need to investigate the failure on CPU and file bugs. +XLA_TEST_P(ExhaustiveF32BinaryTest, DISABLED_ON_CPU(AbsComplex)) { + auto host_abs_complex = [](float x, float y) { + return std::abs(std::complex(x, y)); + }; + auto device_abs_complex = [](XlaOp x, XlaOp y) { return Abs(Complex(x, y)); }; + + Run(device_abs_complex, host_abs_complex); +} + +#if defined(BINARY_TEST_TARGET_F32) + +INSTANTIATE_TEST_SUITE_P( + SpecialValues, ExhaustiveF32BinaryTest, + ::testing::Combine( + ::testing::ValuesIn(CreateFpValuesForBoundaryTest()), + ::testing::ValuesIn(CreateFpValuesForBoundaryTest()))); + +INSTANTIATE_TEST_SUITE_P( + SpecialAndNormalValues, ExhaustiveF32BinaryTest, + ::testing::Combine( + ::testing::ValuesIn(CreateFpValuesForBoundaryTest()), + ::testing::Values(GetNormals(2000)))); + +INSTANTIATE_TEST_SUITE_P( + NormalAndSpecialValues, ExhaustiveF32BinaryTest, + ::testing::Combine( + ::testing::Values(GetNormals(2000)), + ::testing::ValuesIn(CreateFpValuesForBoundaryTest()))); + +INSTANTIATE_TEST_SUITE_P( + NormalAndNormalValues, ExhaustiveF32BinaryTest, + ::testing::Combine(::testing::Values(GetNormals(2000)), + ::testing::Values(GetNormals(2000)))); + +// Tests a total of 40000 ^ 2 inputs, with 2000 ^ 2 inputs in each sub-test. +// Comparing with the unary tests, the binary tests use a smaller set of inputs +// for each sub-test to avoid timeout because the implementation of ExpectNear +// more than 2x slower for binary test. +INSTANTIATE_TEST_SUITE_P( + LargeAndSmallMagnituedNormalValues, ExhaustiveF32BinaryTest, + ::testing::Combine( + ::testing::ValuesIn(GetFpValuesForMagnitudeExtremeNormals(40000, + 2000)), + ::testing::ValuesIn( + GetFpValuesForMagnitudeExtremeNormals(40000, 2000)))); + +#endif + +XLA_TEST_P(ExhaustiveF64BinaryTest, Add) { + auto host_add = [](double x, double y) { return x + y; }; + Run(AddEmptyBroadcastDimension(Add), host_add); +} + +XLA_TEST_P(ExhaustiveF64BinaryTest, Sub) { + auto host_sub = [](double x, double y) { return x - y; }; + Run(AddEmptyBroadcastDimension(Sub), host_sub); +} + +// TODO(bixia): Need to investigate the failure on CPU and file bugs. +XLA_TEST_P(ExhaustiveF64BinaryTest, DISABLED_ON_CPU(Mul)) { + auto host_mul = [](double x, double y) { return x * y; }; + Run(AddEmptyBroadcastDimension(Mul), host_mul); +} + +// TODO(bixia): Need to investigate the failure on CPU and file bugs. +XLA_TEST_P(ExhaustiveF64BinaryTest, DISABLED_ON_CPU(Div)) { + auto host_div = [](double x, double y) { return x / y; }; + Run(AddEmptyBroadcastDimension(Div), host_div); +} + +XLA_TEST_P(ExhaustiveF64BinaryTest, Max) { + Run(AddEmptyBroadcastDimension(Max), ReferenceMax); +} + +XLA_TEST_P(ExhaustiveF64BinaryTest, Min) { + Run(AddEmptyBroadcastDimension(Min), ReferenceMin); +} + +// TODO(bixia): Need to investigate the failure on CPU and file bugs. +XLA_TEST_P(ExhaustiveF64BinaryTest, DISABLED_ON_CPU(AbsComplex)) { + auto host_abs_complex = [](double x, double y) { + return std::abs(std::complex(x, y)); + }; + auto device_abs_complex = [](XlaOp x, XlaOp y) { return Abs(Complex(x, y)); }; + + Run(device_abs_complex, host_abs_complex); +} + +#if defined(BINARY_TEST_TARGET_F64) + +#if !defined(XLA_BACKEND_DOES_NOT_SUPPORT_FLOAT64) +INSTANTIATE_TEST_SUITE_P( + SpecialValues, ExhaustiveF64BinaryTest, + ::testing::Combine( + ::testing::ValuesIn(CreateFpValuesForBoundaryTest()), + ::testing::ValuesIn(CreateFpValuesForBoundaryTest()))); + +INSTANTIATE_TEST_SUITE_P( + SpecialAndNormalValues, ExhaustiveF64BinaryTest, + ::testing::Combine( + ::testing::ValuesIn(CreateFpValuesForBoundaryTest()), + ::testing::Values(GetNormals(1000)))); + +INSTANTIATE_TEST_SUITE_P( + NormalAndSpecialValues, ExhaustiveF64BinaryTest, + ::testing::Combine( + ::testing::Values(GetNormals(1000)), + ::testing::ValuesIn(CreateFpValuesForBoundaryTest()))); + +INSTANTIATE_TEST_SUITE_P( + NormalAndNormalValues, ExhaustiveF64BinaryTest, + ::testing::Combine(::testing::Values(GetNormals(1000)), + ::testing::Values(GetNormals(1000)))); + +// Tests a total of 40000 ^ 2 inputs, with 1000 ^ 2 inputs in each sub-test. +// Similar to ExhaustiveF64BinaryTest, we use a smaller set of inputs for each +// for each sub-test comparing with the unary test to avoid timeout. +INSTANTIATE_TEST_SUITE_P( + LargeAndSmallMagnituedNormalValues, ExhaustiveF64BinaryTest, + ::testing::Combine( + ::testing::ValuesIn( + GetFpValuesForMagnitudeExtremeNormals(40000, 2000)), + ::testing::ValuesIn( + GetFpValuesForMagnitudeExtremeNormals(40000, 2000)))); +#endif + +#endif } // namespace } // namespace xla diff --git a/tensorflow/compiler/xla/tests/exhaustive_op_test_utils.cc b/tensorflow/compiler/xla/tests/exhaustive_op_test_utils.cc index 6a763f3a54c..1d3248fe04c 100644 --- a/tensorflow/compiler/xla/tests/exhaustive_op_test_utils.cc +++ b/tensorflow/compiler/xla/tests/exhaustive_op_test_utils.cc @@ -215,6 +215,18 @@ inline ExhaustiveOpTestBase::ErrorSpec DefaultSpecGenerator( return ExhaustiveOpTestBase::ErrorSpec{0.002, 0.02}; } +template <> +inline ExhaustiveOpTestBase::ErrorSpec DefaultSpecGenerator( + double, double) { + return ExhaustiveOpTestBase::ErrorSpec{0.001, 0.001}; +} + +template <> +inline ExhaustiveOpTestBase::ErrorSpec DefaultSpecGenerator( + float, float) { + return ExhaustiveOpTestBase::ErrorSpec{0.001, 0.001}; +} + template <> inline ExhaustiveOpTestBase::ErrorSpec DefaultSpecGenerator( Eigen::half, Eigen::half) { @@ -242,6 +254,8 @@ template class ExhaustiveOpTestBase; template class ExhaustiveOpTestBase; template class ExhaustiveOpTestBase; +template class ExhaustiveOpTestBase; +template class ExhaustiveOpTestBase; template class ExhaustiveOpTestBase; template class ExhaustiveOpTestBase; diff --git a/tensorflow/compiler/xla/tests/exhaustive_op_test_utils.h b/tensorflow/compiler/xla/tests/exhaustive_op_test_utils.h index 51da80e36fd..d66da60c66c 100644 --- a/tensorflow/compiler/xla/tests/exhaustive_op_test_utils.h +++ b/tensorflow/compiler/xla/tests/exhaustive_op_test_utils.h @@ -1001,6 +1001,7 @@ class FpValues { const FpValues* fp_values_; }; + FpValues() : bit_chunks_(), offsets_() {} FpValues(absl::Span chunks, absl::Span offsets) { CHECK_EQ(chunks.size(), offsets.size() - 1); CHECK_EQ(chunks.size(), kTotalBitChunks);