DenseElementsAttr operand_init;
    if (!matchPattern(reduce_op.getInitValues().front(),
                      m_Constant(&operand_init)))
      return failure();
    if (!IsValueInitValue(operand_init)) return failure();

    // Check that the iota init is zero.
    DenseElementsAttr iota_init;
    if (!matchPattern(reduce_op.getInitValues().back(), m_Constant(&iota_init)))
      return failure();

        GetInvertPermutedTensorType(reduce_type, transpose_perm);
    auto new_reduce_op = rewriter.create<stablehlo::ReduceWindowOp>(
        reduce_op.getLoc(), new_reduce_type, transpose_input,
        reduce_op.getInitValues()[0], new_window_dimensions, new_window_strides,
        new_base_dilations, new_win_dilations, new_padding_attr);
    IRMapping mapping;
    reduce_op.getBody().cloneInto(&new_reduce_op.getBody(), mapping);

    if (rw.getNumResults() != 1 || rw.getInputs().size() != 1 ||
        rw.getInitValues().size() != 1)
      return failure();

    if (failed(MatchBinaryReduceFunction<BinaryOp>(rw.getBody())))
      return failure();

    // Ensure that initial_values are as expected.
    auto const_op = llvm::dyn_cast_or_null<mhlo::ConstantOp>(
        rw.getInitValues()[0].getDefiningOp());
    if (!const_op) return failure();

  const auto &inputs_ty = op.getInputs().getType();
  int n_inputs = inputs_ty.size();
  if (n_inputs < 1) return op.emitOpError() << "No inputs";

  const auto &init_values_ty = op.getInitValues().getType();
  int n_init_values = init_values_ty.size();
  if (n_init_values != n_inputs) {
    return op.emitOpError() << "Number of inputs (" << n_inputs

        [](Type ty) { return mlir::cast<ShapedType>(ty).getElementType(); })};

    // Create the mhlo.reduce op.
    auto reduce_op = rewriter.create<mhlo::ReduceOp>(
        loc, op.getInputs(), op.getInitValues(),
        GetI64ElementsAttr(op.getDimensionsToReduce()), elementTypes);

    // Insert a call to the reducer in the region of the mhlo op.
    BuildBodyWithCall(rewriter, loc, func, func_ty, &reduce_op.getBody());