@triton.jit
def _sgmv_expand_slice_kernel(input_ptr, lora_ptr, out_ptr, N, K,
    b_seq_start_loc, seq_lens, lora_indices, xm_stride, xk_stride,
    l0_stride, lora_k_stride, lora_n_stride, cm_stride, cn_stride,
    slice_offset, BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl
    .constexpr, EVEN_K: tl.constexpr, ADD_INPUTS: tl.constexpr, CAST_TYPE:
    tl.constexpr):
    """

    Similar to the 'sgmv_expand' operator, but with an added parameter 
    'slice_offset'. The reason for not reusing the 'sgmv_expand' operator 
    might be that in the future, we could implement a fusion operator to 
    achieve the current functionality instead of having to call it multiple 
    times.
    """
    pid = tl.program_id(axis=0)
    cur_batch = tl.program_id(axis=1)
    cta_n_num = tl.cdiv(N, BLOCK_N)
    pid_m = pid // cta_n_num
    pid_n = pid % cta_n_num
    M = tl.load(seq_lens + cur_batch)
    if pid_m * BLOCK_M > M:
        return
    lora_index = tl.load(lora_indices + cur_batch)
    if lora_index == -1:
        return
    cur_seq_start = tl.load(b_seq_start_loc + cur_batch)
    offset_m = tl.arange(0, BLOCK_M) + pid_m * BLOCK_M
    offset_n = tl.arange(0, BLOCK_N) + pid_n * BLOCK_N
    offset_k = tl.arange(0, BLOCK_K)
    ram = tl.max_contiguous(tl.multiple_of(offset_m % M, BLOCK_M), BLOCK_M)
    rbn = tl.max_contiguous(tl.multiple_of(offset_n % N, BLOCK_N), BLOCK_N)
    a_ptr = input_ptr + cur_seq_start * xm_stride + ram[:, None
        ] * xm_stride + offset_k[None, :] * xk_stride,
    b_ptr = lora_ptr + l0_stride * lora_index + offset_k[:, None
        ] * lora_n_stride + rbn[None, :] * lora_k_stride
    accumulator = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)
    for k in range(tl.cdiv(K, BLOCK_K)):
        if EVEN_K:
            tiled_a = tl.load(a_ptr)
            tiled_b = tl.load(b_ptr)
        else:
            tiled_a = tl.load(a_ptr, mask=offset_k[None, :] < K - k *
                BLOCK_K, other=0)
            tiled_b = tl.load(b_ptr, mask=offset_k[:, None] < K - k *
                BLOCK_K, other=0)
        if CAST_TYPE:
            tiled_a = tiled_a.to(lora_ptr.dtype.element_ty)
        accumulator += tl.dot(tiled_a, tiled_b)
        a_ptr += BLOCK_K * xk_stride
        b_ptr += BLOCK_K * lora_n_stride
    tiled_c = accumulator.to(lora_ptr.dtype.element_ty)
    offset_cm = cur_seq_start + tl.arange(0, BLOCK_M) + pid_m * BLOCK_M
    offset_cn = tl.arange(0, BLOCK_N) + pid_n * BLOCK_N + slice_offset
    c_ptr = out_ptr + offset_cm[:, None] * cm_stride + offset_cn[None, :
        ] * cn_stride
    M = tl.load(seq_lens + cur_batch)
    c_mask = (offset_cm[:, None] < cur_seq_start + M) & (offset_cn[None, :] <
        slice_offset + N)
    if ADD_INPUTS:
        tiled_out = tl.load(c_ptr, mask=c_mask, other=None)
        tiled_c += tiled_out
    tl.store(c_ptr, tiled_c, mask=c_mask)
