Home | History | Annotate | Download | only in kernels 
      1 /* Copyright 2017 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 // XLA-specific reduction Ops.
     17 
     18 #include "tensorflow/compiler/tf2xla/kernels/reduction_ops.h"
     19 #include "tensorflow/compiler/tf2xla/type_util.h"
     20 #include "tensorflow/compiler/tf2xla/xla_helpers.h"
     21 #include "tensorflow/compiler/tf2xla/xla_op_kernel.h"
     22 #include "tensorflow/compiler/xla/literal_util.h"
     23 #include "tensorflow/core/framework/kernel_def_builder.h"
     24 
     25 namespace tensorflow {
     26 
     27 XlaReductionOp::XlaReductionOp(OpKernelConstruction* ctx) : XlaOpKernel(ctx) {
     28   const DataType dt = BaseType(input_type(0));
     29   OP_REQUIRES_OK(ctx, ctx->MatchSignature({dt, DT_INT32}, {dt}));
     30 
     31   OP_REQUIRES_OK(ctx, ctx->GetAttr("keep_dims", &keep_dims_));
     32 }
     33 
     34 // Return the base case for the reduction. Defaults to zero.
     35 xla::ComputationDataHandle XlaReductionOp::InitialValue(
     36     xla::ComputationBuilder* builder) {
     37   return XlaHelpers::Zero(builder, input_type(0));
     38 }
     39 
     40 // Unless BuildFinalizer is overridden the reduction has no
     41 // finalizer.
     42 xla::ComputationDataHandle XlaReductionOp::BuildFinalizer(
     43     xla::ComputationBuilder* builder,
     44     const xla::ComputationDataHandle& reduce_output,
     45     int64 num_elements_reduced) {
     46   return reduce_output;
     47 }
     48 
     49 void XlaReductionOp::Compile(XlaOpKernelContext* ctx) {
     50   const TensorShape data_shape = ctx->InputShape(0);
     51   const TensorShape axes_tensor_shape = ctx->InputShape(1);
     52   VLOG(1) << "ReductionOp: " << ctx->op_kernel().name();
     53 
     54   if (axes_tensor_shape.num_elements() == 0) {
     55     // The reduction axes is an empty vector, which means there are no
     56     // axes to reduce so just pass the input directly through to the
     57     // output.
     58     ctx->SetOutput(0, ctx->Input(0));
     59     return;
     60   }
     61 
     62   // Evaluate the constant, reshaping to a 1-vector if it is a scalar.
     63   xla::Literal axes_literal;
     64   OP_REQUIRES_OK(ctx,
     65                  ctx->ConstantInputReshaped(
     66                      1, {axes_tensor_shape.num_elements()}, &axes_literal));
     67 
     68   VLOG(1) << "data shape: " << data_shape.DebugString();
     69   VLOG(1) << "axes      : " << axes_literal.ToString();
     70 
     71   gtl::InlinedVector<bool, 4> bitmap(data_shape.dims(), false);
     72   std::vector<int64> xla_axes;
     73   int64 num_elements_reduced = 1LL;
     74   for (int64 i = 0; i < axes_tensor_shape.num_elements(); ++i) {
     75     int32 index = axes_literal.Get<int>({i});
     76     OP_REQUIRES(ctx,
     77                 !(index < -data_shape.dims() || index >= data_shape.dims()),
     78                 errors::InvalidArgument("Invalid reduction dimension (", index,
     79                                         " for input with ", data_shape.dims(),
     80                                         " dimension(s)"));
     81     index = (index + data_shape.dims()) % data_shape.dims();
     82     bitmap[index] = true;
     83     xla_axes.push_back(index);
     84     num_elements_reduced *= data_shape.dim_size(index);
     85   }
     86 
     87   std::vector<int64> final_shape;
     88   for (int i = 0; i < data_shape.dims(); ++i) {
     89     if (!bitmap[i]) {
     90       // If we are not reducing along dimension i.
     91       int64 dim = data_shape.dim_size(i);
     92       final_shape.push_back(dim);
     93     } else if (keep_dims_) {
     94       // We are reducing along dimension i, but we want to keep the
     95       // same number of dimensions, so we set the dimension of i to
     96       // '1'.
     97       final_shape.push_back(1);
     98     }
     99   }
    100 
    101   string desc = ctx->op_kernel().name();
    102 
    103   // Call virtual method to get the initial value.
    104   const xla::ComputationDataHandle initial = InitialValue(ctx->builder());
    105   // Construct the builder for the reduction lambda.
    106   xla::ComputationBuilder r(ctx->builder()->client(),
    107                             strings::StrCat(desc, "-reduction"));
    108   xla::PrimitiveType type;
    109   TF_CHECK_OK(DataTypeToPrimitiveType(input_type(0), &type));
    110   // Make two scalar parameters of the desired type for the lambda.
    111   xla::ComputationDataHandle rx =
    112       r.Parameter(0, xla::ShapeUtil::MakeShape(type, {}), "x");
    113   xla::ComputationDataHandle ry =
    114       r.Parameter(1, xla::ShapeUtil::MakeShape(type, {}), "y");
    115 
    116   auto data = ctx->Input(0);
    117 
    118   // Call virtual method to build the reduction lambda.
    119   BuildReducer(&r, rx, ry);
    120   xla::Computation reduction_computation = r.Build().ConsumeValueOrDie();
    121   xla::ComputationDataHandle reduce =
    122       ctx->builder()->Reduce(data, initial, reduction_computation, xla_axes);
    123 
    124   xla::ComputationDataHandle finalized =
    125       BuildFinalizer(ctx->builder(), reduce, num_elements_reduced);
    126 
    127   xla::ComputationDataHandle result;
    128   if (keep_dims_) {
    129     result = ctx->builder()->Reshape(finalized, final_shape);
    130   } else {
    131     result = finalized;
    132   }
    133   ctx->SetOutput(0, result);
    134 }
    135 
    136 }  // namespace tensorflow
    137