1 /* Copyright 2018 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 #include <algorithm>
17
18 #include "tensorflow/compiler/tf2xla/shape_util.h"
19 #include "tensorflow/compiler/tf2xla/xla_op_kernel.h"
20 #include "tensorflow/compiler/tf2xla/xla_op_registry.h"
21 #include "tensorflow/compiler/xla/client/xla_builder.h"
22 #include "tensorflow/core/framework/op_kernel.h"
23
24 #include "tensorflow/compiler/tf2xla/type_util.h"
25 #include "tensorflow/compiler/tf2xla/xla_helpers.h"
26 #include "tensorflow/core/framework/kernel_def_builder.h"
27
28 namespace tensorflow {
29 namespace {
30
SliceVector(xla::XlaOp input,int64 rank)31 absl::InlinedVector<xla::XlaOp, 4> SliceVector(xla::XlaOp input, int64 rank) {
32 absl::InlinedVector<xla::XlaOp, 4> scalar_indices;
33 scalar_indices.reserve(rank);
34 for (int i = 0; i < rank; i++)
35 scalar_indices.push_back(
36 xla::Reshape(xla::Slice(input, {i}, {i + 1}, {1}), {}));
37 return scalar_indices;
38 }
39
40 class DynamicUpdateSliceOp : public XlaOpKernel {
41 public:
DynamicUpdateSliceOp(OpKernelConstruction * context)42 explicit DynamicUpdateSliceOp(OpKernelConstruction* context)
43 : XlaOpKernel(context) {}
44
Compile(XlaOpKernelContext * ctx)45 void Compile(XlaOpKernelContext* ctx) override {
46 DataType index_type = ctx->InputType("indices");
47 CHECK(index_type == DT_INT32 || index_type == DT_INT64);
48
49 const TensorShape input_shape = ctx->InputShape("input");
50 const TensorShape update_shape = ctx->InputShape("update");
51 const TensorShape index_shape = ctx->InputShape("indices");
52
53 int64 rank = input_shape.dims();
54 OP_REQUIRES(
55 ctx,
56 TensorShapeUtils::IsVector(index_shape) &&
57 index_shape.num_elements() == rank,
58 errors::InvalidArgument("index must be a vector with length equal to "
59 "the number of input dimensions"));
60 OP_REQUIRES(
61 ctx, rank == update_shape.dims(),
62 errors::InvalidArgument("input and update must have the same rank,"
63 " input shape is ",
64 input_shape.DebugString(), "; update shape is ",
65 update_shape.DebugString()));
66
67 xla::XlaOp indices = ctx->Input("indices");
68 xla::XlaOp result = xla::DynamicUpdateSlice(
69 ctx->Input("input"), ctx->Input("update"), SliceVector(indices, rank));
70 ctx->SetOutput(0, result);
71 }
72 };
73
74 REGISTER_XLA_OP(Name("XlaDynamicUpdateSlice"), DynamicUpdateSliceOp);
75
76 class DynamicSliceOp : public XlaOpKernel {
77 public:
DynamicSliceOp(OpKernelConstruction * context)78 explicit DynamicSliceOp(OpKernelConstruction* context)
79 : XlaOpKernel(context) {}
80
Compile(XlaOpKernelContext * ctx)81 void Compile(XlaOpKernelContext* ctx) override {
82 DataType index_type = ctx->InputType("start_indices");
83 CHECK(index_type == DT_INT32 || index_type == DT_INT64);
84 CHECK(index_type == ctx->InputType("size_indices"));
85
86 const TensorShape input_shape = ctx->InputShape("input");
87 const TensorShape start_indices_shape = ctx->InputShape("start_indices");
88 const TensorShape size_indices_shape = ctx->InputShape("size_indices");
89
90 int64 rank = input_shape.dims();
91 OP_REQUIRES(ctx,
92 TensorShapeUtils::IsVector(start_indices_shape) &&
93 start_indices_shape.num_elements() == rank,
94 errors::InvalidArgument(
95 "start_indices must be a vector with length equal to "
96 "input rank, but input rank is ",
97 rank, " and start_indices has shape ",
98 start_indices_shape.DebugString()));
99 OP_REQUIRES(ctx,
100 TensorShapeUtils::IsVector(size_indices_shape) &&
101 size_indices_shape.num_elements() == rank,
102 errors::InvalidArgument(
103 "size_indices must be a vector with length equal to "
104 "input rank, but input rank is ",
105 input_shape.dims(), " and size_indices has shape ",
106 size_indices_shape.DebugString()));
107
108 std::vector<int64> size_indices;
109 OP_REQUIRES_OK(
110 ctx, ctx->ConstantInputAsIntVector("size_indices", &size_indices));
111
112 xla::XlaOp start_indices = ctx->Input("start_indices");
113 xla::XlaOp result = xla::DynamicSlice(
114 ctx->Input("input"), SliceVector(start_indices, rank), size_indices);
115 ctx->SetOutput(0, result);
116 }
117 };
118
119 REGISTER_XLA_OP(
120 Name("XlaDynamicSlice").CompileTimeConstantInput("size_indices"),
121 DynamicSliceOp);
122
123 } // namespace
124 } // namespace tensorflow
125