// CHECK-LABEL: QuantizeCustomOp
// CustomOp-LABEL: QuantizeCustomOp
func.func @QuantizeCustomOp(%arg0: tensor<1x1x1x1xf32>) -> (tensor<*xf32>, tensor<*xf32>, tensor<*xf32>) attributes {tf.entry_function = {inputs = "input", outputs = "custom_op"}} {
  %0 = "quantfork.stats"(%arg0) {layerStats = dense<[0.000000e+00, 2.550000e+02]> : tensor<2xf32>} : (tensor<1x1x1x1xf32>) -> tensor<1x1x1x1xf32>
  %w_1 = arith.constant dense<127.0> : tensor<4096x1x1x1xf32>

func.func @no_qdq_for_non_weight_constant(%arg0: tensor<1x2xf32>, %arg1: tensor<2x3xf32>) -> tensor<1x3xf32> attributes {tf._original_func_name = "main_0"} {
  %cst = "tf.Const"() {value = dense<4.000000e-02> : tensor<3xf32>} : () -> tensor<3xf32>
  %0 = "tf.XlaCallModule"(%arg0, %arg1, %cst) {
    Sout = [#tf_type.shape<1x3>], _entry_function = @composite_dot_general_with_bias_fn,
    _original_entry_function = "composite_dot_general_with_bias_fn",

  if (!quantDenseAttr) {
    return nullptr;
  }

  // Cast from an expressed-type-based type to storage-type-based type,
  // preserving the sparse shape (i.e. tensor<4xf32> -> tensor<4xi8>).
  ShapedType newSparseType = mlir::dyn_cast_or_null<ShapedType>(
      quantizedElementType.castExpressedToStorageType(
          realSparseAttr.getType()));
  if (!newSparseType) {
    return nullptr;

// CHECK-NEXT:      %12 = mhlo.dynamic_reshape %2, %11 : (tensor<4x256x?xf32>, tensor<3xi32>) -> tensor<4x256x?xf32>
// CHECK-NEXT:      %13 = "tfl.batch_matmul"(%3, %12) <{adj_x = false, adj_y = false, asymmetric_quantize_inputs = false}> : (tensor<4x4x256xf32>, tensor<4x256x?xf32>) -> tensor<4x4x?xf32>
// CHECK-NEXT:      %14 = "tfl.shape"(%arg0) : (tensor<4x4x256xf32>) -> tensor<3xi32>

  // CHECK-SAME:      -> tensor<40x40xf32>
  // CHECK:         %[[stat:.*]] = "quantfork.stats"(%arg0) <{layerStats = dense
  // CHECK-SAME:      <[-1.000000e+00, 1.000000e+00]> : tensor<2xf32>}>
  // CHECK-SAME:      : (tensor<40x37xf32>) -> tensor<40x37xf32>
  // CHECK-NEXT:    %[[cst:.*]] = "tfl.pseudo_const"() <{value = dense<
  // CHECK-SAME:      1.000000e+00> : tensor<40xf32>}> : () -> tensor<40xf32>

    tensor type, the shape shouldn't be changed during the conversion.

    Example:

    ```mlir
    %1 = tfr.constant_tensor(%0) : f32 -> tensor<f32>
    %3 = tfr.constant_tensor(%2) : vector<1xf32> -> tensor<1xf32>
    ```
  }];

  let arguments = (ins TFR_AllAttrTypes:$arg);

  let results = (outs TFR_singleTensorType:$out);

  let hasCanonicalizer = 1;

  let hasVerifier = 1;
}

func.func @tf2xla_fallback_op_approx_top_k(%arg0: tensor<16xf32>) -> (tensor<?xf32>, tensor<?xi32>) {
  %0:2 = "tf_device.cluster"() ({
    // CHECK: tf.ApproxTopK
    // CHECK-NOT: _xla_outside_compilation
    %1:2 = "tf.ApproxTopK"(%arg0) {k = 2} : (tensor<16xf32>) -> (tensor<?xf32>, tensor<?xi32>)
    tf_device.return %1#0, %1#1 : tensor<?xf32>, tensor<?xi32>

          %1 = call @then_branch_func(%arg1) : (tensor<*xf32>) -> tensor<*xf32>
          "tf.Yield"(%1) : (tensor<*xf32>) -> ()
        },  {
          %1 = call @else_branch_func(%arg1) : (tensor<*xf32>) -> tensor<*xf32>
          "tf.Yield"(%1) : (tensor<*xf32>) -> ()
        }) {is_stateless = false} : (tensor<i1>) -> tensor<*xf32>
    ```

    will be transformed into this functional operation

///      : (tensor<f32>, tensor<f32>, tensor<f32>) -> tensor<5xf32>
///
/// Output would be:
///   %iota = "mhlo.iota"() {iota_dimension = 0 : i64} : () -> tensor<5xf32>
///   %scaled = "mhlo.multiply"(%iota, %delta)
///       {broadcast_dimensions = dense<[]> : tensor<0xi64>} :
///       (tensor<5xf32>, tensor<f32>) -> tensor<5xf32>
///   %result = "mhlo.add"(%scaled, %offset)

      while (auto cast_op =
                 llvm::dyn_cast_or_null<CastOp>(value.getDefiningOp())) {
        // Consider cast compatibility in case
        //    %cast = "tf.Cast"(%0) : (tensor<2xi64>) -> tensor<2xf32>
        // is skipped.
        if (cast_op.getSrcT() != cast_op.getDstT()) {
          break;
        }
        value = cast_op.getOperand();
      }
      return value;
    };