Value operands[] = {op.getInput(), index_values};
    Value init_values[] = {init_value, index_init_value};
    DenseIntElementsAttr reduction_dimensions =
        GetI64ElementsAttr({axis}, &rewriter);

    auto reduction = rewriter.create<ReduceOp>(
        loc, llvm::ArrayRef<Value>(operands),
        llvm::ArrayRef<Value>(init_values), reduction_dimensions,
        TypeRange({input_element_type, index_element_type}));

  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
                            << ") is different than number of init_values ("
                            << n_init_values << ")";
  }

  auto input_ty_0 = inputs_ty[0].cast<ShapedType>();

    const float zp = q_type.getZeroPoint();

    auto input_values = input_dequant.getValue();

    // mapValues always takes a function returning APInt, even when the output
    // is actually float.
    using DequantizeFuncType = llvm::APInt(const llvm::APInt&);
    auto dequantize_func = [&](const APInt& ap_int_value) -> APInt {
      const int64_t int_value = ap_int_value.getSExtValue();

              "args": {
                "grpc.http2.max_pings_without_data": {
                  "int_value": 0
                },
                "grpc.keepalive_time_ms": {
                  "int_value": 10000
                },
                "grpc.keepalive_timeout_ms": {
                  "int_value": 20000
                }
              }
            }
          },
          "initial_metadata": [

      "variable",
  });

  EXPECT_EQ(names, expected);
}

Status MakeLoop(Scope s, Output init_value, absl::string_view loop_name) {
  s = s.NewSubScope(std::string(loop_name));
  ops::internal::Enter enter(s.WithOpName("init_value"), init_value, loop_name);
  ops::Merge merge(s.WithOpName("merge"), {init_value, init_value});
  Output next_iteration =
      ops::NextIteration(s.WithOpName("next_itr"), merge.output);

    }

    if (input_values.hasRank() && output_values.hasRank() &&
        input_values.getRank() != output_values.getRank()) {
      return func.emitError() << "Input " << kValues << " and output "
                              << kValues << " should have the same rank";
    }
  } else {
    auto input_values = GetInputType(func, kValues);
    if (!RankEquals(input_values, 1) ||

    %1 = "tf_device.cluster"() ( {
      %init_value = "tf.ReadVariableOp"(%resource_handle)
      "tf.AssignAddVariableOp"(%resource_handle, %init_value)
      %new_value = "tf.ReadVariableOp"(%resource_handle)
      tf_device.return %new_value
    })
    ```

    After this pass, the computation would become:

    ```mlir
    %resource_handle = "tf.VarHandleOp"()
    %init_value = "tf.ReadVariableOp"(%resource_handle)

    for (int exponent : asList(2, 3, 4, 9, 15, 16, 17, 24, 25, 30)) {
      int x = 1 << exponent;
      intValues.add(x, x + 1, x - 1);
    }
    intValues.add(9999).add(10000).add(10001).add(1000000); // near powers of 10
    intValues.add(5792).add(5793); // sqrt(2^25) rounded up and down
    POSITIVE_INTEGER_CANDIDATES = intValues.build();
    NEGATIVE_INTEGER_CANDIDATES =
        ImmutableList.copyOf(
            Iterables.concat(

    inputs.push_back(std::move(tensor));
  }

  std::vector<tensorflow::TensorValue> input_values;
  for (tensorflow::Tensor& tensor : inputs) {
    input_values.emplace_back();
    input_values.back().tensor = &tensor;
  }

  tensorflow::OpKernelContext::Params params;
  params.inputs = input_values;
  params.device = runner->device();
  params.op_kernel = runner->op_kernel();

  void rewrite(AddOp op, PatternRewriter& rewriter) const override {
    DeferRhsTransposeForBinaryOp(op, rewriter);
  }
};

// Rewrites the `reduce_window(transpose(%activation), %init_value)` patterns to
// `transpose(reduce_window(%activation), %init_value)`, deferring the transpose
// to the result. The reduce function should be equivalent to
// `stablehlo.maximum`, representing max pooling.