o ÓÙ¾izaã@sRddlZddlZddlmZddlmZddlZddlmZddlm Z ddl m m mm Z ddlmZddlmZeGdd „d eƒƒZeGd d „d eƒƒZejd ejed ejfdd„ƒZej  ddejdejdejdedejd ejedejefdd„ƒZej  ddejdejdejdejd ejedejef dd„ƒZdS)éN)ÚTuple)Ú dataclass)ÚFloat32é)Ú ParamsBase)Ú SeqlenInfoQKc @sbeZdZUeed<ejeed<ej ed<ej ed<dZ ejeed<dZ edBed<e  d%dedejedejededBfd d „ƒZ d&d d „Z d'dejdeeBdBd efdd„Z d'dejdeeBdBd efdd„Zej  d(dej dejedejed ej fdd„ƒZej d)dedeej BdBd ej fdd „ƒZejd!ej d"ej d dfd#d$„ƒZdS)*ÚSoftmaxÚ scale_log2Únum_rowsÚrow_maxÚrow_suméPÚarchNÚ softmax_scalecCs*t |t¡}t |t¡}t||||||ƒS©N)ÚcuteÚmake_rmem_tensorrr)r r rrr r ©rúb/home/ubuntu/.local/lib/python3.10/site-packages/sglang/jit_kernel/flash_attention/cute/softmax.pyÚcreates  zSoftmax.createÚreturncCs |j tj ¡|j d¡dS)Nç)r ÚfillrÚinfr )ÚselfrrrÚreset%sz Softmax.resetÚ acc_S_rowÚinit_valcCótj|||jdS©N)r)ÚutilsÚ fmax_reducer)rrrrrrÚ_compute_row_max)ózSoftmax._compute_row_maxÚ acc_S_row_expcCrr)r Ú fadd_reducer)rr$rrrrÚ_compute_row_sum.r#zSoftmax._compute_row_sumFTÚacc_SÚis_firstÚ check_infcCsht |¡}t |jt¡}|j}|j}|j}|j} t j t  |¡ddD]} || df  ¡} tj | t  | ¡r9|| nd| d} tj| tjjdd} || } | || <t  |¡r_| tj kr]dn| } t  |¡r~| |}t | ||¡}tj|d| d}d|| <n&| |}t | ||¡}t | | |¡|| <tj||| || | d}||| <|| df |¡q"|S) zÝApply online softmax and return the row_scale to rescale O. :param acc_S: acc_S tensor :type acc_S: cute.Tensor :param is_first: is first n_block :type is_first: cutlass.Constexpr T©Ú unroll_fullN)rré©Úwidthrçð?)r Úmake_acc_tensor_mn_viewrÚmake_fragment_liker rr r rÚcutlassÚrangeÚsizeÚloadr!Ú const_exprÚ warp_reduceÚfmaxrÚexp2fr%Ústore)rr'r(r)Úacc_S_mnÚ row_scaler r r rÚrrÚ row_max_curÚ row_max_prevÚrow_max_cur_scaledr$Ú acc_S_row_sumrrrÚonline_softmax3s@ ý   ÿzSoftmax.online_softmaxr/Ú final_scaleÚsink_valc Cs\t |duo t|tjƒ¡rt |¡t |j¡ksJ‚|j}|j}|j}|  t j |  ¡t jdd¡t |t¡}tjt |¡ddD]j}t |du¡rot|tjƒsR|n||}t tj¡} ||t  || |||¡7<||dkp|||||k} tj | s†||nd¡|||<||} t d¡} | s¥|||t  | ¡| ntj ||<qA|S) zAFinalize the online softmax by computing the scale and logsumexp.Nr,r-Tr*rr/g@)r2r6Ú isinstancerÚTensorr4r r r r:r r7r5ÚoperatorÚaddr1rr3ÚmathÚlog2Úer9rÚ rcp_approxÚlogÚlog2fr) rrCrDr r r r<r=Ú sink_val_curÚLOG2_EÚacc_O_mn_row_is_zero_or_nanÚ row_sum_curÚLN2rrrÚfinalizeqs.  &þ ÿýzSoftmax.finalizeÚacc_Or<cCslt |¡}t |¡tj|dgdksJ‚tjt |¡ddD]}||df ||df ¡||¡qdS)zÏScale each row of acc_O by the given scale tensor. :param acc_O: input tensor :type acc_O: cute.Tensor :param row_scale: row_scale tensor :type row_scale: cute.Tensor r©ÚmodeTr*N)r r0rr4r2r3r:r5)rrUr<Úacc_O_mnr=rrrÚ rescale_O”s (ÿzSoftmax.rescale_O)r N)rNr)FT)r/N)Ú__name__Ú __module__Ú __qualname__rÚ__annotations__r2Ú ConstexprÚintrrFrrÚ staticmethodrrÚ TensorSSAÚfloatr"r&ÚjitÚboolrBrTrYrrrrrsv   üÿþý ü ÿÿ ÿ þÿÿ ÿ þüþýü û=ÿÿ ÿ þ" rc@seZdZUdZejeed<e  d!de dejede dBfdd„ƒZ e j de j d ed ee e ffd d „ƒZ d"de j de d ed dfdd„Ze j de jde fdd„ƒZe j   d#de jde jdejedejedejedejef dd„ƒZe j de jde de jfdd „ƒZdS)$Ú SoftmaxSm100rÚrescale_thresholdNr rc Cs6d}d}t |t¡}t |t¡}t|||||||dS)Nréd)rf)rrrre)r rfrr rr r rrrr¦s  ùzSoftmaxSm100.createrr(rcCs°t |¡r| |¡}|tjj kr|nd}d}n7|jd}|j||d}|tjj kr-|nd}|||j}t |¡}t |j dk¡rO||j krO|}|}d}||jd<||fS)Nrr©rr/) r2r6r"rrr r r r9rf)rrr(Ú row_max_newÚ row_max_safeÚ acc_scaleÚ row_max_oldÚ acc_scale_rrrÚupdate_row_maxºs      zSoftmaxSm100.update_row_maxFr$r<cCs6t | ¡r |jd|nd}|j||d|jd<dS)Nrrh)r2r6r r&)rr$r<r(rrrrÚupdate_row_sumÎszSoftmaxSm100.update_row_sumr cCsˆt |j¡ddksJdƒ‚||j}tjdt |j¡dddD]!}t ||||df|j|jf| | f¡\||<||d<q dS)Nérú.acc_S_row must have an even number of elementsTr*r)rr4Úshaper r2r3r Úfma_packed_f32x2)rrr Úrow_max_scaledÚirrrÚscale_subtract_rowmax×s   ýÿz"SoftmaxSm100.scale_subtract_rowmaxér,rÚacc_S_row_convertedÚe2eÚe2e_freqÚe2e_resÚ e2e_frg_limitc Cs¼t |j¡ddksJdƒ‚d}|ddksJ‚t |¡|}t |¡|dks*J‚t |t |¡¡} t |t |¡¡} t |¡D]š} t dtj| dgdd¡D]w} t | ¡rytj  | | | f¡| | | f<tj  | | d| f¡| | d| f<qPt | |||kpˆ| ||k¡r«tj  | | | f¡| | | f<tj  | | d| f¡| | d| f<qPt   | | | f| | d| f¡\| | | f<| | d| f<qP| d| f  | d| f  ¡ |j¡¡qAdS)Nrprrqé rVr)rr4rrÚlogical_divideÚ make_layoutr2Úrange_constexprr6rÚexp2r Úex2_emulation_2r:r5ÚtoÚ element_type) rrrxryrzr{r|Úfrg_tileÚfrg_cntÚ acc_S_row_frgÚacc_S_row_converted_frgÚjÚkrrrÚapply_exp2_convertæs4  ÿ &ÿ&ÿ ÿîzSoftmaxSm100.apply_exp2_convertc Cstt |j¡ddksJdƒ‚| |j}t dt |j¡d¡D]}t ||||df|j|jf||f¡\||<||d<qd}|ddksIJ‚t |¡|}t |¡|dks[J‚t |t  |¡¡}t |t  |¡¡} t |¡D]E} t dtj|dgdd¡D]"} tj   || | f¡|| | f<tj   || d| f¡|| d| f<q| d| f  |d| f  ¡ |j¡¡qrdS)Nrprrqrr}rV)rr4rrr r2r€r rsr~rrrr:r5rƒr„) rrr rxÚminus_row_max_scaledrur…r†r‡rˆr‰rŠrrrÚscale_apply_exp2_converts0  ý ÿ & ÿóz%SoftmaxSm100.scale_apply_exp2_convert)rN)F)Frwr,r)rZr[r\rfr2r^rbr]r`rrrrcrar_rrnrorFrvrdr‹rrrrrre¢sp ýÿþ ý"ÿÿÿÿ þ þ ýùþýüûú ù(þýüreÚqhead_per_kvheadrcCst |dk¡r |S||S)z,Convert q_idx to packed format for Pack-GQA.r)r2r6)Úq_idxrŽrrrÚfloor_if_packedBsrFÚ score_modÚvec_sizeÚ qk_acc_dtypeÚ seqlen_infoÚconstant_q_idxÚtranspose_indicesc( CsÐt | ¡rt d¡}t d¡}n t d¡}t d¡}t t |j¡¡}t ||¡}t |tj¡}t |tj¡  |f¡}t |tj¡}t | dkoK| du¡rTt |tj¡}tj d||ddD]}tj |ddD]}||||||<t | dko|| du¡r˜||||}|| }||| }|| |||<t |duo¡| du¡rÛt | du¡rÅ| \}}t ||||| ƒ}t ||ƒ\}}|||<n| \}}t |||||ƒ\}} | ||<qg| durìt ||||| ƒ||<||||||<qg|  ¡}!|  ¡}"t | du¡r |  ¡}#n | }$t |$tj¡  |f¡}#t | dko#| du¡r+|  ¡}%n t |tj¡  |f¡}%g}&t |du¡rB|}&||!||%|#|"| |&d}'| |'¡tj |ddD] }|||||<qYq]dS)aBShared implementation for applying score modification. Args: score_tensor: The scores to modify (acc_S for flash_fwd, tSrS_t2r for sm100) index_tensor: Index positions (tScS for flash_fwd, tScS_t2r for sm100) score_mod: The score modification function to apply batch_idx: Batch index head_idx: Head index softmax_scale: Scale to apply vec_size: Vector size for processing elements qk_acc_dtype: Data type for accumulator aux_tensors: Optional aux_tensors for FlexAttention fastdiv_mods: Tuple of (seqlen_q_divmod, seqlen_k_divmod) for wrapping seqlen_info: Sequence length info constant_q_idx: If provided, use this constant for all q_idx values If None, compute q_idx per-element qhead_per_kvhead_packgqa: Pack-GQA replication factor. Divide q_idx by this when greater than 1 so score mods see logical heads. transpose_indices: If True, swap q_idx/kv_idx in index_tensor (for bwd kernel where S is transposed) rrNTr*©rÚkv_idxr”Ú aux_tensors)r2r6rr4rrrÚInt32r Ú scalar_to_ssaÚ broadcast_tor3rÚdivmodr5r:)(Ú score_tensorÚ index_tensorr‘Ú batch_idxÚhead_idxrr’r“r™Ú fastdiv_modsr”r•rŽr–Ú q_idx_posÚ kv_idx_posÚn_valsÚ score_vecÚ kv_idx_vecÚ batch_idx_ssaÚ q_idx_vecÚ head_idx_vecrur‰Ú q_idx_packedÚ q_idx_logicalÚ head_offsetÚseqlen_q_divmodÚseqlen_k_divmodÚ q_idx_flooredÚ_Ú q_idx_wrappedÚkv_idx_wrappedÚ score_ssaÚ kv_idx_ssaÚ q_idx_ssaÚ q_idx_constÚ head_idx_ssaÚaux_argsÚpost_mod_scoresrrrÚapply_score_mod_innerMsx (      ÿ    ù ÿÀr»Ú score_mod_bwdc* Csòt |¡rt d¡}t d¡}n t d¡}t d¡}t t |j¡¡}t ||¡}t ||¡}t |tj¡}t |tj¡  |f¡}t |tj¡}t | dkoQ| du¡rZt |tj¡}tj d||ddD]}tj |ddD]—}|||||<||||||<t | dkoŠ| du¡r¦||||}|| }||| }|| |||<t | duo¯| du¡rét | du¡rÓ| \}}t ||||| ƒ}t ||ƒ\} }!|!||<n| \} }t |||||ƒ\} }"|"||<qm| durút ||||| ƒ||<||||||<qm|  ¡}#|  ¡}$|  ¡}%t | du¡r|  ¡}&n t | tj¡  |f¡}&t | dko3| du¡r;|  ¡}'n t |tj¡  |f¡}'g}(t | du¡rR| }(||#|$||'|&|%| |(d})| |)¡tj |ddD] }|||||<qjqcdS)aÅApply backward score modification (joint graph). Args: grad_tensor: in/out: dlogits rewritten in-place with d(scaled_scores) score_tensor: pre-mod scores (unscaled QK tile), scaled by softmax_scale internally index_tensor: Index positions (same as forward) score_mod_bwd: The backward score modification function (joint graph) batch_idx: Batch index head_idx: Head index softmax_scale: Scale to apply to score_tensor vec_size: Vector size for processing elements qk_acc_dtype: Data type for accumulator aux_tensors: Optional aux_tensors for FlexAttention fastdiv_mods: Tuple of (seqlen_q_divmod, seqlen_k_divmod) for wrapping seqlen_info: Sequence length info constant_q_idx: If provided, use this constant for all q_idx values qhead_per_kvhead: Pack-GQA replication factor transpose_indices: If True, swap q_idx/kv_idx in index_tensor rrNTr*r—)r2r6rr4rrÚ make_fragmentršr r›rœr3rrr5r:)*Ú grad_tensorržrŸr¼r r¡rr’r“r™r¢r”r•rŽr–r£r¤r¥Úgrad_vecr¦r§r¨r©rªrur‰r«r¬r­r®r¯r°r±r²r³Úgrad_ssar´rµr¶r¸r¹Ú grad_out_ssarrrÚapply_score_mod_bwd_innerÏs~ (       ÿ    ø ÿÃrÂ)rF)rIrGÚtypingrÚ dataclassesrr2Ú cutlass.cuterrÚ,sglang.jit_kernel.flash_attention.cute.utilsÚ jit_kernelÚflash_attentionr Úcute_dsl_utilsrr”rrrercr^r_rFrrdr»rÂrrrrÚsr       þ ý òýùø õ ô ó òñüø ÷ óòñ