// CHECK: %[[INIT:.*]] = mhlo.constant dense<-2147483648> : tensor<i32>
  // CHECK: "mhlo.reduce_window"(%[[ARG]], %[[INIT]])
  // CHECK: <{window_dimensions = dense<[1, 2, 2, 1]> : tensor<4xi64>, window_strides = dense<[1, 4, 4, 1]> : tensor<4xi64>}>
  // CHECK: mhlo.maximum
  // CHECK: mhlo.return

            /*init_value=*/op.getOperand(1),
            /*window_dimensions=*/
            PermuteI64ArrayAttr(rewriter, op.getWindowDimensions(),
                                kNchwToNhwcPermutation),
            /*window_strides=*/
            PermuteI64ArrayAttr(rewriter, op.getWindowStrides(),
                                kNchwToNhwcPermutation),
            /*base_dilations=*/
            PermuteI64ArrayAttr(rewriter, op.getBaseDilations(),

    output_quantization_min_val = -2147483648 : i64, padding = "SAME",
    rhs_dilation = [1, 1], rhs_quantization_axis = -1 : i64,
    rhs_quantization_max_val = 127 : i64, rhs_quantization_min_val = -128 : i64,
    window_strides = [1, 1]
  } : (tensor<1x9x9x9x!tf_type.qint8>, tensor<3x3x9x10x!tf_type.qint8>,
    tensor<f32>, tensor<i32>, tensor<f32>, tensor<i32>, tensor<f32>, tensor<i32>
  ) -> tensor<1x9x9x10x!tf_type.qint32>

  absl::flat_hash_map<std::string, Attribute (*)(PatternRewriter&,
                                                 llvm::StringMap<Attribute>&)>
      attribute_getter_map;

  attribute_getter_map = {{"window_strides", GetWindowStridesValue},
                          {"lhs_dilation", GetLhsDilationValue},
                          {"rhs_dilation", GetRhsDilationValue},
                          {"padding", GetPaddingValue},

  lhs_contracting_dimensions : [long];
  rhs_contracting_dimensions : [long];
  precision_config : [StablehloPrecisionConfig];
}

table StablehloReduceWindowOptions{
  window_dimensions : [long];
  window_strides : [long];
  base_dilations : [long];
  window_dilations : [long];
  padding : [long];
  body_subgraph_index : int;
}

table StablehloWhileOptions{
  cond_subgraph_index : int;

  %1 = stablehlo.convolution(%arg0, %0) dim_numbers = [b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f], window = {pad = [[1, 1], [1, 1]]} {batch_group_count = 1 : i64, feature_group_count = 1 : i64, window_strides = array<i64: 2, 2>} : (tensor<1x32x32x3x!quant.uniform<i8:f32, 2.000000e+00:0>>, tensor<3x3x3x2x!quant.uniform<i8:f32:3, {3.000000e+00, 3.000000e+00}>>) -> tensor<1x16x16x2x!quant.uniform<i32:f32:3, {6.000000e+00, 6.000000e+00}>>

  { auto _e = window_strides(); if (_e) { _o->window_strides.resize(_e->size()); for (::flatbuffers::uoffset_t _i = 0; _i < _e->size(); _i++) { _o->window_strides[_i] = _e->Get(_i); } } else { _o->window_strides.resize(0); } }

  auto wd = reduce_window_op.getWindowDimensions();
  auto window_dimensions =
      builder_.CreateVector(std::vector<int64_t>(wd.begin(), wd.end()));
  auto window_strides = builder_.CreateVector(
      mlir::GetOptionalVector<int64_t>(reduce_window_op.getWindowStrides()));
  auto base_dilations = builder_.CreateVector(

  // expected-error @below {{'tf.UniformQuantizedConvolutionHybrid' op quantization_axis is -1, scales must have 0 rank.}}
  %0 = "tf.UniformQuantizedConvolutionHybrid"(%lhs, %rhs, %rhs_scales, %rhs_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,

    // CHECK: %[[RESULT:.*]] = "mhlo.select_and_scatter"(%[[INPUT]], %[[GRAD]], %[[ZERO]]) <{
    // CHECK-SAME:  padding = dense<0> : tensor<4x2xi64>, window_dimensions = dense<[1, 2, 2, 1]> : tensor<4xi64>, window_strides = dense<[1, 2, 2, 1]> : tensor<4xi64>
    // CHECK-SAME }> ({
    // CHECK: ^bb0(%[[VALUE_A:.*]]: tensor<f32>, %[[VALUE_B:.*]]: tensor<f32>):