// input activations need to be asymmetrically quantized. Refer to the
        // implementation for fully-connected as an example in
        // tensorflow/lite/kernels/fully_connected.cc. The kernels will handle
        // the asymmetric_quantize_inputs attribute in the builtin option.
        dynamic_range_quantized_user->setAttr(
            kAsymmetricQuantizeInputsAttr,

  %4 = "tfl.dequantize"(%3) : (tensor<32x12x!quant.uniform<u8<1:255>:f32, 0.021826678373682216:151>>) -> tensor<32x12xf32>
  %5 = "tfl.fully_connected"(%2, %4, %cst) {fused_activation_function = "NONE", keep_num_dims = false, weights_format = "DEFAULT"} : (tensor<1x224x224x3xf32>, tensor<32x12xf32>, tensor<32xf32>) -> tensor<1x112x112x32xf32>

      return dequantize.getInput();
    } else {
      return value;
    }
  };

  // Returns true if the TF::BatchMatMul operation can be converted to
  // tfl.fully_connected.
  auto can_convert_to_fully_connected =
      [&](BatchMatMulOpType& batch_matmul_op) {
        Value input_rhs = get_real_input_value(batch_matmul_op.getY());

        DenseElementsAttr constant;

  func.return %0, %1, %2 : tensor<2xi32>, tensor<2xf32>, tensor<i32>
}

// -----

func.func @fully_connected(%arg0: tensor<1x37xf32>, %arg1: tensor<40x37xf32>, %arg2: tensor<40xf32>) -> tensor<1x40xf32> {

// CHECK: %[[ARG:.*]] = "tfl.transpose"(%arg1, %[[CST]]) : (tensor<37x40xf32>, tensor<2xi32>) -> tensor<40x37xf32>
// CHECK: %[[CST_0:.*]] = "tfl.no_value"() <{value}> : () -> none
// CHECK: "tfl.fully_connected"(%arg0, %[[ARG]], %[[CST_0]]) <{fused_activation_function = "NONE", keep_num_dims = false, weights_format = "DEFAULT"}> : (tensor<40x37xf32>, tensor<40x37xf32>, none) -> tensor<40x40xf32>
}

    // [b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f]
    // `stablehlo.dot_general` can take various formats. We only per-channel
    // quantize non-batch ops.
    // `stablehlo.dot_general` legalizable to `tfl.fully_connected` has a
    // filter rank of 2 with the last dimension as the channel dimension.
    const int64_t quantization_dimension =
        mlir::cast<ShapedType>(filter_type).getShape().size() - 1;

  %4 = "tfl.dequantize"(%3) : (tensor<32x12x!quant.uniform<u8<1:255>:f32, 0.021826678373682216:151>>) -> tensor<32x12xf32>
  %5 = "tfl.fully_connected"(%2, %4, %cst) {fused_activation_function = "NONE", keep_num_dims = false, weights_format = "DEFAULT"} : (tensor<1x224x224x3xf32>, tensor<32x12xf32>, tensor<32xf32>) -> tensor<1x112x112x32xf32>

  %5 = "tfl.pseudo_const"() {value = dense<1.000000e+00> : tensor<128x128xf32>} : () -> tensor<128x128xf32>
  %6 = "tfl.pseudo_const"() {value = dense<1.000000e+00> : tensor<128xf32>} : () -> tensor<128xf32>
  %7 = "tfl.fully_connected"(%4, %5, %6) {tac.device = "CPU", tac.inference_type = "FLOAT", fused_activation_function = "NONE", keep_num_dims = false, weights_format = "DEFAULT"} : (tensor<128x128xf32>, tensor<128x128xf32>, tensor<128xf32>) -> tensor<128x128xf32>

    std::vector<int> GetQuantizableOperandIndices() { return {1}; }
  }];
}

// TODO(jpienaar): Update post discussion on semantics of FC OP.
def TFL_FullyConnectedOp : TFL_Op<"fully_connected", [
    Pure, AccumulatorUniformScale<2, 0, 1>,
    AffineQuantizedOpInterface,
    AffineOpCoefficient<0, 1>,
    TFL_SparseOp,
    DeclareOpInterfaceMethods<TFL_ArithmeticCount>,
    QuantizableResult,

    llvm::cl::init(""));

// NOLINTNEXTLINE
opt<bool> unfold_batchmatmul(
    "unfold_batchmatmul",
    llvm::cl::desc(
        "Whether to unfold TF BatchMatMul to a set of TFL FullyConnected ops."),
    llvm::cl::init(false));

// NOLINTNEXTLINE
opt<bool> unfold_large_splat_constant(
    "unfold-large-splat-constant",
    llvm::cl::desc("Whether to unfold large splat constant tensors to reduce "