avoid atomic contention on AsyncValue's refcount.
  }];

  let arguments = (ins
    TFTensorType:$operand
  );

  let results = (outs
    Variadic<TFTensorType>:$results
  );

  let assemblyFormat = "$operand attr-dict `:` `(` type($operand) `)` `->` `(` type($results) `)`";
}

def CreateOp: FallbackAsync_Op<"createop", [CoreRT_TypedAttributeTrait]> {

// -> tfl.batch_matmul(mhlo.transpose(mhlo.reshape(operand)), ...).
// Note:
// 1) Reshape/transpose are inserted because tfl.BatchMatMul requires
// size(contracting_dimensions) = 1 and size(output_dim) = 1, whereas
// mhlo.dot_general has no such restriction.
// 2) Inserted mhlo.reshape/transpose will be legalized to tf.reshape/transpose
// in LegalizeHloToTf (then from tf to tfl later).
// 3) If the operands are dynamic shaped tensors, mhlo.DynamicReshapeOp is

	cgofunc:   "cgo function",
}

// An operand represents an intermediate value during type checking.
// Operands have an (addressing) mode, the expression evaluating to
// the operand, the operand's type, a value for constants, and an id
// for built-in functions.
// The zero value of operand is a ready to use invalid operand.
type operand struct {
	mode operandMode
	expr syntax.Expr
	typ  Type

  return absl::OkStatus();
}

Status MlirAbstractOp::Create(ArrayRef<Value> operands,
                              OperationState** state) {
  state_->operands = llvm::to_vector<4>(operands);
  Builder builder(context_);

  if (current_ods_input_ != op_def_->input_arg_size())
    return InvalidArgument(absl::StrCat("Mismatch in operands number: got ",
                                        current_ods_input_, " expected ",

// commutativity of the AddV2: we want the FusedBatchNormV3 on any side.
// Also we need some native calls to handle the "hasOneUse" aspects and the
// optional extra operands for the AddV2 case.
struct ReluToFusedBatchNorm : public OpRewritePattern<ReluOp> {
  using OpRewritePattern<ReluOp>::OpRewritePattern;

  LogicalResult matchAndRewrite(ReluOp relu_op,

    parallel_execute = mlir::TF::BuildParallelExecuteOp(cluster_func, &builder);

  // Rewrite results before rewriting operands as `tf.TPUPartitionedInputV2`
  // resource handle results is an indicator for a partitioned resource
  // variable. These `tf.TPUPartitionedInputV2` will be removed when rewriting
  // the operands.
  for (Value result : parallel_execute.getExecuteOutputs()) {
    if (!result.hasOneUse()) continue;

to an aliased device name. Otherwise the device already set in tf_device.launch
is used instead.

Operands are replicated inputs and packed inputs.

replicated_inputs: each group of `n` inputs corresponds to an input for a single
individual replica and is mapped to a single region argument. Inside one group
the operands are matching in order the `devices` attribute. Each replicated
input must have compatible shapes and types.

#include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h"

namespace mlir {
namespace quant {

// Copied from tensorflow/compiler/mlir/lite/transforms/prepare_tf.cc.
// By removing identity ops, constant operands with dynamic shapes have static
// shape information which is necessary for correct pattern matching in this
// pass.
struct RemoveIdentity : public OpRewritePattern<TF::IdentityOp> {

    return rewriter.notifyMatchFailure(
        concat, "Concatenate op should have at least two operands");

  auto first = concat.getVal()[0].getDefiningOp<mhlo::SliceOp>();
  auto second = concat.getVal()[1].getDefiningOp<mhlo::SliceOp>();
  if (!first || !second)
    return rewriter.notifyMatchFailure(concat, "operands are not slice ops");
  if (first.getOperand() != second.getOperand())

//
// If the format specifier includes a %w verb with an error operand,
// the returned error will implement an Unwrap method returning the operand.
// If there is more than one %w verb, the returned error will implement an
// Unwrap method returning a []error containing all the %w operands in the
// order they appear in the arguments.
// It is invalid to supply the %w verb with an operand that does not implement