// CHECK: %[[minimum:.*]] = "tf.Minimum"(%[[maximum]]
// CHECK: %[[round:.*]] = "tf.Round"(%[[minimum]]
// CHECK: %[[quant:.*]] = "tf.Cast"(%[[round]]) : (tensor<1x3x4x3xf32>) -> tensor<1x3x4x3xi8>
// CHECK: %[[pad:.*]] = "tf.PadV2"(%[[quant]]
// CHECK: %[[xlaconv:.*]] = "tf.XlaConvV2"(%[[pad]]
// CHECK: %[[sub:.*]] = "tf.Sub"(%[[xlaconv]]
// CHECK: %[[cast:.*]] = "tf.Cast"(%[[sub]]) <{Truncate = false}> : (tensor<1x3x2x2xi32>) -> tensor<1x3x2x2xf32>

    "MatrixDiagV2",
    "MatrixDiagV3",
    "MatrixSetDiag",
    "MatrixSetDiagV2",
    "MatrixSetDiagV3",
    "MirrorPad",
    "OneHot",
    "OnesLike",
    "Pack",
    "Pad",
    "PadV2",
    "ParallelConcat",
    "PlaceholderV2",
    "PlaceholderWithDefault",
    "PreventGradient",
    "QuantizeAndDequantize",
    "QuantizeV2",
    "Rank",
    "Reshape",

  }

// CHECK-LABEL: func @conv_with_filter_larger_than_1GB
// CHECK-DAG: %[[CONST:.*]] = "tf.Const"() <{value = dense<-237772800> : tensor<1x1x1x512xi32>}> : () -> tensor<1x1x1x512xi32>
// CHECK: %[[PADV2_0:.*]] = "tf.PadV2"
// CHECK: %[[XLACONVV2_0:.*]] = "tf.XlaConvV2"(%[[PADV2_0]]
// CHECK: %[[SUB_0:.*]] = "tf.Sub"(%[[XLACONVV2_0]], %[[CONST]])

^bb0(%arg0: tensor<*xf32>, %arg1: tensor<5x3xi32>):
  %cst = arith.constant dense<2.0> : tensor<f32>
  // CHECK: "tfl.padv2"(%arg0, %arg1, %cst)
  %0 = "tfl.padv2"(%arg0, %arg1, %cst) : (tensor<*xf32>, tensor<5x3xi32>, tensor<f32>) -> tensor<? x f32>
  func.return %0#0 : tensor<? x f32>
}

// -----

// test PadV2 with invalid paddings size

  grad_outputs->push_back(NoGradient());
  // PadV2 adds a "constant_values" input.
  if (IsPadV2) {
    grad_outputs->push_back(NoGradient());
  }
  return scope.status();
}
REGISTER_GRADIENT_OP("Pad", PadGrad<false>);
REGISTER_GRADIENT_OP("PadV2", PadGrad<true>);

Status SpaceToBatchGrad(const Scope& scope, const Operation& op,

// CHECK-DAG-SAME{LITERAL}: value = dense<[[[[-22016, -23680]]]]>
// CHECK-DAG: %[[CONST_8:.*]] = "tf.Const"() <{value = dense<[162, 160]> : tensor<2xi32>}> : () -> tensor<2xi32>
// CHECK: %[[PADV2_0:.*]] = "tf.PadV2"({{.*}}, %[[CONST_4]], %[[CONST_5]]) : (tensor<1x3x4x3xi8>, tensor<4x2xi32>, tensor<i8>) -> tensor<1x4x5x3xi8>

  // CHECK-DAG: "tf.PadV2"(%arg0, {{.*}}, {{.*}}) {{.*}} -> tensor<3x?x?x10xf16>
  // CHECK-DAG: return {{.*}}: tensor<?x?x?x10xf16>
  %0 = "tf.SpaceToBatchND"(%input, %block_shape, %paddings) : (tensor<3x5x7x10xf16>, tensor<2xi64>, tensor<2x2xi64>) -> tensor<?x?x?x10xf16>
  func.return %0 : tensor<?x?x?x10xf16>
}

// Verify the result shape for the tf.PadV2 op.
// CHECK-LABEL: const_paddings_space_to_batch_nd

// CHECK: %[[PAD:.+]] = "tfl.padv2"(%[[ARG0]], %[[PADDING]], %[[ARG1]]) : (tensor<2x3x!quant.uniform<i8:f32, 2.000000e+00:-1>>, tensor<2x2xi32>, tensor<!quant.uniform<i8:f32, 2.000000e+00:-1>>) -> tensor<4x5x!quant.uniform<i8:f32, 2.000000e+00:-1>>
// CHECK: return %[[PAD]]

// -----

// Tests that a quantized `stablehlo.pad` with interior padding is converted to
// `tfl.dilate` and `tfl.padv2`.

func.func @pad_with_interior_padding(

  // CHECK-LABEL: padv2
  // CHECK:  "tfl.padv2"(%arg0, %arg1, %cst) : (tensor<2x1x3xf32>, tensor<3x2xi32>, tensor<f32>) -> tensor<?xf32>
  // CHECK:  return
}

func.func @padv2_5D(tensor<2x1x3x1x1xf32>, tensor<5x2xi32>) -> tensor<? x f32> {
^bb0(%arg0: tensor<2x1x3x1x1xf32>, %arg1: tensor<5x2xi32>):
  %cst = "tf.Const"() { value = dense<2.0> : tensor<f32> } : () -> tensor<f32>

  DELEGATE = 51,
  BIDIRECTIONAL_SEQUENCE_LSTM = 52,
  CAST = 53,
  PRELU = 54,
  MAXIMUM = 55,
  ARG_MAX = 56,
  MINIMUM = 57,
  LESS = 58,
  NEG = 59,
  PADV2 = 60,
  GREATER = 61,
  GREATER_EQUAL = 62,
  LESS_EQUAL = 63,
  SELECT = 64,
  SLICE = 65,
  SIN = 66,
  TRANSPOSE_CONV = 67,
  SPARSE_TO_DENSE = 68,
  TILE = 69,
  EXPAND_DIMS = 70,