ó
    KëgjP6  ã                   ó‚  • S SK r S SKJs  Jr  S SKrS SKJr  S SK	J
r
  S SKJr  S SKJr  S SKJr  S SKJr  S SKJrJrJrJrJrJr  \" 5       (       a  SS	/OS
S/r/ SQr\(       a  / SQO/ SQrS\ R:                  4S\ R<                  S\ R<                  S\ R<                  S-  S\ R>                  S\ R<                  4
S jjr SS\ R:                  4S\ R<                  S\ R<                  S\ R<                  S-  S\!S\ R>                  S\ R<                  4S jjr"\RF                  " S S S S.5      \" SF\ V VVs/ s H)  n \  H  nS  H  n\RI                  SU 0XS9PM     M!     M+     snnn SS/S .\D6\RJ                  " S!/S"9S\RL                  S\RL                  S\RL                  S#\RL                  S$\RL                  S%\RL                  S&\RL                  4S' j5       5       5       r'\RF                  " S( S) S* S.5      \" SF\ VVs/ s H  nS  H  n\RI                  0 XS9PM     M     snnSS/S .\D6\RJ                  " S!/S"9S\RL                  S\RL                  S\RL                  S#\RL                  S$\RL                  S%\RL                  S&\RL                  4S+ j5       5       5       r(\
" S,5      SS\ R:                  4S\ R<                  S\ R<                  S\ R<                  S-  S\!S-  S\ R>                  S\ R<                  4S- jj5       r)\
" S,5         SGS\ R<                  S\ R<                  S\ R<                  S-  S.\ R<                  S-  S\!S-  S\*\ R<                  \ R<                  \ R<                  S-  4   4S/ jj5       r+ " S0 S1\ RX                  RZ                  5      r.\
" S,5      \ R^                  R`                  SS\ R:                  4S\ R<                  S\ R<                  S\ R<                  S-  S\!S-  S\ R>                  S\ R<                  \*\ R<                  \ R<                  4   -  4S2 jj5       5       r1\RF                  " S3 S4 S5 S6 S7.5      \" SF\ VVs/ s H  nS8  H  n\RI                  S9U0US:9PM     M!     snn/ S;QS .\D6\RJ                  " S!/S"9S\RL                  S<\RL                  S\RL                  S9\RL                  S=\RL                  S$\RL                  S>\RL                  S?\RL                  S&\RL                  4S@ j5       5       5       r2\\
" S,5      SSS\ RB                  SS4S\ R<                  S\ R<                  SA\3SB\!S\ R<                  S-  SC\ R<                  S-  S\ R>                  S-  SD\ Rh                  S-  S\!S-  S\ R<                  4SE jj5       5       r5gs  snnn f s  snnf s  snnf )Hé    N)Údispatch)Úfla_cache_autotune)Úprepare_chunk_indices)Úexp)Úsoftplus)ÚIS_AMDÚautocast_custom_bwdÚautocast_custom_fwdÚautotune_cache_kwargsÚcheck_shared_memÚinput_guardé    é@   é   )r   r   é€   )é   é   é   r   )r   r   r   r   ÚgÚA_logÚdt_biasÚoutput_dtypeÚreturnc                 ó<  • U R                   SS u  pEU R                  5       n Ub  XR                  US5      -   n UR                  US5      R                  5       R                  5       * [        R
                  " U R                  5       5      -  R                  U5      n U $ )aÅ  
Torch reference implementation for KDA gate computation.

Computes: g = -A_log.exp().unsqueeze(-1) * softplus(g + dt_bias.view(g.shape[-2:]))

Args:
    g (torch.Tensor):
        Input tensor of shape `[..., H, K]`.
    A_log (torch.Tensor):
        Parameter tensor with `H` elements.
    dt_bias (torch.Tensor | None):
        Optional bias tensor added to `g` before activation, shape `[H * K]`.

Returns:
    Output tensor of shape `[..., H, K]` .
