105 lines
3.7 KiB
C++
105 lines
3.7 KiB
C++
/* Copyright 2017 The TensorFlow Authors. All Rights Reserved.
|
|
|
|
Licensed under the Apache License, Version 2.0 (the "License");
|
|
you may not use this file except in compliance with the License.
|
|
You may obtain a copy of the License at
|
|
|
|
http://www.apache.org/licenses/LICENSE-2.0
|
|
|
|
Unless required by applicable law or agreed to in writing, software
|
|
distributed under the License is distributed on an "AS IS" BASIS,
|
|
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
|
See the License for the specific language governing permissions and
|
|
limitations under the License.
|
|
==============================================================================*/
|
|
|
|
// Native XLA implementations of indexing ops.
|
|
|
|
#include "tensorflow/compiler/tf2xla/kernels/index_ops.h"
|
|
|
|
#include "tensorflow/compiler/tf2xla/type_util.h"
|
|
#include "tensorflow/compiler/tf2xla/xla_op_kernel.h"
|
|
#include "tensorflow/compiler/tf2xla/xla_op_registry.h"
|
|
#include "tensorflow/compiler/xla/client/lib/arithmetic.h"
|
|
#include "tensorflow/compiler/xla/shape_util.h"
|
|
#include "tensorflow/core/framework/bounds_check.h"
|
|
#include "tensorflow/core/framework/kernel_def_builder.h"
|
|
#include "tensorflow/core/framework/op_kernel.h"
|
|
#include "tensorflow/core/framework/register_types.h"
|
|
#include "tensorflow/core/framework/tensor.h"
|
|
#include "tensorflow/core/framework/tensor_shape.h"
|
|
|
|
namespace tensorflow {
|
|
XlaArgMinMaxOp::XlaArgMinMaxOp(OpKernelConstruction* ctx, bool is_min)
|
|
: XlaOpKernel(ctx),
|
|
is_min_(is_min),
|
|
is_gpu_(ctx->device_type().type_string() == DEVICE_GPU_XLA_JIT) {}
|
|
|
|
void XlaArgMinMaxOp::Compile(XlaOpKernelContext* ctx) {
|
|
const TensorShape input_shape = ctx->InputShape(0);
|
|
const TensorShape dimension_shape = ctx->InputShape(1);
|
|
|
|
OP_REQUIRES(ctx, TensorShapeUtils::IsScalar(dimension_shape),
|
|
errors::InvalidArgument(
|
|
"dim must be a scalar, but received tensor of shape: ",
|
|
dimension_shape.DebugString()));
|
|
|
|
int64 dim;
|
|
OP_REQUIRES_OK(ctx, ctx->ConstantInputAsIntScalar(1, &dim));
|
|
|
|
const int input_dims = input_shape.dims();
|
|
const int axis = dim < 0 ? dim + input_dims : dim;
|
|
|
|
OP_REQUIRES(
|
|
ctx, axis >= 0 && axis < input_dims,
|
|
errors::InvalidArgument("Expected dimension in the range [", -input_dims,
|
|
", ", input_dims, "), but got ", dim));
|
|
const int64 axis_size = input_shape.dim_size(axis);
|
|
OP_REQUIRES(
|
|
ctx, axis_size > 0,
|
|
errors::InvalidArgument("Reduction axis ", dim, " is empty in shape ",
|
|
input_shape.DebugString()));
|
|
|
|
DataType index_type = output_type(0);
|
|
xla::PrimitiveType index_xla_type;
|
|
OP_REQUIRES_OK(ctx, DataTypeToPrimitiveType(index_type, &index_xla_type));
|
|
|
|
xla::XlaOp input = ctx->Input(0);
|
|
xla::XlaOp output;
|
|
// One pass ArgMin/ArgMax is slow on GPUs.
|
|
if (is_min_) {
|
|
if (is_gpu_) {
|
|
output = xla::ArgMinTwoPass(input, index_xla_type, axis);
|
|
} else {
|
|
output = xla::ArgMin(input, index_xla_type, axis, /*stable=*/true);
|
|
}
|
|
} else {
|
|
if (is_gpu_) {
|
|
output = xla::ArgMaxTwoPass(input, index_xla_type, axis);
|
|
} else {
|
|
output = xla::ArgMax(input, index_xla_type, axis, /*stable=*/true);
|
|
}
|
|
}
|
|
|
|
ctx->SetOutput(0, output);
|
|
}
|
|
|
|
XlaArgMaxOp::XlaArgMaxOp(OpKernelConstruction* ctx)
|
|
: XlaArgMinMaxOp(ctx, /*is_min=*/false) {}
|
|
REGISTER_XLA_OP(Name("ArgMax").CompileTimeConstantInput("dimension"),
|
|
XlaArgMaxOp);
|
|
|
|
namespace {
|
|
|
|
class XlaArgMinOp : public XlaArgMinMaxOp {
|
|
public:
|
|
explicit XlaArgMinOp(OpKernelConstruction* ctx);
|
|
};
|
|
XlaArgMinOp::XlaArgMinOp(OpKernelConstruction* ctx)
|
|
: XlaArgMinMaxOp(ctx, /*is_min=*/true) {}
|
|
REGISTER_XLA_OP(Name("ArgMin").CompileTimeConstantInput("dimension"),
|
|
XlaArgMinOp);
|
|
|
|
} // namespace
|
|
} // namespace tensorflow
|