// TODO(fengliuai): make build_list op only take tensor argument
      llvm::SmallVector<Attribute, 4> list_input_types;
      llvm::SmallVector<Value, 4> list_inputs;
      for (auto list_input : list_op.getOperands()) {
        auto cast_op = dyn_cast_or_null<CastOp>(list_input.getDefiningOp());
        if (!cast_op) return failure();
        list_inputs.push_back(cast_op.getArg());

bool AddAccessedResourceIds(
    Operation* op,
    const mlir::TF::ResourceAliasAnalysis::Info& resource_analysis_info,
    llvm::SmallDenseSet<int64_t>& resource_ids) {
  for (Value operand : TF::filter_resources(op->getOperands())) {
    if (resource_analysis_info.IsUnknownResource(operand)) {
      VLOG(2) << "  unknown access";
      return true;
    }
    const auto& ids = resource_analysis_info.GetResourceUniqueIds(operand);

             tfl::SparseQConstOp, mlir::TFL::NoValueOp,
             mlir::stablehlo::ConstantOp, mlir::vhlo::ConstantOpV1>(op);
}

static bool IsTFResourceOp(Operation* op) {
  for (const auto& operand : op->getOperands()) {
    auto elementType = getElementTypeOrSelf(operand.getType());
    if (mlir::isa<mlir::TF::ResourceType>(elementType)) {
      return true;
    }
  }
  for (const auto& result : op->getResults()) {

  IslandOp const_wrapper = GetDummyConstant(builder, const_type, loc);

  // Get new operand and result types.
  auto new_operands = llvm::to_vector<4>(while_op->getOperands());
  auto new_result_types = llvm::to_vector<4>(while_op->getResultTypes());
  Value const_output = const_wrapper.getOutputs()[0];
  for (int i = 0; i < num_resource_inputs; ++i) {
    new_operands.push_back(const_output);

  auto maybe_replicate =
      llvm::dyn_cast<tf_device::ReplicateOp>(cluster_func->getParentOp());

  llvm::SmallVector<int64_t, 8> transform_input_indices;
  for (const auto& input : llvm::enumerate(cluster_func.getOperands())) {
    if (auto block_arg = mlir::dyn_cast<BlockArgument>(input.value())) {
      if (block_arg.getArgNumber() != arg_num) continue;
      // For a block argument, consider transforms only when it is a replicated

  for (int i = 0; i < num_cores_per_replica; ++i)
    input_list->emplace_back(llvm::SmallVector<mlir::Value, 4>());

  llvm::SmallVector<mlir::Value, 4> cluster_func_inputs(
      cluster_func.getOperands());
  auto sharding_attrs =
      cluster_func.getOperation()->getAttrOfType<mlir::ArrayAttr>(
          kInputShardingAttr);
  // If sharding attribute does not exist, then all inputs are placed on 0th

  mlir::Operation& return_op = region.back().back();
  mlir::Location loc = return_op.getLoc();
  op_builder.setInsertionPointToEnd(&region.back());
  op_builder.create<mlir::stablehlo::ReturnOp>(loc, return_op.getOperands());
  return_op.erase();
}

void InlineVhloOpRegion(mlir::Region& region, mlir::func::FuncOp func) {
  OpBuilder op_builder{region};
  mlir::IRMapping mapper;

      }

      mhlo_op = rewriter.create<DstOpT>(loc, TypeRange(while_result_types),
                                        adaptor.getOperands());
    } else {
      mhlo_op = rewriter.create<DstOpT>(loc, op.getResultTypes(),
                                        adaptor.getOperands());
    }

    int64_t num_regions = op.getNumRegions();
    for (int64_t idx = 0; idx < num_regions; ++idx) {

  // `tf.While` or `tf.If` ops which are not part of the op definition but
  // appear in a variadic input list.
  add_remaining_effects(filter_resources(op->getOperands()));
  add_remaining_effects(filter_resources(op->getResults()));

  if (!found_any_effect) {
    // We haven't collected any side effect but the op is potentially

    Type output_type = op.getResult().getType();
    uint32_t axis = CastI64ToI32(op.getDimension()).value();
    rewriter.replaceOpWithNewOp<TFL::ConcatenationOp>(
        op, output_type, op.getOperands(), axis,
        /*fused_activation_function=*/rewriter.getStringAttr("NONE"));
  }
};

// Rewrites quantized stablehlo.pad to tfl.padv2.
// tfl.dilate is introduced in between when interior padding exists.