o ۾i!"@sddlmZmZddlZddlZddlmZejdej dej dej dej dej dej f d d Z   dd ej d ej dej dee ej fdeej dee dej fddZ dS))OptionalUnionNBLOCK_KIS_SEQLEN_OFFSETS_TENSOR IS_VARLEN INTERLEAVED CONJUGATEBLOCK_Mc-Cstjdd}tjdd}tjdd}|d}|s-||| ||}||| || }n&t||}t||d|}|||||}||| || }|||kr[dS||td|}|sl||}n |t||}td|}td|d}|s||dddf||dddf|}||dddf||dddf}||dddf||dddf}tj||dddf|k|dddf|k@ddtj} tj||dddf|k|dddf|k@ddtj}!tj||dddf|k|dddf|k@ddtj}"tj||||dddf|k|dddf|k@ddtj}#|rF|! }!|"| |#|!}$|"|!|#| }%||dddf| |dddf| }tj||$|dddf|k|dddf|k@dtj||| |%|dddf|k|dddf|k@ddS||dddd}&td|d}'||dddf||dddf|}(||dddf||&dddf|})||dddf||'dddf}||dddf||'dddf}tj||dddf|k|'dddf|k@ddtj} tj||dddf|k|'dddf|k@ddtj}!tj|(|dddf|k|dddf|k@ddtj}"tj|)|dddf|k|&dddf|k@ddtj}#|r|! }!|"| }*|#|!}+t|dddfddk|*|+|*|+},||dddf| |dddf| }tj||,|dddf|k|dddf|k@ddS) Nr)axisg?)maskotherg)r )tl program_idloadarangetofloat32storewhere)-OUTXCOSSIN CU_SEQLENSSEQLEN_OFFSETSseqlen rotary_dim seqlen_rostride_out_batchstride_out_seqlenstride_out_nheadsstride_out_headdimstride_x_batchstride_x_seqlenstride_x_nheadsstride_x_headdimrrrrrr pid_mpid_head pid_batchrotary_dim_half start_idxrmrm_csrkrk_halfcossinx0x1o0o1rk_swap rk_repeatX0X1x0_cosx1_sinoutr>Z/home/ubuntu/.local/lib/python3.10/site-packages/vllm/vllm_flash_attn/ops/triton/rotary.py rotary_kernel s      ,((*** &,4 &  ,,((&&4**,8r@Fxr1r2seqlen_offsets cu_seqlens max_seqlenreturnc s|du} | s|j\} n|dusJd|j\} } |jd} | d||j\} }|j|jks4J|d9}|| ks@Jd| dksHJd| ksPJd |j|jksbJd |jd |j|j|jkstJd |jd |j||}}t|tjr|jfksJ|jtjtjfvsJ|}n|| ksJ|st|n|}|| kr|s|d |df |d |df|dkrdn |dkrdn|dkrdnd}fdd}|rdn|dkrdnd}tj |j j Tt ||||||||| | s|dnd|d|d|d| s|dnd|d|d|d|t|tj| ||||dkr;dnddWd|S1sKwY|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 >= seqlenz*cos and sin must have the same dtype, got z and z0Input and cos/sin must have the same dtype, got . @cst|dfS)Nr )tritoncdiv)METAbatchnheadsrr>r?szapply_rotary..) num_warps)shapedtype contiguous isinstancetorchTensorint32int64 empty_likecopy_cudadeviceindexr@stride)rAr1r2rBrCrD interleavedinplace conjugate is_varlenheaddim total_seqlen batch_p_1rroutputrgridr r>rMr? apply_rotarys      "  rn)rNNFFF)typingrrr[rJtriton.languagelanguagerjit constexprr@r\intrnr>r>r>r?sN  }