// RUN: stablehlo-quant-opt %s -split-input-file -stablehlo-lift-quantizable-spots-as-functions | FileCheck %s

// CHECK-LABEL: @conv_fn(
// CHECK-SAME:          %[[ARG_0:.*]]: tensor<1x3x3x4xf32>
func.func @conv_fn(%arg0: tensor<1x3x3x4xf32>) -> tensor<1x3x3x4xf32> {
  %0 = stablehlo.constant dense<2.000000e+00> : tensor<3x3x4x4xf32>

    quantization_axis = -1 : i64, quantization_min_val = -128 : i64,
    quantization_max_val = 127 : i64
  } : (
    tensor<3x3x10x20x!tf_type.qint8>, tensor<f32>, tensor<i32>
  ) -> tensor<3x3x10x20xf32>
  %0 = "tf.Conv2D"(%input, %filter_new) {
    Tin = "tfdtype$DT_FLOAT", Tout = "tfdtype$DT_FLOAT",
    attr_map = "", batch_group_count = 1 : i64,
    explicit_padding = [], feature_group_count = 1 : i64, lhs_dilation = [1, 1],

  };
  return values;
}

inline const char * const *EnumNamesBuiltinOperator() {
  static const char * const names[208] = {
    "ADD",
    "AVERAGE_POOL_2D",
    "CONCATENATION",
    "CONV_2D",
    "DEPTHWISE_CONV_2D",
    "DEPTH_TO_SPACE",
    "DEQUANTIZE",
    "EMBEDDING_LOOKUP",
    "FLOOR",
    "FULLY_CONNECTED",
    "HASHTABLE_LOOKUP",
    "L2_NORMALIZATION",

      "Bincount",
      "Bucketize",
      "Case",
      "CheckNumerics",
      "Cholesky",
      "ControlTrigger",
      "Conv",
      "Conv2D",
      "Conv2DBackpropFilter",
      "Conv2DBackpropInput",
      "Conv3D",
      "Conv3DBackpropFilterV2",
      "Conv3DBackpropInputV2",
      "Cross",
      "Cumprod",
      "Cumsum",
      "CumulativeLogsumexp",

*   Fixes a `CHECK` fail in `tf.sparse.cross` ([CVE-2022-35997](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2022-35997))
*   Fixes a floating point exception in `Conv2D` ([CVE-2022-35996](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2022-35996))
*   Fixes a `CHECK` fail in `AudioSummaryV2` ([CVE-2022-35995](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2022-35995))

  %0 = "tf.Conv3D"(%arg0, %arg1) {padding = "SAME", strides = [1, 1, 1, 1, 1]} : (tensor<?x?x?x?x?xf32>, tensor<?x?x?x?x?xf32>) -> tensor<?x?x?x?x?xf32>
  func.return %0: tensor<?x?x?x?x?xf32>

  // CHECK-LABEL: conv3d_valid
  // CHECK:  %[[CST:.*]] = "tfl.no_value"() <{value}> : () -> none

        %device_launch = "tf_device.cluster_func"(%input,...) {func = @_func,...)
        return ...
      }
      func @_func(%input: tensor<2x224x224x3xf32>, %filter: tensor<7x7x3x64xf32>) {
        %6 = "tf.Conv2D"(%input, %filter)  {strides = [1, 2, 2, 1]}: (tensor<2x230x230x3xf32>, tensor<7x7x3x64xf32>) -> tensor<2x112x112x64xf32>
      }
    }
    ```

    The program will be transformed into:

    ```mlir

    if (add.getType() != op.getType()) {
      add = rewriter.create<tensor::CastOp>(loc, op.getType(), add);
    }
    rewriter.replaceOp(op, {add});
    return success();
  }
};

// Conterts tf.Conv2D to mhlo.dynamic_conv.
// TODO(disc): To recover static special case's performance with adding folding,
// canonicalization func and removing ConvertConvOp.
template <typename OpT, int num_spatial_dims, bool depthwise_conv = false>

}

type RawSockaddrL2TPIP struct {
	Family  uint16
	Unused  uint16
	Addr    [4]byte /* in_addr */
	Conn_id uint32
	_       [4]uint8
}

type RawSockaddrL2TPIP6 struct {
	Family   uint16
	Unused   uint16
	Flowinfo uint32
	Addr     [16]byte /* in6_addr */
	Scope_id uint32
	Conn_id  uint32
}

type RawSockaddrIUCV struct {
	Family  uint16
	Port    uint16
	Addr    uint32

    return emitOptionalError(location, "invalid padding format provided");
  }

  // Output always have rank 4. All dimensions are initialized to
  // dynamic size and can be partially inferred.
  // TFL's conv2d is always NHWC format & the filter is OHWI.
  SmallVector<int64_t, 4> return_shape(4, ShapedType::kDynamic);
  return_shape[0] = input_ty.getDimSize(0);
  return_shape[3] = filter_ty.getDimSize(0);