Kgj+mSSKJr SSKrSSKJr SSKJs Jr SSKrSSK J r SSK J r SSKJrJrJr SSKJr SSKJr SSKJrJrJr SSKJrJr SS KJr S r \(dS OS r!\RD"S SS.5\RFS\ RHS\ RHS\ RHS\ RH4Sj55r%\"S5S:S\&S-S\&S-S\RNS-4Sjj5r(\RFS\ RHS\ RHS\ RHS\ RHS\ RHS\ RH4 S j5r)\RFS!\ RHS\ RH4S"j5r*\"S5S;S#\RVS$\RXS%\RVS&\RVS'\-S\&S\&S\&S(\-S\.S)\/S*\&S+\/4S,jj5r0\"S5S-\RVS.\RVS/\RVS0\RV4S1j5r1"S2S3\RdRf5r4S;S#\RVS$\RXS%\RVS&\RVS'\-S\&S\&S\&S(\-S\.S)\/S*\&S+\/S4\5\RV\RV44S5jjr6"S6S7\Rn5r8"S8S9\5r9g!\\4a SrGNzf=f)<)partialN) DeviceMesh) ReplicateSharddistribute_module) ParallelStyle)dispatch)explogtanh)IS_AMD input_guard)DTensori cUSSL$)Nscaleargss ڴC:\Users\johnn\AppData\Local\Packages\PythonSoftwareFoundation.Python.3.13_qbz5n2kfra8p0\LocalCache\local-packages\Python313\site-packages\fla/modules/fused_linear_cross_entropy.pyr(s d7m47cUSSL$)N softcappingrrs rrr)sD$7t$Cr) HAS_SCALEHAS_SOFTCAPPINGDBrrc[R"S5R[R5[R"S5R[R5pX-[R"SU5-n X:n [R "XU--U -U [ S5*S9n U(aX-n U(aU[X- 5-n [R"U S5n [[R"[X- 5S55U -n[R"X[R"XE5--U -U5 g)Nrinfmaskother)tl program_idtoint64arangeloadfloatr maxr sumr storecdiv)xzrrrrrri_ni_do_dm_db_xb_mb_zs rlogsumexp_fwd_kernelr:'s}}Q""288,bmmA.>.A.A"((.K 'BIIaO #C 'C ''!Ag+##eEl] CCkD!233 &&a.C bffS^Q' (3 .CHHQrwwq}$ $s *C0rmodulesrrdtypec URnURSUS5nURupV[[R"U5S5n[R "Xg5nUR XX[RS9n [XX4"UU UUUUS9 U RS5R"USS6n Ub%U[R:waU RU5n U $)Nir<)r1r2rrrr) shapeviewmintritonnext_power_of_2r0 new_emptytorchr,r: logsumexpr() r1rrr<r@NrrNDr2s r logsumexp_fwdrJDs GGE r59A 77DA F " "1 %y1A Q B A -A!!   BeCRj)A Uekk1 DDK Hrlabel_smoothing logit_scalelogit_softcapping reductionVBVc  [R"S5R[R5n [R"X5n [R "X,-5nX U -- nX:XaJ[ SX5H9nU[R"SU 5-n[R"UU-SUU :S9 M; g[R "X-5R[R5U-nUb[UU- 5nUU-nSUU-- n[R "X-5nUU- nSnXj- n[R"5 [ SU 5GHnUU -[R"SU 5-n[R "UU-UU :[S5S9R[R5U-nUb[UU- 5nUU-nOUnUS:a6U[R"[R"UU :U*U-S55- n[UU- 5U- U-nUb USWU-- -nU S:XaUU- n[R"UU-UUU :S9 [R"5 GM US:aUS U- -UUU---n[R "X-5nUbWnOSnU S:XaUU- nUUS - U- U-U-- nOUUS - U-U-- n[R"X<-U5 [R"X-U5 g) aP This kernel computes both cross entropy loss and the gradient of the input. We only consider hard label + mean reduction for now. Please refer to https://pytorch.org/docs/stable/generated/torch.nn.CrossEntropyLoss.html for the math. Args: logits: Pointer to logits tensor. lse: Pointer to logsumexp tensor. target: Pointer to target tensor. loss: Pointer to tensor to store the loss. V (int): The number of columns in the input tensor. total (int): The number of non-ignored classes. ignore_index (int): The index to ignore in the target. label_smoothing (float): The amount of smoothing when computing the loss, where 0.0 means no smoothing. reduction (str): The string for the reduction to apply BV (int): The block size for vocab. rr$N?z-infr#meanr!)r&r'r(r)r0r+ranger*r/float32r debug_barrierr,r.wherer )logitslsetargetlosstotal ignore_indexrKrLrMrNrOrPr3NVb_yio_vb_lb_t_yb_softcap_deriv_yb_lseb_lossr9epsivb_logitsb_tb_cappedb_p b_sc_factors rcross_entropy_kernelrp`sX --   bhh 'C B ''&, C AgF q!Abii2&&C HHVc\3S1W 5!  ''&,  " "2:: . 1./359P3PQ ''&, C$'  F% !#u,{:[HH !#{2[@@HHTZ HHV\3rrHc8[R"S5R[R5nXC-[R"SU5-n[R "U5n[R "X-XR:S9n[R "X-Xv-XR:S9 g)am This function multiplies each element of the tensor pointed by x with the value pointed by g. The multiplication is performed in-place on the tensor pointed by x. Parameters: x: Pointer to the input tensor. g: Pointer to the gradient output value. N (int): The number of columns in the input tensor. B (int): The block size for Triton operations. rrSN)r&r'r(r)r*r+r/)r1grHri_xo_xb_gr7s relementwise_mul_kernelrvsp. --   bhh 'C 'BIIaO #C ''!*C ''!' (CHHQWcicg.rr1r\weightbiasr_ num_chunks use_l2warpl2_penalty_factoraccumulate_grad_in_fp32c URn /URQURSP7upn[[[R "U55n[U[R "UU55n[R "[R "UU55n[R "UU5n[R"X S9nU (a[RO URnSnUb#U (a[RO URnUb[R"X-US9OSnUb[R"X=US9OSn[R"X[RS9nURU5R5R5n[!U5GHnUU-[US-U-U5nnUUUn["R$"UX#5nUbURU:waUR'US9nUUUn [)UXg[RS9n!UUUn"U (a[R*"USSS 9un#n$[,URS4"UU!U U"UUUUUU UU[.S 9 U (a[R"U5n%X- n&W#U&-n'U%R1SW$U'5 Ub2[R2"UU%R55R'US9UUS 9 Ub$[R6"UU%RSUS9US 9 [R8"U%U5n(OS n([R8"UU5U(-UUU&Ub2[R2"UUR55R'US9UUS 9 UcGM[R6"UURSUS9US 9 GM UR5nUbUR'U5nUbUR'U5nUUUU4$)Nrdevice)rr<r!r?)rrr<r>T)keepdim) rZr[r\r]r^r_rKrLrMrNrOrP num_warps)inputmat1mat2out)rr%rrR)rr@rBMAX_FUSED_SIZErCrDr0rF zeros_likerWr<zerosr,ner.itemrVFlinearr(rJr-rp STATIC_WARPSscatter_addmmtaddmm))r1r\rwrxr_rKrLrMryrNrzr{r|rrHHrOrPNCCdx grad_dtypebias_grad_dtypedwdbr]r^icstartendc_xc_logitsc_targetc_lsec_lossc_maxxc_ids g_logits_l2 l2_factor penalty_graddx_l2_contributions) r"fused_linear_cross_entropy_forwardrs XXF(qww' Q'GA! ^V33A6 7B ZQ* +Bv{{1b12A Q B   ! +B"9v||JO +B%-- GMFX  &z B^bBIMIY  $_ E_cB ;;qu{{ ;D IIl # ' ' ) . . 0EBi!VS"q&Aq1scl88C.  #))z"9&&z&*C%$hk`e`k`kleC !IIhDAMFE hnnQ/12%+#/"  **84K*-I!I-L  UL 9! $++*+=   +//!?/*SY[\!&+v!> !$ 625GG5   KKZZ\__:_6     IIBhll1Ol&LRT UY\ 88:D ~ UU6] ~ UU4[ R rdorrrc6[R"U[R"SURS95(aURupE[ [ [R"U55n[[R"XE-U54"UUXE-U[S9 Ub=URupu[[R"Xu-U54"UUXu-U[S9 Ub9URSn[[R"Xv54"UUUU[S9 XU4$)NrTr~)r1rrrHrrr) rFrtensorrr@rBrrCrDrvr0r)rrrrrHrrrOs r#fused_linear_cross_entropy_backwardrs xxELLRYY788xx  6 6q 9 : AE1 578c"   >88DA "FKKq$9#; <#&   > A "FKK$5#7 8&   2:rc\rSrSr\\SS\RS\RS\RS\RS\ S\ S \ S \ S \ S \ S \ S\ S\ 4Sjj55r \\S55rSrg)FusedLinearCrossEntropyFunctioniNr1r\rwrxr_rKrLrMryrNrzr{r|c[UUUUUUUUU U U U U 5 upnnURUR5UbUR5OSUbUR55 U$S5 U$)aJ Fusing the last linear layer with cross-entropy loss Reference: https://github.com/mgmalek/efficient_cross_entropy Handle the forward and backward pass of the final linear layer via cross-entropy loss by avoiding the materialization of the large logits tensor. Since Cross Entropy Loss is the last layer, we can compute the gradient at the forward pass. By doing so, we don't have to store the x and target for the backward pass. x (torch.Tensor): [batch_size * seq_len, hidden_size] target (torch.LongTensor): [batch_size * seq_len] where each value is in [0, vocab_size). weight (torch.Tensor): [vocab_size, hidden_size] where `vocab_size` is the number of classes. bias (Optional[torch.Tensor]): [vocab_size] where `vocab_size` is the number of classes. ignore_index: the index to ignore in the target. label_smoothing: the amount of smoothing when computing the loss, where 0.0 means no smoothing. logit_scale: float = 1.0, A scaling factor applied to the logits. Default: 1.0 logit_softcapping: float = None, If > 0, apply logit softcapping: logits = softcap * tanh(logits / softcap). Default: 0.0 num_chunks: int The number of chunks to split the input tensor into for processing. This can help optimize memory usage and computation speed. Default: 8 reduction: Specifies the reduction to apply to the output: 'mean' | 'sum'. 'mean': the weighted mean of the output is taken, 'sum': the output will be summed. Default: 'mean'. use_l2warp: bool = False, Whether to use L2 regularization on the logits to prevent overconfidence. Default: False l2_penalty_factor: float = 1e-4, The L2Warp penalty factor. Default: 1e-4 accumulate_grad_in_fp32: bool = True, Whether to accumulate weight and bias gradients in fp32 before casting them back to the parameter dtype. Default: True N)rsave_for_backwarddetach)ctxr1r\rwrxr_rKrLrMryrNrzr{r|r]rrrs rforward'FusedLinearCrossEntropyFunction.forwardsz>            # "b  IIK!-BIIK4+BIIK  26  rc ZURup#n[XX45up#nUSX4SSSSSSSSS4 $N) saved_tensorsr)rrrrrs rbackward(FusedLinearCrossEntropyFunction.backward sC&& 8H 4tT4tT4QUUUrr NrRrTNrUF-C6?T)__name__ __module__ __qualname____firstlineno__ staticmethodrrFTensor LongTensorintr,strboolrr__static_attributes__rrrrrs " !$ #' #'(,P <<P  P P ll P  PPP!PPPP!P"&PPdVVrrreturnc D[RUUUUUUUUUU U U U 5 $)a6 Args: x (torch.Tensor): [batch_size * seq_len, hidden_size] target (torch.LongTensor): [batch_size * seq_len] where each value is in [0, vocab_size). weight (torch.Tensor): [vocab_size, hidden_size] where `vocab_size` is the number of classes. bias (Optional[torch.Tensor]): [vocab_size] where `vocab_size` is the number of classes. ignore_index: int. If target == ignore_index, the loss is set to 0.0. label_smoothing: float logit_scale: float A scaling factor applied to the logits. Default: 1.0 logit_softcapping: float If > 0, apply logit softcapping: logits = softcap * tanh(logits / softcap). Default: 0.0 num_chunks: int The number of chunks to split the input tensor into for processing. This can help optimize memory usage and computation speed. Default: 8 reduction: Specifies the reduction to apply to the output: 'mean' | 'sum'. 'mean': the weighted mean of the output is taken, 'sum': the output will be summed. Default: 'mean'. use_l2warp: Whether to add the L2Warp logit regularization gradient. The penalty is normalized by the full number of input tokens, matching `fla.modules.l2warp.l2_warp`. Default: `False`. l2_penalty_factor: The L2Warp penalty factor. Default: `1e-4`. accumulate_grad_in_fp32: Whether to accumulate weight and bias gradients in fp32 before casting them back to the parameter dtype. Default: `True`. Returns: losses: [batch,], float )rapply) r1r\rwrxr_rKrLrMryrNrzr{r|s rfused_linear_cross_entropy_lossr(s?j + 0 0   rc^\rSrSrSS\S\S\S\S\S\S \S \S \4U4S jjjr\ RRSS \ RS\ RS\ RS\ RS-4Sjj5rSrU=r$)FusedLinearCrossEntropyLossinNr_rKrLrMryrNrzr{r|c >[T U]5 US;d SUS35eXlX lX0lX@lXPlX`lXplXl Xl g)a Args: ignore_index: int. If target == ignore_index, the loss is set to 0.0. label_smoothing: float logit_scale: float A scaling factor applied to the logits. Default: 1.0 logit_softcapping: float If > 0, apply logit softcapping: logits = softcap * tanh(logits / softcap). Default: 0.0 num_chunks: int The number of chunks to split the input tensor into for processing. This can help optimize memory usage and computation speed. Default: 8 reduction: Specifies the reduction to apply to the output: 'mean' | 'sum'. 'mean': the weighted mean of the output is taken, 'sum': the output will be summed. Default: 'mean'. use_l2warp: Whether to add the L2Warp logit regularization gradient. The penalty is normalized by the full number of input tokens, matching `fla.modules.l2warp.l2_warp`. Default: `False`. l2_penalty_factor: The L2Warp penalty factor. Default: `1e-4`. accumulate_grad_in_fp32: Whether to accumulate weight and bias gradients in fp32 before casting them back to the parameter dtype. Default: `True`. )rUr.z reduction: z is not supportedN) super__init__r_rKrLrMryrNrzr{r|) selfr_rKrLrMryrNrzr{r| __class__s rr$FusedLinearCrossEntropyLoss.__init__psaR O+W{9+EV-WW+(.&!2$"$!2'>$rr1r\rwrxc<[URSURS5URS5UUURURUR UR URURURURURS9 nU$)a} Args: x (torch.Tensor): [batch_size, seq_len, hidden_size] target (torch.LongTensor): [batch_size, seq_len] where each value is in [0, V). weight (torch.Tensor): [vocab_size, hidden_size] where `vocab_size` is the number of classes. bias (Optional[torch.Tensor]): [vocab_size] where `vocab_size` is the number of classes. Returns: loss r>) rwrxr_rKrLrMryrNrzr{r|) rrAr@r_rKrLrMryrNrzr{r|)rr1r\rwrxr]s rr#FusedLinearCrossEntropyLoss.forwards(/ FF2qwwr{ # KKO** 00(("44nn"44$($@$@  r) r|r_r{rKrLrMryrNrz) rrRrTNrrUFrTr)rrrrrr,rrrrFcompilerdisablerrrr __classcell__rs@rrrns!!$ #' #'(,5?5?5? 5? ! 5?  5?5?5?!5?"&5?5?n ^^ %) " <<"  " " llT! ""rrc^\rSrSrSSS.S\S\4U4Sjjjr\S5r\S 5r S \ RS \ S \ R4S jr SrU=r$)LinearLossParallelir!F) sequence_dimuse_local_outputrrcP>[TU]5 [U54UlX lgr)rrrsequence_shardingr)rrrrs rrLinearLossParallel.__init__s& "' "5!7 0rc~UupEpg[U[5(d[R"XCU5nURU:waUR USS9n[U[5(d [R"XS[ 5/5nURU:waUR USS9n[U[5(d [R"Xc[ 5/5nUR[ 5/:waUR [ 5/SS9nUb5[U[5(d [R"Xs[ 5/5nUb2UR[ 5/:waUR [ 5/SS9nUR 5UR 5UR 5UbUR 54$U4$)NT) placementsasync_op) isinstancer from_localr redistributerto_local)rmodinputs device_meshr1r\rwrxs r_prepare_input_fn$LinearLossParallel._prepare_input_fnsq"(6!W%%""13DEA <<, ,*;dKA&'**''ik]KF    1 1((4EPT(UF&'**''ik]KF    -((Y[MD(QF  JtW$=$=%%d)+GD  IK= @$$ $MDzz|V__.0AVZVf4==?pplppprc4U(aUR5$U$r)r)rroutputsrs r_prepare_output_fn%LinearLossParallel._prepare_output_fns%5w!B7Brmodulerrc [UUS[URUR5[URUR 5S9$)N) partition_fninput_fn output_fn)rrrrrr)rrrs r_applyLinearLossParallel._applysD  T33T5K5KLd55t7L7LM   r)rr)rrrrrrrrrrnnModulerrrrrs@rrrs|!& 1 1 1 1qq4CC RYY Z BII  rr)NNNr): functoolsrrFtorch.nnrtorch.nn.functional functionalrrCtriton.languagelanguager&torch.distributedrtorch.distributed.tensorrrr!torch.distributed.tensor.parallelrfla.modules.backendsr fla.ops.utils.opr r r fla.utilsr rr ImportErrorAttributeErrorrr heuristicsjit constexprr:r,r<rJrprvrrrrrrrautogradFunctionrtuplerrrrrrrrs  (HH;)++)0rR 7C1  || 1  || 1||1\\1  10 ) $ $  4<   ;;    6P \\P P ||P ||P  ||P  P  P f/ ||/ || / /> )  ##$(E ||E   E LLE ,, E  E  EEEEEEE"EEP )) ) )  )  ))X[Venn&=&=[VD ##$(C ||C   C LLC ,, C  C  CCCCCCC"C 5<< %&CL\"))\~2 2 c ^$GsK KK