key: "narrow_range"
    value {
      b: true
    }
  }
  attr {
    key: "num_bits"
    value {
      i: 4
    }
  }
}
node {
  name: "BoxPredictor_4/ClassPredictor/Conv2D"
  op: "Conv2D"
  input: "input"
  input: "BoxPredictor_4/ClassPredictor/weights_quant/FakeQuantWithMinMaxVarsPerChannel"
  attr {
    key: "T"
    value {
      type: DT_FLOAT
    }
  }
  attr {

  %6 = "tfl.var_handle"() {container = "", shared_name = "conv_variable/state"} : () -> tensor<!tf_type.resource>
  %7 = "tfl.read_variable"(%6) : (tensor<!tf_type.resource>) -> tensor<1x3x1x1xf32>

func.func @testConv2D(tensor<256x32x32x3xf32>, tensor<16x3x3x3xf32>, tensor<16xf32>) -> tensor<256x32x32x16xf32> {
^bb0(%arg0: tensor<256x32x32x3xf32>, %arg1: tensor<16x3x3x3xf32>, %arg2: tensor<16xf32>):
  // CHECK: "tfl.conv_2d"(%arg0, %arg1, %arg2)

            FindUserOfType<TFL::QuantizeOp>(op) != nullptr);
  }
};

// Rewrites `stablehlo.convolution` into fused `tfl.conv_2d`.
// If available, fuse bias and activation adjacent to `stablehlo.convolution`.
// This RewritePattern rewrites both the following into `tfl.conv_2d` op:
//
// StableHLO Quantizer output:
//   * input: per-tensor qi8
//   * filter: per-channel qi8 (`quantization_dimension` = 3)

// CHECK:           %[[VAL_5:.*]] = "tfl.conv_2d"(%[[VAL_0]], %[[VAL_1]], %[[VAL_2]]) <{dilation_h_factor = 1 : i32, dilation_w_factor = 1 : i32, fused_activation_function = "NONE", padding = "VALID", stride_h = 1 : i32, stride_w = 1 : i32}> {tac.device = "GPU", tac.inference_type = "FLOAT"} : (tensor<256x32x32x3xf32>, tensor<16x3x3x3xf32>, tensor<16xf32>) -> tensor<256x30x30x16xf32>

  rewriter.replaceOp(splitv_op, slice_outputs);

  return success();
}

// ================== conv_2d ========================

LogicalResult EnsureBiasForConv2d::matchAndRewrite(
    TFL::Conv2DOp conv_op, PatternRewriter& rewriter) const {
  return EnsureBias(conv_op, 2, rewriter);
}

// ================== slice ============================

// CHECK:           %[[VAL_12:.*]] = "tfl.reshape"(%[[VAL_10]], %[[VAL_7]]) {tac.device = "GPU", tac.inference_type = "FLOAT"} : (tensor<128x512xf32>, tensor<4xi32>) -> tensor<128x1x1x512xf32>

// object containing configuration parameters, builtins have a predetermined
// set of acceptable options.
// LINT.IfChange
enum BuiltinOperator : int32 {
  ADD = 0,
  AVERAGE_POOL_2D = 1,
  CONCATENATION = 2,
  CONV_2D = 3,
  DEPTHWISE_CONV_2D = 4,
  DEPTH_TO_SPACE = 5,
  DEQUANTIZE = 6,
  EMBEDDING_LOOKUP = 7,
  FLOOR = 8,
  FULLY_CONNECTED = 9,
  HASHTABLE_LOOKUP = 10,
  L2_NORMALIZATION = 11,
  L2_POOL_2D = 12,

// object containing configuration parameters, builtins have a predetermined
// set of acceptable options.
// LINT.IfChange
enum BuiltinOperator : byte {
  ADD = 0,
  AVERAGE_POOL_2D = 1,
  CONCATENATION = 2,
  CONV_2D = 3,
  DEPTHWISE_CONV_2D = 4,
  DEPTH_TO_SPACE = 5,
  DEQUANTIZE = 6,
  EMBEDDING_LOOKUP = 7,
  FLOOR = 8,
  FULLY_CONNECTED = 9,
  HASHTABLE_LOOKUP = 10,
  L2_NORMALIZATION = 11,
  L2_POOL_2D = 12,

  fused_activation_function:ActivationFunctionType;
  dilation_w_factor:int = 1;
  dilation_h_factor:int = 1;
  // Parameters for Conv2D version 8 or above.
  // When set, quantized_bias_type defines the dtype for both bias and accumulator.
  quantized_bias_type: TensorType;
}

// Options for both Conv3D and Conv3DTranspose.
table Conv3DOptions {
  padding:Padding;
  stride_d:int;
  stride_w:int;
  stride_h:int;