diff options
Diffstat (limited to 'tensorflow/compiler/xla/shape_util.cc')
-rw-r--r-- | tensorflow/compiler/xla/shape_util.cc | 140 |
1 files changed, 90 insertions, 50 deletions
diff --git a/tensorflow/compiler/xla/shape_util.cc b/tensorflow/compiler/xla/shape_util.cc index 2166c34358..ec901af1e2 100644 --- a/tensorflow/compiler/xla/shape_util.cc +++ b/tensorflow/compiler/xla/shape_util.cc @@ -46,28 +46,14 @@ namespace xla { using ::tensorflow::strings::StrAppend; using ::tensorflow::strings::StrCat; -string ShapeIndex::ToString() const { - return StrCat("{", tensorflow::str_util::Join(indices_, ","), "}"); -} +string ShapeIndex::ToString() const { return ShapeIndexView(*this).ToString(); } string ShapeIndexView::ToString() const { - return StrCat("{", - tensorflow::str_util::Join( - tensorflow::gtl::make_range(begin_, end_), ","), - "}"); + return StrCat("{", tensorflow::str_util::Join(indices_, ","), "}"); } bool ShapeIndexView::operator==(const ShapeIndexView& other) const { - if (size() != other.size()) { - return false; - } - for (auto it = begin(), other_it = other.begin(); it != end(); - ++it, ++other_it) { - if (*it != *other_it) { - return false; - } - } - return true; + return indices_ == other.indices_; } bool ShapeIndexView::operator!=(const ShapeIndexView& other) const { @@ -696,7 +682,7 @@ StatusOr<Shape> ParseShapeStringInternal(tensorflow::StringPiece* s) { CompatibleIgnoringElementType); } else { // Opaque, token, etc types are vacuously compatible. - return true; + return lhs.element_type() == rhs.element_type(); } } @@ -711,7 +697,7 @@ StatusOr<Shape> ParseShapeStringInternal(tensorflow::StringPiece* s) { CompatibleIgnoringFpPrecision); } else { // Opaque, token, etc types are vacuously compatible. - return true; + return lhs.element_type() == rhs.element_type(); } } @@ -891,44 +877,62 @@ StatusOr<Shape> ParseShapeStringInternal(tensorflow::StringPiece* s) { /* static */ Status ShapeUtil::ValidateShapeSize(const Shape& shape) { VLOG(3) << "Validating shape size: " << ShapeUtil::HumanString(shape); - auto invalid_argument = - InvalidArgument("Shape %s size may overflow int64.", - ShapeUtil::HumanString(shape).c_str()); + if (!IsArray(shape)) { return Status::OK(); } - int64 shape_size; - if (LayoutUtil::IsSparseArray(shape)) { - shape_size = LayoutUtil::MaxSparseElements(shape.layout()); - if (shape_size < 0) { - return invalid_argument; - } - shape_size = MultiplyWithoutOverflow(shape_size, ShapeUtil::Rank(shape)); - if (shape_size < 0) { - return invalid_argument; + + int64 shape_size = [&shape]() { + if (LayoutUtil::IsSparseArray(shape)) { + int64 max_sparse_elements = LayoutUtil::MaxSparseElements(shape.layout()); + if (max_sparse_elements < 0) { + return max_sparse_elements; + } + int64 sparse_elements_size = MultiplyWithoutOverflow( + max_sparse_elements, ByteSizeOfPrimitiveType(shape.element_type())); + if (sparse_elements_size < 0) { + return sparse_elements_size; + } + int64 sparse_indices_size = + MultiplyWithoutOverflow(max_sparse_elements, ShapeUtil::Rank(shape)); + if (sparse_indices_size < 0) { + return sparse_indices_size; + } + sparse_indices_size = + MultiplyWithoutOverflow(sparse_indices_size, sizeof(int64)); + if (sparse_indices_size < 0) { + return sparse_indices_size; + } + // At this point, both sparse_indices_size and sparse_elements_size are + // non-negative, so we can easily check if adding them wraps. + if (static_cast<uint64>(sparse_elements_size) + + static_cast<uint64>(sparse_indices_size) > + INT64_MAX) { + return static_cast<int64>(-1); + } } - shape_size = MultiplyWithoutOverflow(shape_size, sizeof(int64)); - if (shape_size < 0) { - return invalid_argument; + + // This is intentionally unconditional: even if the shape is sparse, we want + // to verify the densified version has a reasonable size. + int64 dense_shape_size = 1; + if (shape.dimensions().empty()) { + return dense_shape_size; } - } - // This is intentionally unconditional: even if the shape is sparse, we want - // to verify the densified version has a reasonable size. - if (shape.dimensions().empty()) { - return Status::OK(); - } - shape_size = 1; - for (int64 dim : shape.dimensions()) { - shape_size = MultiplyWithoutOverflow(shape_size, dim); - if (shape_size < 0) { - return invalid_argument; + for (int64 dim : shape.dimensions()) { + dense_shape_size = MultiplyWithoutOverflow(dense_shape_size, dim); + if (dense_shape_size < 0) { + return dense_shape_size; + } } - } - shape_size = MultiplyWithoutOverflow( - shape_size, ByteSizeOfPrimitiveType(shape.element_type())); + dense_shape_size = MultiplyWithoutOverflow( + dense_shape_size, ByteSizeOfPrimitiveType(shape.element_type())); + return dense_shape_size; + }(); + if (shape_size < 0) { - return invalid_argument; + return InvalidArgument("Shape %s size may overflow int64.", + ShapeUtil::HumanString(shape).c_str()); } VLOG(3) << "Shape size is valid: " << shape_size; @@ -1119,12 +1123,41 @@ Status ForEachMutableSubshapeHelper( for (auto dim : Permute(permutation, shape.dimensions())) { new_shape.add_dimensions(dim); } + + // If `shape` has a layout, by contract we choose a new layout such that the + // transpose defined by this permutation is a bitcast. + // + // Some formalism helps to understand the correct way to do this. We're going + // to do algebra in the group of permutations of the dimensions of `shape`. + // + // Since the order of `shape`'s dimensions is not permuted relative to itself, + // `shape`'s list of dimensions is isomorphic to the identity I. + // + // Let `shape`'s layout be L. A layout is a permutation which maps a + // minor-to-major physical layout to the order of a shape's logical dims. + // Therefore inverse of a layout maps from logical to physical dims, and so + // the physical layout of I is simply L'.I = L', where L' is the inverse of L. + // + // Let the argument `permutation` be P. This is a permutation over `shape`'s + // dimensions, so our return value will be a shape with dims P.I = P. Our + // goal is to construct a layout permutation L* that we can apply to P such + // that that the physical dimension ordering of the returned shape is the same + // as that of the original shape, namely L'. + // + // Our returned shape has dims P and layout L*, so its in-memory layout is + // L*'.P. Setting this equal to L' and solving for L*, we get: + // + // L*'.P = L' => + // L*' = L'P' => + // L* = P.L + // if (shape.has_layout()) { CHECK(LayoutUtil::IsDenseArray(shape)); Layout* new_layout = new_shape.mutable_layout(); new_layout->set_format(DENSE); new_layout->clear_minor_to_major(); - for (auto index : Permute(permutation, shape.layout().minor_to_major())) { + for (auto index : ComposePermutations( + permutation, AsInt64Slice(shape.layout().minor_to_major()))) { new_layout->add_minor_to_major(index); } if (shape.layout().padded_dimensions_size() > 0) { @@ -1134,6 +1167,13 @@ Status ForEachMutableSubshapeHelper( new_layout->add_padded_dimensions(dim); } } + // The permutation accepted by TransposeIsBitcast is the inverse of the + // permutation here. + CHECK(TransposeIsBitcast(shape, new_shape, InversePermutation(permutation))) + << "shape=" << HumanStringWithLayout(shape) + << ", new_shape=" << HumanStringWithLayout(new_shape) + << ", permutation={" << tensorflow::str_util::Join(permutation, ",") + << "}"; } return new_shape; } |