ó
    KëgjV  ã                   óÆ  • S SK r S SK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JrJr  / SQr\(       a  / SQO/ SQr\R"                  " S#\ V s/ s H  n \R%                  0 U S9PM     sn S	/S
.\D6\R&                  S\R(                  4S j5       5       r\R"                  " S#\ V s/ s H  n \R%                  0 U S9PM     sn S	/S
.\D6\R&                  S\R(                  4S j5       5       r\
" S#S V Vs/ s H  n \  H  n\R%                  SU0U S9PM     M!     snn S	S/S
.\D6\R&                  " S/S9S	\R(                  S\R(                  S\R(                  S\R(                  4S j5       5       r\
" S#S V Vs/ s H  n \  H  n\R%                  SU0U S9PM     M!     snn S	S/S
.\D6\R&                  " S/S9S	\R(                  S\R(                  S\R(                  S\R(                  4S j5       5       r\" S5        S$S\ R2                  S\S\ R6                  S-  4S jj5       r\" S5       S%S\ R2                  S\ R2                  S\ R2                  S\4S jj5       r " S S\ R<                  R>                  5      r   S$S\ R2                  S\S\ R6                  S-  S\ R2                  4S  jjr!\!r" " S! S"\RF                  5      r$gs  sn f s  sn f s  snn f s  snn f )&é    N)Údispatch)Úfla_cache_autotune)ÚIS_AMDÚautotune_cache_kwargsÚinput_guard)é   é   é    é@   é€   )é   é   é   r   r	   )r   r   r   r   r	   r
   )Ú	num_warpsÚD)ÚconfigsÚkeyÚBDc                 ó¸  • [         R                  " S5      nXU-  -  n XU-  -  n[         R                  " SU5      nXt:  n[         R                  " X-   USS9R	                  [         R
                  5      n	S[         R                  " [         R                  " X™-  5      U-   5      -  n
Xš-  n[         R                  " X-   X¸S9  [         R                  " X&-   U
5        g ©Nr   ç        ©ÚmaskÚotherr   ©r   )	ÚtlÚ
program_idÚarangeÚloadÚtoÚfloat32ÚsqrtÚsumÚstore)ÚxÚyÚrstdÚepsr   r   Úi_tÚcolsr   Úb_xÚb_rstdÚb_ys               Ú C:\Users\johnn\AppData\Local\Packages\PythonSoftwareFoundation.Python.3.13_qbz5n2kfra8p0\LocalCache\local-packages\Python313\site-packages\fla/modules/l2norm.pyÚl2norm_fwd_kernel1r/      s¬   € ô �-Š-˜Ó
€CØˆq‰�L€AØˆq‰�L€Aä�9Š9�Q˜Ó€DØ‰8€Dä
�'Š'�!‘( ¨SÑ
1×
4Ñ
4´R·Z±ZÓ
@€CØ”—’œŸš ¡	Ó*¨SÑ0Ó1Ñ1€FØ
‰,€CÜ‡H‚HˆQ‰X�sÒ&Ü‡H‚HˆT‰Z˜Õ ó    c                 óB  • [         R                  " S5      nXU-  -  n X7U-  -  nX'U-  -  n[         R                  " SU5      nX…:  n	[         R                  " X-   U	SS9R	                  [         R
                  5      n
[         R                  " X-   5      R	                  [         R
                  5      n[         R                  " X(-   U	SS9R	                  [         R
                  5      nXË-  [         R                  " XÊ-  5      U
-  U-  -
  n[         R                  " X8-   XÙS9  g )Nr   r   r   r   )r   r   r   r   r    r!   r#   r$   )r&   r'   ÚdyÚdxr(   r   r   r)   r*   r   r-   r,   Úb_dyÚb_dxs                 r.   Úl2norm_bwd_kernel1r6   1   sà   € ô �-Š-˜Ó
€CØˆq‰�L€AØ�‰'�M€BØ�‰'�M€Bä�9Š9�Q˜Ó€DØ‰8€DÜ
�'Š'�!‘( ¨SÑ
1×
4Ñ
4´R·Z±ZÓ
@€CÜ�WŠW�T‘ZÓ ×#Ñ#¤B§J¡JÓ/€FÜ�7Š7�2‘9 4¨sÑ3×6Ñ6´r·z±zÓB€DØ‰=œ2Ÿ6š6 $¡*Ó-°Ñ3°fÑ<Ñ<€DÜ‡H‚HˆR‰Y˜Ó(r0   ÚBTÚNBÚT)Údo_not_specializec	                 ó4  • [         R                  " S5      R                  [         R                  5      n	X˜-  [         R                  " SU5      -   n
[         R                  " SU5      nX¤:  nUS S 2S 4   US S S 24   U:  -  nX
S S 2S 4   U-  -   US S S 24   -   nXS S 2S 4   U-  -   US S S 24   -   nX*-   n[         R
                  " XíSS9R                  [         R                  5      nS[         R                  " [         R                  " UU-  S5      U-   5      -  nUUS S 2S 4   -  n[         R                  " UUR                  UR                  R                  5      US9  [         R                  " UUR                  UR                  R                  5      US9  g r   )r   r   r    Úint64r   r   r!   r"   r#   r$   ÚdtypeÚ
element_ty)r%   r&   r'   r(   r9   r   r   r8   r7   r)   Úo_tÚo_dÚm_tÚm_xÚp_xÚp_yÚp_rstdr+   r,   r-   s                       r.   Úl2norm_fwd_kernelrF   N   sW  € ô" �-Š-˜Ó
×
Ñ
œbŸh™hÓ
'€CØ
‰(”R—Y’Y˜q "Ó%Ñ
%€CÜ
�)Š)�A�rÓ
€CØ
‰'€CØ
Ša�ˆg‰,˜#˜d¢A˜g™,¨Ñ*Ñ
+€CØ
’!�T�'‰l˜QÑÑ
  Tª1 W¡Ñ
-€CØ
’!�T�'‰l˜QÑÑ
  Tª1 W¡Ñ
-€CØ‰Z€Fä
�'Š'�# sÑ
+×
.Ñ
.¬r¯z©zÓ
:€CØ”—’œŸš  c¡	¨1Ó-°Ñ3Ó4Ñ4€FØ
�’q˜$�w‘Ñ
€Cä‡H‚HˆS�#—&‘&˜Ÿ™×-Ñ-Ó.°SÒ9Ü‡H‚HˆV�V—Y‘Y˜vŸ|™|×6Ñ6Ó7¸cÓBr0   c
                 ó´  • [         R                  " S5      R                  [         R                  5      n
X©-  [         R                  " SU	5      -   n[         R                  " SU5      nXµ:  nUS S 2S 4   US S S 24   U:  -  nXS S 2S 4   U-  -   US S S 24   -   nX-   nX+S S 2S 4   U-  -   US S S 24   -   nX;S S 2S 4   U-  -   US S S 24   -   n[         R
                  " XþSS9R                  [         R                  5      n[         R
                  " UUSS9R                  [         R                  5      n[         R
                  " UUSS9R                  [         R                  5      nUUS S 2S 4   -  [         R                  " UU-  S5      S S 2S 4   U-  US S 2S 4   -  -
  n[         R                  " UUR                  UR                  R                  5      US9  g r   )r   r   r    r<   r   r   r!   r#   r$   r=   r>   )r&   r'   r2   r3   r(   r9   r   r   r8   r7   r)   r?   r@   rA   rB   rD   rE   Úp_dyÚp_dxr-   r,   r4   r5   s                          r.   Úl2norm_bwd_kernelrJ   p   s£  € ô$ �-Š-˜Ó
×
Ñ
œbŸh™hÓ
'€CØ
‰(”R—Y’Y˜q "Ó%Ñ
%€CÜ
�)Š)�A�rÓ
€CØ
‰'€CØ
Ša�ˆg‰,˜#˜d¢A˜g™,¨Ñ*Ñ
+€CØ
’!�T�'‰l˜QÑÑ
  Tª1 W¡Ñ
-€CØ‰Z€FØ’A�t�G‘˜qÑ Ñ  3 tªQ w¡<Ñ/€DØ’A�t�G‘˜qÑ Ñ  3 tªQ w¡<Ñ/€Dä
�'Š'�# sÑ
+×
.Ñ
.¬r¯z©zÓ
:€CÜ�WŠW�V #¨SÑ1×4Ñ4´R·Z±ZÓ@€FÜ�7Š7�4˜c¨Ñ-×0Ñ0´·±Ó<€DØ�&š˜D˜‘/Ñ!¤B§F¢F¨4°#©:°qÓ$9º!¸T¸'Ñ$BÀSÑ$HÈ6ÒRSÐUYÐRYÉ?Ñ$ZÑZ€DÜ‡H‚HˆT�4—7‘7˜4Ÿ:™:×0Ñ0Ó1¸Ó<r0   Úmodulesr%   r(   Úoutput_dtypec                 óæ  ^• U R                   nU R                  SU R                   S   5      n Uc  [        R                  " U 5      nO[        R                  " XS9nUR	                  S5      S:X  d   eU R                   S   U R                   S   smnSU R                  5       -  n[        U[        R                  " U5      5      nXW:”  a  [        S5      e[        R                  " T4[        R                  U R                  S9nUS::  a2  [        R                  " TS5      n	U4S	 jn
[        U
   " U UUUTUUU	S
9  O[        T4   " U UUUUUS9  UR                  U5      UR                  US S 5      4$ )Néÿÿÿÿ)r=   r   r   é   z/This layer doesn't support feature dim >= 64KB.)r=   Údeviceé   c                 ó:   >• [         R                  " TU S   5      4$ ©Nr7   ©ÚtritonÚcdiv©Úmetar9   s    €r.   ÚgridÚl2norm_fwd.<locals>.grid¯   ó   ø€ Ü—K’K  4¨¡:Ó.Ð0Ð0r0   )r%   r&   r'   r(   r9   r   r   r8   )r%   r&   r'   r(   r   r   )ÚshapeÚviewÚtorchÚ
empty_likeÚstrideÚelement_sizeÚminrU   Únext_power_of_2ÚRuntimeErrorÚemptyr!   rP   rV   rF   r/   )r%   r(   rL   Ú
x_shape_ogr&   r   ÚMAX_FUSED_SIZEr   r'   r8   rY   r9   s              @r.   Ú
l2norm_fwdrh   “   s_  ø€ ð —‘€JØ	�‰ˆr�1—7‘7˜2‘;Ó€AàÑÜ×Ò˜QÓ‰ä×Ò˜QÑ3ˆØ�8‰8�B‹<˜1ÓÐÐØ�7‰7�1‰:�q—w‘w˜r‘{€D€A€qà˜aŸn™nÓ.Ñ.€NÜ	ˆ^œV×3Ò3°AÓ6Ó	7€BØƒvÜÐLÓMÐMä�;Š;˜�t¤5§=¡=¸¿¹ÑB€DØˆCƒxô �[Š[˜˜IÓ&ˆõ	1ô 	˜$ÒØØØØØØØØó		
ô 	˜A˜4Ò ØØØØØØò	
ð �6‰6�*Ó˜tŸy™y¨°C°R¨Ó9Ð9Ð9r0   r&   r'   r2   c                 ó|  ^• U R                   nU R                  SUR                   S   5      n UR                  SUR                   S   5      nUR                   U R                   :X  d   e[        R                  " U 5      nU R                   S   U R                   S   smnSU R	                  5       -  n[        U[        R                  " U5      5      nXh:”  a  [        S5      eUS::  a3  [        R                  " TS5      n	U4S jn
[        U
   " U UUUUTUUU	S9	  O[        T4   " U UUUUUUS9  UR                  U5      $ )	NrN   r   rO   z4This layer norm doesn't support feature dim >= 64KB.rQ   c                 ó:   >• [         R                  " TU S   5      4$ rS   rT   rW   s    €r.   rY   Úl2norm_bwd.<locals>.gridâ   r[   r0   )	r&   r'   r2   r3   r(   r9   r   r   r8   )r&   r'   r2   r3   r(   r   r   )r\   r]   r^   r_   ra   rb   rU   rc   rd   rV   rJ   r6   )r&   r'   r2   r(   Ú
y_shape_ogr3   r   rg   r   r8   rY   r9   s              @r.   Ú
l2norm_bwdrm   È   s7  ø€ ð —‘€JØ	�‰ˆr�2—8‘8˜B‘<Ó €AØ	�‰��R—X‘X˜b‘\Ó	"€BØ�8‰8�q—w‘wÓÐÐä	×	Ò	˜!Ó	€BØ�7‰7�1‰:�q—w‘w˜r‘{€D€A€qà˜aŸn™nÓ.Ñ.€NÜ	ˆ^œV×3Ò3°AÓ6Ó	7€BØƒvÜÐQÓRÐRàˆCƒxô �[Š[˜˜IÓ&ˆõ	1ô 	˜$ÒØØØØØØØØØó
	
ô 	˜A˜4Ò ØØØØØØØò	
ð �7‰7�:ÓÐr0   c                   óP   • \ rS rSr\\  SS j5       5       r\\S 5       5       rSrg)ÚL2NormFunctionéþ   Nc                 ór   • [        XU5      u  pEX l        UR                  U l        U R	                  XE5        U$ ©N)rh   r(   r=   Úx_dtypeÚsave_for_backward)Úctxr%   r(   rL   r&   r'   s         r.   ÚforwardÚL2NormFunction.forwardÿ   s5   € ô ˜Q \Ó2‰ˆØŒØ—g‘gˆŒØ×Ñ˜aÔ&Øˆr0   c                 óT   • U R                   u  p#[        X#XR                  5      nUS S 4$ rr   )Úsaved_tensorsrm   r(   )ru   r2   r&   r'   r3   s        r.   ÚbackwardÚL2NormFunction.backward  s-   € ð ×#Ñ#‰ˆÜ˜ §W¡WÓ-ˆØ�4˜ˆ~Ðr0   © ©ç�íµ ÷Æ°>N)	Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__Ústaticmethodr   rv   rz   Ú__static_attributes__r|   r0   r.   ro   ro   þ   s@   † ØØð Øó	
ó ó ð
ð Øñó ó ór0   ro   Úreturnc                 ó.   • [         R                  XU5      $ rr   )ro   Úapply)r%   r(   rL   s      r.   Úl2normrˆ     s   € ô
 ×Ñ ¨Ó5Ð5r0   c                   ó�   ^ • \ rS rSr  S
S\S\R                  S-  4U 4S jjjrS\R                  S\R                  4S jr	S	r
U =r$ )ÚL2Normi   Nr(   rL   c                 ó:   >• [         TU ]  5         Xl        X l        g rr   )ÚsuperÚ__init__r(   rL   )Úselfr(   rL   Ú	__class__s      €r.   r�   ÚL2Norm.__init__!  s   ø€ ô
 	‰ÑÔØŒØ(Õr0   r%   r…   c                 óB   • [        XR                  U R                  5      $ rr   )rˆ   r(   rL   )rŽ   r%   s     r.   rv   ÚL2Norm.forward*  s   € Ü�aŸ™ 4×#4Ñ#4Ó5Ð5r0   )r(   rL   r}   )r   r€   r�   r‚   Úfloatr^   r=   r�   ÚTensorrv   r„   Ú__classcell__)r�   s   @r.   rŠ   rŠ      sP   ø† ð Ø+/ñ)àð)ð —k‘k DÑ(÷)ð )ð6˜Ÿ™ð 6¨%¯,©,÷ 6ò 6r0   rŠ   r|   r}   )r~   )%r^   Útorch.nnÚnnrU   Útriton.languageÚlanguager   Úfla.modules.backendsr   Úfla.ops.utils.cacher   Ú	fla.utilsr   r   r   ÚBT_LISTÚNUM_WARPS_AUTOTUNEÚautotuneÚConfigÚjitÚ	constexprr/   r6   rF   rJ   r”   r“   r=   rh   rm   ÚautogradÚFunctionro   rˆ   Úl2_normÚModulerŠ   )r   r7   s   00r.   Ú<module>r§      s�  ðó Ý Û Ý å )Ý 2ß @Ñ @â
€Þ)/Ó%Ò5IÐ ð ‡‚ð ÙEWÓXÒEW¸	ˆV�]‰]˜2¨ˆ]Ó3ÑEWÑXØ	ˆñð ñð
 ‡�ð!ð 	�‰ó!ó óð!ð, ‡‚ð ÙEWÓXÒEW¸	ˆV�]‰]˜2¨ˆ]Ó3ÑEWÑXØ	ˆñð ñð
 ‡�ð)ð 	�‰ó)ó óð)ñ. ð ÙM]ÔpÒM]À	ÔhoÐbdˆV�]‰]˜D "˜:°ˆ]Ó;ÑhoÑ;ÑM]ÒpØ	ˆdˆñð ñð
 ‡‚˜s˜eÑ$ðCð 
‡|�|ðCð 	�‰ðCð 	�‰ðCð 	�‰óCó %óðCñ8 ð ÙM]ÔpÒM]À	ÔhoÐbdˆV�]‰]˜D "˜:°ˆ]Ó;ÑhoÑ;ÑM]ÒpØ	ˆdˆñð ñð
 ‡‚˜s˜eÑ$ð=ð 
‡|�|ð=ð 	�‰ð=ð 	�‰ð=ð 	�‰ó=ó %óð=ñ: 
ˆ)Óð Ø'+ñ1:Ø‡|�|ð1:à	ð1:ð —+‘+ Ñ$ô1:ó ð1:ñh 
ˆ)Óð
 ñ	2Ø‡|�|ð2à
�,‰,ð2ð 	�‰ð2ð 
ô	2ó ð2ôj�U—^‘^×,Ñ,ô ð2 Ø'+ñ6Ø‡|�|ð6à	ð6ð —+‘+ Ñ$ð6ð ‡\�\õ	6ð €ô6ˆR�Y‰Yõ 6ùòU Yùò8 Yùó: qùóD qs   ÁK
Â'K
Ã6&KÅ<&K