// Dynamic start indices should be clamped with upper limit of
      // `shape(operand) - slice_sizes)` as per semantics of
      // `stablehlo.dynamic_slice`.
      // (https://github.com/openxla/stablehlo/blob/main/docs/spec.md#dynamic_slice)
      start_indices[i] =
          rewriter.create<TFL::MinimumOp>(op->getLoc(), cast, upper_limit_cst);
    }

    Value concatenated = start_indices[0];

// CHECK-NOT: tfl.gather

// -----

// Tests that a quantized `stablehlo.dynamic_slice` is converted to `tfl.slice`.

// CHECK-LABEL: func @dynamic_slice
// CHECK-SAME: %[[ARG0:.+]]: tensor<4x4x!quant.uniform<i8:f32, 3.000000e-01:-5>>, %[[ARG1:.+]]: tensor<i64>, %[[ARG2:.+]]: tensor<i64>
func.func @dynamic_slice(
    %arg0: tensor<4x4x!quant.uniform<i8:f32, 3.000000e-01:-5>>,
    %arg1: tensor<i64>,

    Type start_indices_element_type =
        mlir::cast<ShapedType>(op.getStartIndices().front().getType())
            .getElementType();

    // The mhlo dynamic_slice's start_indices can be either signed/unsigned
    // int32/int64. However, TF only takes in either i32 or i64 types for begin,
    // so we will always put a cast.
    Type signed_start_indices_element_type;