%i8_min = "tf.Const"() {value = dense<-128> : tensor<i8>} : () -> tensor<i8>
    %i8_max = "tf.Const"() {value = dense<127> : tensor<i8>} : () -> tensor<i8>

    %clip_min = "tf.PartitionedCall"(%i8_min, %out_scale, %out_zp) {
        config = "", config_proto = "", executor_type = "", f=@internal_dequantize_i8_fn
      } : (tensor<i8>, tensor<*xf32>, tensor<*xi32>) -> tensor<*xf32>

// CHECK-NO-UNPACK-DAG: %[[CONST:.+]] = stablehlo.constant() <{value = dense<{{.*}}> : tensor<1024x3xi8>}> : () -> tensor<1024x3x!quant.uniform<i8<-127:127>:f32:1, {{.*}}>>
// CHECK-NO-UNPACK: %[[QUANTIZE_0:.+]] = stablehlo.uniform_quantize %[[ARG_0]] : (tensor<1x1024xf32>) -> tensor<1x1024x!quant.uniform<i8:f32, {{.*}}>>
// CHECK-NO-UNPACK: %[[DOT:.+]] = stablehlo.dot_general %[[QUANTIZE_0]], %[[CONST]]

  let arguments = (
    ins TFL_TensorOf<[F32, I8]>:$input,

    // Forward LSTM Weights
    TFL_TensorOfOrNone<[F32, I8]>:$fw_input_to_input_weights,
    TFL_TensorOf<[F32, I8]>:$fw_input_to_forget_weights,
    TFL_TensorOf<[F32, I8]>:$fw_input_to_cell_weights,
    TFL_TensorOf<[F32, I8]>:$fw_input_to_output_weights,

    // Forward Recurrent weights
    TFL_TensorOfOrNone<[F32, I8]>:$fw_recurrent_to_input_weights,

    %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>

	}},
	{func(i1, i2, i3, i4, i5, i6, i7, i8 uintptr) uintptr {
		return i1 + i2 + i3 + i4 + i5 + i6 + i7 + i8
	}},
	{func(i1, i2, i3, i4, i5, i6, i7, i8, i9 uintptr) uintptr {
		return i1 + i2 + i3 + i4 + i5 + i6 + i7 + i8 + i9
	}},

	// Non-uintptr parameters.
	{func(i1, i2, i3, i4, i5, i6, i7, i8, i9 uint8) uintptr {
		return uintptr(i1 + i2 + i3 + i4 + i5 + i6 + i7 + i8 + i9)
	}},

  %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: %[[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>

// `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 "