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 "mlir/IR/Attributes.h" // from @llvm-project
22 #include "mlir/IR/Builders.h" // from @llvm-project
23 #include "mlir/IR/BuiltinTypes.h" // from @llvm-project
24 #include "tensorflow/compiler/mlir/hlo/include/mlir-hlo/Dialect/mhlo/IR/hlo_ops.h"
25 #include "tensorflow/compiler/mlir/hlo/include/mlir-hlo/utils/convert_op_folder.h"
26 #include "tensorflow/compiler/xla/service/hlo_instruction.h"
27
28 namespace xla {
29
30 StatusOr<mlir::DenseElementsAttr> CreateDenseElementsAttrFromLiteral(
31 const LiteralBase& literal, mlir::Builder builder);
32
33 Status CopyDenseElementsDataToXlaFormat(mlir::DenseElementsAttr data,
34 std::vector<uint8>* output);
35
36 StatusOr<int> GetElementTypeBytes(mlir::Type type);
37
38 // Creates an DenseIntElementsAttr using the elements of the vector and the
39 // optional shape.
40 mlir::DenseIntElementsAttr CreateDenseIntElementsAttrFromVector(
41 const llvm::ArrayRef<int64> vector, mlir::Builder builder,
42 llvm::ArrayRef<int64_t> shape = {});
43
44 StatusOr<mlir::Type> ConvertPrimitiveTypeToMLIRType(PrimitiveType element_type,
45 mlir::Builder builder);
46
47 mlir::mhlo::GatherDimensionNumbers CreateGatherDimensionNumbers(
48 const GatherDimensionNumbers& input, mlir::Builder builder);
49
50 template <typename TypeT>
ConvertTensorShapeToType(const Shape & shape,mlir::Builder builder)51 static StatusOr<TypeT> ConvertTensorShapeToType(const Shape& shape,
52 mlir::Builder builder) {
53 auto element_type_or =
54 ConvertPrimitiveTypeToMLIRType(shape.element_type(), builder);
55 if (!element_type_or.ok()) return element_type_or.status();
56
57 auto dimensions = shape.dimensions();
58 llvm::SmallVector<int64_t, 4> array(dimensions.begin(), dimensions.end());
59 return TypeT::get(array, element_type_or.ValueOrDie());
60 }
61
62 StatusOr<mlir::MemRefType> ConvertTensorShapeToMemRefType(
63 const Shape& shape, mlir::Builder builder);
64
65 template <>
ConvertTensorShapeToType(const Shape & shape,mlir::Builder builder)66 inline StatusOr<mlir::MemRefType> ConvertTensorShapeToType(
67 const Shape& shape, mlir::Builder builder) {
68 return ConvertTensorShapeToMemRefType(shape, builder);
69 }
70
71 template <typename TypeT>
ConvertShapeToType(const Shape & shape,mlir::Builder builder)72 static StatusOr<mlir::Type> ConvertShapeToType(const Shape& shape,
73 mlir::Builder builder) {
74 if (shape.IsTuple()) {
75 llvm::SmallVector<mlir::Type, 4> contents;
76 contents.reserve(shape.tuple_shapes_size());
77 for (const auto& subtype : shape.tuple_shapes()) {
78 TF_ASSIGN_OR_RETURN(auto mlir_subtype,
79 ConvertShapeToType<TypeT>(subtype, builder));
80 contents.push_back(mlir_subtype);
81 }
82 return builder.getTupleType(contents);
83 }
84 if (shape.IsToken()) {
85 return mlir::mhlo::TokenType::get(builder.getContext());
86 }
87 return ConvertTensorShapeToType<TypeT>(shape, builder);
88 }
89
90 ::xla::StatusOr<::xla::HloOpcode> MhloToHloOpcode(mlir::Operation* op);
91
92 } // namespace xla
93
94 #endif // TENSORFLOW_COMPILER_MLIR_XLA_UTILS_H_
95