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

func.func @main(tensor<4x5xbf16>) -> tensor<4x5xbf16> {
^bb0(%arg0: tensor<4x5xbf16>):

// CHECK:  {
// CHECK-NEXT:      version: 3,
// CHECK-NEXT:      operator_codes: [ {
// CHECK-NEXT:        deprecated_builtin_code: 53,
// CHECK-NEXT:        version: 7,
// CHECK-NEXT:        builtin_code: CAST
// CHECK-NEXT:      } ],

func.func @main(%arg0: tensor<i32>, %arg1: tensor<1xf32>) -> tensor<1xf32> {
  %0:2 = "tfl.while"(%arg0, %arg1) ({
  ^bb0(%arg2: tensor<*xi32>, %arg3: tensor<*xf32>):
    %1 = func.call @cond(%arg2, %arg3) : (tensor<*xi32>, tensor<*xf32>) -> tensor<i1>
    "tfl.yield"(%1) : (tensor<i1>) -> ()
  },  {
  ^bb0(%arg2: tensor<*xi32>, %arg3: tensor<*xf32>):

// RUN: flatbuffer_translate -mlir-to-tflite-flatbuffer %s -o - | flatbuffer_to_string - | FileCheck %s

func.func @main(tensor<3x!quant.uniform<i8:f32, 0.1>>) -> tensor<3x!quant.uniform<i8:f32, 0.1>> {
^bb0(%arg0: tensor<3x!quant.uniform<i8:f32, 0.1>>):
  // CHECK:      {
  // CHECK-NEXT:  version: 3,
  // CHECK-NEXT:  operator_codes: [ {
  // CHECK-NEXT:    deprecated_builtin_code: 18,
  // CHECK-NEXT:    version: 2,
  // CHECK-NEXT:    builtin_code: MUL

// RUN: flatbuffer_translate -mlir-to-tflite-flatbuffer %s -o - | flatbuffer_to_string - | FileCheck %s

func.func @main(tensor<3x!quant.uniform<i8:f32, 1.0>>) -> tensor<3x!quant.uniform<i8:f32, 1.0>> {
^bb0(%arg0: tensor<3x!quant.uniform<i8:f32, 1.0>>):
  // CHECK:      {
  // CHECK-NEXT:  version: 3,
  // CHECK-NEXT:  operator_codes: [ {
  // CHECK-NEXT:    deprecated_builtin_code: 18,
  // CHECK-NEXT:    version: 3,
  // CHECK-NEXT:    builtin_code: MUL

// RUN: flatbuffer_translate -mlir-to-tflite-flatbuffer %s -o - | flatbuffer_to_string - | FileCheck %s

func.func @main(tensor<3x2xf32>) -> tensor<3x2xi32> {
^bb0(%arg0: tensor<3x2xf32>):
  // CHECK:      {
  // CHECK-NEXT:     version: 3,
  // CHECK-NEXT:     operator_codes: [ {
  // CHECK-NEXT:       deprecated_builtin_code: 127,
  // CHECK-NEXT:       version: 1,
  // CHECK-NEXT:       builtin_code: BUCKETIZE
  // CHECK-NEXT:     } ],

    %10 = "quantfork.dcast"(%9) : (tensor<2x3x1x3x!quant.uniform<i8:f32, 3.000000e-01:1>>) -> tensor<2x3x1x3xf32>
    %11 = "stablehlo.reduce_window"(%10, %3) ({
    ^bb0(%arg1: tensor<f32>, %arg2: tensor<f32>):
      %14 = stablehlo.maximum %arg1, %arg2 : tensor<f32>
      stablehlo.return %14 : tensor<f32>

  %tmp1 = "tfl.add"(%const, %const) { fused_activation_function = "NONE" } : (tensor<1xf32>,tensor<1xf32>) -> (tensor<1xf32>)
  %tmp2 = "tfl.custom_tf"(%tmp1, %tmp1) ({
    ^bb0(%arg0: tensor<1xf32>, %arg1: tensor<1xf32>) :
      %wrap_result = "tf.AddV2"(%arg0, %arg1) { device = "" } : (tensor<1xf32>, tensor<1xf32>) -> tensor<1xf32>
      "tfl.yield"(%wrap_result): (tensor<1xf32>)->()

  // CHECK: tf.WhileRegion
  %while = "tf.WhileRegion"(%max_size) ({
    // CHECK: ^bb0(%[[BARG0:.*]]: tensor<i32>
    ^bb0(%barg0: tensor<i32>):
     // CHECK: "tf._SomeOp"(%[[BARG0]])
     %pred = "tf._SomeOp"(%barg0) : (tensor<i32>) -> tensor<i1>
    "tf.Yield"(%pred) : (tensor<i1>) -> ()
  }, {
    // CHECK: ^bb0(%[[BARG0:.*]]: tensor<i32>
    ^bb0(%barg0: tensor<i32>):

}

func.func @main(%arg0: tensor<i32>, %arg1: tensor<1xf32>) -> tensor<1xf32> {
  %0:2 = "tfl.while"(%arg0, %arg1) ({
  ^bb0(%arg2: tensor<*xi32>, %arg3: tensor<*xf32>):
    %1 = func.call @WhileOp_cond(%arg2, %arg3) : (tensor<*xi32>, tensor<*xf32>) -> tensor<i1>
    "tfl.yield"(%1) : (tensor<i1>) -> ()
  },  {
  ^bb0(%arg2: tensor<*xi32>, %arg3: tensor<*xf32>):

// RUN: flatbuffer_translate -mlir-to-tflite-flatbuffer %s -o - | flatbuffer_to_string - | FileCheck %s

func.func @main(tensor<4xi1>) -> tensor<4xi1> {
^bb0(%arg0: tensor<4xi1>):
  // CHECK:      {
  // CHECK-NEXT:   version: 3,
  // CHECK-NEXT:   operator_codes: [ {
  // CHECK-NEXT:     deprecated_builtin_code: 84,
  // CHECK-NEXT:     version: 1,
  // CHECK-NEXT:     builtin_code: LOGICAL_OR
  // CHECK-NEXT:   }, {