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/mlir/hlo/include/mlir-hlo/Dialect/mhlo/IR/hlo_ops.h"
26 #include "tensorflow/compiler/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>* 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> 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::GatherDimensionNumbers CreateGatherDimensionNumbers(
50 const GatherDimensionNumbers& input, mlir::Builder builder);
51
52 // Converts the given XLA shape for tensors to the template MLIR type. Note that
53 // any dynamic bounds in the input shape is lost and those dimensions are fully
54 // dynamic in the MLIR type.
55 template <typename TypeT>
ConvertTensorShapeToType(const Shape & shape,mlir::Builder builder)56 static StatusOr<TypeT> ConvertTensorShapeToType(const Shape& shape,
57 mlir::Builder builder) {
58 auto element_type_or =
59 ConvertPrimitiveTypeToMLIRType(shape.element_type(), builder);
60 if (!element_type_or.ok()) return element_type_or.status();
61
62 auto dimensions = shape.dimensions();
63 llvm::SmallVector<int64_t, 4> array(dimensions.begin(), dimensions.end());
64 for (auto element : llvm::enumerate(shape.dynamic_dimensions())) {
65 bool is_dynamic = element.value();
66 if (is_dynamic) {
67 array[element.index()] = mlir::ShapedType::kDynamicSize;
68 }
69 }
70 return TypeT::get(array, element_type_or.ValueOrDie());
71 }
72
73 StatusOr<mlir::MemRefType> ConvertTensorShapeToMemRefType(
74 const Shape& shape, mlir::Builder builder);
75
76 template <>
ConvertTensorShapeToType(const Shape & shape,mlir::Builder builder)77 inline StatusOr<mlir::MemRefType> ConvertTensorShapeToType(
78 const Shape& shape, mlir::Builder builder) {
79 if (shape.is_dynamic()) {
80 return tensorflow::errors::FailedPrecondition(
81 "MemRefType don't support dynamic shapes");
82 }
83 return ConvertTensorShapeToMemRefType(shape, builder);
84 }
85
86 // Converts the given XLA shape to the template MLIR type. Note that any
87 // dynamic bounds in the input shape is lost and those dimensions are fully
88 // dynamic in the MLIR type.
89 template <typename TypeT>
ConvertShapeToType(const Shape & shape,mlir::Builder builder)90 static StatusOr<mlir::Type> ConvertShapeToType(const Shape& shape,
91 mlir::Builder builder) {
92 if (shape.IsTuple()) {
93 llvm::SmallVector<mlir::Type, 4> contents;
94 contents.reserve(shape.tuple_shapes_size());
95 for (const auto& subtype : shape.tuple_shapes()) {
96 TF_ASSIGN_OR_RETURN(auto mlir_subtype,
97 ConvertShapeToType<TypeT>(subtype, builder));
98 contents.push_back(mlir_subtype);
99 }
100 return builder.getTupleType(contents);
101 }
102 if (shape.IsToken()) {
103 return mlir::mhlo::TokenType::get(builder.getContext());
104 }
105 return ConvertTensorShapeToType<TypeT>(shape, builder);
106 }
107
108 ::xla::StatusOr<::xla::HloOpcode> MhloToHloOpcode(mlir::Operation* op);
109
110 } // namespace xla
111
112 #endif // TENSORFLOW_COMPILER_MLIR_XLA_UTILS_H_
113