let arguments = (ins
     TFL_TensorOf<[I8, I16, I32, I64, UI8, UI16, UI32, UI64, F32, F64]>:$input,
     TFL_TensorOf<[I32]>:$dilations,
     TFL_0DTensorOf<[I8, I16, I32, I64, UI8, UI16, UI32, UI64, F32, F64]>:$padding_value
  );

  let results = (outs TFL_TensorOf<[I8, I16, I32, I64, UI8, UI16, UI32, UI64, F32, F64]>:$output);
}

def TFL_AddOp : TFL_Op<"add", [

    if (!type || type.getElementType().isUnsignedInteger(64)) {
      return rewriter.notifyMatchFailure(dynamic_iota_op,
                                         "TF::RangeOp doesn't support UI64");
    }
    // Only support 1D for now.
    if (type.getRank() > 1 || dynamic_iota_op.getIotaDimension() != 0) {
      return rewriter.notifyMatchFailure(
          dynamic_iota_op, [&](::mlir::Diagnostic& diag) {

      lhs_quantization_axis = -1 : i64,
      lhs_quantization_min_val = -128 : i64,
      lhs_quantization_max_val = 127 : i64,
      rhs_quantization_axis = -1 : i64,
      rhs_quantization_min_val = -128 : i64,
      rhs_quantization_max_val = 127 : i64,
      output_quantization_axis = -1 : i64,
      output_quantization_min_val = -2147483648 : i64,
      output_quantization_max_val = 2147483647 : i64} : (

  func.return %0 : tensor<1x1x5x6x7xf32>
  // CHECK-LABEL: strided_slice_big_dims

  %2 = "tf.Shape"(%arg0) : (tensor<?x1x2x?xf32>) -> tensor<4xi32>
  %3 = "tf.StridedSlice"(%2, %1, %0, %1) {begin_mask = 0 : i64, ellipsis_mask = 0 : i64, end_mask = 0 : i64, new_axis_mask = 0 : i64, shrink_axis_mask = 0 : i64} : (tensor<4xi32>, tensor<1xi32>, tensor<1xi32>, tensor<1xi32>) -> tensor<2xi32>
  func.return %3 : tensor<2xi32>

  func.func @map_outside_compilation_must_be_ranked() -> () {
    "tf_device.cluster"() ({
      %0 = "tf.OpA"() : () -> tensor<*xi64>
      // expected-error @+1 {{must be ranked}}
      %1 = "tf.OpB"(%0) {_xla_map_outside_compilation = "0", _xla_outside_compilation = "from_launch"} : (tensor<*xi64>) -> tensor<*xi64>
      tf_device.return

    %one = "tf.Const"() {value = dense<1> : tensor<4xi32>} : () -> tensor<4xi32>
    %0 = "tf.StridedSlice"(%input, %zero, %one, %one) {begin_mask = 15 : i64, device = "", ellipsis_mask = 0 : i64, end_mask = 0 : i64, new_axis_mask = 0 : i64, shrink_axis_mask = 0 : i64} : (tensor<1024x1024x1024x1024xf32>, tensor<4xi32>, tensor<4xi32>, tensor<4xi32>) -> tensor<*xf32>

    // CHECK: tensor<1x1x1x1xf32>
    func.return %0 : tensor<*xf32>

func.func @dynamic_slice(
    %arg0: tensor<4x4x!quant.uniform<i8:f32, 3.000000e-01:-5>>,
    %arg1: tensor<i64>,
    %arg2: tensor<i64>
  ) -> tensor<2x1x!quant.uniform<i8:f32, 3.000000e-01:-5>> {
  %0 = "stablehlo.dynamic_slice"(%arg0, %arg1, %arg2) {
    slice_sizes = array<i64: 2, 1>
  } : (

// CHECK:           [[VAL_24:%.*]] = "tf.StridedSlice"([[VAL_20]], [[VAL_21]], [[VAL_22]], [[VAL_23]]) <{begin_mask = 6 : i64, ellipsis_mask = 0 : i64, end_mask = 6 : i64, new_axis_mask = 0 : i64, shrink_axis_mask = 1 : i64}> : (tensor<?x8x10xf32>, tensor<3xi32>, tensor<3xi32>, tensor<3xi32>) -> tensor<8x10xf32>

      // expected-remark@above {{ID: 0}}
      %1 = "tf.RandomUniform"(%0) {device = "", seed = 3 : i64, seed2 = 5 : i64} : (tensor<1xi32>) -> tensor<3xf32>
      // expected-remark@above {{ID: 1}}
      // expected-remark@above {{Successors: {4}}}
      %2 = "tf.RandomUniform"(%0) {device = "", seed = 3 : i64, seed2 = 5 : i64} : (tensor<1xi32>) -> tensor<3xf32>
      // expected-remark@above {{ID: 2}}