// ops of the same device, under a `tf_device.launch` op.
  bool form_clusters = false;
  // If `unfold_batch_matmul` is true, the tf.BatchMatMul is unfolded to a set
  // of tfl.fully_connected ops.
  bool unfold_batch_matmul = true;
  // Whether to outline WhileOp at the end of the pipeline.
  bool outline_tf_while = false;
  // Whether to do shape inference.
  bool shape_inference = true;

          .getRank();

  // To quantize rhs per-channel, we currently only consider the case where
  // `stablehlo.dot_general` is legalizable to `tfl.fully_connected`.
  const bool is_per_axis_quantizable =
      IsDotGeneralFullyConnected(dot_general_op).value();
  if (!is_per_axis_quantizable) return std::nullopt;
  return filter_rank - 1;
}

  %w = arith.constant dense<127.0> : tensor<512x12xf32>
  %b = arith.constant dense<0.0> : tensor<512xf32>
  %fc = "tfl.fully_connected"(%0, %w, %b) {fused_activation_function = "NONE", keep_num_dims = false, weights_format = "DEFAULT"} : (tensor<1x224x224x3xf32>, tensor<512x12xf32>, tensor<512xf32>) -> tensor<1x112x112x512xf32>

  %0 = "tfl.pseudo_const"() {value = dense<0.0> : tensor<40xf32>} : () -> tensor<40xf32>
  %1 = "tfl.fully_connected"(%arg0, %arg1, %0) {fused_activation_function = "NONE", keep_num_dims = false, weights_format = "DEFAULT"} : (tensor<1x37xf32>, tensor<40x37xf32>, tensor<40xf32>) -> tensor<1x40xf32>
// CHECK: "tfl.fully_connected"
// CHECK-SAME: (tensor<1x37xf32>, tensor<40x37xf32>, none) -> tensor<1x40xf32>
  func.return %1 : tensor<1x40xf32>

// filter value from [i, o] -> [o, i]. This is because we assume `[i, o]`
// format for `stablehlo.dot_general` (i.e. contracting dimension == 1)
// whereas `tfl.fully_connected` accepts an OI format.
TFL::QConstOp CreateTransposedTflConstOpForFilter(
    stablehlo::ConstantOp filter_constant_op, PatternRewriter& rewriter,
    bool is_per_channel) {

// Remove Reshape before FullyConnected when `keep_num_dims=false` and Reshape
// does not alter the last dimension as FullyConnected will collapse all other
// dimensions into a single dimension. For example,
//
//   %shape = arith.constant dense<[1, 128, 64]> : tensor<3xi32>
//   %reshape = tfl.reshape(%input, %shape) // %input: tensor<128x64xf32>
//   %fc = tfl.fully_connected(%reshape, %filter, %bias)

// is quantized.
bool IsHybridQuantizedOp(Operation* op);

// Returns whether a given `stablehlo.dot_general` can be legalizable to
// `tfl.fully_connected`.
absl::StatusOr<bool> IsDotGeneralFullyConnected(
    ::mlir::stablehlo::DotGeneralOp dot_general_op);

// Returns the quantization dimension for a given `stablehlo.dot_general` op,

}

// An implementation of TensorFlow fully_connected (a.k.a Dense) layer.
table FullyConnectedOptions {
  // Parameters for FullyConnected version 1 or above.
  fused_activation_function:ActivationFunctionType;

  // Parameters for FullyConnected version 2 or above.
  weights_format:FullyConnectedOptionsWeightsFormat = DEFAULT;

  // Parameters for FullyConnected version 5 or above.

  rewriter.replaceOp(slice_op, output_reshape_op.getResult());

  return success();
}

// ================== fully_connected ========================

// TFL fully_connected basically does:
// Weight * Input + bias.
// Input layout is : [..., depth]
// Weight layout is : [output, depth]
// Bias is [output].
//
// While conv2d is:
// Filter: [NHWC]

// CHECK-NOT: "tfl.reshape"
// CHECK: "tfl.slice"

// -----

func.func @fullyConnectedToConv(%arg0: tensor<384x384xf32>, %arg1: tensor<512x384xf32>, %arg2: tensor<512xf32>) -> tensor<384x512xf32> {
  %0 = "tfl.fully_connected"(%arg0, %arg1, %arg2) {fused_activation_function = "NONE", keep_num_dims = false, weights_format = "DEFAULT"} : (tensor<384x384xf32>, tensor<512x384xf32>, tensor<512xf32>) -> tensor<384x512xf32>
  func.return %0: tensor<384x512xf32>