éþÿÿÿNéÿÿÿÿé   )ÚshapeÚfloatÚviewr   ÚFr   Úto)r   r   r   r   ÚHÚ_s         ÚžC:\Users\johnn\AppData\Local\Packages\PythonSoftwareFoundation.Python.3.13_qbz5n2kfra8p0\LocalCache\local-packages\Python313\site-packages\fla/ops/kda/gate.pyÚnaive_kda_gater&      s‚   € ð, �7‰7�2�3ˆ<�D€AØ	�‰‹	€AØÑØ—‘˜Q Ó#Ñ#ˆà
�*‰*�Q˜Ó
×
 Ñ
 Ó
"×
&Ñ
&Ó
(Ð	(¬1¯:ª:°a·g±g³iÓ+@Ñ	@×DÑDÀ\ÓR€AØ€Hó    g      ÀÚlower_boundc                 ó  • U R                   SS  u  pVU R                  5       n Ub  XR                  US5      -   n U[        R                  " UR                  US5      R                  5       U -  5      -  n U R                  U5      $ )Nr   r   r   )r   r   r    r!   Úsigmoidr   r"   )r   r   r   r(   r   r#   r$   s          r%   Únaive_kda_lowerbound_gater+   :   st   € ð �7‰7�2�3ˆ<�D€AØ	�‰‹	€AØÑØ—‘˜Q Ó#Ñ#ˆØ”a—i’i §
¡
¨1¨aÓ 0× 4Ñ 4Ó 6¸Ñ :Ó;Ñ;€AØ�4‰4�ÓÐr'   c                 ó   • U S   S L$ ©Nr   © ©Úargss    r%   Ú<lambda>r1   J   ó   € ˜T )™_°DÑ8r'   c                 ó   • U S   S L$ ©NÚbetar.   r/   s    r%   r1   r1   K   ó   € ˜T &™\°Ñ5r'   c                 ó   • U S   S L$ ©Nr(   r.   r/   s    r%   r1   r1   L   ó   €  D¨Ñ$7¸tÑ$Cr'   )ÚHAS_BIASÚHAS_BETAÚUSE_LOWER_BOUND)r   é   ÚBT)Ú	num_warpsÚ
num_stagesr#   ÚD)ÚconfigsÚkeyÚT)Údo_not_specializeÚBDr:   r;   r<   c                 ó„  • [         R                  " S5      R                  [         R                  5      [         R                  " S5      nn[         R                  " UU-   5      R                  [         R
                  5      nXú-  [         R                  " SU
5      -   n[         R                  " SU5      nUU:  nUS S 2S 4   US S S 24   U	:  -  nU UU	-  -   US S 2S 4   X‰-  -  -   US S S 24   -   nUUU	-  -   US S 2S 4   X‰-  -  -   US S S 24   -   n[         R                  " UUSS9R                  [         R
                  5      nU(       a[  UU	-  [         R                  " SU5      -   nU[         R                  " UU-   UX‰-  :  SS9R                  [         R
                  5      -   nU(       d  [        U5      * [        U5      -  nO%U[         R                  " [        U5      U-  5      -  n[         R                  " UUR                  UR                  R                  5      US9  U(       a—  UU-   UU-  -   nUU-   UU-  -   n[         R                  " [         R                  " UUSS9R                  [         R
                  5      5      n[         R                  " UUR                  UR                  R                  5      US9  g g )Nr   r   ç        ©ÚmaskÚother©rJ   )ÚtlÚ
program_idr"   Úint64ÚloadÚfloat32Úaranger   r   r*   ÚstoreÚdtypeÚ
element_ty)r   r   r   r5   ÚygÚybr(   rD   r#   rA   r>   rF   r:   r;   r<   Úi_tÚi_hÚb_AÚo_tÚo_dÚm_tÚm_gÚp_gÚp_ygÚb_gÚo_bÚb_ygÚp_bÚp_ybÚb_ybs                                 r%   Úkda_gate_fwd_kernelrg   I   s6  € ôB �}Š}˜QÓ×"Ñ"¤2§8¡8Ó,¬b¯mªm¸AÓ.>ˆ€Cä
�'Š'�%˜#‘+Ó
×
!Ñ
!¤"§*¡*Ó
-€Cà
‰(”R—Y’Y˜q "Ó%Ñ
%€CÜ
�)Š)�A�rÓ
€CØ
�‰'€CØ
Ša�ˆg‰,˜#˜d¢A˜g™,¨Ñ*Ñ
+€CØ
ˆc�A‰g‰+˜šA˜t˜G™¨©Ñ.Ñ
.°°Tº1°W±Ñ
=€CØ��a‘‰<˜#ša ˜g™,¨!©%Ñ0Ñ0°3°tºQ°w±<Ñ?€Dä
�'Š'�#˜C sÑ
+×
.Ñ
.¬r¯z©zÓ
:€CÞØ�A‰gœŸ	š	 ! RÓ(Ñ(ˆØ”B—G’G˜G c™M°°a±e±À3ÑG×JÑJÌ2Ï:É:ÓVÑVˆÞÜ�C“ˆyœ8 C›=Ñ(‰àœRŸZšZ¬¨C«°3©Ó7Ñ7ˆÜ‡H‚HˆT�4—7‘7˜4Ÿ:™:×0Ñ0Ó1¸Ò<æØ�S‰j˜3 ™7Ñ"ˆØ�C‰x˜# ™'Ñ!ˆÜ�zŠzœ"Ÿ'š' #¨C°sÑ;×>Ñ>¼r¿z¹zÓJÓKˆÜ
�Š��t—w‘w˜tŸz™z×4Ñ4Ó5¸CÓ@ð	 r'   c                 ó   • U S   S L$ r-   r.   r/   s    r%   r1   r1   ‡   r2   r'   c                 ó   • U S   S L$ r4   r.   r/   s    r%   r1   r1   ˆ   r6   r'   c                 ó   • U S   S L$ r8   r.   r/   s    r%   r1   r1   ‰   r9   r'   c                 óø  • [         R                  " S5      R                  [         R                  5      [         R                  " S5      nn[         R                  " UU-   5      R                  [         R
                  5      nUU-  [         R                  " SU5      -   n[         R                  " SU5      nUU
:  nUS S 2S 4   US S S 24   U:  -  nU UU-  -   US S 2S 4   X¼-  -  -   US S S 24   -   nUUU-  -   US S 2S 4   X¼-  -  -   US S S 24   -   nUUU-  -   US S 2S 4   X¼-  -  -   US S S 24   -   n[         R                  " UUSS9R                  [         R
                  5      n[         R                  " UUSS9R                  [         R
                  5      nU(       a[  UU-  [         R                  " SU5      -   nU[         R                  " UU-   UX¼-  :  SS9R                  [         R
                  5      -   nU(       df  [        U5      * nU[        U5      -  nUU[         R                  " U5      -  -  n [         R                  " [         R                  " UU-  S5      S5      n!Oj[        U5      nUU-  n"[         R                  " U"5      n#U#SU#-
  -  n$UU	U$-  -  n%U%U-  n [         R                  " [         R                  " U U-  S5      S5      n![         R                  " UU R                  UR                  R                  5      US9  [         R                  " UUU-  -   U-   U!5        U(       aÊ  UU-   UU-  -   n&UU-   UU-  -   n'UU-   UU-  -   n([         R                  " U&USS9R                  [         R
                  5      n)[         R                  " U(USS9R                  [         R
                  5      U)-  SU)-
  -  n*[         R                  " U'U*R                  U'R                  R                  5      US9  g g )Nr   r   rH   rI   g      ð?rL   )rM   rN   r"   rO   rP   rQ   rR   r   r   r*   ÚsumrS   rT   rU   )+r   r   r   r5   ÚdygÚdybÚdgÚdAÚdbetar(   rD   r#   rA   r>   rF   r:   r;   r<   rX   rY   rZ   r[   r\   r]   r^   r_   Úp_dgÚp_dygra   Úb_dygrb   rc   Úb_dgÚb_dAÚb_innerÚb_sigÚb_dsigÚb_d_inner_termrd   Úp_dbÚp_dybÚb_bÚb_dbs+                                              r%   Úkda_gate_bwd_kernelr   †   se  € ôF �}Š}˜QÓ×"Ñ"¤2§8¡8Ó,¬b¯mªm¸AÓ.>ˆ€Cä
�'Š'�%˜#‘+Ó
×
!Ñ
!¤"§*¡*Ó
-€Cà
�‰(”R—Y’Y˜q "Ó%Ñ
%€CÜ
�)Š)�A�rÓ
€CØ
�‰'€CØ
Ša�ˆg‰,˜#˜d¢A˜g™,¨Ñ*Ñ
+€CØ
ˆc�A‰g‰+˜šA˜t˜G™¨©Ñ.Ñ
.°°Tº1°W±Ñ
=€CØ��a‘‰<˜#ša ˜g™,¨!©%Ñ0Ñ0°3°tºQ°w±<Ñ?€DØ�#˜‘'‰M˜C¢ 4 ™L¨A©EÑ2Ñ2°S¸ºq¸±\ÑA€Eô �'Š'�#˜C sÑ
+×
.Ñ
.¬r¯z©zÓ
:€CÜ�GŠG�E ¨3Ñ/×2Ñ2´2·:±:Ó>€EæØ�A‰gœŸ	š	 ! RÓ(Ñ(ˆØ”B—G’G˜G c™M°°a±e±À3ÑG×JÑJÌ2Ï:É:ÓVÑVˆö Ü�3‹xˆiˆØ”X˜c“]Ñ"ˆØ�eœbŸjšj¨›oÑ-Ñ.ˆÜ�vŠv”b—f’f˜U T™\¨1Ó-¨qÓ1‰ä�#‹hˆØ˜‘)ˆÜ—
’
˜7Ó#ˆØ˜# ™+Ñ&ˆà +°Ñ"6Ñ7ˆà Ñ#ˆÜ�vŠv”b—f’f˜T C™Z¨Ó+¨QÓ/ˆä‡H‚HˆT�4—7‘7˜4Ÿ:™:×0Ñ0Ó1¸Ò<Ü‡H‚HˆR�#˜‘'‰\˜CÑ Ô&æØ�S‰j˜3 ™7Ñ"ˆØ�s‰{˜S 1™WÑ$ˆØ�c‘	˜C !™GÑ#ˆä�gŠg�c ¨3Ñ/×2Ñ2´2·:±:Ó>ˆÜ�wŠw�u 3¨cÑ2×5Ñ5´b·j±jÓAÀCÑGÈ3ÐQTÉ9ÑUˆÜ
�Š��t—w‘w˜tŸz™z×4Ñ4Ó5¸CÓ@ð r'   Úkdac                 óæ   ^^	• U R                   SS  u  mnU R                  5       TU-  -  m	[        R                  " XS9nUU	4S jn[        U   " U UUS US T	TU[
        R                  " U5      US9  U$ )Nr   ©rT   c                 ó<   >• [         R                  " TU S   5      T4$ )Nr>   ©ÚtritonÚcdiv)Úmetar#   rD   s    €€r%   ÚgridÚkda_gate_fwd.<locals>.gridè   s   ø€ Ü—’˜A˜t D™zÓ*¨AÐ.Ð.r'   )r   r   r   r5   rV   rW   rD   r#   rA   rF   r(   )r   ÚnumelÚtorchÚ
empty_likerg   r…   Únext_power_of_2)
r   r   r   r(   r   ÚKrV   rˆ   r#   rD   s
           @@r%   Úkda_gate_fwdr�   Û   sƒ   ù€ ð �7‰7�2�3ˆ<�D€A€qØ	�‰‹	�a˜!‘eÑ€Aä	×	Ò	˜!Ñ	0€Bö/ô ˜ÒØ
ØØØØØØ
Ø
Ø
Ü×!Ò! !Ó$Øòð €Ir'   rm   c                 ó¢  • U R                   SS  u  pVU R                  5       XV-  -  nSn[        R                  " Xx5      n	[        R
                  " U [        R                  S9n
UR                  X•[        R                  S9n[        R                  " Xx5      U4n[        U   " U UUS US U
US UUUU[        R                  " U5      US9  U
R                  U 5      R                  U 5      n
UR                  S5      R                  U5      R                  U5      nUb2  U
R                  SXV-  5      R                  S5      R                  U5      OS nX«U4$ )Nr   r   r‚   )r   r   r   r5   rm   rn   ro   rp   rq   rD   r#   rA   r>   rF   r(   r   r   )r   rŠ   r…   r†   r‹   rŒ   rQ   Ú	new_emptyr   r�   Úview_asÚtype_asrl   r    r"   )r   r   r   rm   r(   r#   rŽ   rD   r>   ÚNTro   rp   rˆ   Údbiass                 r%   Úkda_gate_bwdr–   û   s4  € ð �7‰7�2�3ˆ<�D€AØ	�‰‹	�a‘eÑ€AØ	€BÜ	�Š�QÓ	€Bä	×	Ò	˜!¤5§=¡=Ñ	1€BØ	�‰˜¤e§m¡mˆÐ	4€Bä�KŠK˜Ó Ð"€DÜ˜ÒØ
ØØØØØØØØØ
Ø
Ø
ØÜ×!Ò! !Ó$Øòð$ 
�‰�A‹×	Ñ	˜qÓ	!€BØ	�‰�‹×	Ñ	˜5Ó	!×	)Ñ	)¨%Ó	0€BØ5<Ñ5HˆB�G‰G�B˜™Ó×"Ñ" 1Ó%×(Ñ(¨Ô1Èd€Eà�5ˆ=Ðr'   c                   ó"  • \ rS rSr\\\SS\R                  4S\R                  S\R                  S\R                  S-  S\
S-  S\R                  S\R                  4S	 jj5       5       5       r\\\S
\R                  4S j5       5       5       rSrg)ÚKDAGateFunctioni%  Nr   r   r   r(   r   r   c                 óP   • [        UUUUUS9nU R                  XU5        X@l        U$ )N©r   r   r   r(   r   )r�   Úsave_for_backwardr(   )Úctxr   r   r   r(   r   rV   s          r%   ÚforwardÚKDAGateFunction.forward&  s9   € ô ØØØØ#Ø%ñ
ˆð 	×Ñ˜a¨Ô0Ø%ŒØˆ	r'   rm   c                 ó`   • U R                   u  p#n[        UUUUU R                  S9u  pVnXVUS S 4$ )N)r   r   r   rm   r(   )Úsaved_tensorsr–   r(   )rœ   rm   r   r   r   ro   rp   r•   s           r%   ÚbackwardÚKDAGateFunction.backward<  sF   € ð  ×-Ñ-Ñˆ�'Ü$ØØØØØŸ™ñ
‰ˆ�ð �u˜d DÐ(Ð(r'   r.   )Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__Ústaticmethodr   r
   r‹   rQ   ÚTensorr   rT   r�   r	   r¡   Ú__static_attributes__r.   r'   r%   r˜   r˜   %  s¹   † ØØØð
 (,Ø$(Ø$)§M¡Mñà�<‰<ðð �|‰|ðð —‘ Ñ$ð	ð
 ˜T‘\ðð —k‘kðð 
�‰ôó ó ó ðð& ØØð	)˜5Ÿ<™<ó 	)ó ó ó ó	)r'   r˜   c                 ó0   • [         R                  XX#U5      $ )a½  
Fused KDA gate computation with autograd support.

Computes: g = -A_log.exp().unsqueeze(-1) * softplus(g + dt_bias.view(g.shape[-2:]))

Args:
    g (torch.Tensor):
        Input tensor of shape `[..., H, K]`.
    A_log (torch.Tensor):
        Parameter tensor with `H` elements.
    dt_bias (torch.Tensor | None):
        Optional bias tensor added to `g` before activation, shape `[H * K]`.

Returns:
    Output tensor of shape `[..., H, K]`.
)r˜   Úapplyrš   s        r%   Úfused_kda_gater¬   K  s   € ô2 × Ñ  ¨7ÀÓNÐNr'   c                 ó   • U S   S L$ r-   r.   r/   s    r%   r1   r1   h  r2   r'   c                 ó   • U S   S L$ )NÚscaler.   r/   s    r%   r1   r1   i  s   € ˜d 7™m°4Ñ7r'   c                 ó   • U S   S L$ )NÚ
cu_seqlensr.   r/   s    r%   r1   r1   j  s   € ˜d <Ñ0¸Ñ<r'   c                 ó   • U S   S L$ r8   r.   r/   s    r%   r1   r1   k  r9   r'   )r:   Ú	HAS_SCALEÚ	IS_VARLENr<   )r   r   r   ÚBS)r?   )r#   ÚSr>   r´   ÚREVERSEr¶   r·   r³   r´   c                 ó  • [         R                  " S5      [         R                  " S5      R                  [         R                  5      [         R                  " S5      nnnUU	-  UU	-  nnU(       aê  [         R                  " UUS-  -   5      R                  [         R
                  5      [         R                  " UUS-  -   S-   5      R                  [         R                  5      nn[         R                  " UU-   5      R                  [         R
                  5      [         R                  " UU-   S-   5      R                  [         R
                  5      nnUU-
  nOUU-  UU-  U-   nnUU-  [         R                  " SU5      -   nUU-  [         R                  " SU5      -   nUS S 2S 4   U:  US S S 24   U
:  -  nU UU	-  U-   U
-  -   US S 2S 4   Xš-  -  -   US S S 24   -   nUUU	-  U-   U
-  -   US S 2S 4   Xš-  -  -   US S S 24   -   n[         R                  " UUSS9R                  [         R                  5      nU(       aK  [         R                  " UUU
-  -   U-   UU
:  SS9R                  [         R                  5      n UU S S S 24   -   n[         R                  " UU-   5      R                  [         R                  5      n!U(       d  [        U!5      * [        U5      -  n"O%U[         R                  " [        U!5      U-  5      -  n"U(       a  [         R                  " U"SSS9n#O[         R                  " U"SS9n#U(       a  U#U-  n#[         R                  " UU#R                  UR                  R                  5      US	9  g )
Nr   r   r   rH   rI   T)ÚaxisÚreverse)r¹   rL   )rM   rN   r"   rO   rP   Úint32rR   rQ   r   r   r*   ÚcumsumrS   rT   rU   )$Úsr   r   Úor¯   r±   Úchunk_indicesr(   rD   r#   r¶   r>   rµ   r·   r:   r³   r´   r<   Úi_srX   Úi_bhÚi_brY   Úi_nÚbosÚeosr[   Úo_sÚm_sÚp_sÚp_oÚb_sÚb_biasrZ   Úb_gateÚb_os$                                       r%   Ú#kda_gate_chunk_cumsum_vector_kernelrÎ   g  sí  € ôH —]’] 1Ó%¤r§}¢}°QÓ'7×':Ñ':¼2¿8¹8Ó'DÄbÇmÂmÐTUÓFVˆdˆ€CØ�q‰y˜$ ™(ˆ€CÞÜ—7’7˜=¨3°©7Ñ2Ó3×6Ñ6´r·x±xÓ@Ä"Ç'Â'È-ÐZ]Ð`aÑZaÑJaÐdeÑJeÓBf×BiÑBiÔjl×jrÑjrÓBsˆSˆÜ—7’7˜:¨Ñ+Ó,×/Ñ/´·±Ó9¼2¿7º7À:ÐPSÑCSÐVWÑCWÓ;X×;[Ñ;[Ô\^×\dÑ\dÓ;eˆSˆØ�#‰I‰à˜‘7˜C !™G a™KˆSˆà
�‰(”R—Y’Y˜q "Ó%Ñ
%€CØ
�‰(”R—Y’Y˜q "Ó%Ñ
%€CØŠq�$ˆw‰<˜!Ñ  Dª! G¡¨qÑ 0Ñ
1€CØ
ˆs�Q‰w˜‰} Ñ!Ñ
! Cª¨4¨¡L°A±CÑ$8Ñ
8¸3¸tÂQ¸w¹<Ñ
G€CØ
ˆs�Q‰w˜‰} Ñ!Ñ
! Cª¨4¨¡L°A±CÑ$8Ñ
8¸3¸tÂQ¸w¹<Ñ
G€Cä
�'Š'�#˜C sÑ
+×
.Ñ
.¬r¯z©zÓ
:€Cö Ü—’˜ 3¨¡7Ñ*¨SÑ0°s¸Q±wÀcÑJ×MÑMÌbÏjÉjÓYˆØ�F˜4¢˜7‘OÑ#ˆä
�'Š'�%˜#‘+Ó
×
!Ñ
!¤"§*¡*Ó
-€CÞä�c“(�œX c›]Ñ*‰àœrŸzšz¬#¨c«(°S©.Ó9Ñ9ˆö Ü�iŠi˜ Q°Ñ5‰ä�iŠi˜ QÑ'ˆæØˆu‰ˆÜ‡H‚HˆS�#—&‘&˜Ÿ™×-Ñ-Ó.°SÓ9r'   Ú
chunk_sizer¯   r±   r¿   c	                 óà  ^^^• Ub  U R                   S   S:X  d   S5       e[        U R                   5      S:X  d   eU R                   u  mn
mnUnUc  Ub  [        X\5      nUc  [        R                  " X¬5      O
[        U5      mUSUR                  5       S-
  -  :X  d   S5       eU [        R                  " X=(       d    U R                  S9pUUU4S jn[        U   " UUUU UUUUU
TUUS	S
9  U $ )Nr   r   z;Only batch size 1 is supported when cu_seqlens are providedr   r   zchunk_size must be a power of 2r‚   c                 óJ   >• [         R                  " U S   U S   5      TTT-  4$ )Nr¶   rµ   r„   )r‡   ÚBr#   r”   s    €€€r%   rˆ   Ú#kda_gate_chunk_cumsum.<locals>.gridÌ  s&   ø€ œFŸKšK¨¨S©	°4¸±:Ó>ÀÀAÈÁEÐJÐJr'   F)r½   r   r   r¾   r¯   r±   r¿   r(   rD   r#   r¶   r>   r·   )
r   Úlenr   r…   r†   Ú
bit_lengthr‹   rŒ   rT   rÎ   )r   r   rÏ   r¯   r   r±   r   r¿   r(   ÚkwargsrD   r¶   r>   Úg_orgrˆ   rÒ   r#   r”   s                  @@@r%   Úkda_gate_chunk_cumsumrØ   ³  s   ú€ ð ÑØ�w‰w�q‰z˜Q‹Ð]Ð ]Ó]ˆÜˆq�w‰w‹<˜1ÓÐÐØ—‘�J€A€qˆ!ˆQØ	€BØÑ Ñ!7Ü-¨jÓ=ˆØ)Ñ1Œ�Š�QÔ	´s¸=Ó7I€BØ˜˜Z×2Ñ2Ó4°QÑ6Ñ7Ó7ÐZÐ9ZÓZÐ7à”%×"Ò" 1×,C¸A¿G¹GÑDˆ1ßJÜ'¨Ò-Ø
ØØØ
ØØØ#ØØ
Ø
Ø
ØØòð €Hr'   r.   )NNN)6r‹   Útorch.nn.functionalÚnnÚ
functionalr!   r…   Útriton.languageÚlanguagerM   Úfla.ops.backendsr   Úfla.ops.utils.cacher   Úfla.ops.utils.indexr   Úfla.ops.utils.opr   Úfla.ops.utils.softplusr   Ú	fla.utilsr   r	   r
   r   r   r   ÚBS_LISTÚBT_LIST_AUTOTUNEÚNUM_WARPS_AUTOTUNErQ   r¨   rT   r&   r   r+   Ú
heuristicsÚConfigÚjitÚ	constexprrg   r   r�   Útupler–   ÚautogradÚFunctionr˜   ÚcompilerÚdisabler¬   rÎ   ÚintÚ
LongTensorrØ   )r>   r?   r@   rµ   s   0000r%   Ú<module>rò      s   ðó ß Ð Û Ý å %Ý 2Ý 5Ý  Ý +ß |× |á&×(Ñ(ˆ2ˆr‰(¨r°2¨h€Ú Ð Þ&,“]².Ð ð $(Ø %§¡ñ	Ø‡|�|ðà�<‰<ðð �\‰\˜DÑ ðð —+‘+ð	ð
 ‡\�\õðD $(ØØ %§¡ñØ‡|�|ðà�<‰<ðð �\‰\˜DÑ ðð ð	ð
 —+‘+ðð ‡\�\õð ×ÒÙ8Ù5ÙCñó ñ
 ð 	ñ #õâ"ˆBÜ+ˆIÜ ˆJð 	�‰�t˜R�j¨IˆÓMñ !ñ 	Ná+ñ 	NÙ"óð 
ˆcˆ
ñ	ð ñ	ð ‡‚˜s˜eÑ$ð*Að 
‡|�|ð*Að 
‡|�|ð*Að 	�‰ð*Að 	�‰ð*Að �l‰lð*Að �l‰lð*Að —\‘\ó*Aó %ó	óð *AðZ ×ÒÙ8Ù5ÙCñó ñ
 ð ñ ,ôâ+ˆIÜ ˆJð 	�‰�b IˆÓEá ñ 	FÙ+òð
 
ˆcˆ
ñð ñð ‡‚˜s˜eÑ$ðCAð 
‡|�|ðCAð 
‡|�|ðCAð 	�‰ðCAð 	�‰ðCAð  �l‰lð!CAð" �l‰lð#CAð$ —\‘\ó%CAó %óóðCAñL 
ˆ%ƒð $(Ø $Ø %§¡ñØ‡|�|ðà�<‰<ðð �\‰\˜DÑ ðð ˜‘ð	ð
 —+‘+ðð ‡\�\ôó ðñ> 
ˆ%ƒð $(Ø#Ø $ñ&Ø‡|�|ð&à�<‰<ð&ð �\‰\˜DÑ ð&ð 
�‰˜Ñ	ð	&ð
 ˜‘ð&ð ˆ5�<‰<˜Ÿ™ u§|¡|°dÑ':Ð:Ñ;ô&ó ð&ôR#)�e—n‘n×-Ñ-ô #)ñL 
ˆ%ƒØ‡�×Ñð $(Ø $Ø %§¡ñOØ‡|�|ðOà�<‰<ðOð �\‰\˜DÑ ðOð ˜‘ð	Oð
 —+‘+ðOð ‡\�\�E˜%Ÿ,™,¨¯©Ð4Ñ5Ñ5ôOó ó ðOð4 ×ÒÙ8Ù7Ù<ÙCñ	ó ñ ð ñ ôâˆBÜ"ˆIð 	�‰�t˜R�j¨IˆÓ6á"ñ 	7Ùòò
 	1ñð ñð ‡‚˜s˜eÑ$ð9:ð 
‡|�|ð9:ð 
‡|�|ð9:ð 	�‰ð9:ð 	�‰ð9:ð �\‰\ð9:ð �l‰lð9:ð  �|‰|ð!9:ð" �|‰|ð#9:ð$ —\‘\ó%9:ó %óóð 9:ðx Ù	ˆ%ƒð
 Ø#'Ø&*Ø',§{¡{Ø-1Ø $ñ'Ø‡|�|ð'à�<‰<ð'ð ð'ð ð	'ð
 �\‰\˜DÑ ð'ð —‘˜tÑ#ð'ð —+‘+ Ñ$ð'ð ×#Ñ# dÑ*ð'ð ˜‘ð'ð ‡\�\ô'ó ó ñ'ùôMùózùóDs   Ä50V.È$V5Ñ&V;