// Create new ConstantOp-ConvertOp-Operation sequences. At this moment,
      // old ConstantOp is guaranteed to have one F32->F16 convert op regardless
      // of its number of users.
      rewriter.setInsertionPointAfter(op);
      // create new F16 constant op in that location
      ConstantOp new_const = rewriter.create<ConstantOp>(
          op->getLoc(), new_result_type, new_value_attr);

  %5 = "tfl.mul"(%0, %1) {fused_activation_function = "NONE"} : (tensor<  f16>, tensor<  f16>) -> tensor<  f16>
  %6 = "tfl.mul"(%0, %3) {fused_activation_function = "NONE"} : (tensor<  f16>, tensor<4xf16>) -> tensor<4xf16>
  %7 = "tfl.mul"(%2, %1) {fused_activation_function = "NONE"} : (tensor<4xf16>, tensor<  f16>) -> tensor<4xf16>

  func.return %0 : tensor<4xcomplex<f64>>
}

func.func @f16() -> tensor<4xf16> {
  // CHECK-LABEL: @f16
  // CHECK: value = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf16>
  %0 = "tfl.pseudo_const"() { value = dense<[1.0, 2.0, 3.0, 4.0]> : tensor<4xf16> } : () -> tensor<4xf16>

  func.return %0 : tensor<4xcomplex<f64>>
}

func.func @f16() -> tensor<4xf16> {
  // CHECK-LABEL: @f16
  // CHECK: value = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf16>
  %0 = "tfl.pseudo_const"() { value = dense<[1.0, 2.0, 3.0, 4.0]> : tensor<4xf16> } : () -> tensor<4xf16>

include "mhlo/IR/hlo_ops.td"

def SignedIntTensor : TensorOf<[I1, I8, I16, I32, I64]>;
def UnsignedIntTensor : TensorOf<[UI8, UI16, UI32, UI64]>;

// IEEE compliant floating point tensors.
def IEEEFloatTensor : TensorOf<[F16, F32, F64]>;

//===----------------------------------------------------------------------===//
// BatchNorm op patterns.
//===----------------------------------------------------------------------===//

        tensor<256xf64>) -> tensor<4x256xf64>
  func.return %0 : tensor<4x256xf64>
}

// CHECK-LABEL: @batchNormInference_f16
// Validate that epsilon is properly down to f16
// CHECK-DAG: %[[EPS:.+]] = mhlo.constant dense<1.000000e+00> : tensor<256xf16>
func.func @batchNormInference_f16(
    %x: tensor<4x256xf16>, %scale: tensor<256xf16>, %offset: tensor<256xf16>,

    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.

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

      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",