TF_RegisterOpDefinition(op_builder, status);
  CHECK_EQ(TF_GetCode(status), TF_OK)
      << "Bitcast op registration failed: " << TF_Message(status);
  TF_DeleteStatus(status);
}

TF_ATTRIBUTE_UNUSED static bool IsBitcastOpRegistered = []() {
  if ((&TF_NewStatus != nullptr) && SHOULD_REGISTER_OP("Bitcast")) {
    RegisterBitcastOp();
  }
  return true;

func.func @bitcast_convert_i32_f32(%arg0: tensor<1x256128xi32>) -> tensor<1x256128xf32> {
  // CHECK: %[[BITCAST:.*]] = stablehlo.bitcast_convert %arg0 : (tensor<1x256128xi32>) -> tensor<1x256128xf32>
  // CHECK: %[[CONVERT:.*]] = stablehlo.convert %[[BITCAST]] : (tensor<1x256128xf32>) -> tensor<1x256128xbf16>
  // CHECK: return %[[CONVERT]] : tensor<1x256128xbf16>

        "//tensorflow/c:tf_tensor",
        "//tensorflow/core:framework",
        "//tensorflow/core:lib",
        "//tensorflow/core:protos_all_cc",
    ],
)

tf_gen_op_libs(
    op_lib_names = ["bitcast"],
    deps = [
        "//tensorflow/c:ops",
        "//tensorflow/c:tf_datatype",
        "//tensorflow/c:tf_status",
        "//tensorflow/c:tf_tensor",
        "//tensorflow/core:lib",
    ],
)

// CHECK:  return %[[RES0]] : tensor<8xf32>
}

func.func @bitcast(%arg0: tensor<8xi32>) -> tensor<8xui32> {
  %0 = "tf.Bitcast"(%arg0) : (tensor<8xi32>) -> tensor<8xui32>
  func.return %0 : tensor<8xui32>

// CHECK-LABEL: bitcast
// CHECK: %[[RES0:.*]] = "tfl.bitcast"(%arg0) : (tensor<8xi32>) -> tensor<8xui32>
// CHECK:  return %[[RES0]] : tensor<8xui32>
}

func.func @testDoubleBitcast(%arg0: tensor<8x16x32x64xf32>) -> tensor<8x16x32x64xi32> {
  %0 = "tf.Bitcast"(%arg0) : (tensor<8x16x32x64xf32>) -> tensor<8x16x32x64x2xi16>
  %1 = "tf.Bitcast"(%0) {device = "/job:localhost/replica:0/task:0/device:GPU:0"} : (tensor<8x16x32x64x2xi16>) -> tensor<8x16x32x64xi32>
  func.return %1: tensor<8x16x32x64xi32>

// CHECK: %[[BITCAST1:.+]] = "tfl.bitcast"(%[[ARG1]]) : (tensor<i64>) -> tensor<1xi64>
// CHECK: %[[MIN1:.+]] = "tfl.minimum"(%[[BITCAST1]], %[[MAX1]]) : (tensor<1xi64>, tensor<1xi64>) -> tensor<1xi64>
// CHECK: %[[BITCAST2:.+]] = "tfl.bitcast"(%[[ARG2]]) : (tensor<i64>) -> tensor<1xi64>
// CHECK: %[[MIN2:.+]] = "tfl.minimum"(%[[BITCAST2]], %[[MAX2]]) : (tensor<1xi64>, tensor<1xi64>) -> tensor<1xi64>

  func.return %0 : tensor<?xf32>
}

// -----

// CHECK-LABEL: func @bitcast
func.func @bitcast(%arg0: tensor<2xf32>) -> tensor<2xf32> {
  // CHECK:  mhlo.bitcast_convert %arg0 : (tensor<2xf32>) -> tensor<2xf32>
  %0 = "tf.Bitcast"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
  func.return %0 : tensor<2xf32>
}

// -----

  DYNAMIC_UPDATE_SLICE = 151,
  RELU_0_TO_1 = 152,
  UNSORTED_SEGMENT_PROD = 153,
  UNSORTED_SEGMENT_MAX = 154,
  UNSORTED_SEGMENT_SUM = 155,
  ATAN2 = 156,
  UNSORTED_SEGMENT_MIN = 157,
  SIGN = 158,
  BITCAST = 159,
  BITWISE_XOR = 160,
  RIGHT_SHIFT = 161,
  // All Operators start with STABLEHLO_ prefixes are subject to change
  // Many of the ops below can not be executed by TFlite runtime

  let hasFolder = 1;
}

def TFL_BitcastOp : TFL_Op<"bitcast", [Pure]> {
  let summary = "Bitcast operator";

  let description = [{
    Bitcasts a tensor from one type to another.
  }];

  let arguments = (ins AnyTensor:$input);

  let results = (outs AnyTensor:$output);

  // TFLite's bitcast bitop does not utilize options, instead derives types
  // from the TfLiteTensors.

  func.return %0 : tensor<?xf32>
}

// CHECK-LABEL:   func @bitcast(
// CHECK-SAME:                  %[[VAL_0:.*]]: tensor<2xf32>) -> tensor<2xf32> {
// CHECK:           %[[VAL_1:.*]] = "tf.Bitcast"(%[[VAL_0]]) : (tensor<2xf32>) -> tensor<2xf32>
// CHECK:           return %[[VAL_1]] : tensor<2xf32>
// CHECK:         }
func.func @bitcast(%arg0: tensor<2xf32>) -> tensor<2xf32> {