}

func.func @addV2I16(%arg0: tensor<1xi16>, %arg1: tensor<1xi16>) -> tensor<1xi16> {
  %0 = "tf.AddV2"(%arg0, %arg1) : (tensor<1xi16>, tensor<1xi16>) -> tensor<1xi16>
  func.return %0 : tensor<1xi16>

// CHECK-LABEL: addV2I16
// CHECK:  tfl.add %arg0, %arg1 {fused_activation_function = "NONE"} : tensor<1xi16>
}

  let arguments = (ins
     TFL_TensorOf<[I8, I16, I32, I64, UI8, UI16, UI32, UI64, F32, F64]>:$input,
     TFL_TensorOf<[I32]>:$dilations,
     TFL_0DTensorOf<[I8, I16, I32, I64, UI8, UI16, UI32, UI64, F32, F64]>:$padding_value
  );

  let results = (outs TFL_TensorOf<[I8, I16, I32, I64, UI8, UI16, UI32, UI64, F32, F64]>:$output);
}

def TFL_AddOp : TFL_Op<"add", [

}

// -----

// CHECK-LABEL: func @testMul
func.func @testMul(%arg0: tensor<2xui16>) -> (tensor<2xui16>) {
  %0 = "tf.Mul"(%arg0, %arg0) {T = "tfdtype$DT_UINT16", device = "/device:CPU:0", name = "Mul"} : (tensor<2xui16>, tensor<2xui16>) -> tensor<2xui16>
  func.return %0 : tensor<2xui16>
}

// -----

// Test error message for incompatible element types.

// 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.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();