| 88 | } |
| 89 | |
| 90 | void Compile(XlaOpKernelContext* ctx) override { |
| 91 | const TensorShape out_backprop_shape = ctx->InputShape(0); |
| 92 | |
| 93 | OP_REQUIRES(ctx, TensorShapeUtils::IsMatrixOrHigher(out_backprop_shape), |
| 94 | errors::InvalidArgument("Input tensor must be at least 2D: ", |
| 95 | out_backprop_shape.DebugString())); |
| 96 | |
| 97 | // feature_dim is the channel (C) dimension of the data. |
| 98 | int feature_dim = (data_format_ == FORMAT_NHWC) |
| 99 | ? out_backprop_shape.dims() - 1 |
| 100 | : /*data_format == FORMAT_NCHW*/ 1; |
| 101 | OP_REQUIRES( |
| 102 | ctx, feature_dim >= 0, |
| 103 | errors::InvalidArgument("Input tensor does not have enough dimensions " |
| 104 | "to contain the feature dimension")); |
| 105 | |
| 106 | std::vector<int64> reduce_dims(out_backprop_shape.dims() - 1); |
| 107 | std::iota(reduce_dims.begin(), reduce_dims.begin() + feature_dim, 0); |
| 108 | std::iota(reduce_dims.begin() + feature_dim, reduce_dims.end(), |
| 109 | feature_dim + 1); |
| 110 | xla::XlaBuilder* const b = ctx->builder(); |
| 111 | const DataType accumulation_type = |
| 112 | XlaHelpers::SumAccumulationType(input_type(0)); |
| 113 | auto converted = |
| 114 | XlaHelpers::ConvertElementType(ctx->Input(0), accumulation_type); |
| 115 | auto reduce = |
| 116 | xla::Reduce(converted, XlaHelpers::Zero(b, accumulation_type), |
| 117 | *ctx->GetOrCreateAdd(accumulation_type), reduce_dims); |
| 118 | ctx->SetOutput(0, XlaHelpers::ConvertElementType(reduce, input_type(0))); |
| 119 | } |
| 120 | |
| 121 | private: |
| 122 | TensorFormat data_format_; |
nothing calls this directly
no test coverage detected