return true;
  };
  return just_call(while_op.getBody()) && just_call(while_op.getCond());
}

bool IsCompatibleTypeWithTFLCastOp(Type type) {
  auto elemType = getElementTypeOrSelf(type);
  // F16, F32, F64, BF16 types are allowed.
  if (elemType.isBF16() || elemType.isF16() || elemType.isF32() ||
      elemType.isF64())
    return true;

  // I1, I4, I8, I16, I32, I64 types are allowed.

func.func @testAddNWrongOperandResultType(tensor<? x f16>, tensor<? x f16>, tensor<? x f16>) -> tensor<? x f16> {
^bb0(%arg0: tensor<? x f16>, %arg1: tensor<? x f16>, %arg2: tensor<? x f16>):
  // expected-error @+1 {{'tfl.add_n' op operand #0 must be variadic of tensor of 32-bit float or 32-bit signless integer}}
  %0 = "tfl.add_n"(%arg0, %arg1, %arg2): (tensor<? x f16>, tensor<? x f16>, tensor<? x f16>) -> tensor<? x f16>
  func.return %0 : tensor<? x f16>
}

  // The bytes of floats are stored little-endian.
  switch (elem_type.getIntOrFloatBitWidth()) {
    case 16: {
      assert(bytes_len % 2 == 0);
      // Supports both BF16 and F16.
      assert(elem_type.isF16() || elem_type.isBF16());
      int elem_count = bytes_len / 2;

      if (elem_type.isF16()) {
        std::vector<Eigen::half> values;
        values.reserve(elem_count);

          new_result_type, ArrayRef<Eigen::half>(new_values));

      // Create new ConstantOp-Dequantize-Operation sequences. At this moment,
      // old ConstantOp is guaranteed to have one F32->F16 cast regardless of
      // its number of users.
      rewriter.setInsertionPointAfter(op);
      auto new_const = rewriter.create<arith::ConstantOp>(
          op->getLoc(), new_result_type, new_value_attr);

  if (!input_ty) return getDataFormat();

  // For f16 data type on devices with Tensor Cores support NHWC data format
  // is up to ~2x faster.
  const bool is_f16 = input_ty.getElementType().isF16();
  if (is_f16 && CanUseTensorCores(devices)) return "NHWC";

  // For f32/f16 data type decision depends on the filter size in spatial

      Option<"unfold_batch_matmul_", "unfold_batchmatmul",
             "bool", "true",
             "Unfold BatchMatMul into individual MatMul ops.">,
      Option<"allow_bf16_and_f16_type_legalization_", "allow-bf16-and-f16-type-legalization",
             "bool", "false",
             "Allow bf16 type legalization.">,
      Option<"use_fake_quant_num_bits_", "use-fake-quant-num-bits",
             "bool", "false",

  let description = [{
    Fill the tensor with given value.
  }];

  let arguments = (ins TFL_I32OrI64Tensor:$dims,
                   TFL_TensorOf<[F32, F16, I32, I64, I1, QI8, QI16, TFL_Str]>:$input);

  let results = (outs TFL_TensorOf<[F32, F16, I32, I64, I1, QI8, QI16, TFL_Str]>:$result);

  let hasOptions = 0;
}

def TFL_FloorOp: TFL_Op<"floor", [
    Pure,

	"R31", // REGTMP

	"F0",
	"F1",
	"F2",
	"F3",
	"F4",
	"F5",
	"F6",
	"F7",
	"F8",
	"F9",
	"F10",
	"F11",
	"F12",
	"F13",
	"F14",
	"F15",
	"F16",
	"F17",
	"F18",
	"F19",
	"F20",
	"F21",
	"F22",
	"F23",
	"F24",
	"F25",
	"F26",
	"F27",
	"F28",
	"F29",
	"F30",

	MOVD	-16(R12), g
	MOVD	-8(R12), R31
	RET
TEXT runtime·elf_savefpr(SB),NOSPLIT|NOFRAME,$0
	// R0 holds the LR of the caller's caller, R1 holds save location
	FMOVD	F14, -144(R1)
	FMOVD	F15, -136(R1)
	FMOVD	F16, -128(R1)
	FMOVD	F17, -120(R1)
	FMOVD	F18, -112(R1)
	FMOVD	F19, -104(R1)
	FMOVD	F20, -96(R1)
	FMOVD	F21, -88(R1)
	FMOVD	F22, -80(R1)
	FMOVD	F23, -72(R1)
	FMOVD	F24, -64(R1)
	FMOVD	F25, -56(R1)

  // CHECK: %[[CAST:.*]] = mhlo.convert %arg0 : tensor<4x8xf16>
  // CHECK: %[[INITIAL:.*]] = mhlo.constant dense<0xFC00> : tensor<f16>
  // CHECK: %[[REDUCED:.*]] = mhlo.reduce(%[[CAST]] init: %[[INITIAL]]) applies mhlo.maximum across dimensions = [1] : (tensor<4x8xf16>, tensor<f16>) -> tensor<4xf16>
  // CHECK: %[[CAST_BACK:.*]] = mhlo.convert %[[REDUCED]] : tensor<4xf16>