@triton.jit
def triton_jagged_to_dense(jagged_value_ptr, jagged_offsets_ptr,
    jagged_value_row_stride, output_dense_ptr, dense_indices_ptr,
    dense_col_stride, dense_row_stride, dense_matrix_stride, JAGGED_DIM: tl
    .constexpr, thread_block_row_size: tl.constexpr, thread_block_col_size:
    tl.constexpr, operation_function: tl.constexpr, operation_dense) ->None:
    pid = tl.program_id(0)
    begin = tl.load(jagged_offsets_ptr + pid)
    end = tl.load(jagged_offsets_ptr + (pid + 1))
    jagged_value_ptr += begin * jagged_value_row_stride
    if JAGGED_DIM > 2:
        dense_indice = tl.load(dense_indices_ptr + pid)
        if dense_indice == -1:
            return
        output_dense_ptr += dense_indice
        if operation_function is not None:
            operation_dense += dense_indice
    else:
        output_dense_ptr += pid * dense_matrix_stride
        if operation_function is not None:
            operation_dense += pid * dense_matrix_stride
    offset_row = tl.arange(0, thread_block_row_size)
    N = tl.minimum(dense_row_stride, jagged_value_row_stride)
    M = tl.minimum(dense_matrix_stride // dense_row_stride, end - begin)
    for _i in range(begin, end, thread_block_row_size):
        offset_col = tl.arange(0, thread_block_col_size)
        block_offset = offset_row[:, None] * dense_row_stride + offset_col[
            None, :] * dense_col_stride
        for _j in range(0, N, thread_block_col_size):
            mask = (offset_row[:, None] < M) & (offset_col[None, :] < N)
            jagged_val = tl.load(jagged_value_ptr + block_offset, mask=mask,
                other=0)
            if operation_function is not None:
                val1 = jagged_val
                val2 = tl.load(operation_dense + block_offset, mask=mask,
                    other=0)
                if operation_function == 'add':
                    jagged_val = tensor_elementwise_add(val1, val2)
                else:
                    jagged_val = tensor_elementwise_mul(val1, val2)
            tl.store(output_dense_ptr + block_offset, jagged_val, mask=mask)
            offset_col += thread_block_col_size
            block_offset += thread_block_col_size
        offset_row += thread_block_row_size
