XlaCompileOnDemandOp);                               \
  REGISTER_KERNEL_BUILDER(Name("XlaReduceWindow")                              \
                              .HostMemory("window_dimensions")                 \
                              .HostMemory("window_strides")                    \
                              .HostMemory("base_dilations")                    \

// CHECK: <{window_dimensions = array<i64: 1, 2, 2, 1>, window_strides = array<i64: 1, 2, 2, 1>}>
// CHECK: ^bb0(%[[REDUCE_ARG_0:.+]]: tensor<f32>, %[[REDUCE_ARG_1:.+]]: tensor<f32>):
// CHECK: %[[MAX:.+]] = stablehlo.maximum %[[REDUCE_ARG_0]], %[[REDUCE_ARG_1]]
// CHECK: stablehlo.return %[[MAX]]

// Check that the attributes window_dimensions & window_strides are also

// CHECK-SAME:          %[[PAD_VAL:.*]]: tensor<f32>) -> tensor<1x64x56x56xf32> {
// CHECK:           %[[REDUCE:.*]] = "stablehlo.reduce_window"(%[[INPUT]], %[[PAD_VAL]])
// CHECK:               <{window_dimensions = array<i64: 1, 3, 3, 1>,
// CHECK:                 window_strides = array<i64: 1, 2, 2, 1>}> ({
// CHECK:           ^bb0(%[[ARG0:.*]]: tensor<f32>, %[[ARG1:.*]]: tensor<f32>):

    // CHECK: %[[REDUCE:.*]] = "stablehlo.reduce_window"(%[[CALL]], %[[Q0]])
    // CHECK{LITERAL}: padding = dense<[[0, 0], [1, 1], [1, 1], [0, 0]]> : tensor<4x2xi64>
    // CHECK-SAME: window_dimensions = array<i64: 1, 3, 3, 1>
    // CHECK: %[[ARG1:.*]]: tensor<!quant.uniform<i8:f32, 3.000000e-01:1>>, %[[ARG2:.*]]: tensor<!quant.uniform<i8:f32, 3.000000e-01:1>>

    return true;
  if (op_name == "stablehlo.reduce_window" &&
      (field_name == "window_dimensions" || field_name == "window_strides" ||
       field_name == "base_dilations" || field_name == "window_dilations"))
    return true;
  if (op_name == "stablehlo.select_and_scatter" &&
      (field_name == "window_dimensions" || field_name == "window_strides"))
    return true;
  if (op_name == "stablehlo.slice" &&

    return;
  }
  if (const auto* op = op_union.AsStablehloReduceWindowOptions()) {
    if (!op->window_dimensions.empty()) {
      attributes.emplace_back(builder.getNamedAttr(
          "window_dimensions",
          BuildVhloTensorV1Attr(
              {static_cast<int64_t>(op->window_dimensions.size())},
              op->window_dimensions, builder)));
    }
    if (!op->window_strides.empty()) {

    ^bb0(%arg1: tensor<f32>, %arg2: tensor<f32>):
      %2 = stablehlo.maximum %arg1, %arg2 : tensor<f32>
      stablehlo.return %2 : tensor<f32>
  }) {padding = dense<[[0, 0], [1, 1], [1, 1], [0, 0]]> : tensor<4x2xi64>, window_dimensions = array<i64: 1, 3, 3, 1>} : (tensor<2x3x1x3xf32>, tensor<f32>) -> tensor<2x3x1x3xf32>
  return %1 : tensor<2x3x1x3xf32>
}

// -----

  %4 = "stablehlo.reduce_window"(%3, %5) ({  // max pool
  ^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>
}

          << x_name
          << "\nNumber of window dimensions: " << window_dimensions.size()
          << "\nNumber of " << x_name << ": " << x << "\n";
      return false;
    }
  };

  if (!(verify_size(window_dimensions.size(), "window_dimensions") &&
        verify_size(window_strides.size(), "window strides") &&
        verify_size(padding.size(), "padding entries") &&

    mhlo.return %7 : tensor<f32>
  }) {window_dimensions = dense<[1, 1, 3, 3]> : tensor<4xi64>} : (tensor<1x3x6x6xf32>, tensor<f32>) -> tensor<1x3x4x4xf32>
  %4 = "mhlo.reduce_window"(%1, %2) ({
  ^bb0(%arg1: tensor<f32>, %arg2: tensor<f32>):
    %7 = mhlo.add %arg1, %arg2 : tensor<f32>
    mhlo.return %7 : tensor<f32>
  }) {window_dimensions = dense<3> : tensor<2xi64>} : (tensor<6x6xf32>, tensor<f32>) -> tensor<4x4xf32>