%0 = stablehlo.constant() : () -> tensor<...x!quant.uniform<i8:f32:3 ...>>
    %1 = stablehlo.uniform_quantize %2 : (tensor<...xf32>) -> tensor<...x!quant.uniform<i8:f32 ...>>
    %3 = stablehlo.my_quantized_op(%1, %0) : (tensor<...x!quant.uniform<i8:f32 ...>>, tensor<...x!quant.uniform<i8:f32:3 ...>>) -> tensor<...x!quant.uniform<i8:f32 ...>>
    %4 = stablehlo.dequantize %3 : (tensor<...x!quant.uniform<i8:f32 ...>>) -> tensor<...xf32>
    ```
  }];

// CHECK-DAG: %[[CONST:.*]] = arith.constant dense<0.000000e+00> : tensor<2x1024xf32>
// CHECK: %0 = "quantfork.qcast"(%[[CONST]]) : (tensor<2x1024xf32>) -> tensor<2x1024x!quant.uniform<i8<-127:127>:f32, 3.9370078740157481E-9>>
// CHECK: %1 = "quantfork.dcast"(%0) : (tensor<2x1024x!quant.uniform<i8<-127:127>:f32, 3.9370078740157481E-9>>) -> tensor<2x1024xf32>

  // CHECK: %[[cst:.*]] = stablehlo.constant
  // CHECK: %[[q1:.*]] = "quantfork.qcast"(%[[cst]])
  // CHECK-SAME: quant.uniform<i8:f32, 0.0040316890267764818:127>
  // CHECK: %[[dq1:.*]] = "quantfork.dcast"(%[[q1]])
  // CHECK-SAME: quant.uniform<i8:f32, 0.0040316890267764818:127>
  %cst = stablehlo.constant dense<[[-0.960978984, -0.390246302], [-0.790828585, -0.601039409], [-1.0280807, -1.02731466]]> : tensor<3x2xf32>

  %1 = "quantfork.dcast"(%0) : (tensor<1x3x4x3x!quant.uniform<i8:f32, 0.0011764706057660721:-43>>) -> tensor<1x3x4x3xf32>
  %q_w = "quantfork.qcast"(%cst) : (tensor<2x3x3x2xf32>) -> tensor<2x3x3x2x!quant.uniform<i8:f32, 0.0125:-24>>
  %dq_w = "quantfork.dcast"(%q_w) : (tensor<2x3x3x2x!quant.uniform<i8:f32, 0.0125:-24>>) -> tensor<2x3x3x2xf32>

}

// CHECK: %[[q1:.*]] = "quantfork.qcast"(%arg0)
// CHECK-SAME: tensor<3x!quant.uniform<i8:f32, 0.0076314610593459188:-3>>
// CHECK: %[[dq1:.*]] = "quantfork.dcast"(%[[q1]])
// CHECK: %[[q2:.*]] = "quantfork.qcast"(%arg1)
// CHECK-SAME: tensor<4x3x!quant.uniform<i8<-127:127>:f32:1, {0.003937007874015748,0.0039370079913477263:-25,0.003937007874015748:51}>>
// CHECK: %[[dq2:.*]] = "quantfork.dcast"(%[[q2]])

      }
    }
  }
}
versions {
  producer: 216
}

# CHECK-NO-SUBTYPE: func @main(%arg0: tensor<i32>, %arg1: tensor<!tf_type.variant>, %arg2: tensor<i8>, %arg3: tensor<f64>, %arg4: tensor<!tf_type.resource>) -> (tensor<i32>, tensor<!tf_type.variant>, tensor<i8>, tensor<f64>, tensor<!tf_type.resource>)

// CHECK: %[[out:.*]] = "tf.PartitionedCall"([[ARG0:%arg[0-9]+]], %[[q_cst]]) <{config = "", config_proto = "", executor_type = "", f = @composite_matmul_fn}> {_tfl_quant_trait = "fully_quantizable"} : (tensor<1x2x2x3xf32>, tensor<2x1024x!quant.uniform<i8<-127:127>:f32, 3.9370078740157481E-9>>) -> tensor<*xf32>

// CHECK: %[[q1:.*]] = "quantfork.qcast"(%[[cst_1]]) {volatile}
// CHECK-SAME: tensor<2x3x3x2x!quant.uniform<i8<-127:127>:f32:3, {0.075176584439014829,0.072960192762960605}>>
// CHECK: %[[dq1:.*]] = "quantfork.dcast"(%[[q1]])

// CHECK: %[[q2:.*]] = "quantfork.qcast"(%arg0)
// CHECK-SAME: tensor<1x3x4x3x!quant.uniform<i8:f32, 0.58754816990272674:-128>>
// CHECK: %[[dq2:.*]] = "quantfork.dcast"(%[[q2]])

// `tfl.quantize` or `tfl.dequantize` ops. ui8, i8 and i16 are supported.
bool IsSupportedByTfliteQuantizeOrDequantizeOps(IntegerType storage_type) {
  if (storage_type.getWidth() == 8 ||
      (storage_type.isSigned() && storage_type.getWidth() == 16)) {
    return true;
  }
  LLVM_DEBUG(llvm::dbgs()
             << "Uniform quantize / dequantize op only supports ui8, i8 or "

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

func.func @main(tensor<4 x f32>, tensor<4 x i8>, tensor<4 x f32>, tensor<4 x f32>) -> tensor<4 x f32> {
// CHECK:      {
// CHECK-NEXT:     version: 3,
// CHECK-NEXT:     operator_codes: [ {
// CHECK-NEXT:       deprecated_builtin_code: 27,
// CHECK-NEXT:       version: 2,
// CHECK-NEXT:       builtin_code: SVDF
// CHECK-NEXT:     } ],