aboutsummaryrefslogtreecommitdiffhomepage
diff options
context:
space:
mode:
Diffstat
-rw-r--r--tensorflow/compiler/tests/BUILD2
-rw-r--r--tensorflow/compiler/tests/binary_ops_test.py106
-rw-r--r--tensorflow/compiler/tests/xla_test.py11
-rw-r--r--tensorflow/compiler/xla/service/hlo_instruction.cc1
-rw-r--r--tensorflow/compiler/xla/service/hlo_verifier.cc7
-rw-r--r--tensorflow/compiler/xla/service/shape_inference.cc5
-rw-r--r--tensorflow/compiler/xla/tests/dynamic_ops_test.cc8
7 files changed, 121 insertions, 19 deletions
diff --git a/tensorflow/compiler/tests/BUILD b/tensorflow/compiler/tests/BUILD
index 4143aa1f80..85a2adab28 100644
--- a/tensorflow/compiler/tests/BUILD
+++ b/tensorflow/compiler/tests/BUILD
@@ -98,7 +98,7 @@ tf_xla_py_test(
tf_xla_py_test(
name = "binary_ops_test",
- size = "small",
+ size = "medium",
srcs = ["binary_ops_test.py"],
shard_count = 5,
tags = [
diff --git a/tensorflow/compiler/tests/binary_ops_test.py b/tensorflow/compiler/tests/binary_ops_test.py
index 6bcfed7b69..ba7b9bacd2 100644
--- a/tensorflow/compiler/tests/binary_ops_test.py
+++ b/tensorflow/compiler/tests/binary_ops_test.py
@@ -232,11 +232,16 @@ class BinaryOpsTest(XLATestCase):
expected=np.right_shift(lhs, rhs))
if dtype in [np.int8, np.int16, np.int32, np.int64]:
- lhs = np.array([-1, -5, -3, -14], dtype=dtype)
- rhs = np.array([5, 0, 1, 11], dtype=dtype)
- self._testBinary(
- bitwise_ops.right_shift, lhs, rhs,
- expected=np.right_shift(lhs, rhs))
+ lhs = np.array([-1, -5, -3, -14, -2], dtype=dtype)
+ rhs = np.array([5, 0, 1, 11, 36], dtype=dtype)
+ # HLO has saturating shift behavior.
+ bits = np.ceil(
+ np.log(np.iinfo(dtype).max - np.iinfo(dtype).min) / np.log(2))
+ expected = [
+ np.right_shift(l, r) if r < bits else np.sign(l)
+ for l, r in zip(lhs, rhs)
+ ]
+ self._testBinary(bitwise_ops.right_shift, lhs, rhs, expected=expected)
def testNumericOps(self):
for dtype in self.numeric_types:
@@ -255,12 +260,18 @@ class BinaryOpsTest(XLATestCase):
np.array([[1], [2]], dtype=dtype),
dtype(7),
expected=np.array([[8], [9]], dtype=dtype))
+ self._testBinary(
+ math_ops.add,
+ np.array([0xffffffff, 0xfffffffff, 1, 1], dtype=np.int64),
+ np.array([1, 1, 0xffffffff, 0xfffffffff], dtype=np.int64),
+ expected=np.array(
+ [1 << 32, 1 << 36, 1 << 32, 1 << 36], dtype=np.int64))
self._testBinary(
math_ops.subtract,
- np.array([1, 2], dtype=dtype),
- np.array([10, 20], dtype=dtype),
- expected=np.array([-9, -18], dtype=dtype))
+ np.array([1, 2, 100], dtype=dtype),
+ np.array([10, 20, -1], dtype=dtype),
+ expected=np.array([-9, -18, 101], dtype=dtype))
self._testBinary(
math_ops.subtract,
dtype(5),
@@ -668,6 +679,11 @@ class BinaryOpsTest(XLATestCase):
np.array([[10], [7], [2]], dtype=np.float32),
np.float32(7),
expected=np.array([[False], [False], [True]], dtype=np.bool))
+ self._testBinary(
+ less_op,
+ np.array([[10], [7], [2], [-1]], dtype=np.int64),
+ np.int64(7),
+ expected=np.array([[False], [False], [True], [True]], dtype=np.bool))
for less_equal_op in [math_ops.less_equal, (lambda x, y: x <= y)]:
self._testBinary(
@@ -686,6 +702,80 @@ class BinaryOpsTest(XLATestCase):
np.float32(7),
expected=np.array([[False], [True], [True]], dtype=np.bool))
+ def testS64Comparisons(self):
+ for op in [(lambda x, y: x < y), (lambda x, y: x <= y),
+ (lambda x, y: x >= y), (lambda x, y: x > y)]:
+ lhs = np.array(
+ [
+ np.int64(0x000000007FFFFFFF),
+ np.int64(0x000000007FFFFFFF),
+ np.int64(0x0000000080000000),
+ np.int64(0x0000000080000000),
+ np.int64(0x0000000080000001),
+ np.int64(0x00000000FFFF0000),
+ np.int64(0x00000000FFFF0000),
+ np.int64(0x00000000FFFFFFFE),
+ np.int64(0x00000000FFFFFFFF),
+ np.int64(0x00000000FFFFFFFF),
+ np.int64(0x0000000100000000),
+ np.int64(0x0000000200000002),
+ np.int64(0x0000000200000002),
+ np.int64(0x0000000200000002),
+ np.int64(0x0000000200000002),
+ np.int64(0x0000000200000002),
+ np.int64(0x0000000200000002),
+ np.int64(0x0000000200000002),
+ np.int64(0x0000000200000002),
+ np.int64(0x0000000200000002),
+ np.int64(-0x7FFFFFFF00000002),
+ np.int64(-0x7FFFFFFF00000002),
+ np.int64(-0x7FFFFFFF00000001),
+ np.int64(-0x7FFFFFFF00000001),
+ np.int64(-0x7FFFFFFF00000001),
+ np.int64(-0x7FFFFFFF00000001),
+ np.int64(0x7ffffffefff00010),
+ np.int64(0x7ffffffefff00010),
+ np.int64(-1),
+ np.int64(-1)
+ ],
+ dtype=np.int64)
+ rhs = np.array(
+ [
+ np.int64(0x000000007FFFFFFE),
+ np.int64(0x000000007FFFFFFF),
+ np.int64(0x000000007FFFFFFF),
+ np.int64(0x0000000080000000),
+ np.int64(0x0000000080000001),
+ np.int64(0x00000000FFFF0000),
+ np.int64(0x00000000FFFF0001),
+ np.int64(0x00000000FFFFFFFF),
+ np.int64(0x00000000FFFFFFFE),
+ np.int64(0x00000000FFFFFFFF),
+ np.int64(0x00000000FFFFFFFF),
+ np.int64(0x0000000100000001),
+ np.int64(0x0000000100000002),
+ np.int64(0x0000000100000003),
+ np.int64(0x0000000200000001),
+ np.int64(0x0000000200000002),
+ np.int64(0x0000000200000003),
+ np.int64(0x0000000300000001),
+ np.int64(0x0000000300000002),
+ np.int64(0x0000000300000003),
+ np.int64(0x00000000FFFFFFFF),
+ np.int64(-0x7FFFFFFF00000001),
+ np.int64(0x00000000FFFFFFFE),
+ np.int64(0x00000000FFFFFFFF),
+ np.int64(-0x7FFFFFFF00000002),
+ np.int64(-0x7FFFFFFF00000001),
+ np.int64(0x00000000FFFFFFFF),
+ np.int64(-0x7FFFFFFF00000001),
+ np.int64(-2),
+ np.int64(-1)
+ ],
+ dtype=np.int64)
+ expected = np.array([op(l, r) for l, r in zip(lhs, rhs)], dtype=np.bool)
+ self._testBinary(op, lhs, rhs, expected=expected)
+
def testBroadcasting(self):
"""Tests broadcasting behavior of an operator."""
diff --git a/tensorflow/compiler/tests/xla_test.py b/tensorflow/compiler/tests/xla_test.py
index 7e1f5c76ed..cc778f1c3c 100644
--- a/tensorflow/compiler/tests/xla_test.py
+++ b/tensorflow/compiler/tests/xla_test.py
@@ -71,14 +71,14 @@ class XLATestCase(test.TestCase):
self._all_types = set(
[dtype.as_numpy_dtype for dtype in self._all_tf_types])
- self.int_types = set([dtype.as_numpy_dtype for dtype in self.int_tf_types])
+ self._int_types = set([dtype.as_numpy_dtype for dtype in self.int_tf_types])
self._float_types = set(
[dtype.as_numpy_dtype for dtype in self._float_tf_types])
self.complex_types = set([
dtype.as_numpy_dtype for dtype in self.complex_tf_types
])
- self._numeric_types = set(
- self.int_types | self._float_types | self.complex_types)
+ self._numeric_types = set(self._int_types | self._float_types
+ | self.complex_types)
# Parse the manifest file, if any, into a regex identifying tests to
# disable
@@ -131,6 +131,11 @@ class XLATestCase(test.TestCase):
return self._float_tf_types - self._method_types_filter.get(name, set())
@property
+ def int_types(self):
+ name = '{}.{}'.format(type(self).__name__, self._testMethodName)
+ return self._int_types - self._method_types_filter.get(name, set())
+
+ @property
def numeric_tf_types(self):
name = '{}.{}'.format(type(self).__name__, self._testMethodName)
tf_types = set([dtypes.as_dtype(t)
diff --git a/tensorflow/compiler/xla/service/hlo_instruction.cc b/tensorflow/compiler/xla/service/hlo_instruction.cc
index af9d772b00..d33add23d0 100644
--- a/tensorflow/compiler/xla/service/hlo_instruction.cc
+++ b/tensorflow/compiler/xla/service/hlo_instruction.cc
@@ -182,6 +182,7 @@ StatusOr<std::unique_ptr<HloInstruction>> HloInstruction::CreateFromProto(
/* static */ std::unique_ptr<HloInstruction>
HloInstruction::CreateGetTupleElement(const Shape& shape,
HloInstruction* operand, int64 index) {
+ CHECK(ShapeUtil::IsTuple(operand->shape()));
auto instruction =
WrapUnique(new HloInstruction(HloOpcode::kGetTupleElement, shape));
instruction->tuple_index_ = index;
diff --git a/tensorflow/compiler/xla/service/hlo_verifier.cc b/tensorflow/compiler/xla/service/hlo_verifier.cc
index b1fd068115..8c875698eb 100644
--- a/tensorflow/compiler/xla/service/hlo_verifier.cc
+++ b/tensorflow/compiler/xla/service/hlo_verifier.cc
@@ -762,11 +762,14 @@ StatusOr<bool> HloVerifier::Run(HloModule* module) {
} else if (instruction->opcode() == HloOpcode::kBroadcast) {
// If you see this failure then someone has confused the difference
// between the HLO broadcast op, and the UserComputation broadcast
- // op. See https://groups.google.com/forum/#!topic/xla-dev/9LqijHmTt_I
+ // op. See https://groups.google.com/forum/#!topic/xla-dev/9LqijHmTt_I
// or ComputationLowerer::Visit()
TF_RET_CHECK(instruction->dimensions().size() ==
ShapeUtil::Rank(instruction->operand(0)->shape()))
- << "Broadcast HLO has invalid number of dimensions.";
+ << "Broadcast HLO (" << instruction->ToShortString()
+ << ") has invalid number of dimensions: "
+ << instruction->dimensions().size()
+ << " != " << ShapeUtil::Rank(instruction->operand(0)->shape());
} else if (instruction->opcode() == HloOpcode::kWhile) {
auto* while_cond = instruction->while_condition();
auto* while_body = instruction->while_body();
diff --git a/tensorflow/compiler/xla/service/shape_inference.cc b/tensorflow/compiler/xla/service/shape_inference.cc
index 2ff7ae97b7..74f744a62b 100644
--- a/tensorflow/compiler/xla/service/shape_inference.cc
+++ b/tensorflow/compiler/xla/service/shape_inference.cc
@@ -193,7 +193,10 @@ tensorflow::Status VerifyReducerShape(const ProgramShape& reducer_shape,
const Shape& accumulator_shape = reducer_shape.result();
if (ShapeUtil::Rank(accumulator_shape) != 0) {
- return InvalidArgument("Reduction function must have rank 0.");
+ return InvalidArgument(
+ "Reduction function must have rank 0 (rank %lld reduction function "
+ "given).",
+ ShapeUtil::Rank(accumulator_shape));
}
// Check that the accumulator can be passed in as the first argument.
diff --git a/tensorflow/compiler/xla/tests/dynamic_ops_test.cc b/tensorflow/compiler/xla/tests/dynamic_ops_test.cc
index 877dc7db0e..4f354e6aef 100644
--- a/tensorflow/compiler/xla/tests/dynamic_ops_test.cc
+++ b/tensorflow/compiler/xla/tests/dynamic_ops_test.cc
@@ -206,19 +206,19 @@ XLA_TEST_F(DynamicSliceTest, Int32R1BF16) { TestR1<int32, bfloat16>(); }
XLA_TEST_F(DynamicSliceTest, Int32R1) { TestR1<int32, int32>(); }
XLA_TEST_F(DynamicSliceTest, Int32R1Wrap) { TestR1Wrap<int32, int32>(); }
XLA_TEST_F(DynamicSliceTest, Int64R1) { TestR1<int64, float>(); }
-XLA_TEST_F(DynamicSliceTest, UInt64R1) { TestR1<uint64, double>(); }
+XLA_TEST_F(DynamicSliceTest, UInt64R1) { TestR1<uint64, float>(); }
XLA_TEST_F(DynamicSliceTest, Int32R2BF16) { TestR2<int32, bfloat16>(); }
XLA_TEST_F(DynamicSliceTest, Int32R2) { TestR2<int32, int32>(); }
XLA_TEST_F(DynamicSliceTest, Int32R2Wrap) { TestR2Wrap<int32, int32>(); }
-XLA_TEST_F(DynamicSliceTest, Int64R2) { TestR2<int64, double>(); }
+XLA_TEST_F(DynamicSliceTest, Int64R2) { TestR2<int64, float>(); }
XLA_TEST_F(DynamicSliceTest, UInt64R2) { TestR2<uint64, int32>(); }
XLA_TEST_F(DynamicSliceTest, Int32R3BF16) { TestR3<int32, bfloat16>(); }
XLA_TEST_F(DynamicSliceTest, Int32R3) { TestR3<int32, float>(); }
XLA_TEST_F(DynamicSliceTest, Int32R3Wrap) { TestR3Wrap<int32, float>(); }
XLA_TEST_F(DynamicSliceTest, Int64R3) { TestR3<int64, float>(); }
-XLA_TEST_F(DynamicSliceTest, UInt64R3) { TestR3<uint64, double>(); }
+XLA_TEST_F(DynamicSliceTest, UInt64R3) { TestR3<uint64, float>(); }
XLA_TEST_F(DynamicSliceTest, Int32R1Pred) {
// Slice at dimension start.
@@ -506,7 +506,7 @@ XLA_TEST_F(DynamicUpdateSliceTest, DISABLED_ON_CPU_PARALLEL(Int32R1BF16)) {
}
XLA_TEST_F(DynamicUpdateSliceTest, Int32R1) { TestR1<int32, float>(); }
XLA_TEST_F(DynamicUpdateSliceTest, Int64R1) { TestR1<int64, float>(); }
-XLA_TEST_F(DynamicUpdateSliceTest, UInt64R1) { TestR1<uint64, double>(); }
+XLA_TEST_F(DynamicUpdateSliceTest, UInt64R1) { TestR1<uint64, float>(); }
// TODO(b/71820067): The CPU parallel backend failed for this on 2018-01-10.
XLA_TEST_F(DynamicUpdateSliceTest, DISABLED_ON_CPU_PARALLEL(Int32R2BF16)) {
generated by cgit v1.2.3 (git 2.39.1) at 2024-06-26 06:09:51 +0000