"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)

  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.

  }

  void Run(std::initializer_list<int> input_values) {
    Run(outputs_, input_values);
  }

  void Run(const std::vector<TF_Output>& run_outputs,
           std::initializer_list<int> input_values) {
    DCHECK_EQ(inputs_.size(), input_values.size());
    std::vector<std::pair<TF_Operation*, TF_Tensor*>> inputs(inputs_.size());
    int i = 0;
    for (int v : input_values) {
      inputs[i] = {inputs_[i].oper, Int32Tensor(v)};

    } else {
      llvm::SmallVector<Value> input_values(node_output_tensors.size(),
                                            scalar_one);
      for (const auto& [output_index, tensor_index, tensor_value] :
           node_output_tensors) {
        input_values[tensor_index] = tensor_value;
      }
      identity_op = builder.create<TF::IdentityNOp>(
          new_loc, TypeRange(ValueRange(input_values)), input_values);
    }

      Type current_input_type = mlir::cast<TypeAttr>(input_types[i]).getValue();
      if (current_input_type != target_input_type) {
        input_values[i] =
            rewriter.create<TF::CastOp>(loc, result_type, input_values[i]);
      }
    }
  }

  const SymbolTable& symbol_table_;
  const bool materialize_derived_attrs_;
  const llvm::SmallDenseSet<StringRef, 4> fixed_elt_type_attrs_{"i32_", "i64_",

// CHECK: %[[ADD:.+]] = stablehlo.add %[[TRANSPOSE_1]], %[[ARG_1]] : tensor<1x4x5x5xf32>
// CHECK: return %[[ADD]]

// -----

// Tests that a `reduce_window{max}(add(convolution(%activation, %weight), %bias), %init_value)`
// with the activation tensor of NCHW format is converted to NHWC convolution +
// add + reduce_window (with max) operation. Transpose ops are inserted to

  TF_Output output{dequeue_op, 0};
  TF_Tensor* ret;
  TF_SessionRun(session, /*run_options*/ nullptr,
                // input related parameters
                /*inputs*/ nullptr, /*input_values*/ nullptr, /*ninputs*/ 0,
                // output related parameters
                /*outputs*/ &output, /*output_values*/ &ret,
                /*noutputs*/ 1,
                /*targets*/ nullptr, /*ntargets*/ 0,