o ½Ù¾iIaã@sFddlZddlZddlmZddlmZddlZddlmZddlm Z ddl mm Z ddl m Z ddlmZeGdd„de ƒƒZeGd d „d eƒƒZejd ejed ejfd d„ƒ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)rr r rr r ©rúR/home/ubuntu/.local/lib/python3.10/site-packages/flash_attn_origin/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  |¡rW| tj krUdn| } t  |¡rv| |} 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 rr ÚcutlassÚrangeÚsizeÚloadr Ú const_exprÚ warp_reduceÚfmaxrÚexp2fr$Ústore)rr&r'r(Úacc_S_mnÚ row_scaler r rr ÚrrÚ row_max_curÚrow_max_cur_scaledr#Ú acc_S_row_sumÚ row_max_prevrrrÚ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@)r1r5Ú isinstancerÚTensorr3r r rr9rr6r4ÚoperatorÚaddr0rr2ÚmathÚlog2Úer8r Ú rcp_approxÚlogÚlog2fr) rrBrCr r rr;r<Ú sink_val_curÚLOG2_EÚacc_O_mn_row_is_zero_or_nanÚ row_sum_curÚLN2rrrÚfinalizeos.  &þ ÿý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)rr/rr3r1r2r9r4)rrTr;Úacc_O_mnr<rrrÚ rescale_O’s (ÿzSoftmax.rescale_O)r N)rNr)FT)r.N)Ú__name__Ú __module__Ú __qualname__rÚ__annotations__r1Ú ConstexprÚintrrEr rÚ staticmethodrrÚ TensorSSAÚfloatr!r%ÚjitÚboolrArSrXrrrrrsv   üÿþý ü ÿÿ ÿ þÿÿ ÿ þüþýü û;ÿÿ ÿ þ" 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_thresholdNrrc Cs6d}d}t |t¡}t |t¡}t|||||||dS)Nééd)re)rrrrd)rrerr r r 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.) r1r5r!rrr rrr8re)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)r1r5r 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)rf)rr3Úshaperr1r2rÚfma_packed_f32x2)rrr Úrow_max_scaledÚirrrÚscale_subtract_rowmaxÕs   ýÿz"SoftmaxSm100.scale_subtract_rowmaxér+rfÚ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é rUrf)rr3rrÚlogical_divideÚ make_layoutr1Úrange_constexprr5r Úexp2rÚex2_emulation_2r9r4Ú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)Nrprrqrfr}rU)rr3rrrr1r€rrsr~rr rr9r4rƒr„) rrr rxÚminus_row_max_scaledrur…r†r‡rˆr‰rŠrrrÚscale_apply_exp2_convert s0  ý ÿ & ÿóz%SoftmaxSm100.scale_apply_exp2_convert)rN)F)Frwr+rf)rYrZr[rer1r]rar\r_rrrrbr`r^rrnrorErvrcr‹rrrrrrd sp ýÿþ ý"ÿÿÿÿ þ þ ýùþýüûú ù(þýürdÚqhead_per_kvheadrcCst |dk¡r |S||S)z,Convert q_idx to packed format for Pack-GQA.rf)r1r5)Úq_idxrŽrrrÚfloor_if_packed@srrfFÚ 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) rfrNTr)©rÚkv_idxr”Ú aux_tensors)r1r5rr3rrrÚInt32rÚ scalar_to_ssaÚ broadcast_tor2rÚdivmodr4r9)(Ú 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_innerKsx (      ÿ    ù ÿÀ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 rfrNTr)r—)r1r5rr3rrÚ make_fragmentršrr›rœr2rrr4r9)*Ú 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Â)rfF)rHrFÚtypingrÚ dataclassesrr1Ú cutlass.cuterrÚflash_attn_origin.cute.utilsrÚ%flash_attn_origin.cute.cute_dsl_utilsrÚ"flash_attn_origin.cute.seqlen_inforrrdrbr]r^rErrcr»rÂrrrrÚsr       þ ý òýùø õ ô ó òñüø ÷ óòñ