Support Position Ids in Rope

[fabianlim]

This could be very possible by just providing the correct sin and cos values adjusted according to position ids. This can be done outside of the kernel and then passed in:

def _rope_embedding(
    Q,     Q_row_stride,
    cos, cos_row_stride,
    sin, sin_row_stride,
    seqlen,
    head_dim      : tl.constexpr,
    n_heads       : tl.constexpr,
    BACKWARD_PASS : tl.constexpr,
    BLOCK_SIZE    : tl.constexpr,
):

[fabianlim]

actually just realized that actually the position_ids arg is depracted huggingface/transformers#29143 and the sin, cos are handled outside of the function. So closing this