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

			for _, i := range s.i {
				for _, j := range s.i {
					fmt.Fprintf(w, "itd%s{a: %d, b: %d, add: %d, sub: %d, mul: %d", s.sn, i, j, icast(i+j, s), icast(i-j, s), icast(i*j, s))
					if j != 0 {
						fmt.Fprintf(w, ", div: %d, mod: %d", icast(i/j, s), icast(i%j, s))
					}
					fmt.Fprint(w, "},\n")
				}
			}
			fmt.Fprintf(w, "}\n")
		}
	}

// CHECK-SAME: (tensor<*xf32>) -> tensor<*xf32>
// CHECK: %[[round:.*]] = "tf.Round"(%[[avgpool_f32]])
// CHECK: %[[icast:.*]] = "tf.Cast"(%[[round]]) <{Truncate = false}> : (tensor<*xf32>) -> tensor<*xi8>
// CHECK: %[[sc2:.*]] = "quantfork.scast"(%[[icast]])
// CHECK: %[[dq:.*]] = "quantfork.dcast"(%[[sc2]]) : (tensor<*x!quant.uniform<i8:f32, 5.000000e-02:-10>>)

    auto dequantize_func = [&](const APInt& ap_int_value) -> APInt {
      const int64_t int_value = ap_int_value.getSExtValue();

      const float real = (int_value - zp) * scale;

      auto real_int = absl::bit_cast<int32_t>(real);
      return APInt(/*numBits=*/32, real_int);
    };

    auto dequant_values =
        mlir::cast<DenseIntOrFPElementsAttr>(input_values)