// CHECK: %[[R0:.*]] = "tf.Transpose"(%arg0, %[[CST]])
  // CHECK: %[[R1:.*]] = "tf.MaxPool"(%[[R0]]) <{data_format = "NHWC", explicit_paddings = [], ksize = [1, 3, 3, 1], padding = "SAME", strides = [1, 2, 2, 1]}>
  // CHECK: %[[CST_0:.*]] = "tf.Const"() <{value = dense<[0, 3, 1, 2]> : tensor<4xi64>}>
  // CHECK: "tf.Transpose"(%[[R1]], %[[CST_0]])
  %0 = "tf.MaxPool"(%arg0)
       {
         data_format = "NCHW",
         ksize = [1, 1, 3, 3],

 BiasAdd 1
 BiasAddGrad 1
 Cast 3
 Const 357
 Conv2D 53
 Conv2DBackpropFilter 53
 Conv2DBackpropInput 52
 DivNoNan 1
 Equal 1
 FusedBatchNorm 53
 FusedBatchNormGrad 53
 Identity 2
 MatMul 3
 MaxPool 1
 MaxPoolGrad 1
 Mean 1
 Mul 164
 Pad 1
 ReadVariableOp 646
 Relu 49
 ReluGrad 49
 Reshape 2
 ResourceApplyKerasMomentum 161
 ShapeN 50
 Softmax 1
 SparseSoftmaxCrossEntropyWithLogits 1

  %0 = "quantfork.qcast"(%arg0) : (tensor<*xf32>) -> tensor<*x!quant.uniform<i8:f32, 5.000000e-02:-10>>
  %1 = "quantfork.dcast"(%0) : (tensor<*x!quant.uniform<i8:f32, 5.000000e-02:-10>>) -> tensor<*xf32>
  %2 = "tf.MaxPool"(%1) {data_format = "NHWC", device = "", explicit_paddings = [], ksize = [1, 2, 2, 1], padding = "VALID", strides = [1, 2, 2, 1]} : (tensor<*xf32>) -> tensor<*xf32>

 AssignAddVariableOp 1
 BiasAdd 1
 BiasAddGrad 1
 Cast 115
 Const 407
 Conv2D 53
 Conv2DBackpropFilter 53
 Conv2DBackpropInput 52
 Equal 1
 FusedBatchNormGradV2 53
 FusedBatchNormV2 53
 MatMul 3
 MaxPool 1
 MaxPoolGrad 1
 Mean 1
 Mul 218
 Pad 2
 ReadVariableOp 538
 Relu 49
 ReluGrad 49
 Reshape 2
 ResourceApplyKerasMomentum 161
 Slice 1
 Softmax 1
 SparseSoftmaxCrossEntropyWithLogits 1

  %0 = "quantfork.qcast"(%arg0) : (tensor<*xf32>) -> tensor<*x!quant.uniform<i8:f32, 5.000000e-02:-10>>
  %1 = "quantfork.dcast"(%0) : (tensor<*x!quant.uniform<i8:f32, 5.000000e-02:-10>>) -> tensor<*xf32>
  %2 = "tf.MaxPool"(%1) {data_format = "NHWC", device = "", explicit_paddings = [], ksize = [1, 2, 2, 1], padding = "VALID", strides = [1, 2, 2, 1]} : (tensor<*xf32>) -> tensor<*xf32>

  let description = [{
    Operation supports folding operand(s) transposes into the operation itself.

    (1) Operation might have layout dependent operands and results...

      Example:  MaxPool(Transpose($arg, $perm))
                  -> Transpose(MaxPool($arg, $perm))

    (2) ... or it might have only layout dependent operands:

      Example: Mean(Transpose($arg, $reduction_dims))
                 -> Mean($arg, Transpose($reduction_dims))

// CHECK: %[[maxpool:.*]] = "tf.MaxPool"(%[[conv_quant]]) <{data_format = "NHWC", ksize = [1, 2, 2, 1], padding = "VALID", strides = [1, 1, 1, 1]}> : (tensor<*xi8>) -> tensor<*xi8>
// CHECK: %[[dequantize:.*]] = "tf.PartitionedCall"(%[[maxpool]]
// CHECK-SAME: f = @dequantize_i8
// CHECK: return %[[dequantize]]

// -----

// CHECK-LABEL: maxpool_explicit_padding
func.func @maxpool_explicit_padding(%arg0: tensor<2x12x20x7xi32>) -> tensor<2x3x5x7xi32> {
  // CHECK: tf.MaxPool
  // TODO(b/165938852): need to support explicit padding in max_pool.

  %0 = "tf.MaxPool"(%arg0) {data_format = "NHWC", ksize = [1, 2, 2, 1], padding = "EXPLICIT", strides = [1, 4, 4, 1]} : (tensor<2x12x20x7xi32>) -> tensor<2x3x5x7xi32>

    ConvertAvgPoolGradOp<TF::AvgPool3DGradOp, /*num_dims=*/5>;

// Converts MaxPool op to HLO ReduceWindow op by setting appropriate window
// dimensions with max as the reduction function.
//
// Sample result for VALID padding mode:
//
//   %init = arith.constant dense<...> : tensor<i32>
//   %max_pool = "mhlo.reduce"(%inp, %init) ["mhlo.maximum"]

def _composite_max_pool(input_, stride_w, stride_h, filter_width, filter_height,
                        padding):
  ksize = [1, filter_width, filter_height, 1]
  strides = [1, stride_w, stride_h, 1]
  return tf.raw_ops.MaxPool(
      input=input_, ksize=ksize, strides=strides, padding=padding)


@tf.RegisterGradient('NewMaxPool')
def _max_pool_grad(op: ops.Operation, grad):
  filter_width = op.get_attr('filter_width')