// CHECK-SAME: rhs_dilation = [2, 2]
// CHECK-SAME: rhs_quantization_axis = -1 : i64
// CHECK-SAME: rhs_quantization_max_val = 127 : i64
// CHECK-SAME: rhs_quantization_min_val = -128 : i64
// CHECK-SAME: window_strides = [1, 2]
// CHECK-SAME: (tensor<1x2x2x3xf32>, tensor<2x3x3x2x!tf_type.qint8>, tensor<f32>, tensor<i32>) -> tensor<*xf32>

// CHECK-LABEL: func private @quantized_conv2d_fn_1

    dnums.add_kernel_spatial_dimensions(i - 1);
    dnums.add_output_spatial_dimensions(i);
  }

  Value padding, window_strides, lhs_dilation, rhs_dilation,
      feature_group_count;
  PrepareXlaConvParams(builder, loc, strides, dilations, feature_group_cnt,
                       /*window_strides=*/window_strides,
                       /*lhs_dilation=*/lhs_dilation,
                       /*rhs_dilation=*/rhs_dilation,

  xla::Window window;
  for (size_t i = 0; i < window_dimensions.size(); i++) {
    auto dim = window.add_dimensions();
    dim->set_size(window_dimensions[i]);
    if (!window_strides.empty()) {
      dim->set_stride(window_strides[i]);
    } else {
      dim->set_stride(1);
    }
    if (!padding.empty()) {
      dim->set_padding_low(padding[i].first);
      dim->set_padding_high(padding[i].second);

  ^bb0(%arg1: tensor<f32>, %arg2: tensor<f32>):
      %6 = stablehlo.maximum %arg1, %arg2 : tensor<f32>
      stablehlo.return %6 : tensor<f32>
  }) {
    window_dimensions = array<i64: 1, 1, 2, 2>,
    window_strides = array<i64: 1, 1, 2, 2>
  } : (tensor<1x4x5x5xf32>, tensor<f32>) -> tensor<1x4x2x2xf32>
  return %4 : tensor<1x4x2x2xf32>
}
// CHECK-DAG: %[[WEIGHT_CONST:.+]] = stablehlo.constant {{.*}} : tensor<3x3x2x4xf32>

    }
    auto rhs_dilations_attr = rewriter.getNamedAttr(
        "rhs_dilation", GetI64ElementsAttr(rhs_dilations, &rewriter));

    auto window_strides_attr = rewriter.getNamedAttr(
        "window_strides", GetI64ElementsAttr(window_strides, &rewriter));

    auto dimension_numbers_attr = GetConvDimensionNumbersAttr(
        spatial_dim_indices, data_format, &rewriter);

    const int64_t input_channels =

    // Use NHWC format.
    int32_t stride_h = 1;
    int32_t stride_w = 1;
    std::optional<ArrayRef<int64_t>> window_strides = op.getWindowStrides();
    if (window_strides.has_value()) {
      stride_h = CastI64ToI32((*window_strides)[1]).value();
      stride_w = CastI64ToI32((*window_strides)[2]).value();
    }
    auto stride_h_attr = IntegerAttr::get(rewriter.getI32Type(), stride_h);

                                             window_reversal = #vhlo.tensor_v1<dense<false> : tensor<2xi1>>,
                                             window_strides = #vhlo.tensor_v1<dense<1> : tensor<2xi64>>}>
                                              : (tensor<1x1x1600x32xf32>, tensor<1x13x1x32xf32>) -> tensor<1x1x1600x32xf32>
  return %0 : tensor<1x1x1600x32xf32>
}

                                %input_scale, %input_zp, %filter_scale, %filter_zp, %out_scale, %out_zp) {
        Tin = "tfdtype$DT_QINT8",
        Tout = "tfdtype$DT_QINT32",
        window_strides = [1, 1],
        padding = "SAME",
        explicit_padding = [],
        lhs_dilation = [],
        rhs_dilation = [],
        batch_group_count = 1,
        feature_group_count = 1,

  // CHECK-SAME: (tensor<1x6x6x3xf32>, tensor<2x3x3x2x!quant.uniform<i8:f32, 1.000000e+00:3>>) -> tensor<1x4x1x2xf32>

  %0 = "tf.UniformQuantizedConvolutionHybrid"(%input, %weight, %weight_scales, %weight_zps) {
    window_strides = [1, 2],
    padding = "VALID",
    explicit_padding = [],
    lhs_dilation = [1, 1],
    rhs_dilation = [2, 2],
    batch_group_count = 1 : i64,
    feature_group_count = 1 : i64,

    %7 = stablehlo.maximum %arg1, %arg2 : tensor<f32>
    stablehlo.return %7 : tensor<f32>
  }) {padding = dense<[[0, 0], [0, 1], [0, 1], [0, 0]]> : tensor<4x2xi64>, window_dimensions = array<i64: 1, 3, 3, 1>, window_strides = array<i64: 1, 2, 2, 1>} : (tensor<?x112x112x64xf32>, tensor<f32>) -> tensor<?x56x56x64xf32>
  return %6 : tensor<?x56x56x64xf32>