| 30 | explicit PadOp(OpKernelConstruction* ctx) : XlaOpKernel(ctx) {} |
| 31 | |
| 32 | void Compile(XlaOpKernelContext* ctx) override { |
| 33 | const TensorShape input_shape = ctx->InputShape("input"); |
| 34 | const TensorShape pad_shape = ctx->InputShape("paddings"); |
| 35 | const int dims = input_shape.dims(); |
| 36 | OP_REQUIRES( |
| 37 | ctx, |
| 38 | TensorShapeUtils::IsMatrix(pad_shape) && pad_shape.dim_size(1) == 2, |
| 39 | errors::InvalidArgument("paddings must be a matrix with 2 columns: ", |
| 40 | pad_shape.DebugString())); |
| 41 | OP_REQUIRES( |
| 42 | ctx, dims == pad_shape.dim_size(0), |
| 43 | errors::InvalidArgument( |
| 44 | "The first dimension of paddings must be the rank of inputs", |
| 45 | pad_shape.DebugString(), " ", input_shape.DebugString())); |
| 46 | |
| 47 | xla::XlaOp input = ctx->Input("input"); |
| 48 | if (dims == 0) { |
| 49 | // Tensor is rank 0. Return it unchanged. |
| 50 | ctx->SetOutput(0, input); |
| 51 | return; |
| 52 | } |
| 53 | |
| 54 | xla::Literal pad_literal; |
| 55 | OP_REQUIRES_OK(ctx, |
| 56 | ctx->ConstantInputAsInt64Literal("paddings", &pad_literal)); |
| 57 | |
| 58 | xla::PaddingConfig config; |
| 59 | for (int i = 0; i < dims; ++i) { |
| 60 | auto* dim = config.add_dimensions(); |
| 61 | int before = pad_literal.Get<int64>({i, 0}); |
| 62 | int after = pad_literal.Get<int64>({i, 1}); |
| 63 | OP_REQUIRES(ctx, before >= 0 && after >= 0, |
| 64 | errors::InvalidArgument( |
| 65 | "Paddings must be non-negative: ", before, " ", after)); |
| 66 | dim->set_edge_padding_low(before); |
| 67 | dim->set_edge_padding_high(after); |
| 68 | } |
| 69 | |
| 70 | // PadV2 added a "constant_values" input that indicates the pad value. |
| 71 | xla::XlaOp constant_values; |
| 72 | if (ctx->num_inputs() == 3) { |
| 73 | OP_REQUIRES( |
| 74 | ctx, TensorShapeUtils::IsScalar(ctx->InputShape("constant_values")), |
| 75 | errors::InvalidArgument("constant_values must be a scalar.")); |
| 76 | ctx->SetOutput(0, xla::Pad(input, ctx->Input("constant_values"), config)); |
| 77 | } else { |
| 78 | auto zero = XlaHelpers::Zero(ctx->builder(), input_type(0)); |
| 79 | ctx->SetOutput(0, xla::Pad(input, zero, config)); |
| 80 | } |
| 81 | } |
| 82 | }; |
| 83 | |
| 84 | REGISTER_XLA_OP(Name("Pad").CompileTimeConstantInput("paddings"), PadOp); |
nothing calls this directly
no test coverage detected