0`i!">ddlmZmZddlZddlZddlmZejdej dej dej dej dej dej f d Z dd ej d ej d ej dee ej fdeej dee dej fdZ dS))OptionalUnionNBLOCK_KIS_SEQLEN_OFFSETS_TENSOR IS_VARLEN INTERLEAVED CONJUGATEBLOCK_Mc 8 tjd}tjd}tjd}|dz}|s||| zz||zz}||| zz|| zz}nPtj||z}tj||zdz|z }|||zz||zz}||| zz|| zz}||z|krdS||ztjd|z}|s||z}n|tj||zz}tjd|}tjd|dz}|sc||dddf|z|dddf|zzz}||dddf|z|dddfzz}||dddf|z|dddfzz}tj||dddf|k|dddf|kzdtj} tj||dddf|k|dddf|kzdtj}!tj||dddf|k|dddf|kzdtj}"tj|||zz|dddf|k|dddf|kzdtj}#|r|! }!|"| z|#|!zz }$|"|!z|#| zz}%||dddf| z|dddf| zzz}tj||$|dddf|k|dddf|kztj||| zz|%|dddf|k|dddf|kzdS||dzdzdzzdz }&tjd|dz}'||dddf|z|dddf|zzz}(||dddf|z|&dddf|zzz})||dddf|z|'dddfzz}||dddf|z|'dddfzz}tj||dddf|k|'dddf|kzdtj} tj||dddf|k|'dddf|kzdtj}!tj|(|dddf|k|dddf|kzdtj}"tj|)|dddf|k|&dddf|kzdtj}#|r|! }!|"| z}*|#|!z}+tj|dddfdzdk|*|+z |*|+z},||dddf| z|dddf| zzz}tj||,|dddf|k|dddf|kzdS) 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_headdimrrrrr r pid_mpid_head pid_batchrotary_dim_half start_idxrmrm_csrkrk_halfcossinx0x1o0o1rk_swap rk_repeatX0X1x0_cosx1_sinouts- z/home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/vllm/vllm_flash_attn/ops/triton/rotary.py rotary_kernelrA s: Mq ! ! !E}!$$$H 1%%%I AoO Q N* *X-G GI 0008>O3OOGJ233 i/!344y@ O+ +h.H HI 111H?P4PP w&   29Q00 0B #9^#RW^i7888 1g  Bi7a<((G BU AAAtG.qqq1ADT1TT UU111d7^o5aaa8HHIU111d7^o5aaa8HHIg uQQQW~ 1gdAAAg6F6XYad   "RZ.. g uQQQW~ 1gdAAAg6F6XYad   "RZ.. W R4[6)gdAAAg.>.PQY\   "RZ.. W "22 2QQQW+&747+;o+MN    "RZ..  $C #XS  #XS R4[#44wtQQQw7GJ\7\\] b111d7 f 4qqq9IO9[\]]]]  /$66 6 QQQW+&747+;o+MN      a1 ))A-Ia))Q. "QQQW+/"T111W+@P2PP Q "QQQW+/'$'2BEU2UU VU111d7^o5 $'8JJKU111d7^o5 $'8JJKg 4.9,471Co1UV    "RZ.. g 4.9,471Co1UV    "RZ.. WRr!!!T'{V347 j8PQY\ ] ] ] ` ` J  W bDkF*wtQQQw/?*/LMUX   "RZ..   $Ccchr$'{Q!+Vf_fvoNNR4[#44r$'{EW7WWX cAAAtGv!5"T111W+ :R STTTTTTFxr3r4seqlen_offsets cu_seqlens max_seqlenreturnc @|du} | s |j\} n+| Jd|j\} } |jd} | dz ||j\} }|j|jksJ|dz}|| ks Jd| dks Jd| ks Jd |j|jksJd |jd |j|j|jksJd |jd |j||}}t|tjrD|jfksJ|jtjtjfvsJ|}n |z| ksJ|st j|n|}|| kr+|s)|d |df |d |df|dkrdn|dkrdn |dkrdnd}fd}|rdn |dkrdnd}tj |j j 5t||||||||| | s|dnd|d|d|d| s|dnd|d|d|d|t|tj| ||||dkrdnddddn #1swxYwY|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 rzrotary_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 . @c@tj|dfS)Nr )tritoncdiv)METAbatchnheadsrs r@zapply_rotary..sVT)_==vuMrB) num_warps)shapedtype contiguous isinstancetorchTensorint32int64 empty_likecopy_cudadeviceindexrAstride)rCr3r4rDrErF interleavedinplace conjugate is_varlenheaddim total_seqlen batch_p_1r!r outputrgridr rQrRrs @@@r@ apply_rotaryrqs.$&I )*&vvww%%'c%%%()% fg$Q' A IIz 9 ! ! ! !!OJ  "A c>>>8>>>     =    SYOCIOOCIOO  39S!'SS SS ~~!1!1C.%,//4#x////# U['AAAAA'2244&)3333(/ 6U a QFGGsJKK &&qjkk)9':;;;    "$$bb*2C2C33 N M M M M MDCaa**;*;QQG   18> * *  d        $- 4FMM!   1 MM"   MM"   MM"  ( /AHHQKKKa HHRLL HHRLL HHRLL  ~u| 4 4    %++aa1                   6 MsC.LLL)rNNFFF)typingrrr^rNtriton.languagelanguagerjit constexprrAr_intrqrBr@rysh#""""""" vU*\+vU,!l-vU.|/vU01vU2|3vU4\5vUvUvU vUz01)- $ __ |_ _ _#u|+, _ & _  _ \______rB