// When begin/end values are dynamic, the ellipsis mask, if set, must refer
    // to the last dimension.
    int ellipsis_mask = op.getEllipsisMask();
    if (!(ellipsis_mask == 0 || ellipsis_mask == (1 << last_dim)))
      return rewriter.notifyMatchFailure(
          op,
          "requires that ellipsis_mask, if set, refer to the last dimension of "
          "input (when begin/end values are dynamic)");

      builder->create<TF::ConstOp>(loc, builder->getI32TensorAttr({idx + 1})),
      /*strides=*/
      builder->create<TF::ConstOp>(loc, builder->getI32TensorAttr({1})),
      /*begin_mask=*/0, /*end_mask=*/0, /*ellipsis_mask=*/0,
      /*new_axis_mask=*/0, /*shrink_axis_mask=*/1);
}

// Return true if the passed quantized type is unsigned.
bool QuantizedTypeIsUnsigned(Type type) {
  return TypeSwitch<Type, bool>(type)

  %0 = "tfl.dequantize"(%arg0) : (tensor<12x2x2x5x!quant.uniform<u8:f32, 0.1>>) -> tensor<12x2x2x5xf32>
  %1 = "tfl.strided_slice"(%0, %arg1, %arg2, %arg3) {begin_mask = 0 : i32, ellipsis_mask = 0 : i32, end_mask = 0 : i32, new_axis_mask = 0 : i32, shrink_axis_mask = 0 : i32, offset = false} : (tensor<12x2x2x5xf32>, tensor<1xi32>, tensor<1xi32>, tensor<1xi32>) -> tensor<1x2x2x5xf32>
  func.return %1 : tensor<1x2x2x5xf32>

  // The kernel will reshape the input tensor with new axis, it only supports
  // this reshaped tensor up to 5D.
  uint32_t ellipsis_mask = op.getEllipsisMask();
  uint32_t new_axis_mask = op.getNewAxisMask();
  int num_added_axis = 0;
  for (int i = 0; i < 8; ++i) {
    if (!((1 << i) & ellipsis_mask) && ((1 << i) & new_axis_mask)) {
      num_added_axis++;
    }
  }

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

    rewriter.replaceOpWithNewOp<TFL::StridedSliceOp>(
        op, output_type, op.getOperand(), start_idx, slice_size, stride,
        /*begin_mask=*/0, /*end_mask=*/0,
        /*ellipsis_mask=*/0, /*new_axis_mask=*/0, /*shrink_axis_mask=*/0,
        /*offset=*/false);
  }
};

// Rewrites quantized stablehlo.broadcast_in_dim to tfl.broadcast_to.

  auto _end_mask = _o->end_mask;
  auto _ellipsis_mask = _o->ellipsis_mask;
  auto _new_axis_mask = _o->new_axis_mask;
  auto _shrink_axis_mask = _o->shrink_axis_mask;
  auto _offset = _o->offset;
  return tflite::CreateStridedSliceOptions(
      _fbb,
      _begin_mask,
      _end_mask,
      _ellipsis_mask,
      _new_axis_mask,
      _shrink_axis_mask,
      _offset);

    TF_Tensor:$value,

    DefaultValuedOptionalAttr<I64Attr, "0">:$begin_mask,
    DefaultValuedOptionalAttr<I64Attr, "0">:$end_mask,
    DefaultValuedOptionalAttr<I64Attr, "0">:$ellipsis_mask,
    DefaultValuedOptionalAttr<I64Attr, "0">:$new_axis_mask,
    DefaultValuedOptionalAttr<I64Attr, "0">:$shrink_axis_mask
  );

  let results = (outs);

  // expected-error @+1 {{cannot have multiple ellipses}}
  %0 = "tf.StridedSlice"(%input, %begin, %end, %strides) {ellipsis_mask = 3}: (tensor<4x8xf32>, tensor<2xi64>, tensor<2xi64>, tensor<2xi64>) -> tensor<?x?xf32>
  func.return %0 : tensor<?x?xf32>
}

// -----

// CHECK: %[[STRIDE:.+]] = arith.constant
// CHECK{LITERAL}: dense<[2, 3]> : tensor<2xi32>