Output a = ops::Const(scope.WithOpName("a"), 10.0f, {});
    Output read_var;
    if (use_resource) {
      Output var =
          ops::VarHandleOp(scope.WithOpName("var"), DataType::DT_FLOAT, {});
      if (use_identity) {
        Output identity = ops::Identity(scope.WithOpName("identity"), var);
        read_var =
            ops::ReadVariableOp(scope.WithOpName("var/Read/ReadVariableOp"),

    assign_var.erase();
  }

  for (OpOperand& operand : cluster_func->getOpOperands()) {
    auto read_var = llvm::dyn_cast_or_null<TF::ReadVariableOp>(
        operand.get().getDefiningOp());
    if (!read_var) continue;
    auto partitioned_input =
        llvm::dyn_cast_or_null<TF::TPUPartitionedInputV2Op>(
            read_var.getResource().getDefiningOp());
    if (!partitioned_input || !AllResourceTypesHaveSubtypes(

// CHECK: tf_device.replicate([%[[ARG_0]], %[[ARG_1]]] as %[[RI:[a-z0-9]*]]: tensor<*x!tf_type.resource>)
// CHECK:   %[[READ_VAR:[0-9]*]] = "tf.ReadVariableOp"(%[[RI]])
// CHECK:   %[[OP_A:[0-9]*]] = "tf.opA"(%[[READ_VAR]], %[[VAR_SHAPE]], %[[RI]])
// CHECK:   tf_device.return %[[READ_VAR]], %[[VAR_SHAPE]], %[[OP_A]]

// CHECK-LABEL: func @replicate_arg_shape_with_packed

	REG_A2
	REG_A3
	REG_A4
	REG_A5
	REG_A6
	REG_A7

	REG_MSR
	REG_FPSCR
	REG_CR

	REG_SPECIAL = REG_CR0

	REG_CRBIT0 = REG_CR0LT // An alias for a Condition Register bit 0

	REG_SPR0 = obj.RBasePPC64 + 1024 // first of 1024 registers

	REG_XER = REG_SPR0 + 1
	REG_LR  = REG_SPR0 + 8
	REG_CTR = REG_SPR0 + 9

	REGZERO = REG_R0 /* set to zero */
	REGSP   = REG_R1
	REGSB   = REG_R2

		register[obj.Rconv(i)] = int16(i)
	}
	for i := ppc64.REG_MSR; i <= ppc64.REG_CR; i++ {
		register[obj.Rconv(i)] = int16(i)
	}
	for i := ppc64.REG_CR0LT; i <= ppc64.REG_CR7SO; i++ {
		register[obj.Rconv(i)] = int16(i)
	}
	register["CR"] = ppc64.REG_CR
	register["XER"] = ppc64.REG_XER
	register["LR"] = ppc64.REG_LR
	register["CTR"] = ppc64.REG_CTR
	register["FPSCR"] = ppc64.REG_FPSCR
	register["MSR"] = ppc64.REG_MSR

)

gentbl_cc_library(
    name = "tensorflow_reduce_patterns_inc_gen",
    tbl_outs = [
        (
            ["-gen-rewriters"],
            "reducer/tf_reduce_patterns.inc",
        ),
    ],
    tblgen = "@llvm-project//mlir:mlir-tblgen",
    td_file = "reducer/tf_mlir_reduce_patterns.td",
    deps = [
        "//tensorflow/compiler/mlir/tensorflow:tensorflow_ops_td_files",
    ],
)

cc_library(

      EraseFuncOpArguments(func, args_to_erase);

      func.eraseResults(result_indices_to_erase);
    } else if (auto read_var = dyn_cast<TF::ReadVariableOp>(user_op)) {
      // Read variables was already replaced by constant op. Just remove the op.
      read_var->erase();
    } else if (auto while_op = dyn_cast<TF::WhileOp>(user_op)) {
      GetUpdatedWhileOp<TF::WhileOp>(

  %2 = mhlo.constant dense<0> : tensor<i32>
  %3:2 = mhlo.reduce(%arg0 init: %1), (%0 init: %2) across dimensions = [0] : (tensor<2xi1>, tensor<2xi32>, tensor<i1>, tensor<i32>) -> (tensor<i1>, tensor<i32>)
    reducer(%arg1: tensor<i1>, %arg3: tensor<i1>) (%arg2: tensor<i32>, %arg4: tensor<i32>)  {
    %4 = mhlo.compare  GT, %arg1, %arg3 : (tensor<i1>, tensor<i1>) -> tensor<i1>
    %5 = mhlo.select %4, %arg1, %arg3 : tensor<i1>, tensor<i1>

	Ac_minflt                 uint64
	Ac_majflt                 uint64
	Coremem                   uint64
	Virtmem                   uint64
	Hiwater_rss               uint64
	Hiwater_vm                uint64
	Read_char                 uint64
	Write_char                uint64
	Read_syscalls             uint64
	Write_syscalls            uint64
	Read_bytes                uint64
	Write_bytes               uint64

	Ac_minflt                 uint64
	Ac_majflt                 uint64
	Coremem                   uint64
	Virtmem                   uint64
	Hiwater_rss               uint64
	Hiwater_vm                uint64
	Read_char                 uint64
	Write_char                uint64
	Read_syscalls             uint64
	Write_syscalls            uint64
	Read_bytes                uint64
	Write_bytes               uint64