From e0374aaae8d1d841bd85c2361fe5f7ab8dbbd79f Mon Sep 17 00:00:00 2001 From: Adrian Kuegel Date: Thu, 6 Feb 2020 07:03:15 -0800 Subject: [PATCH] Split up exhaustive_binary_test.cc into two files. This also allows us to get rid of #define flags, and we can fix the issue that required the usage of the GTEST_ALLOW_UNINSTANTIATED_PARAMETERIZED_TEST macro. PiperOrigin-RevId: 293584602 Change-Id: Idd6723a3304200c8df5b6066a08159b5c45cb470 --- tensorflow/compiler/xla/tests/BUILD | 54 +----- .../tests/exhaustive_binary_16_bit_test.cc | 143 ++++++++++++++ ...t.cc => exhaustive_binary_test_f32_f64.cc} | 181 +----------------- .../xla/tests/exhaustive_op_test_utils.h | 42 ++++ 4 files changed, 193 insertions(+), 227 deletions(-) create mode 100644 tensorflow/compiler/xla/tests/exhaustive_binary_16_bit_test.cc rename tensorflow/compiler/xla/tests/{exhaustive_binary_test.cc => exhaustive_binary_test_f32_f64.cc} (59%) diff --git a/tensorflow/compiler/xla/tests/BUILD b/tensorflow/compiler/xla/tests/BUILD index 7aa18642bfa..bf2a1d64476 100644 --- a/tensorflow/compiler/xla/tests/BUILD +++ b/tensorflow/compiler/xla/tests/BUILD @@ -836,14 +836,12 @@ xla_test( ) xla_test( - name = "exhaustive_binary_test_f16", - srcs = ["exhaustive_binary_test.cc"], + name = "exhaustive_binary_16_bit_test", + srcs = ["exhaustive_binary_16_bit_test.cc"], backends = [ "gpu", "cpu", ], - copts = ["-DBINARY_TEST_TARGET_F16"], - real_hardware_only = True, # Very slow on the interpreter. shard_count = 48, tags = [ "optonly", @@ -857,56 +855,12 @@ xla_test( ) xla_test( - name = "exhaustive_binary_test_bf16", - srcs = ["exhaustive_binary_test.cc"], + name = "exhaustive_binary_test_f32_f64", + srcs = ["exhaustive_binary_test_f32_f64.cc"], backends = [ "gpu", "cpu", ], - copts = ["-DBINARY_TEST_TARGET_BF16"], - 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", - "no_pip", - ], - deps = [ - ":exhaustive_op_test_utils", - ], -) - -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", - "no_pip", - ], - 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", diff --git a/tensorflow/compiler/xla/tests/exhaustive_binary_16_bit_test.cc b/tensorflow/compiler/xla/tests/exhaustive_binary_16_bit_test.cc new file mode 100644 index 00000000000..4682319b5e7 --- /dev/null +++ b/tensorflow/compiler/xla/tests/exhaustive_binary_16_bit_test.cc @@ -0,0 +1,143 @@ +/* Copyright 2020 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/tests/exhaustive_op_test_utils.h" + +#ifdef __FAST_MATH__ +#error("Can't be compiled with fast math on"); +#endif + +namespace xla { +namespace exhaustive_op_test { +namespace { + +// Exhaustive test for binary operations for 16 bit floating point types, +// including float16 and bfloat. +// +// Test parameter is a pair of (begin, end) for range under test. +template +class Exhaustive16BitBinaryTest + : public ExhaustiveBinaryTest, + public ::testing::WithParamInterface> { + public: + int64 GetInputSize() override { + int64 begin, end; + std::tie(begin, end) = GetParam(); + return end - begin; + } + + // Given a range of uint64 representation, uses bits 0..15 and bits 16..31 for + // the values of src0 and src1 for a 16 bit binary operation being tested, + // and generates the cartesian product of the two sets as the two inputs for + // the test. + void FillInput(std::array* input_literals) override { + int64 input_size = GetInputSize(); + CHECK_EQ(input_size, (*input_literals)[0].element_count()); + CHECK_EQ(input_size, (*input_literals)[1].element_count()); + + int64 begin, end; + std::tie(begin, end) = GetParam(); + VLOG(2) << "Checking range [" << begin << ", " << end << "]"; + + absl::Span input_arr_0 = (*input_literals)[0].data(); + absl::Span input_arr_1 = (*input_literals)[1].data(); + for (int64 i = 0; i < input_size; i++) { + uint32 input_val = i + begin; + // Convert the lower 16 bits to the NativeT and replaced known incorrect + // input values with 0. + input_arr_0[i] = ConvertAndReplaceKnownIncorrectValueWith(input_val, 0); + input_arr_1[i] = + ConvertAndReplaceKnownIncorrectValueWith(input_val >> 16, 0); + } + } + + protected: + using typename ExhaustiveBinaryTest::NativeT; + using ExhaustiveBinaryTest::ConvertAndReplaceKnownIncorrectValueWith; +}; + +#if !defined(XLA_BACKEND_DOES_NOT_SUPPORT_FLOAT16) +using ExhaustiveF16BinaryTest = Exhaustive16BitBinaryTest; +#define BINARY_TEST_F16(test_name, ...) \ + XLA_TEST_P(ExhaustiveF16BinaryTest, test_name) \ + __VA_ARGS__ +#else +#define BINARY_TEST_F16(test_name, ...) +#endif + +#if defined(XLA_BACKEND_SUPPORTS_BFLOAT16) +using ExhaustiveBF16BinaryTest = Exhaustive16BitBinaryTest; +#define BINARY_TEST_BF16(test_name, ...) \ + XLA_TEST_P(ExhaustiveBF16BinaryTest, test_name) \ + __VA_ARGS__ +#else +#define BINARY_TEST_BF16(test_name, ...) +#endif + +#define BINARY_TEST_16BIT(test_name, ...) \ + BINARY_TEST_F16(test_name, __VA_ARGS__) \ + BINARY_TEST_BF16(test_name, __VA_ARGS__) + +BINARY_TEST_16BIT(Add, { + auto host_add = [](float x, float y) { return x + y; }; + Run(AddEmptyBroadcastDimension(Add), host_add); +}) + +BINARY_TEST_16BIT(Sub, { + auto host_sub = [](float x, float y) { return x - y; }; + Run(AddEmptyBroadcastDimension(Sub), host_sub); +}) + +// TODO(bixia): Mul fails with bfloat16 on CPU. +BINARY_TEST_16BIT(DISABLED_ON_CPU(Mul), { + auto host_mul = [](float x, float y) { return x * y; }; + Run(AddEmptyBroadcastDimension(Mul), host_mul); +}) + +// TODO(bixia): Div fails with bfloat16 on CPU. +BINARY_TEST_16BIT(DISABLED_ON_CPU(Div), { + auto host_div = [](float x, float y) { return x / y; }; + Run(AddEmptyBroadcastDimension(Div), host_div); +}) + +BINARY_TEST_16BIT(Max, { + Run(AddEmptyBroadcastDimension(Max), ReferenceMax); +}) + +BINARY_TEST_16BIT(Min, { + Run(AddEmptyBroadcastDimension(Min), ReferenceMin); +}) + +// TODO(bixia): Pow fails with bfloat16 on CPU. +BINARY_TEST_16BIT(DISABLED_ON_CPU(Pow), + { Run(AddEmptyBroadcastDimension(Pow), std::powf); }) + +// TODO(bixia): Atan2 fails with bfloat16 on CPU. +BINARY_TEST_16BIT(DISABLED_ON_CPU(Atan2), + { Run(AddEmptyBroadcastDimension(Atan2), std::atan2f); }) + +#if !defined(XLA_BACKEND_DOES_NOT_SUPPORT_FLOAT16) +INSTANTIATE_TEST_SUITE_P(F16, ExhaustiveF16BinaryTest, + ::testing::ValuesIn(CreateExhaustiveF32Ranges())); +#endif + +#if defined(XLA_BACKEND_SUPPORTS_BFLOAT16) +INSTANTIATE_TEST_SUITE_P(BF16, ExhaustiveBF16BinaryTest, + ::testing::ValuesIn(CreateExhaustiveF32Ranges())); +#endif + +} // namespace +} // namespace exhaustive_op_test +} // namespace xla diff --git a/tensorflow/compiler/xla/tests/exhaustive_binary_test.cc b/tensorflow/compiler/xla/tests/exhaustive_binary_test_f32_f64.cc similarity index 59% rename from tensorflow/compiler/xla/tests/exhaustive_binary_test.cc rename to tensorflow/compiler/xla/tests/exhaustive_binary_test_f32_f64.cc index 9c67ee4e01a..14d3b343b6c 100644 --- a/tensorflow/compiler/xla/tests/exhaustive_binary_test.cc +++ b/tensorflow/compiler/xla/tests/exhaustive_binary_test_f32_f64.cc @@ -23,171 +23,6 @@ namespace xla { namespace exhaustive_op_test { namespace { -template -using ExhaustiveBinaryTest = ExhaustiveOpTestBase; - -// Exhaustive test for binary operations for 16 bit floating point types, -// including float16 and bfloat. -// -// Test parameter is a pair of (begin, end) for range under test. -template < - PrimitiveType T, - typename std::enable_if< - std::is_same::type, - half>::value || - std::is_same::type, - bfloat16>::value>::type* = nullptr> -class Exhaustive16BitBinaryTest - : public ExhaustiveBinaryTest, - public ::testing::WithParamInterface> { - public: - int64 GetInputSize() override { - int64 begin, end; - std::tie(begin, end) = GetParam(); - return end - begin; - } - - // Given a range of uint64 representation, uses bits 0..15 and bits 16..31 for - // the values of src0 and src1 for a 16 bit binary operation being tested, - // and generates the cartesian product of the two sets as the two inputs for - // the test. - void FillInput(std::array* input_literals) override { - int64 input_size = GetInputSize(); - CHECK_EQ(input_size, (*input_literals)[0].element_count()); - CHECK_EQ(input_size, (*input_literals)[1].element_count()); - - int64 begin, end; - std::tie(begin, end) = GetParam(); - VLOG(2) << "Checking range [" << begin << ", " << end << "]"; - - absl::Span input_arr_0 = (*input_literals)[0].data(); - absl::Span input_arr_1 = (*input_literals)[1].data(); - for (int64 i = 0; i < input_size; i++) { - uint32 input_val = i + begin; - // Convert the lower 16 bits to the NativeT and replaced known incorrect - // input values with 0. - input_arr_0[i] = ConvertAndReplaceKnownIncorrectValueWith(input_val, 0); - input_arr_1[i] = - ConvertAndReplaceKnownIncorrectValueWith(input_val >> 16, 0); - } - } - - protected: - using typename ExhaustiveBinaryTest::NativeT; - using ExhaustiveBinaryTest::ConvertAndReplaceKnownIncorrectValueWith; -}; - -using ExhaustiveF16BinaryTest = Exhaustive16BitBinaryTest; -using ExhaustiveBF16BinaryTest = Exhaustive16BitBinaryTest; - -// Returns a wrapper of the given build method, which build an HLO operation -// with an empty broadcast dimension. -inline std::function AddEmptyBroadcastDimension( - std::function)> build_method) { - return [&](XlaOp src0, XlaOp src1) -> XlaOp { - return build_method(src0, src1, {}); - }; -} - -#if defined(BINARY_TEST_TARGET_F16) && defined(BINARY_TEST_TARGET_BF16) -#error "Can't define both BINARY_TEST_TARGET_F16 and BINARY_TEST_TARGET_BF16" -#endif - -#if defined(BINARY_TEST_TARGET_F16) && \ - !defined(XLA_BACKEND_DOES_NOT_SUPPORT_FLOAT16) -#define BINARY_TEST_16BIT(test_name, ...) \ - XLA_TEST_P(ExhaustiveF16BinaryTest, test_name) \ - __VA_ARGS__ -#elif defined(BINARY_TEST_TARGET_BF16) && defined(XLA_BACKEND_SUPPORTS_BFLOAT16) -#define BINARY_TEST_16BIT(test_name, ...) \ - XLA_TEST_P(ExhaustiveBF16BinaryTest, test_name) \ - __VA_ARGS__ -#else -#define BINARY_TEST_16BIT(test_name, ...) -#endif - -BINARY_TEST_16BIT(Add, { - auto host_add = [](float x, float y) { return x + y; }; - Run(AddEmptyBroadcastDimension(Add), host_add); -}) - -BINARY_TEST_16BIT(Sub, { - auto host_sub = [](float x, float y) { return x - y; }; - Run(AddEmptyBroadcastDimension(Sub), host_sub); -}) - -// TODO(bixia): Mul fails with bfloat16 on CPU. -BINARY_TEST_16BIT(DISABLED_ON_CPU(Mul), { - auto host_mul = [](float x, float y) { return x * y; }; - Run(AddEmptyBroadcastDimension(Mul), host_mul); -}) - -// TODO(bixia): Div fails with bfloat16 on CPU. -BINARY_TEST_16BIT(DISABLED_ON_CPU(Div), { - auto host_div = [](float x, float y) { return x / y; }; - Run(AddEmptyBroadcastDimension(Div), host_div); -}) - -template ::value || - std::is_same::value>::type* = nullptr> -T ReferenceMax(T x, T y) { - // We need to propagate NAN here because std::max may not propagate NAN. - if (std::fpclassify(x) == FP_NAN) { - return x; - } - if (std::fpclassify(y) == FP_NAN) { - return y; - } - - return std::max(x, y); -} - -template ::value || - std::is_same::value>::type* = nullptr> -T ReferenceMin(T x, T y) { - // We need to propagate NAN here because std::max may not propagate NAN. - if (std::fpclassify(x) == FP_NAN) { - return x; - } - if (std::fpclassify(y) == FP_NAN) { - return y; - } - - return std::min(x, y); -} - -BINARY_TEST_16BIT(Max, { - Run(AddEmptyBroadcastDimension(Max), ReferenceMax); -}) - -BINARY_TEST_16BIT(Min, { - Run(AddEmptyBroadcastDimension(Min), ReferenceMin); -}) - -// TODO(bixia): Pow fails with bfloat16 on CPU. -BINARY_TEST_16BIT(DISABLED_ON_CPU(Pow), - { Run(AddEmptyBroadcastDimension(Pow), std::powf); }) - -// TODO(bixia): Atan2 fails with bfloat16 on CPU. -BINARY_TEST_16BIT(DISABLED_ON_CPU(Atan2), - { Run(AddEmptyBroadcastDimension(Atan2), std::atan2f); }) - -#if defined(BINARY_TEST_TARGET_F16) -#if !defined(XLA_BACKEND_DOES_NOT_SUPPORT_FLOAT16) -INSTANTIATE_TEST_SUITE_P(F16, ExhaustiveF16BinaryTest, - ::testing::ValuesIn(CreateExhaustiveF32Ranges())); -#endif -#endif - -#if defined(BINARY_TEST_TARGET_BF16) -#if defined(XLA_BACKEND_SUPPORTS_BFLOAT16) -INSTANTIATE_TEST_SUITE_P(BF16, ExhaustiveBF16BinaryTest, - ::testing::ValuesIn(CreateExhaustiveF32Ranges())); -#endif -#endif - // Exhaustive test for binary operations for float and double. // // Test parameter is a tuple of (FpValues, FpValues) describing the possible @@ -236,20 +71,10 @@ class Exhaustive32BitOrMoreBinaryTest }; using ExhaustiveF32BinaryTest = Exhaustive32BitOrMoreBinaryTest; -GTEST_ALLOW_UNINSTANTIATED_PARAMETERIZED_TEST( - ExhaustiveF32BinaryTest); // TODO(b/139702016) go/are-your-tests-running -using ExhaustiveF64BinaryTest = Exhaustive32BitOrMoreBinaryTest; -GTEST_ALLOW_UNINSTANTIATED_PARAMETERIZED_TEST( - ExhaustiveF64BinaryTest); // TODO(b/139702016) go/are-your-tests-running - -#if defined(BINARY_TEST_TARGET_F32) #define BINARY_TEST_FLOAT_32(test_name, ...) \ XLA_TEST_P(ExhaustiveF32BinaryTest, test_name) \ __VA_ARGS__ -#else -#define BINARY_TEST_FLOAT_32(test_name, ...) -#endif BINARY_TEST_FLOAT_32(Add, { auto host_add = [](float x, float y) { return x + y; }; @@ -332,8 +157,8 @@ INSTANTIATE_TEST_SUITE_P( ::testing::ValuesIn( GetFpValuesForMagnitudeExtremeNormals(40000, 2000)))); -#if defined(BINARY_TEST_TARGET_F64) && \ - !defined(XLA_BACKEND_DOES_NOT_SUPPORT_FLOAT64) +#if !defined(XLA_BACKEND_DOES_NOT_SUPPORT_FLOAT64) +using ExhaustiveF64BinaryTest = Exhaustive32BitOrMoreBinaryTest; #define BINARY_TEST_FLOAT_64(test_name, ...) \ XLA_TEST_P(ExhaustiveF64BinaryTest, test_name) \ __VA_ARGS__ @@ -381,6 +206,7 @@ BINARY_TEST_FLOAT_64(DISABLED_ON_CPU(AbsComplex), { Run(device_abs_complex, host_abs_complex); }) +#if !defined(XLA_BACKEND_DOES_NOT_SUPPORT_FLOAT64) INSTANTIATE_TEST_SUITE_P( SpecialValues, ExhaustiveF64BinaryTest, ::testing::Combine( @@ -414,6 +240,7 @@ INSTANTIATE_TEST_SUITE_P( GetFpValuesForMagnitudeExtremeNormals(40000, 2000)), ::testing::ValuesIn( GetFpValuesForMagnitudeExtremeNormals(40000, 2000)))); +#endif // !defined(XLA_BACKEND_DOES_NOT_SUPPORT_FLOAT64) } // namespace } // namespace exhaustive_op_test diff --git a/tensorflow/compiler/xla/tests/exhaustive_op_test_utils.h b/tensorflow/compiler/xla/tests/exhaustive_op_test_utils.h index e86a3ca55ac..f0a4687eac7 100644 --- a/tensorflow/compiler/xla/tests/exhaustive_op_test_utils.h +++ b/tensorflow/compiler/xla/tests/exhaustive_op_test_utils.h @@ -1004,6 +1004,45 @@ typename ErrorSpecGenWrapper::type GetDefaultSpecGenerator() { return DefaultSpecGenerator; } +template ::value || + std::is_same::value>::type* = nullptr> +T ReferenceMax(T x, T y) { + // We need to propagate NAN here because std::max may not propagate NAN. + if (std::fpclassify(x) == FP_NAN) { + return x; + } + if (std::fpclassify(y) == FP_NAN) { + return y; + } + + return std::max(x, y); +} + +template ::value || + std::is_same::value>::type* = nullptr> +T ReferenceMin(T x, T y) { + // We need to propagate NAN here because std::max may not propagate NAN. + if (std::fpclassify(x) == FP_NAN) { + return x; + } + if (std::fpclassify(y) == FP_NAN) { + return y; + } + + return std::min(x, y); +} + +// Returns a wrapper of the given build method, which build an HLO operation +// with an empty broadcast dimension. +inline std::function AddEmptyBroadcastDimension( + std::function)> build_method) { + return [&](XlaOp src0, XlaOp src1) -> XlaOp { + return build_method(src0, src1, {}); + }; +} + template class ExhaustiveUnaryTest : public ExhaustiveOpTestBase { public: @@ -1013,6 +1052,9 @@ class ExhaustiveUnaryTest : public ExhaustiveOpTestBase { } }; +template +using ExhaustiveBinaryTest = ExhaustiveOpTestBase; + } // namespace exhaustive_op_test } // namespace xla #endif // TENSORFLOW_COMPILER_XLA_TESTS_EXHAUSTIVE_OP_TEST_UTILS_H_