func.func @TfAdd(%arg0: tensor<4xf64>, %arg1: tensor<4xf64>) -> tensor<4xf64> {
  %0 = "tfl.custom"(%arg0, %arg1) {
    custom_code = "FlexAdd",
    custom_option = #tfl<const_bytes : "0x03416464001412034164641A001A002A070A015412023002320000021B171414042801">
  } : (tensor<4xf64>, tensor<4xf64>) -> tensor<4xf64>

// CHECK: "tf.Add"(%arg0, %arg1) {T = f64}  : (tensor<4xf64>, tensor<4xf64>) -> tensor<4xf64>
  func.return %0 : tensor<4xf64>
}

// RUN: flatbuffer_translate -mlir-to-tflite-flatbuffer %s -emit-select-tf-ops -o - | flatbuffer_to_string - | FileCheck %s

func.func @main(tensor<4xf64>, tensor<4xf64>) -> tensor<4xf64> {
^bb0(%arg0: tensor<4xf64>, %arg1: tensor<4xf64>):
// CHECK:  {
// CHECK-NEXT:  version: 3,
// CHECK-NEXT:  operator_codes: [ {
// CHECK-NEXT:    deprecated_builtin_code: 32,
// CHECK-NEXT:    custom_code: "FlexAdd",
// CHECK-NEXT:    builtin_code: CUSTOM

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

func.func @f64() -> tensor<4xf64> {
  // CHECK-LABEL: @f64
  // CHECK: value = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf64>
  %0 = "tfl.pseudo_const"() { value = dense<[1.0, 2.0, 3.0, 4.0]> : tensor<4xf64> } : () -> tensor<4xf64>
  func.return %0 : tensor<4xf64>
}

func.func @i8() -> tensor<4xi8> {
  // CHECK-LABEL: @i8

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

func.func @f64() -> tensor<4xf64> {
  // CHECK-LABEL: @f64
  // CHECK: value = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf64>
  %0 = "tfl.pseudo_const"() { value = dense<[1.0, 2.0, 3.0, 4.0]> : tensor<4xf64> } : () -> tensor<4xf64>
  func.return %0 : tensor<4xf64>
}

func.func @i8() -> tensor<4xi8> {
  // CHECK-LABEL: @i8

    // CHECK-NOT: tf.BesselI0e
    %0 = "tf.BesselI0e"(%arg0) : (tensor<3xf16>) -> (tensor<3xf16>)
    %1 = "tf.BesselI0e"(%arg1) : (tensor<3xf32>) -> (tensor<3xf32>)
    %2 = "tf.BesselI0e"(%arg2) : (tensor<3xf64>) -> (tensor<3xf64>)
    func.return %0, %1, %2 : tensor<3xf16>, tensor<3xf32>, tensor<3xf64>
  }

  // CHECK-LABEL: bessel_i1e

}

// CHECK-LABEL: func @testEmptyFloat
func.func @testEmptyFloat() -> (tensor<5xf64>) {
  %0 = "tf.Const"() { value = dense<5> : tensor<i32> } : () -> tensor<i32>

  // CHECK: [[VAL:%.+]] = "tf.Const"() <{value =  dense<0.000000e+00> : tensor<5xf64>}>
  // CHECK: return [[VAL]]
  %1 = "tf.Empty"(%0) : (tensor<i32>) -> (tensor<5xf64>)
  func.return %1 : tensor<5xf64>
}

// CHECK-LABEL: func @testEmptyf16

  %cst_2 = arith.constant dense<[8, 6, 6, 18]> : tensor<4xi64>
  %98 = "tf.Reshape"(%97, %cst_2) : (tensor<1x8x6x1x6x1x1x18xbf16>, tensor<4xi64>) -> tensor<8x6x6x18xbf16>
  func.return %98 : tensor<8x6x6x18xbf16>

  // CHECK-DAG: %[[CST:.*]] = "tf.Const"() <{value = dense<[8, 1, 1, 18]> : tensor<4xi64>}> : () -> tensor<4xi64>
  // CHECK-DAG: %[[CST1:.*]] =  "tf.Const"() <{value = dense<[8, 6, 6, 18]> : tensor<4xi64>}> : () -> tensor<4xi64>

  // CHECK: %[[CST:.*]] = "tf.Const"() <{value = dense<[0, 2, 3, 1]> : tensor<4xi64>}>
  // CHECK: %[[R0:.*]] = "tf.Transpose"(%[[ARG0]], %[[CST]])
  // CHECK: %[[R1:.*]] = "tf.BiasAdd"(%[[R0]], %[[ARG1]]) <{data_format = "NHWC"}> {device = ""}
  // CHECK: %[[CST_0:.*]] = "tf.Const"() <{value = dense<[0, 3, 1, 2]> : tensor<4xi64>}>
  // CHECK: "tf.Transpose"(%[[R1]], %[[CST_0]])

// -----

func.func @foldDivF64() -> tensor<2xf64> {
  %0 = stablehlo.constant dense<[2.0, 3.0]> : tensor<2xf64>
  %1 = stablehlo.constant dense<[4.0, 6.0]> : tensor<2xf64>
  %2 = stablehlo.divide %0, %1 : tensor<2xf64>
  return %2 : tensor<2xf64>
}

// CHECK-LABEL: foldDivF64
// CHECK: stablehlo.constant dense<5.000000e-01> : tensor<2xf64>

// -----

func.func @foldDivI32() -> tensor<2xi32> {

// CHECK-LABEL: func @transposeConv2D
func.func @transposeConv2D(%input: tensor<1x32x32x3xf32>, %filter: tensor<1x1x3x8xf32>) -> tensor<1x7x7x8xf32> {

  // CHECK: %[[ARG_PERM:.*]] = "tf.Const"() <{value = dense<[0, 3, 1, 2]> : tensor<4xi64>}>
  // CHECK: %[[ARG_TRANSPOSE:[0-9]*]] = "tf.Transpose"(%arg0, %[[ARG_PERM]])

  // CHECK: %[[CONV2D:[0-9]*]] = "tf.Conv2D"(%[[ARG_TRANSPOSE]], %arg1)
  // CHECK-SAME: data_format = "NCHW"