o iN2@s,ddlmZmZmZddlZddlmmZddl Z ddl m Z de ejvr.ejjej_e dddie jde jde jde jd e jfd d Ze dd die jde jde jfd dZGdddejjZ       ddejdejdeejdededededeejejffddZdS) )TupleOptionalUnionNall_gather_into_tensor HAS_SMOOTHINGcC |ddkSN smoothingargsr r Y/home/ubuntu/vllm_env/lib/python3.10/site-packages/flash_attn/ops/triton/cross_entropy.py r BLOCK_SIZESPLITPRECOMPUTED_LSEc Cstd}||| tj}d}|s~td }d}td| | D]M}|td| }tj|||| ktd dtj|}|rO|t t || k|d7}t |t |}t |||t t ||}|}q t||}t|||nt||}t||}||krd}d}nS|| 8}|dkr|| krt|||}|r|s|nd||| d||}n|s|nd|}n|r||s|nd|| }nd}|s|||}||7}nd}t||||st|||dSdS)Nrr infmaskother)tl program_idtoint64floatrangearangeloadfloat32sumwheremaximummaxexplogstore)loss_ptrlse_ptr z_loss_ptr logits_ptr labels_ptrr logit_scalelse_square_scale ignore_index total_classesclass_start_idxn_colslogits_row_striderrrrrow_idx sum_logitsm_il_i col_offsetcolslogitsm_i_newlse label_idxlossz_loss logits_labelr r rcross_entropy_fwd_kernels^  &     rBcCrrr r r r rrdrcCsJtd}td}||| tj}||| tj}||td|}t||}||kr=t|||}nd}tj|||| ktd dtj|}t||}t||}|d|||7}|| 8}|rd|}|| }t ||k||||}n t ||k|d|}tj ||||||| kddS) Nrrr rrg@?)r) rrrrrr rr!r&r#r() dlogits_ptr dloss_ptrr,r*r-r r.r/r0r1r2r3r4dlogits_row_stridedloss_row_striderrr5 col_block_idx col_offsetsr>dlossr;r=probssmooth_positivesmooth_negativer r rcross_entropy_bwd_kernelbs0  $rNc@s4eZdZe       d ddZedd ZdS) CrossEntropyLossNr rCFc Cs|jtjkr$|ddkr$t|ddddf}|ddks$J|dks*J|j\} } |j| fks7J| dur=dntj| } | | } | durMdntj | }|| }|duob|dkob|dk}| ddkrn| }d }t t | |}|d kr~d n |d krd n|dkrdnd}tj| tj|jd}|r|j| fksJ| }n tj| tj|jd}tj| tj|jd}tj|jj%t| f|||||||||| || | d|| dk||dWdn1swY| dkrM| dkr$tj| | |j|jd}tjj||| dtjj|tjjj| dd}tj|dd}|||7}|dkr@||}|||kd||7}nt|}|||kd||||||||_ ||_!||_"||_#| |_$||_%||_&||fS)Nr)rr.r rrCi@ i )dtypedevice)rrr num_warps)groupT)opr[async_op)dim)'rXtorchlongdata_ptrFpadshape distributedget_world_sizeget_rankstride contiguousmintritonnext_power_of_2emptyrrYcudaindexrBr all_reduceReduceOpSUM logsumexpwaitsquare masked_fill_ zeros_likesave_for_backwardmark_non_differentiabler r.r/r0r1r2inplace_backward)ctxr;labelsprecomputed_lser r.r/r0rz process_groupn_rowsr3 world_sizer1rankr2use_precomputed_lseMAX_BLOCK_SIZErrZlossesr=z_losses lse_allgather handle_lossesr r rforwards          zCrossEntropyLoss.forwardc s~|j\}}}|jr |nt|}|j\ttd}|dkr$dn|dkr*dnd}fdd} tj |j j .t | ||||||j |j |j|j|j|j|d |d |d ||d Wdn1slwY|dddddddddf S) NirSrTrUrVrQcst|dfS)Nr)rkcdiv)METAr3rr rr sz+CrossEntropyLoss.backward..r)rrZ) saved_tensorsrzr_ empty_likerdrjrkrlrnrYrorNr r.r/r0r1r2rh) r{ grad_losses grad_z_lossesr;r=r|dlogitsrrZgridr rrbackwards:  zCrossEntropyLoss.backwardNr rCr rPFN)__name__ __module__ __qualname__ staticmethodrrr r r rrOs irOr rCrPFr;r|r}label_smoothingr.r/rzreturnc Cst||||||||| S)a* Arguments: logits: (batch, vocab_size) labels: (batch,) label_smoothing: float logit_scale: float. Multiply logits by this scale before calculating the loss. lse_square_scale: float. If > 0, we add lse_square_scale * lse(logits) ^ 2 to the loss. This is also referred to as "z-loss". ignore_index: int. If labels == ignore_index, the loss is set to 0.0. inplace_backward: bool. If True, we do the backward pass in-place by modifying the logits. This saves memory. process_group: if not None, we're doing Tensor Parallel: each process is responsible for one part of the vocab. The loss will be aggregated across processes. Returns: losses: (batch,), float z_losses: (batch,), float )rOapply) r;r|r}rr.r/r0rzr~r r rcross_entropy_loss$srr)typingrrrr_torch.nn.functionalnn functionalrbrktriton.languagelanguagerdirre_all_gather_baser heuristicsjit constexprrBrNautogradFunctionrOTensorrboolrr r r rsr  I-