/*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>

// CHECK: %0 = "tfl.dequantize"(%arg0)

  // 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++;
    }
  }
  if (num_input_dims + num_added_axis > 5) return failure();
  return success();
}

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

  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);
}

    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.

        GetI64ElementsAttr(hlo_end_indices, &rewriter),
        GetI64ElementsAttr(hlo_strides, &rewriter));

    // Reshape slice result so that the shape is updated depending on
    // 'new_axis_mask' or 'shrink_axis_mask' attributes.
    rewriter.replaceOpWithNewOp<ReshapeOp>(op, op.getType(), sliced);
    return success();
  }

  LogicalResult rewriteWithUnknownBegin(TF::StridedSliceOp op,

    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);

  TF_DerivedOperandTypeAttr Index = TF_DerivedOperandTypeAttr<1>;

// CHECK{LITERAL}: dense<[2, 3]> : tensor<2xi32>
// CHECK: %[[SLICE:.+]] = "tfl.strided_slice"(%[[ARG0]], %[[START]], %[[SIZE]], %[[STRIDE]]) <{begin_mask = 0 : i32, ellipsis_mask = 0 : i32, end_mask = 0 : i32, new_axis_mask = 0 : i32, offset = false, shrink_axis_mask = 0 : i32}> : (tensor<3x6x!quant.uniform<i8:f32, 2.000000e+00:-1>>, tensor<2xi32>, tensor<2xi32>, tensor<2xi32>) -> tensor<2x2x!quant.uniform<i8:f32, 2.000000e+00:-1>>

    TFL_I32Tensor:$begin,
    TFL_I32Tensor:$end,
    TFL_I32Tensor:$strides,

    I32Attr:$begin_mask,
    I32Attr:$end_mask,
    I32Attr:$ellipsis_mask,
    I32Attr:$new_axis_mask,
    I32Attr:$shrink_axis_mask,
    BoolAttr:$offset
  );

  let results = (outs
    TFL_TensorOf<[F32, I32, I64, I8, UI8, UI32, QI8, QUI8, I1, I16, QI16, TFL_Quint8, TFL_Str]>:$output
  );