// operand's tensor types.
def QI4 : QuantizedType<"Uniform", [4], 1>;
def QUI8 : QuantizedType<"Uniform", [8], 0>;
def QI8 : QuantizedType<"Uniform", [8], 1>;
def QUI16 : QuantizedType<"Uniform", [16], 0>;
def QI16 : QuantizedType<"Uniform", [16], 1>;
def QUI32 : QuantizedType<"Uniform", [32], 0>;
def QI32 : QuantizedType<"Uniform", [32], 1>;

  // CHECK:  %cst = arith.constant dense<0> : tensor<1xi32>
  // CHECK:  %4 = "tfl.reduce_min"(%arg0, %cst) <{keep_dims = false}> : (tensor<2xi16>, tensor<1xi32>) -> tensor<i16>
  // CHECK:  %5 = "tfl.arg_min"(%arg0, %cst) : (tensor<2xi16>, tensor<1xi32>) -> tensor<i32>
  // CHECK:  return %4, %5 : tensor<i16>, tensor<i32>
}


// CHECK-LABEL: func @convert_argmin_constant

// CHECK-SAME:                        %[[VAL_0:.*]]: tensor<1xi1>,
// CHECK-SAME:                        %[[VAL_1:.*]]: tensor<1x2xi1>) -> tensor<1x2xi1> {
// CHECK:           %[[VAL_2:.*]] = "tf.LogicalAnd"(%[[VAL_0]], %[[VAL_1]]) : (tensor<1xi1>, tensor<1x2xi1>) -> tensor<1x2xi1>
// CHECK:           return %[[VAL_2]] : tensor<1x2xi1>
// CHECK:         }
func.func @and_broadcast(%arg0: tensor<1xi1>, %arg1: tensor<1x2xi1>) -> tensor<1x2xi1> {

	ui12 = ui3 / ui0 /* ERROR "division by zero" */
	ui13 = 1 % 0 /* ERROR "division by zero" */
	ui14 = ui1 % 0 /* ERROR "division by zero" */
	ui15 = ui3 % ui0 /* ERROR "division by zero" */

	ui16 = ui2 & ui3
	ui17 = ui2 | ui3
	ui18 = ui2 ^ ui3
	ui19 = 1 /* ERROR "invalid operation" */ % 1.0

	// floating point values
	uf0 = 0.
	uf1 = 1.
	uf2 = 4.2e1

// is ui16. ui16 storage type for quantized type is not compatible with
// `tfl.quantize`.

func.func @uniform_quantize_op_uint16_output(%arg: tensor<2x2xf32>) -> tensor<2x2x!quant.uniform<ui16:f32, 3.000000e+0:127>> {
  %0 = stablehlo.uniform_quantize %arg : (tensor<2x2xf32>) -> tensor<2x2x!quant.uniform<ui16:f32, 3.000000e+0:127>>
  return %0 : tensor<2x2x!quant.uniform<ui16:f32, 3.000000e+0:127>>
}

}

func.func @broadcast_to_i16_low_dim(%input: tensor<3xi16>, %shape: tensor<2xi32>) -> tensor<3x3xi16> {
  %0 = "tf.BroadcastTo"(%input, %shape) : (tensor<3xi16>, tensor<2xi32>) -> tensor<3x3xi16>
  func.return %0: tensor<3x3xi16>

// CHECK-LABEL: broadcast_to_i16_low_dim
// CHECK:    %0 = "tf.BroadcastTo"(%arg0, %arg1) : (tensor<3xi16>, tensor<2xi32>) -> tensor<3x3xi16>
// CHECK:    return %0 : tensor<3x3xi16>
}

}

func TestEmptyMapWithInterfaceKey(t *testing.T) {
	var (
		b    bool
		i    int
		i8   int8
		i16  int16
		i32  int32
		i64  int64
		ui   uint
		ui8  uint8
		ui16 uint16
		ui32 uint32
		ui64 uint64
		uipt uintptr
		f32  float32
		f64  float64
		c64  complex64
		c128 complex128
		a    [4]string
		s    string
		p    *int
		up   unsafe.Pointer

include "stablehlo/dialect/ChloOps.td"
include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.td"
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.

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

// -----

// CHECK-LABEL: testGatherNdI16Indices
func.func @testGatherNdI16Indices(%arg0 : tensor<?xf32>, %arg1 : tensor<?xi16>) -> tensor<?xf32> {
  %0 = "tfl.gather_nd"(%arg0, %arg1) {axis = 1 : i32}: (tensor<?xf32>,tensor<?xi16>) -> tensor<?xf32>
  func.return %0 : tensor<?xf32>
}

// -----

// CHECK-LABEL: testAbs
func.func @testAbs(tensor<? x f32>) -> tensor<? x f32> {

      ConversionPatternRewriter& rewriter) const final {
    RankedTensorType type =
        mlir::dyn_cast_or_null<RankedTensorType>(iota_op.getType());
    // TF::RangeOp doesn't support UI16.
    if (!type || type.getElementType().isUnsignedInteger(16)) return failure();

    const uint64_t dimension = iota_op.getIotaDimension();
    Type element_type = type.getElementType();