CompileTimeConstantOperand<TF::XlaPadOp, 2, 3, 4>,
      // $window_dimensions, $window_strides, $base_dilations,
      // $window_dilations, $padding
      CompileTimeConstantOperand<TF::XlaReduceWindowOp, 2, 3, 4, 5, 6>,
      // $dim_index
      CompileTimeConstantOperand<TF::XlaRemoveDynamicDimensionSizeOp, 1>,
      // $window_dimensions, $window_strides, $padding

  ^bb0(%arg1: tensor<f32>, %arg2: tensor<f32>):
    %7 = stablehlo.maximum %arg1, %arg2 : tensor<f32>
    stablehlo.return %7 : tensor<f32>
  }) {padding = dense<[[0, 0], [0, 1], [0, 1], [0, 0]]> : tensor<4x2xi64>, window_dimensions = array<i64: 1, 3, 3, 1>, window_strides = array<i64: 1, 2, 2, 1>} : (tensor<?x112x112x64xf32>, tensor<f32>) -> tensor<?x56x56x64xf32>
  return %6 : tensor<?x56x56x64xf32>

func.func @maxpool_valid_padding(%arg0: tensor<2x12x20x7xi32>) -> tensor<2x3x5x7xi32> {
  // CHECK: %[[INIT:.*]] = mhlo.constant dense<-2147483648> : tensor<i32>
  // CHECK: "mhlo.reduce_window"(%[[ARG]], %[[INIT]])
  // CHECK: <{window_dimensions = dense<[1, 2, 2, 1]> : tensor<4xi64>, window_strides = dense<[1, 4, 4, 1]> : tensor<4xi64>}>
  // CHECK: mhlo.maximum
  // CHECK: mhlo.return

    auto new_reduce_window_op =
        rewriter.create<mlir::stablehlo::ReduceWindowOp>(
            op.getLoc(), new_result_type, transpose_op.getOperand(),
            /*init_value=*/op.getOperand(1),
            /*window_dimensions=*/
            PermuteI64ArrayAttr(rewriter, op.getWindowDimensions(),
                                kNchwToNhwcPermutation),
            /*window_strides=*/

      "mhlo.return"(%5) : (tensor<f32>) -> ()
    }) {
    base_dilations = dense<1> : tensor<4xi64>,
    padding = dense<0> : tensor<4x2xi64>,
    window_dilations = dense<1> : tensor<4xi64>,
    window_dimensions = dense<[1, 3, 3, 1]> : tensor<4xi64>,
    window_strides = dense<[1, 2, 2, 1]> : tensor<4xi64>} : (tensor<4x16x16x8xf32>, tensor<f32>) -> tensor<4x7x7x8xf32>
  %3 = mhlo.divide %2, %1 : tensor<4x7x7x8xf32>

        padding = #vhlo.tensor_v1<dense<[[0, 0], [159, 0], [0, 0]]> : tensor<3x2xi64>>,
        window_dilations = #vhlo.tensor_v1<dense<1> : tensor<3xi64>>,
        window_dimensions = #vhlo.tensor_v1<dense<[1, 160, 1]> : tensor<3xi64>>,
        window_strides = #vhlo.tensor_v1<dense<1> : tensor<3xi64>>}> ({
    ^bb0(%arg23: tensor<f32>, %arg24: tensor<f32>):

    for (int64_t i = 0, e = window_dimensions.size(); i < e; ++i) {
      int64_t window_dimension = window_dimensions.getValues<int64_t>()[i];
      if (window_dimension == 1) continue;
      // Cumulative axis already set.
      if (cumulative_axis != -1) return failure();
      // Potential cumulative axis is not the right size.
      if (window_dimension != operand_type.getShape()[i]) return failure();

                                PatternRewriter &rewriter) const override {
    DenseElementsAttr window_dimensions, window_strides, base_dilations,
        window_dilations, padding;
    if (!(matchPattern(op.getWindowDimensions(),
                       m_Constant(&window_dimensions)) &&
          matchPattern(op.getWindowStrides(), m_Constant(&window_strides)) &&

    %1 = stablehlo.maximum %arg2, %arg3 : tensor<!quant.uniform<i8:f32, 3.000000e-01:-5>>
    stablehlo.return %1 : tensor<!quant.uniform<i8:f32, 3.000000e-01:-5>>
  }) {window_dimensions = array<i64: 1, 3, 4, 1>, window_strides = array<i64: 1, 2, 3, 1>} : (tensor<2x9x10x3x!quant.uniform<i8:f32, 3.000000e-01:-5>>, tensor<!quant.uniform<i8:f32, 3.000000e-01:-5>>) -> tensor<2x4x3x3x!quant.uniform<i8:f32, 3.000000e-01:-5>>

  uint32_t opcode_index =
      GetOpcodeIndex(op_name, tflite::BuiltinOperator_STABLEHLO_REDUCE_WINDOW);

  auto wd = reduce_window_op.getWindowDimensions();
  auto window_dimensions =
      builder_.CreateVector(std::vector<int64_t>(wd.begin(), wd.end()));
  auto window_strides = builder_.CreateVector(
      mlir::GetOptionalVector<int64_t>(reduce_window_op.getWindowStrides()));