o پiV@sddlZddlZddlmZddlmZejdejdejdejdejdejdejf d d Z ejdejdejdejdejdejd ejd ejd ejdejdejfddZ ejdejdejdejdejdejd ejd ejfddZ ejdejdejdejdejdejd ejdejdejfddZ Gdddej jZejZd*ddZejdejdejfddZ d+dejd ejd!ejd"ejd#ejd$ejded%ejfd&d'ZGd(d)d)ZdS),N) rearrangebhdeBLOCKCBLOCKc .Cs.td} | | }| | }td}||}|||}||| }||| }|| }||}|| }|| }|| }||}|| }||||td| dddf|td|dddf}|||td| dddf|td|dddf}|||td| dddf| td| dddf}||||td| dddf| td| dddf}||} t| }!|}"td| |"| }#tj|||#dddf|kddtj}$tj| | gtjd}%t|"dD]z}&td| |&| }'|#dddf|'dddf}(|!|(})t|(dk|) t d})t |)}*tj|||'dddf|kddtj}+tj|||'dddf|kddtj},t |$|+|*}-|%t |-|,7}%|| |7}|| | 7}qtj ||%|j j||#dddf|kddS)Nrmaskotherdtypez-infr )tl program_idarangeloadtofloat32zerosrangewherefloatexpdotstorer element_ty).QKVOutSrrnrrr NUM_BLOCKroffoff_bh off_block off_cblockoff_h qk_offsetv_offseto_offset block_offsetqk_block_offsetv_block_offseto_block_offset cblock_offsetq_cblock_offseto_cblock_offset Q_block_ptrK_trans_block_ptr V_block_ptr O_block_ptr S_block_ptrsiq_indexqqkvjkv_indexdiffs_indexdecayk_transvqkrGe/home/ubuntu/.local/lib/python3.10/site-packages/sglang/srt/layers/attention/linear/lightning_attn.py_fwd_diag_kernel s       $     rID_FBLOCKE_FBLOCK NUM_FBLOCK NUM_CBLOCKc(Cstd}td}||}|| }||}||}|||}|||}|||}|| ||}|||td|dddf|td| dddf}|||td|dddf|td| dddf}|||td| dddf|td| dddf}||| td|dddf}td|}tj| | gtjd}|| dkr|| d| } n| } t| ||| }!tt| ||}"|||"|7}t|"D]O}#d|#|!}$tj||!||dddf|$kdd}%tj||!||dddf|$kdd}&t|}'|t |%|'|&7}|||7}|||7}||7}qt || |j j dS)Nrr rr r )rrrrrcdivminrrrrrrr)(r r!K_decayKVrrr$rrrr%rJrKrLrrMr'r(r*r.k_block_offsetr0kv_block_offsetk_offsetr, kv_offsetr6r7 KV_block_ptr k_decay_ptrr@kvsplit_n left_shift num_blocksr? left_boundrDrEk_decayrGrGrH_fwd_kv_parallels     $          r^c CsRtd} | |} | | ||}||td| dddf|td| dddf}|| }t|}| ||}||td| dddf|td| dddf}t|tj}t| D]7}t||||}t|tj |}t|tj}t |||j j |||}|||7}qit ||dS)Nr) rrrrrrrrOrrrr)r#rQ KV_HISTORYrrr$rrrr%rJrKr'r*rUrVs_ptrsr:kv_history_offsetKV_HISTORY_block_ptrkv_prer; block_size block_decaykv_currGrGrH_fwd_kv_reduces>     rgc(Cs`td}||}td}|| }|| }td}|| }|| }|| }||}||||||}|||||||}|| ||||||}||td| dddf|td|dddf}||td| dddf|td| dddf}||td|dddf|td| dddf}||}t|}td| } t|tj}!|td| }"tj||"dddf|kddtj}#t|tj || | dddf}$t|#|!|$}%tj||"dddf|kddtj}&|&|%}'tj||'|j j |"dddf|kddS)Nrr r r r) rrrrrrrrrrr)(rr"r#rQrrr$rrrr%rKrrMr'r*off_ncoff_noff_coff_en_offsetc_offsete_offsetr.q_offsetr-rUr5r8rVr9r:c_arrayrXr<r=q_decay qkv_none_diagqkv_diagr>rGrGrH_fwd_none_diag_kernel7sL    ::  (,(  ruc@seZdZeddZdS) _attentioncCsR|}|}|}|}tj}|ddkr tdd|j\}}} } |jd} tj||| | f|j|jd} d} t | | }d}| |}| |dksRJd tj d| |jd d }t | | | d d}||||f}t||||| |||| | | | ||d d }| |}| |dksJ| |}| |dksJd }| |}| |dksJd tj|||| | ftj|jd}|||f}t|||||||| | | | ||||||d|||f}t||||||| | | | |||d ||||f}t||| ||||| | | | ||||d||||||| |_| tj||dgddfS)Nrz(Flash attention currently only supportedzfor compute capability >= 80rdevice z"BLOCK must be a multiple of CBLOCK)rzr )rr%r@)rr%rJrKrLrrM)rr%rJrK)rr%rKrrMrh)dim) contiguoustorchcudaget_device_capability RuntimeErrorshapeemptyrrztritonrNrrreshaperIrr^rgrusave_for_backwardrcat unsqueeze)ctxr=krEr: kv_history capabilityrrr$rrorr%rrMarrayr]gridrLrJrKrXrGrGrHforwards      z_attention.forwardN)__name__ __module__ __qualname__ staticmethodrrGrGrGrHrvsrvr{csJ|jd}|jd}|dkr|dddd}|dkrdnd|dks1Jd|ddfd d t|dD}|d|krK||t|} d} |d urjtj|jd|jd||ftj|j d }n| }t| dD]*} || } || d}|d | |f} |d | |f}t | ||||\}}| |} qv| |fS)a$ Apply lightning attention algorithm to compute attention efficiently. Args: q: Query tensor of shape [batch, heads, seq_len, dim] k: Key tensor of shape [batch, heads, seq_len, dim] v: Value tensor of shape [batch, heads, seq_len, dim_v] ed: Decay rate tensor of shape [heads] block_size: Size of blocks for block-sparse attention kv_history: Optional key-value history from previous computations Returns: output: Attention output kv: Updated key-value history rxr r}rz Dimension d (z) must be divisible by m ()csg|]}|qSrGrG).0r;mrGrH ,sz'lightning_attention..Nry.) rr~viewrappendlenrrrrzclonerlightning_attention_)r=rrEedrdrrrarrr$outputr;r:q1k1rrXrGrrHlightning_attentions0   "     rD BLOCK_SIZEc'Cstd}td}td}t||}|dkrdS|}|}t||}td|}td|||}|dddf| |dddf| }|||| }|||| }|||| }|| || }||k}||k}tj||||dd}tj||||dd} tj||||dd}!| dddf|!dddf}"|dddf|dddf@}#t| }|||}$tj|$|#dd}%|"||%}"|dddftj|"}&tj|&dd}&tj|$|"|#d tj||||&|d dS) z Kernel for linear attention decoding with KV cache. This kernel computes attention for a single token using the KV cache. rr rhrxNr r )axisr) rrrrrrrsumr)'q_ptrk_ptrv_ptr kv_cache_ptr slope_rateslot_idx output_ptrr qkv_b_stride qkv_h_stridecache_b_stridecache_h_stridecache_d0_stridecache_d1_striderpid_bpid_hpid_dslot_idbatch_idhead_idratio qk_d_offsets v_d_offsetscache_d_offsetsrprTr, cache_offsetqk_maskv_maskr=rrEkv_outerkv_maskkv_ptr kv_cache_oldrrGrGrH_linear_attn_decode_kernelEs@    &     rr|r=rrE kv_cachesrrreturncCs|j\}}} } |j||d| fksJ|j||d| fksJt|} ||| |f} |d} |d}|d}|d}|d}|d}t| ||||||| | | ||||||dt| d} | dS)a Perform linear attention decoding using Triton kernels. Args: q: Query tensor of shape [B, H, 1, D] k: Key tensor of shape [B, H, 1, D] v: Value tensor of shape [B, H, 1, D] kv_caches: Key-value cache tensor slope_rate: Decay rate tensor slot_idx: Slot indices for batches BLOCK_SIZE: Size of blocks for processing Returns: output: Attention output tensor r rrh)rzb h n d -> b n (h d))rr empty_likestriderrsqueezer)r=rrErrrrBH_rrrrrrrrrrGrGrHlinear_decode_forward_tritons<        rc@sNeZdZdZe d dejdejdejdejdejded ed ejfd d ZdS)BailingLinearKernela Linear attention kernel implementation for Bailing models. This class is adapted from MiniMaxText01LinearKernel in vllm: https://github.com/vllm-project/vllm/blob/a9138e85b14047e06300685b48e3485b995425fb/vllm/model_executor/models/minimax_text_01.py#L289 The implementation maintains the same functionality while being renamed to match our Bailing model naming convention. Nr=rrErrrd layer_idxrc Ks|tj}|dk}|r|d}|d}|d}|j\} } } } | } |d| | | }t||||||d\}}| |dddddddddf| | | |jddks`Jd| d dd| | | gS)Nrrr )rdrrxzbatch size must be 1) rrrr~rrrrrcopy_r transpose)r=rrErrrdrkwargsshould_pad_dimrrr$rrrrrGrGrHjit_linear_forward_prefixs      2$z-BailingLinearKernel.jit_linear_forward_prefix)N) rrr__doc__rrTensorintrrGrGrGrHrs*  r)r{N)r|)rrtriton.languagelanguagereinopsrjit constexprrIr^rgruautogradFunctionrvapplyrrrrrrrrGrGrGrHs           i   @     R 3 X C