// Only rank size four input will be only available by the tf.Conv2D
    // operator verification.
    if (!input_type || input_type.isDynamicDim(3)) {
      return failure();
    }
    // Check if the given op is based on grouped convolution.
    // Dim size zero will be verified by the tf.Conv2D operator verification.
    if (input_type.getDimSize(3) % filter_type.getDimSize(2) != 0) {

func countrunes(string) int

// Convert non-interface type to the data word of a (empty or nonempty) interface.
func convT(typ *byte, elem *any) unsafe.Pointer

// Same as convT, for types with no pointers in them.
func convTnoptr(typ *byte, elem *any) unsafe.Pointer

// Specialized versions of convT for specific types.
// These functions take concrete types in the runtime. But they may

  // Input: [N, H, W, C] for Conv2D or [N, D, H, W, C] for Conv3D.
  dnums.set_input_batch_dimension(0);
  dnums.set_input_feature_dimension(num_dims - 1);
  // Kernel: [K, K, I, O] for Conv2D or [K, K, K, I, O] for Conv3D.
  dnums.set_kernel_input_feature_dimension(num_dims - 2);
  dnums.set_kernel_output_feature_dimension(num_dims - 1);
  // Output: [N, H, W, C] for Conv2D or [N, D, H, W, C] for Conv3D.

              (UpdateShapeWithAxis<-1> $qtype, $old_value))),
          [(CanUpdateShapeWithAxis<-1> $qtype, $old_value)]>;

// The axis is set to 0 because the transpose is from the legalization of
// tf.conv2d and the new channel axis is the first dimension.
def ReorderTransposeDequantQuantUsedByConv :
      Pat<(TF_TransposeOp:$old_value
              (TFL_DequantizeOp (TFL_QuantizeOp $input, $qtype)), $perm),

    } else if (function_name.contains("conv2d")) {
      // For Conv2D, the channel dimension must be static to calculate the
      // feature group count.
      if (!HasStaticShapeAtDims(call_op->getOperand(0), /*dims=*/3)) {
        return absl::InternalError(
            "The channel dimension of Conv2D is required to be static.");
      }
    } else if (function_name.contains("conv3d")) {

      if (!input_type || !input_type.hasStaticShape()) {
        return false;
      }
      total_count += input_type.getNumElements();
    }
    *count = total_count;
    return true;
  }

  // For conv2d/depthwise_conv/fully_connected ops.
  // This algorithm actually comes from TOCO tooling_util.cc
  static bool GetArithmeticCountForConvAndFullyconnectedOp(mlir::Operation* op,

// CHECK: %[[CONV:.+]] = stablehlo.convolution(%[[ARG0]], %[[DQ]])
// CHECK{LITERAL}: dim_numbers = [b, 0, 1, f]x[o, 0, 1, i]->[b, 0, 1, f], window = {pad = [[1, 1], [1, 1]]} {batch_group_count = 1 : i64, feature_group_count = 1 : i64}
// CHECK-SAME: (tensor<1x3x3x4xf32>, tensor<2x3x3x4xf32>) -> tensor<1x3x3x2xf32>
// CHECK: return %[[CONV]]

// -----

// has a zero MTIME. This is a requirement for the Debian maintainers
// to be able to have deterministic packages.
//
// To patch a .gz file, use the following command:
//
//	$ dd if=/dev/zero bs=1 seek=4 count=4 conv=notrunc of=filename.gz
//
// See https://golang.org/issue/14937.
func TestGZIPFilesHaveZeroMTimes(t *testing.T) {
	// To avoid spurious false positives due to untracked GZIP files that

  %conv = "tf.Conv2D"(%dq_input, %dq_weight) {attr_map = "0:strides,1:use_cudnn_on_gpu,2:padding,3:explicit_paddings,4:dilations", data_format = "NHWC", device = "", dilations = [1, 1, 1, 1], explicit_paddings = [], padding = "VALID", strides = [1, 1, 2, 1], use_cudnn_on_gpu = true} : (tensor<1x3x4x3xf32>, tensor<2x3x3x2xf32>) -> tensor<*xf32>

      mlir::mhlo::createLegalizeDotToDotGeneralPass());
  // Unfuse mhlo BatchNorm to primitive ops.
  pm.addNestedPass<mlir::func::FuncOp>(mlir::odml::createUnfuseBatchNormPass());
  // Fuse Conv + Mul to Conv.
  pm.addNestedPass<mlir::func::FuncOp>(mlir::odml::createFuseConvolutionPass());
  // Fold broadcast_in_dim + Mul.
  pm.addNestedPass<mlir::func::FuncOp>(mlir::odml::createFoldBroadcastPass());