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

// -----

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

  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

	IsNarrowerThan(mask BitMask) bool
	IsLessThan(mask BitMask) bool
	IsGreaterThan(mask BitMask) bool
	String() string
	Count() int
	GetBits() []int
}

type bitMask uint64

// NewEmptyBitMask creates a new, empty BitMask
func NewEmptyBitMask() BitMask {
	s := bitMask(0)
	return &s
}

// NewBitMask creates a new BitMask
func NewBitMask(bits ...int) (BitMask, error) {
	s := bitMask(0)
	err := (&s).Add(bits...)

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

            "Less", "LessEqual", "SigmoidGrad", "SoftplusGrad", "SoftsignGrad",
            "TanhGrad", "Pow", "SquaredDifference", "ApproximateEqual",
            // Others
            "AddN", "Bitcast", "Cast", "ClipByValue", "Const", "Empty",
            "Identity", "IdentityN", "Relu", "Relu6", "ReluGrad", "Relu6Grad",
            "LeakyReluGrad", "Elu", "EluGrad", "Selu", "SeluGrad", "Select",

[..., sizeof(`type`)/sizeof(`T`)] to [...].

tf.bitcast() and tf.cast() work differently when real dtype is casted as a complex dtype
(e.g. tf.complex64 or tf.complex128) as tf.cast() make imaginary part 0 while tf.bitcast()
gives module error.
For example,

Example 1:

>>> a = [1., 2., 3.]
>>> equality_bitcast = tf.bitcast(a, tf.complex128)
Traceback (most recent call last):
...

  // CHECK: return %0 : tensor<8xf32>
}


// -----

// CHECK-LABEL: testBitcast
func.func @testBitcast(%arg0: tensor<8xui32>) -> tensor<8xi32> {
  // CHECK: "tfl.bitcast"(%arg0)
  %0 = "tfl.bitcast"(%arg0) : (tensor<8xui32>) -> tensor<8xi32>
  func.return %0 : tensor<8xi32>
  // CHECK: return %0 : tensor<8xi32>
}

// -----

// CHECK-LABEL: testBitwiseXor

		NUMADistances: distances,
	}

	return numaInfo, nil
}

func (n *NUMAInfo) Narrowest(m1 bitmask.BitMask, m2 bitmask.BitMask) bitmask.BitMask {
	if m1.IsNarrowerThan(m2) {
		return m1
	}
	return m2
}

func (n *NUMAInfo) Closest(m1 bitmask.BitMask, m2 bitmask.BitMask) bitmask.BitMask {
	// If the length of both bitmasks aren't the same, choose the one that is narrowest.
	if m1.Count() != m2.Count() {