1 /* Copyright 2019 The TensorFlow Authors. All Rights Reserved.
2
3 Licensed under the Apache License, Version 2.0 (the "License");
4 you may not use this file except in compliance with the License.
5 You may obtain a copy of the License at
6
7 http://www.apache.org/licenses/LICENSE-2.0
8
9 Unless required by applicable law or agreed to in writing, software
10 distributed under the License is distributed on an "AS IS" BASIS,
11 WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12 See the License for the specific language governing permissions and
13 limitations under the License.
14 ==============================================================================*/
15
16 // This file defines helpers useful when creating or manipulating lhlo/hlo.
17
18 #ifndef TENSORFLOW_COMPILER_MLIR_XLA_UTILS_H_
19 #define TENSORFLOW_COMPILER_MLIR_XLA_UTILS_H_
20
21 #include "llvm/ADT/STLExtras.h"
22 #include "mlir/IR/Attributes.h" // from @llvm-project
23 #include "mlir/IR/Builders.h" // from @llvm-project
24 #include "mlir/IR/BuiltinTypes.h" // from @llvm-project
25 #include "tensorflow/compiler/xla/mlir_hlo/include/mlir-hlo/Dialect/mhlo/IR/hlo_ops.h"
26 #include "tensorflow/compiler/xla/mlir_hlo/include/mlir-hlo/utils/convert_op_folder.h"
27 #include "tensorflow/compiler/xla/service/hlo_instruction.h"
28 #include "tensorflow/core/platform/errors.h"
29
30 namespace xla {
31
32 StatusOr<mlir::DenseElementsAttr> CreateDenseElementsAttrFromLiteral(
33 const LiteralBase& literal, mlir::Builder builder);
34
35 Status CopyDenseElementsDataToXlaFormat(mlir::DenseElementsAttr data,
36 std::vector<uint8_t>* output);
37
38 StatusOr<int> GetElementTypeBytes(mlir::Type type);
39
40 // Creates an DenseIntElementsAttr using the elements of the vector and the
41 // optional shape.
42 mlir::DenseIntElementsAttr CreateDenseIntElementsAttrFromVector(
43 const llvm::ArrayRef<int64_t> vector, mlir::Builder builder,
44 llvm::ArrayRef<int64_t> shape = {});
45
46 StatusOr<mlir::Type> ConvertPrimitiveTypeToMLIRType(PrimitiveType element_type,
47 mlir::Builder builder);
48
49 mlir::mhlo::GatherDimensionNumbersAttr CreateGatherDimensionNumbers(
50 const GatherDimensionNumbers& input, mlir::Builder builder);
51
52 // Converts the given XLA shape for tensors to the template MLIR type.
53 template <typename TypeT>
ConvertTensorShapeToType(const Shape & xla_ty,mlir::Builder builder)54 static StatusOr<TypeT> ConvertTensorShapeToType(const Shape& xla_ty,
55 mlir::Builder builder) {
56 auto element_type_or =
57 ConvertPrimitiveTypeToMLIRType(xla_ty.element_type(), builder);
58 if (!element_type_or.ok()) return element_type_or.status();
59
60 bool is_dynamic = false;
61 int64_t rank = xla_ty.rank();
62 llvm::SmallVector<int64_t, 4> shape(rank, mlir::ShapedType::kDynamicSize);
63 llvm::SmallVector<int64_t, 4> bounds(rank, mlir::ShapedType::kDynamicSize);
64 for (int64_t dim = 0; dim < rank; ++dim) {
65 int64_t dim_size = xla_ty.dimensions(dim);
66 if (xla_ty.is_dynamic_dimension(dim)) {
67 bounds[dim] = dim_size;
68 is_dynamic = true;
69 } else {
70 shape[dim] = dim_size;
71 }
72 }
73 using mlir::mhlo::TypeExtensionsAttr;
74 TypeExtensionsAttr extensions;
75 if (is_dynamic) {
76 extensions = TypeExtensionsAttr::get(builder.getContext(), bounds);
77 }
78 return TypeT::get(shape, element_type_or.ValueOrDie(), extensions);
79 }
80
81 StatusOr<mlir::MemRefType> ConvertTensorShapeToMemRefType(
82 const Shape& shape, mlir::Builder builder);
83
84 template <>
ConvertTensorShapeToType(const Shape & shape,mlir::Builder builder)85 inline StatusOr<mlir::MemRefType> ConvertTensorShapeToType(
86 const Shape& shape, mlir::Builder builder) {
87 if (shape.is_dynamic()) {
88 return tensorflow::errors::FailedPrecondition(
89 "MemRefType don't support dynamic shapes");
90 }
91 return ConvertTensorShapeToMemRefType(shape, builder);
92 }
93
94 // Converts the given XLA shape to the template MLIR type.
95 template <typename TypeT>
ConvertShapeToType(const Shape & shape,mlir::Builder builder)96 static StatusOr<mlir::Type> ConvertShapeToType(const Shape& shape,
97 mlir::Builder builder) {
98 if (shape.IsTuple()) {
99 llvm::SmallVector<mlir::Type, 4> contents;
100 contents.reserve(shape.tuple_shapes_size());
101 for (const auto& subtype : shape.tuple_shapes()) {
102 TF_ASSIGN_OR_RETURN(auto mlir_subtype,
103 ConvertShapeToType<TypeT>(subtype, builder));
104 contents.push_back(mlir_subtype);
105 }
106 return builder.getTupleType(contents);
107 }
108 if (shape.IsToken()) {
109 return mlir::mhlo::TokenType::get(builder.getContext());
110 }
111 return ConvertTensorShapeToType<TypeT>(shape, builder);
112 }
113
114 ::xla::StatusOr<::xla::HloOpcode> MhloToHloOpcode(mlir::Operation* op);
115
116 } // namespace xla
117
118 #endif // TENSORFLOW_COMPILER_MLIR_XLA_UTILS_H_
119