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>

  }
};

void TestBitcastOp(Tensor* input_tensor, DataType out_type,
                   TensorShape expected_shape, error::Code expected_code) {
  Status status;
  NodeDef def;
  def.set_op("Bitcast");
  def.set_device(DEVICE_CPU);

  AttrValue typeAttr;
  SetAttrValue(input_tensor->dtype(), &typeAttr);

  AttrValue outTypeAttr;
  SetAttrValue(out_type, &outTypeAttr);

  {
    auto* builder = TF_NewKernelBuilder("Bitcast", tensorflow::DEVICE_CPU,
                                        &BitcastOp_Create, &BitcastOp_Compute,
                                        &BitcastOp_Delete);
    TF_RegisterKernelBuilder("BitcastOp", builder, status);
    CHECK_EQ(TF_OK, TF_GetCode(status))
        << "Error while registering bitcast kernel";
  }

#if GOOGLE_CUDA || TENSORFLOW_USE_ROCM
  {

        "//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",
    ],
)

                   ],
        )

op_names = [
    "BatchMatrixBandPart",
    "BatchMatrixDiag",
    "BatchMatrixDiagPart",
    "BatchMatrixSetDiag",
    "BatchToSpace",
    "BatchToSpaceND",
    "Bitcast",
    "BroadcastArgs",
    "BroadcastTo",
    "CheckNumerics",
    "ConcatV2",
    "ConjugateTranspose",
    "DebugGradientIdentity",
    "DeepCopy",
    "DepthToSpace",
    "Dequantize",

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

  (TF_BiasAddOp:$dest $arg0, $arg1,
    ConstantAttr<TF_ConvnetDataFormatAttr, "\"NHWC\"">),
  [], [(CopyAttrs $src, $dest)]>;

//===----------------------------------------------------------------------===//
// Bitcast op patterns.
//===----------------------------------------------------------------------===//

// Parameter attributes are part of function attributes, which is not affected

func.func @testBitcast(%arg0: tensor<3x4xui16>) -> tensor<3x4x!tf_type.quint16> {
  %0 = "tf.Bitcast"(%arg0) : (tensor<3x4xui16>) -> tensor<3x4x!tf_type.quint16>
  func.return %0 : tensor<3x4x!tf_type.quint16>
}

// -----

// CHECK-LABEL: func @testBitcast_v2
func.func @testBitcast_v2(%arg0: tensor<3x2xi16>) -> tensor<3xi32> {
  %0 = "tf.Bitcast"(%arg0) : (tensor<3x2xi16>) -> tensor<3xi32>
  func.return %0 : tensor<3xi32>
}

// -----