// RUN: tf-opt %s -tfl-prepare-quantize -tfl-quantize  | FileCheck %s
// RUN: tf-opt %s -tfl-quantize="legacy-quantize=true" | FileCheck --check-prefix=LEGACY %s
// RUN: tf-opt %s -tfl-prepare-quantize -tfl-quantize="ops-blocklist=tfl.fully_connected,tfl.softmax locs-blocklist=Block,NullBlock" | FileCheck --check-prefix=BLOCK %s

// CHECK-LABEL: QuantizeFloatConst
func.func @QuantizeFloatConst() -> tensor<2x2x!quant.uniform<u8:f32, 7.8431372549019615E-4:128>> {

// RUN: tf-opt %s -tfl-prepare-quantize-dynamic-range -tfl-quantize="enable-dynamic-range-quantization=true" | FileCheck %s
// RUN: tf-opt %s -tfl-prepare-quantize-dynamic-range -tfl-quantize="enable-dynamic-range-quantization=true enable-weight-only-quantization=true" | FileCheck --check-prefix=PerChannelWeightOnly %s

    const auto float_graph = model_->subgraphs()->Get(subgraph_idx);
    // The output graph should have an extra tensor from the added dequantize
    // op.
    ASSERT_EQ(quantized_graph->tensors()->size(),
              float_graph->tensors()->size() + 1);
    // Check that a dequantize op exists.
    int32_t dequant_input_idx = -1;
    int32_t dequant_output_idx = -1;

        "MLIR dump file name.">,
    Option<"merge_fusion_with_dequantize_",
        "merge-fusion-with-dequantize",
        "bool", /*default=*/"false",
        "Whether to merge quantized conv/dot_general fusion with subsequent dequantize.">,
  ];
  let dependentDialects = [
    "mlir::arith::ArithDialect",
    "mlir::stablehlo::StablehloDialect",
    "mlir::quant::QuantizationDialect",

  // CHECK: %[[DEQUANTIZE:.*]] = mhlo.uniform_dequantize %[[CONVERT_2]] : (tensor<2x!quant.uniform<i8:f32, 1.000000e+00:3>>) -> tensor<2xf32>
  // CHECK: return %[[DEQUANTIZE]] : tensor<2xf32>

  %0 = "tf.UniformQuantize"(%arg0, %scales, %zps) {
    quantization_axis = -1 : i64, quantization_min_val = -128 : i64, quantization_max_val = 127 : i64

  // If `value` is produced by tf.Dequantize op, then return the Dequantize op's
  // input. Otherwise return `value`.
  auto get_real_input_value = [](Value value) -> Value {
    Operation* defining_op = value.getDefiningOp();
    if (auto dequantize = dyn_cast_or_null<TF::DequantizeOp>(defining_op)) {
      return dequantize.getInput();
    } else if (auto dequantize =

        per_axis_type.getStorageTypeMin(), per_axis_type.getStorageTypeMax());
  }

  auto quantize = builder.create<quantfork::QuantizeCastOp>(
      q_op.getLoc(), new_value_type.clone(new_qtype), new_value);
  auto dequantize = builder.create<quantfork::DequantizeCastOp>(
      dq_op.getLoc(), new_value_type, quantize.getResult());
  return dequantize.getResult();
}

  // hardware performs better with integer ops.
  // Default value: true
  optional bool unpack_quantized_types = 1;

  // When set to True, requantize op in the quantized fusion will merge with the
  // subsequent dequantize op if present.
  // Default value: false
  // TODO: b/321729008 - re-consider default value after testing on prod model.
  bool merge_fusion_with_dequantize = 2;
}

    bool enable_legacy_weight_only = false,
    std::optional<const absl::string_view> mlir_dump_file_prefix =
        std::nullopt);

// Converts dequantize-(quantizable) call-quantize pattern to a single call op
// that has quantized input and output types. It is expected for this pass to
// emit illegal IR with unsupported quantized input and output types. The

       (TF_ConstOp (GetI64ScalarElementsAttr<-1>)))),
  (TF_SoftmaxCrossEntropyWithLogitsOp $features, $adjusted_labels)]>;

//===----------------------------------------------------------------------===//
// Dequantize op patterns.
//===----------------------------------------------------------------------===//

def DequantizeHalfRange : NativeCodeCall<
  "DequantizeHalfRange(&$_builder, $0)">;