}

  // Collects all candidate ops for quantization, which is the operand of
  // `quantize_op`. If successful, this always returns one element which is the
  // operand of `quantize_op`.
  FailureOr<SmallVector<Operation*>> CollectCandidateOps(
      QuantizeOpT quantize_op) const {
    Value operand = quantize_op->getOperand(0);
    if (QuantizedType::getQuantizedElementType(operand.getType())) {

    (BinaryOp:$result (TFL_TileOp $input, (Arith_ConstantOp $tile)),
     $operand, $act_func),
    (BinaryOp $input, $operand, $act_func),
  [(OperandsBroadcastToOutputType $input, $operand, $result),
   (HasRankAtMost<4> $input),
   (HasRankAtMost<4> $operand)]>;

  def FuseTileBroadcastToBinaryOp2#BinaryOp : Pat<
    (BinaryOp:$result $operand,
      (TFL_TileOp $input, (Arith_ConstantOp $tile)), $act_func),

// from the spec
var (
	s uint    = 33
	u         = 1.0 << s // ERROR "invalid operation|shift of non-integer operand"
	v float32 = 1 << s   // ERROR "invalid"
)

// non-constant shift expressions
var (
	e1       = g(2.0 << s) // ERROR "invalid|shift of non-integer operand"
	f1       = h(2 << s)   // ERROR "invalid"
	g1 int64 = 1.1 << s    // ERROR "truncated|must be integer"
)

    return failure();
  }

  MutableArrayRef<BlockArgument> operands =
      entry_func_op.getBody().getArguments();
  // Function must have input, filter, and optionally bias.
  if (operands.size() != 2 && operands.size() != 3) {
    LLVM_DEBUG(llvm::dbgs() << GemmStyleOp::getOperationName()
                            << " op function should have 2 or 3 operands.\n");
    return failure();
  }
  return success();
}

// -> 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

        decomposed_partitioned_call_callees) {
  for (OpOperand& operand : op.getOpOperands()) {
    if (getElementTypeOrSelf(operand.get().getType()).isa<TF::ResourceType>()) {
      return op.emitOpError()
             << "found unexpected type " << operand.get().getType()
             << " of operand #" << operand.getOperandNumber()
             << ", resource type operands are expected to have been "

    new_types.back() = rewriter.getType<tf_executor::ControlType>();

    // Control operand is attached on tf_executor::IslandOp.
    llvm::SmallVector<Value> island_control_operands;
    llvm::SmallVector<Value> inner_op_operands;

    for (Value value : operands) {
      // Because of the property of graph region, the control operands may
      // not have been converted to tf_executor::ControlType.

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

    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]> {

}

// Identical reports whether x and y are identical types.
// Receivers of [Signature] types are ignored.
//
// Predicates such as [Identical], [Implements], and
// [Satisfies] assume that both operands belong to a
// consistent collection of symbols ([Object] values).
// For example, two [Named] types can be identical only if their
// [Named.Obj] methods return the same [TypeName] symbol.