%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>

    appended at the end).
  }];

  let arguments = (ins
    TFL_TensorOf<[I32, I64]>:$indices,
    TFL_I32Tensor:$depth,
    TFL_TensorOf<[F32, I32, I64, I1, I8, UI8]>:$on_value,
    TFL_TensorOf<[F32, I32, I64, I1, I8, UI8]>:$off_value,

    I32Attr:$axis
  );

  let results = (outs
    TFL_TensorOf<[F32, I32, I64, I1, I8, UI8]>:$output
  );

  let hasOptions = 1;
}

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

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>

          "tf.Yield"(%6) : (tensor<i64>) -> ()
        }, {
          "tf.Yield"(%arg3) : (tensor<i64>) -> ()
        }) {_lower_using_switch_merge = true, is_stateless = false} : (tensor<i1>) -> tensor<i64>
        "tf.Yield"(%3, %4) : (tensor<i1>, tensor<i64>) -> ()

}

// CHECK-LABEL: testSubInt64
func.func @testSubInt64(tensor<? x i64>, tensor<? x i64>) -> tensor<? x i64> {
^bb0(%arg0: tensor<? x i64>, %arg1: tensor<? x i64>):
  // CHECK: tfl.sub %arg0, %arg1 {fused_activation_function = "RELU6"}
  %0 = tfl.sub %arg0, %arg1 {fused_activation_function = "RELU6"} : tensor<? x i64>
  func.return %0#0 : tensor<? x i64>
}

// CHECK-LABEL: testMul

    %read1 = "tf.ReadVariableOp"(%arg1) : (tensor<!tf_type.resource<tensor<i32>>>) -> tensor<i32>
    // CHECK-NOT: tf.TPUPartitionedInputV2
    %partitioned_input = "tf.TPUPartitionedInputV2"(%read0, %read1) {N = 2 : i64, partition_dims = []} : (tensor<i32>, tensor<i32>) -> tensor<i32>
    // CHECK: %[[COMPILE_OUTPUT:[0-9]*]]:3 = "tf_device.launch"
    // CHECK-NEXT: "tf._TPUCompileMlir"()
    // CHECK: "tf_device.launch"

  return rewriter.create<TFL::QConstOp>(
      loc, /*output=*/TypeAttr::get(bias_type), /*value=*/bias_value);
}

// Casts the given op shapes from i64 to i32 to fit TFLite spec requirement.
arith::ConstantOp CreateI32ShapeConstantOp(const TensorType op_type,
                                           const Location loc,

// for each dim i, the output tensor is identical to `input`.
bool CanOptimizeIdentitySliceOp(Value input, Attribute begin, Attribute size) {
  // Checks if `begin` and `size` are i32 or i64.
  auto begin_attr = mlir::dyn_cast<DenseIntElementsAttr>(begin);
  auto size_attr = mlir::dyn_cast<DenseIntElementsAttr>(size);
  if (!begin_attr || !size_attr) {
    return false;
  }