o i@sddlmZmZddlZddlZddlmZejdej dej dej dej dej dej d ej fd d Z   dd ej dej dej dee ej fdeej dee dej fddZ dS))OptionalUnionN ROTARY_DIMIS_SEQLEN_OFFSETS_TENSOR IS_VARLEN INTERLEAVED CONJUGATEBLOCK_HBLOCK_Mc- Csdt|}|d}tjdd}tjdd}tjdd}|s*||| }||| }nt||}t||d|}|||}||| }|||krPdS||td|}||td|}|sk||} n |t||} td|d}!|| dddf||!dddf}|| dddf||!dddf}| dddf|k|!dddf|k@}"tj||"ddtj}#tj||"ddtj}$|r|$ }$|s||ddddf||ddddf||!ddddf|}||ddddf| |ddddf| |!ddddf| }|ddddf|k|ddddf|k@|!ddddf|k@}%tj||%ddtj}&tj||||%ddtj}'|&|#|'|$}(|&|$|'|#})tj||(|%dtj||| |)|%ddStd|}*||ddddf||ddddf||*ddddf|}||ddddf| |ddddf| |*ddddf| }|ddddf|k|ddddf|k@|*ddddf|k@}%tj||%ddtj}+t t |+|||ddg\}&}'|&|#|'|$}(|&|$|'|#})t t |(|)|||g},tj||,|%ddS) Nr)axisg?)maskotherg)r) tritonnext_power_of_2tl program_idloadarangetofloat32storesplitreshapejoin)-OUTXCOSSIN CU_SEQLENSSEQLEN_OFFSETSseqlennheads seqlen_rostride_out_batchstride_out_seqlenstride_out_nheadsstride_out_headdimstride_x_batchstride_x_seqlenstride_x_nheadsstride_x_headdimrrrrrr r BLOCK_KROTARY_DIM_HALFpid_headpid_m pid_batch start_idxrhrmrm_csrk_halfmask_cscossinrx0x1o0o1rkxorAR/home/ubuntu/vllm_env/lib/python3.10/site-packages/flash_attn/ops/triton/rotary.py rotary_kernel s^         (((FFB  FFB"rCFr?r8r9seqlen_offsets cu_seqlens max_seqlenreturnc s"|du} | s|j\} n|dusJd|j\} } |jd} | d||j\} }|j|jks4J|d9}|| ks@Jd| dksHJd| ksPJd ||}}t|tjrw|jfksgJ|jtjtjfvsrJ|}n|| ksJ|st|n|}|| kr|s|d |df |d |dffd d }|d krdnd}tj |j j Otj t|||||||| | s|dnd|d|d|d| s|dnd|d|d|d|t|tj| |||ddWd|S1s wY|S)a Arguments: x: (batch, seqlen, nheads, headdim) if cu_seqlens is None else (total_seqlen, nheads, headdim). cos: (seqlen_ro, rotary_dim / 2) sin: (seqlen_ro, rotary_dim / 2) seqlen_offsets: integer or integer tensor of size (batch,) cu_seqlens: (batch + 1,) or None max_seqlen: int Returns: y: (batch, seqlen, nheads, headdim) Nz:If cu_seqlens is passed in, then max_seqlen must be passedrr r zrotary_dim must be <= headdimzOnly support headdim <= 256zseqlen_ro must be >= seqlen.cs"t|dt|dfS)Nr r )rcdiv)METAbatchr#r"rArBs"zapply_rotary..)r r )shape contiguous isinstancetorchTensordtypeint32int64 empty_likecopy_cudadeviceindexlibrary wrap_tritonrCstride)r?r8r9rDrErF interleavedinplace conjugate is_varlenheaddim total_seqlen batch_p_1r$ rotary_dimoutputgridr rArKrB apply_rotaryfsp      "  rn)rNNFFF)typingrrrWrtriton.languagelanguagerjit constexprrCrXintrnrArArArBsR  ]