ó
    IC˜jÈ‹  ã                   ó$  • S SK r S SKJrJrJrJr  S SKJr  \ R                  " \S5      r	\ R                  " \S5      r
\ R                  " \=(       d    \=(       d    \S5      r\ R                  " \S5      r\" 5       (       Ga  S SKrS SKJr  S SKrS	\\R$                     4S
 jr\R(                    SNS j5       r\R(                    SNS j5       r\R(                      SOS j5       r\R(                      SOS j5       r\R2                  " \R$                  " SS0SSS9\R$                  " SS0SSS9\R$                  " SS0SSS9\R$                  " SS0SSS9// S9\R(                    SNS j5       5       r\R2                  " \R$                  " SS0SSS9\R$                  " SS0SSS9\R$                  " SS0SSS9\R$                  " SS0SSS9// S9\R(                    SNS j5       5       r\R2                  " \R$                  " SS0SSS9// S9\R(                    SNS j5       5       r\R2                  " \R$                  " SSS.SSS9\R$                  " SSS.SSS9\R$                  " SSS.SSS9\R$                  " SSS.SSS9// S9\R(                      SPS j5       5       rS r\R2                  " \R$                  " SS0SSS9\R$                  " SS0SSS9// SS S!\0S"9\R(                    SNS# j5       5       r\R(                    SNS$ j5       r \R(                    SNS% j5       r!\R(                      SPS& j5       r"\R(                    SNS' j5       r#\R(                      SPS( j5       r$\R(                    SNS) j5       r%\R(                    SNS* j5       r&\R(                    SNS+ j5       r'\R(                    SNS, j5       r(\R(                  S- 5       r)\R(                      SQS. j5       r*\R(                      SRS/ j5       r+\R(                      SSS0 j5       r,\R(                      SSS1 j5       r-\R(                  S\R\                  4S2 j5       r/\R(                  S\R\                  4S3 j5       r0S S4K1J2r2J3r3  \R(                    SNS5 j5       r4\R(                    SNS6 j5       r5\R(                    SNS7 j5       r6\R(                    SNS8 j5       r7\R(                    SNS9 j5       r8\R2                  " \R$                  " SSSSS:.SSS9\R$                  " SSS;SS:.SSS9// S<QS9\R(                  S=\R\                  S>\R\                  S?\R\                  S@\R\                  4SA j5       5       r9\R(                  S\R\                  4SB j5       r:\R(                  S\R\                  4SC j5       r;\R(                  S\R\                  4SD j5       r<\R(                  S\R\                  4SE j5       r=\R(                  S\R\                  4SF j5       r>\R(                  S\R\                  4SG j5       r?\R(                  S\R\                  4SH j5       r@\R(                    SNSI j5       rA\R(                    SNSJ j5       rB STSK\\C   SL\D4SM jjrEgg)Ué    N)ÚHAS_CUDA_AND_TRITONÚHAS_GPUÚHAS_MTIA_AND_TRITONÚHAS_XPU_AND_TRITON)Ú
has_tritonzrequires cuda and tritonzrequires xpu and tritonzrequires gpu and tritonzrequires gpu)ÚlanguageÚreturnc                  ó,  • [         R                  R                  (       a>  [        R                  " SSSSSSSS.SSS9[        R                  " SSSSSSSS.SSS9/n U $ [        R                  " SSSSS	.SSS9[        R                  " SSS
SS	.SSS9/n U $ )Né   é   é   é   )ÚBLOCK_SIZE_MÚBLOCK_SIZE_NÚBLOCK_SIZE_KÚGROUP_SIZE_MÚmatrix_instr_nonkdimÚwaves_per_euÚkpack©Ú
num_stagesÚ	num_warpsé€   é@   ©r   r   r   r   é    é   )ÚtorchÚversionÚhipÚtritonÚConfig)Úconfigss    Ú²C:\Users\johnn\AppData\Local\Packages\PythonSoftwareFoundation.Python.3.13_qbz5n2kfra8p0\LocalCache\local-packages\Python313\site-packages\torch/testing/_internal/triton_utils.pyÚ_get_strange_configsr%       sØ   € Ü�=‰=××ä—’à(*Ø(*Ø(*Ø()Ø02Ø()Ø!"ñð  !Øñô —’à(+Ø(*Ø(*Ø()Ø02Ø()Ø!"ñð  !ØñðˆGðf ˆô+ —’à(*Ø(*Ø(*Ø()ñ	ð  !Øñ	ô —’à(+Ø(*Ø(*Ø()ñ	ð  !Øñ	ðˆGð, ˆó    Ú
BLOCK_SIZEc                 ó  • [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXs:  n[         R                  " X-   US9n	[         R                  " X-   US9n
Xš-   n[         R                  " X'-   X¸S9  g ©Nr   ©Úaxis©Úmask©ÚtlÚ
program_idÚarangeÚloadÚstore©Úin_ptr0Úin_ptr1Úout_ptrÚ
n_elementsr'   ÚpidÚblock_startÚoffsetsr-   ÚxÚyÚoutputs               r$   Ú
add_kernelr?   X   óu   € ô �mŠm Ñ#ˆØÑ&ˆØ¤§	¢	¨!¨ZÓ 8Ñ8ˆØÑ#ˆÜ�GŠG�GÑ%¨DÑ1ˆÜ�GŠG�GÑ%¨DÑ1ˆØ‘ˆÜ
�Š�Ñ" FÓ6r&   c                 ó  • [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXs:  n[         R                  " X-   US9n	[         R                  " X-   US9n
Xš-
  n[         R                  " X'-   X¸S9  g r)   r.   r4   s               r$   Ú
sub_kernelrB   i   r@   r&   c                 ó  • [         R                  " SS9nXe-  nU[         R                  " SU5      -   nXƒ:  n	[         R                  " X-   U	S9n
US:X  a  [         R                  " X-   U	S9nX«-   nOU
n[         R                  " X(-   XÉS9  g ©Nr   r*   r,   Útwor.   )r5   r6   r7   r8   ÚARGS_PASSEDr'   r9   r:   r;   r-   r<   r=   r>   s                r$   Úadd_kernel_with_optional_paramrG   z   s‚   € ô �mŠm Ñ#ˆØÑ&ˆØ¤§	¢	¨!¨ZÓ 8Ñ8ˆØÑ#ˆÜ�GŠG�GÑ%¨DÑ1ˆØ˜%ÓÜ—’˜Ñ)°Ñ5ˆAØ‘U‰FàˆFÜ
�Š�Ñ" FÓ6r&   c                 ó   • [         R                  " SS9nXv-  nU[         R                  " SU5      -   n	X“:  n
[         R                  " X	U-  -   U
S9nUS:X  a  [         R                  " X-   U
S9nX¼-   nOUn[         R                  " X)U-  -   XÚS9  g rD   r.   )r5   r6   r7   r8   ÚstriderF   r'   r9   r:   r;   r-   r<   r=   r>   s                 r$   Ú-add_kernel_with_none_param_and_equal_to_1_argrJ   �   sŒ   € ô �mŠm Ñ#ˆØÑ&ˆØ¤§	¢	¨!¨ZÓ 8Ñ8ˆØÑ#ˆÜ�GŠG�G¨Ñ.Ñ.°TÑ:ˆØ˜%ÓÜ—’˜Ñ)°Ñ5ˆAØ‘U‰FàˆFÜ
�Š� VÑ+Ñ+¨VÓ?r&   r   r   r   r   r   r   )r#   Úkeyc                 ó  • [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXs:  n[         R                  " X-   US9n	[         R                  " X-   US9n
Xš-   n[         R                  " X'-   X¸S9  g r)   r.   r4   s               r$   Úadd_kernel_autotunedrM   ¥   óu   € ô" �mŠm Ñ#ˆØÑ&ˆØ¤§	¢	¨!¨ZÓ 8Ñ8ˆØÑ#ˆÜ�GŠG�GÑ%¨DÑ1ˆÜ�GŠG�GÑ%¨DÑ1ˆØ‘ˆÜ
�Š�Ñ" FÓ6r&   c                 ó  • [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXs:  n[         R                  " X-   US9n	[         R                  " X-   US9n
Xš-
  n[         R                  " X'-   X¸S9  g r)   r.   r4   s               r$   Úsub_kernel_autotunedrP   ¿   rN   r&   r   r   c                 ó  • [         R                  " SS9nXS-  nU[         R                  " SU5      -   nXr:  n[         R                  " X-   US9n	[         R                  " X-   US9n
Xš-   n[         R                  " XG-   X¸S9  g r)   r.   )r5   r6   r8   r'   r7   r9   r:   r;   r-   r<   r=   r>   s               r$   Ú&add_kernel_autotuned_weird_param_orderrR   Ù   su   € ô  �mŠm Ñ#ˆØÑ&ˆØ¤§	¢	¨!¨ZÓ 8Ñ8ˆØÑ#ˆÜ�GŠG�GÑ%¨DÑ1ˆÜ�GŠG�GÑ%¨DÑ1ˆØ‘ˆÜ
�Š�Ñ" FÓ6r&   )ÚBLOCK_SIZE_XÚBLOCK_SIZE_Yc                 óÐ  • [         R                  " S5      U-  nU[         R                  " SU5      S S 2S 4   -   nXƒ:  n	[         R                  " S5      U-  n
U
[         R                  " SU5      S S S 24   -   nX´:  nUnUn[         R                  " XX>-  -   -   Xœ-  5      n[         R                  " XXM-  -   -   Xœ-  5      nUU-   n[         R                  " X-X>-  -   -   UXœ-  5        g )Nr   é   r.   )r5   r6   r7   Ú
x_elementsÚ
y_elementsrS   rT   ÚxoffsetÚxindexÚxmaskÚyoffsetÚyindexÚymaskÚx1Úy0Útmp0Útmp1Útmp2s                     r$   Úadd_kernel_2d_autotunedrd   ò   sÝ   € ô6 —-’- Ó" \Ñ1ˆØœ2Ÿ9š9 Q¨Ó5²a¸°gÑ>Ñ>ˆØÑ#ˆÜ—-’- Ó" \Ñ1ˆØœ2Ÿ9š9 Q¨Ó5°dºA°gÑ>Ñ>ˆØÑ#ˆØˆØˆÜ�wŠw�w¨
©Ñ"8Ñ9¸5¹=ÓIˆÜ�wŠw�w¨
©Ñ"8Ñ9¸5¹=ÓIˆØ�d‰{ˆÜ
�Š� *¡/Ñ2Ñ3°T¸5¹=ÕIr&   c                 ó   • U $ )N© )r#   Ú_Ú__s      r$   Ú_dummy_early_config_pruneri     s   € Øˆr&   é
   é   Úearly_config_prune)r#   rK   ÚwarmupÚrepÚprune_configs_byc                 ó  • [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXs:  n[         R                  " X-   US9n	[         R                  " X-   US9n
Xš-   n[         R                  " X'-   X¸S9  g r)   r.   r4   s               r$   Ú*add_kernel_autotuned_with_unsupported_argsrq     su   € ô$ �mŠm Ñ#ˆØÑ&ˆØ¤§	¢	¨!¨ZÓ 8Ñ8ˆØÑ#ˆÜ�GŠG�GÑ%¨DÑ1ˆÜ�GŠG�GÑ%¨DÑ1ˆØ‘ˆÜ
�Š�Ñ" FÓ6r&   c                 ó  • [         R                  " SS9nXe-  nU[         R                  " SU5      -   nXƒ:  n	[         R                  " X-   U	S9n
[         R                  " X-   U	S9nX«-   U-  n[         R                  " X(-   XÉS9  g r)   r.   )r5   r6   r7   r8   Úscaling_factorr'   r9   r:   r;   r-   r<   r=   r>   s                r$   Úadd_kernel_with_scalingrt   8  sz   € ô �mŠm Ñ#ˆØÑ&ˆØ¤§	¢	¨!¨ZÓ 8Ñ8ˆØÑ#ˆÜ�GŠG�GÑ%¨DÑ1ˆÜ�GŠG�GÑ%¨DÑ1ˆØ‘%˜>Ñ)ˆÜ
�Š�Ñ" FÓ6r&   c                 ó  • [         R                  " SS9nXC-  n[         R                  " U U/U/[         R                  5      n[         R                  " UU/U/[         R                  5      nXg-   n[         R                  " UUU/5        g ©Nr   r*   ©r/   r0   Ú_experimental_descriptor_loadÚfloat32Ú_experimental_descriptor_store©	Úin_desc_ptr0Úin_desc_ptr1Úout_desc_ptrr'   r9   ÚoffsetÚaÚbr>   s	            r$   Úadd_kernel_with_tma_1d_old_apir‚   J  s†   € ô �mŠm Ñ#ˆØÑ!ˆä×,Ò,ØØˆHØˆLÜ�J‰Jó	
ˆô ×,Ò,ØØˆHØˆLÜ�J‰Jó	
ˆð ‘ˆä
×)Ò)ØØØˆHõ	
r&   c                 óB  • [         R                  " SS9n[         R                  " SS9nXS-  nXd-  n[         R                  " U Xx/X4/[         R                  5      n	[         R                  " UXx/X4/[         R                  5      n
Xš-   n[         R                  " UUXx/5        g ©Nr   r*   rV   rw   ©r|   r}   r~   rS   rT   Úpid_xÚpid_yÚoffset_xÚoffset_yr<   r=   r>   s               r$   Úadd_kernel_with_tma_2d_old_apirŠ   i  s¡   € ô —’ 1Ñ%ˆÜ—’ 1Ñ%ˆØÑ'ˆØÑ'ˆä×,Ò,ØØÐ ØÐ(Ü�J‰Jó	
ˆô ×,Ò,ØØÐ ØÐ(Ü�J‰Jó	
ˆð ‘ˆä
×)Ò)ØØØÐ õ	
r&   c                 óÎ   • [         R                  " SS9nXC-  n[         R                  " U U/5      n[         R                  " UU/5      nXg-   n[         R                  " UU/U5        g rv   ©r/   r0   Úload_tensor_descriptorÚstore_tensor_descriptorr{   s	            r$   Úadd_kernel_with_tma_1d_new_apir�   ‹  sn   € ô �mŠm Ñ#ˆØÑ!ˆä×%Ò%ØØˆHó
ˆô ×%Ò%ØØˆHó
ˆð
 ‘ˆä
×"Ò"ØØˆHØõ	
r&   c                 óþ   • [         R                  " SS9n[         R                  " SS9nXS-  nXd-  n[         R                  " U Xx/5      n	[         R                  " UXx/5      n
Xš-   n[         R                  " UXx/U5        g r„   rŒ   r…   s               r$   Úadd_kernel_with_tma_2d_new_apir‘   ¦  s‡   € ô —’ 1Ñ%ˆÜ—’ 1Ñ%ˆØÑ'ˆØÑ'ˆä×%Ò%ØØÐ ó
ˆô ×%Ò%ØØÐ ó
ˆð
 ‘ˆä
×"Ò"ØØÐ Øõ	
r&   c                 óÖ  • UnUS-   nUS-   n	[         R                  R                  R                  UU Xf/X4/U R                  R
                  S9  [         R                  R                  R                  UUXf/X4/UR                  R
                  S9  [         R                  R                  R                  U	UXf/X4/UR                  R
                  S9  [         R                  R                  R                  U5        [         R                  R                  R                  U5        [         R                  R                  R                  U	5        [         R                  " SS9n
[         R                  " SS9nX¦-  nX¶-  n[         R                  " UXÍ/Xf/[         R                  5      n[         R                  " UXÍ/Xf/[         R                  5      nXï-   n[         R                  " U	UXÍ/5        g )Nr   é   )Údesc_ptrÚglobal_addressÚ	load_sizeÚglobal_sizeÚ
element_tyr   r*   rV   )r/   ÚextraÚcudaÚ&experimental_device_tensormap_create2dÚdtyper˜   Ú)experimental_tensormap_fenceproxy_acquirer0   rx   ry   rz   )Úa_ptrÚb_ptrÚc_ptrÚmÚnÚ	workspacer'   Ú
a_desc_ptrÚ
b_desc_ptrÚ
c_desc_ptrr†   r‡   rˆ   r‰   r€   r�   r>   s                    r$   Ú add_kernel_on_device_tma_old_apir§   Ä  s¯  € ð ˆ
Ø ‘_ˆ
Ø ‘_ˆ
Ü
�‰�‰×<Ñ<ØØ Ø!Ð.Ø˜Ø—{‘{×-Ñ-ð 	=ñ 	
ô 	�‰�‰×<Ñ<ØØ Ø!Ð.Ø˜Ø—{‘{×-Ñ-ð 	=ñ 	
ô 	�‰�‰×<Ñ<ØØ Ø!Ð.Ø˜Ø—{‘{×-Ñ-ð 	=ñ 	
ô 	�‰�‰×?Ñ?À
ÔKÜ
�‰�‰×?Ñ?À
ÔKÜ
�‰�‰×?Ñ?À
ÔKä—’ 1Ñ%ˆÜ—’ 1Ñ%ˆØÑ%ˆØÑ%ˆô ×,Ò,ØØÐ ØÐ$Ü�J‰Jó	
ˆô ×,Ò,ØØÐ ØÐ$Ü�J‰Jó	
ˆð ‘ˆô 	×)Ò)ØØØÐ õ	
r&   c                 ó   • [         R                  " U X4/US/Xf/S9n[         R                  " UX4/US/Xf/S9n[         R                  " UX4/US/Xf/S9n	[         R                  " SS9n
[         R                  " SS9nX¦-  nX¶-  n[         R                  " UXÍ/5      n[         R                  " UXÍ/5      nXï-   n[         R                  " U	XÍ/U5        g )NrV   )ÚbaseÚshapeÚstridesÚblock_shaper   r*   )r/   Úmake_tensor_descriptorr0   r�   rŽ   )rž   rŸ   r    r¡   r¢   r£   r'   Úa_descÚb_descÚc_descr†   r‡   rˆ   r‰   r€   r�   r>   s                    r$   Ú add_kernel_on_device_tma_new_apir±     s   € ô ×*Ò*ØØ�&Ø˜�FØ#Ð0ñ	
ˆô ×*Ò*ØØ�&Ø˜�FØ#Ð0ñ	
ˆô ×*Ò*ØØ�&Ø˜�FØ#Ð0ñ	
ˆô —’ 1Ñ%ˆÜ—’ 1Ñ%ˆØÑ%ˆØÑ%ˆô ×%Ò%ØØÐ ó
ˆô ×%Ò%ØØÐ ó
ˆð ‘ˆô 	×"Ò"ØØÐ Øõ	
r&   c                 óÖ   • [         R                  " SS9nXC-  nU[         R                  " SU5      -   nXb:  n[         R                  " X-   US9nSU-  n	[         R                  " X-   X—S9  g ©Nr   r*   r,   r   r.   )
r5   r7   r8   r'   r9   r:   r;   r-   r<   r>   s
             r$   Úmul2_kernelr´   ?  sd   € ô �mŠm Ñ#ˆØÑ&ˆØ¤§	¢	¨!¨ZÓ 8Ñ8ˆØÑ#ˆÜ�GŠG�GÑ%¨DÑ1ˆØ�Q‘ˆÜ
�Š�Ñ" FÓ6r&   c                 óÖ   • [         R                  " SS9nX2-  nU[         R                  " SU5      -   nXQ:  n[         R                  " X-   US9nSU-  n[         R                  " X-   X†S9  g r³   r.   )	Úptrr8   r'   r9   r:   r;   r-   r<   r>   s	            r$   Úmul2_inplace_kernelr·   N  sb   € ô �mŠm Ñ#ˆØÑ&ˆØ¤§	¢	¨!¨ZÓ 8Ñ8ˆØÑ#ˆÜ�GŠG�C‘M¨Ñ-ˆØ�Q‘ˆÜ
�Š�‘ Ó2r&   c                 ó8   • [         R                  " U S:¬  U S5      $ )Nr   )r/   Úwhere)r<   s    r$   Ú	zero_negsrº   \  s   € ä�xŠx˜˜Q™  1Ó%Ð%r&   c                 ó  • [         R                  " SS9nXS-  nU[         R                  " SU5      -   nXr:  nUS:X  a  [        XUS9  OUS:X  a  [	        X XUS9  [         R
                  " X-   US9n	[         R                  " X-   X˜S9  g )Nr   r*   r·   )r'   r?   r,   )r/   r0   r1   r·   r?   r2   r3   )
r5   r7   r8   r'   Ú
ACTIVATIONr9   r:   r;   r-   r<   s
             r$   Úindirection_kernelr½   `  s‚   € ô �mŠm Ñ#ˆØÑ&ˆØ¤§	¢	¨!¨ZÓ 8Ñ8ˆØÑ#ˆØÐ.Ó.Ü À
ÓKØ˜<Ó'Ü�w¨ÈÒTÜ�GŠG�GÑ%¨DÑ1ˆÜ
�Š�Ñ" AÓ1r&   c                 óŒ  • [         R                  " SS9n[         R                  " SS9nXd-  nXu-  n	U[         R                  " SU5      -   n
U	[         R                  " SU5      -   nUS S 2S 4   U-  U
S S S 24   -   nUS S 2S 4   U-  U
S S S 24   -   n[         R                  " X-   5      n[         R                  " X-   US-  5        g )Nr   r*   rV   g       @r.   )Úin_ptrr7   Úin_y_strideÚout_y_strideÚX_BLOCK_SIZEÚY_BLOCK_SIZEÚxidÚyidÚx_startÚy_startÚ	x_offsetsÚ	y_offsetsÚsrc_offsetsÚdst_offsetsÚsrcs                  r$   Údouble_strided_kernelrÍ   s  sÂ   € ô �mŠm Ñ#ˆÜ�mŠm Ñ#ˆØÑ$ˆØÑ$ˆØœbŸiši¨¨<Ó8Ñ8ˆ	ØœbŸiši¨¨<Ó8Ñ8ˆ	Ø¢ 4 Ñ(¨;Ñ6¸À4ÊÀ7Ñ9KÑKˆØ¢ 4 Ñ(¨<Ñ7¸)ÀDÊ!ÀGÑ:LÑLˆÜ�gŠg�fÑ*Ó+ˆÜ
�Š�Ñ&¨¨c©	Õ2r&   c           	      óº  • [         R                  " U [         R                  " SU5      -   5      n[         R                  " U[         R                  " SU5      -   5      n[         R                  " U/U[         R                  5      n[         R
                  " SSXVU/[         R                  SSS9n[         R                  " U[         R                  " SU5      -   U5        g )Nr   úshf.l.wrap.b32 $0, $1, $2, $3;ú
=r,r, r, rTrV   ©rœ   Úis_pureÚpack©r/   r2   r1   ÚfullÚint32Úinline_asm_elementwiser3   ©	ÚXÚYÚZr¢   ÚBLOCKr<   r=   ÚsÚzs	            r$   Úinline_asm_kernel_is_pure_truerß   ‡  s¢   € ô �GŠG�AœŸ	š	 ! UÓ+Ñ+Ó,ˆÜ�GŠG�AœŸ	š	 ! UÓ+Ñ+Ó,ˆÜ�GŠG�U�G˜Q¤§¡Ó)ˆÜ×%Ò%Ø,ØØ�1ˆIÜ—(‘(ØØñ
ˆô 	�Š�”R—Y’Y˜q %Ó(Ñ(¨!Õ,r&   c           	      óº  • [         R                  " U [         R                  " SU5      -   5      n[         R                  " U[         R                  " SU5      -   5      n[         R                  " U/U[         R                  5      n[         R
                  " SSXVU/[         R                  SSS9n[         R                  " U[         R                  " SU5      -   U5        g )Nr   rÏ   rÐ   FrV   rÑ   rÔ   rØ   s	            r$   Úinline_asm_kernel_is_pure_falserá   ˜  s¢   € ô �GŠG�AœŸ	š	 ! UÓ+Ñ+Ó,ˆÜ�GŠG�AœŸ	š	 ! UÓ+Ñ+Ó,ˆÜ�GŠG�U�G˜Q¤§¡Ó)ˆÜ×%Ò%Ø,ØØ�1ˆIÜ—(‘(ØØñ
ˆô 	�Š�”R—Y’Y˜q %Ó(Ñ(¨!Õ,r&   c                 ój  • [         R                  " SS9nXT-  n[         R                  " [         R                  " U U/S/U/U/S/S9S/S9n[         R                  " [         R                  " UU/S/U/U/S/S9S/S9nXx-   n	[         R                  " [         R                  " UU/S/U/U/S/S9U	S/S9  g ©Nr   r*   rV   )r©   rª   r«   r;   r¬   Úorder)Úboundary_check©r/   r0   r2   Úmake_block_ptrr3   )
Úx_ptrÚy_ptrÚ
output_ptrr8   r'   r9   r:   r<   r=   r>   s
             r$   Úadd_kernel_with_block_ptrrë   ©  sã   € ô �mŠm Ñ#ˆØÑ&ˆÜ�GŠGÜ×ÒØØ!�lØ˜Ø$˜Ø'˜LØ�cñð ˜3ñ

ˆô �GŠGÜ×ÒØØ!�lØ˜Ø$˜Ø'˜LØ�cñð ˜3ñ

ˆð ‘ˆÜ
�ŠÜ×ÒØØ!�lØ˜Ø$˜Ø'˜LØ�cñð Ø˜3ó	
r&   c                 ó  • [         R                  " SS9nXC-  n[         R                  " [         R                  " U US/SS/US/US/SS/S9S/S9nUn[         R                  " [         R                  " UUS/SS/US/US/SS/S9US/S9  g rã   ræ   )rè   rê   r8   r'   r9   r:   r<   r>   s           r$   Úkernel_with_block_ptr_2drí   ×  s¶   € ô �mŠm Ñ#ˆØÑ&ˆÜ�GŠGÜ×ÒØØ! 1�oØ˜A˜Ø$ aÐ(Ø'¨˜OØ˜!�fñð ˜3ñ

ˆð ˆÜ
�ŠÜ×ÒØØ! 1�oØ˜A˜Ø$ aÐ(Ø'¨˜OØ˜!�fñð Ø˜3ó	
r&   )r2   r3   c                 óÀ   • [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXs:  n[        X-   US9n	[        X-   US9n
Xš-   n[	        X'-   X¸S9  g r)   r.   r4   s               r$   Úadd_kernel_with_importrï   û  si   € ô �mŠm Ñ#ˆØÑ&ˆØ¤§	¢	¨!¨ZÓ 8Ñ8ˆØÑ#ˆÜ�Ñ"¨Ñ.ˆÜ�Ñ"¨Ñ.ˆØ‘ˆÜˆgÑ Ó3r&   c                 ó@  • [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXs:  n[         R                  " X-   US9n	[         R                  " X-   US9n
[         R                  " S5      S:X  a  Xš-   nOXš-  n[         R                  " X'-   X¸S9  g r)   r.   r4   s               r$   Úcond_op_kernelrñ     s�   € ô �mŠm Ñ#ˆØÑ&ˆØ¤§	¢	¨!¨ZÓ 8Ñ8ˆØÑ#ˆÜ�GŠG�GÑ%¨DÑ1ˆÜ�GŠG�GÑ%¨DÑ1ˆÜ�=Š=˜Ó˜qÓ Ø‘U‰Fà‘UˆFÜ
�Š�Ñ" FÓ6r&   c                 ó  • [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXs:  n[         R                  " X-   US9n	[         R                  " X-   US9n
Xš-   n[         R                  " X'-   X¸S9  g r)   )r/   r0   r1   r2   Ú
atomic_addr4   s               r$   Úatomic_add_kernelrô      su   € ô �mŠm Ñ#ˆØÑ&ˆØ¤§	¢	¨!¨ZÓ 8Ñ8ˆØÑ#ˆÜ�GŠG�GÑ%¨DÑ1ˆÜ�GŠG�GÑ%¨DÑ1ˆØ‘ˆÜ
�Š�gÑ'¨Ó;r&   c                 óˆ  • [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXs:  n[         R                  " X-   US9n	[         R                  " X-   US9n
[	        S5       H  nXš-   n[         R
                  " X'-   XÈS9  M      SnUS:”  a)  US-  nXš-   n[         R
                  " X'-   XÈS9  US:”  a  M(  g g )Nr   r*   r,   r   rV   )r/   r0   r1   r2   Úranger3   )r5   r6   r7   r8   r'   r9   r:   r;   r-   r<   r=   rg   r>   Úis                 r$   Úadd_4_times_kernelrø   1  s»   € ô �mŠm Ñ#ˆØÑ&ˆØ¤§	¢	¨!¨ZÓ 8Ñ8ˆØÑ#ˆÜ�GŠG�GÑ%¨DÑ1ˆÜ�GŠG�GÑ%¨DÑ1ˆÜ�q–ˆAØ‘UˆFÜ�HŠH�WÑ&¨Ô:ñ ð ˆØ�!‹eØ�‰FˆAØ‘UˆFÜ�HŠH�WÑ&¨Ò:ð �!�er&   c                 ó  • [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXr:  n[         R                  " X-   US9n	[         R                  " X-   US9n
Xš-   n[         R                  " X7-   X¸S9  g r)   r.   )r5   r6   r8   r7   r'   r9   r:   r;   r-   r<   r=   r>   s               r$   Úadd_kernel_out_of_order_fn2rú   H  r@   r&   r   r   )ÚM_ptrÚNÚKr   r   r   r   c
                 óæ  • [         R                  " SS9n
[         R                  " U5      nUS:X  a	  US:”  a  SnOUS:X  a  g [         R                  " X¶5      n[         R                  " XG5      nX�-  nX®-  nXù-  n[	        UU-
  U	5      nUX®-  U-  -   nX®-  U-  nUU-  [         R
                  " SU5      -   U-  nUU-  [         R
                  " SU5      -   U-  n[         R
                  " SU5      nU US S 2S 4   US S S 24   -   -   nUUS S 2S 4   US S S 24   -   -   n[         R                  " Xg4[         R                  S9n[        [         R                  " XX5      5       Hq  n[         R                  " UUS S S 24   UUU-  -
  :  SS9n[         R                  " UUS S 2S 4   UUU-  -
  :  SS9n[         R                  " UUU5      nUU-  nUU-  nMs     UR                  [         R                  5      nUU-  [         R
                  " SU5      -   nUU-  [         R
                  " SU5      -   nUUS S 2S 4   -   US S S 24   -   n US S 2S 4   U:  US S S 24   U:  -  n![         R                  " U UU!S9  g )	Nr   r*   r   i   ©rœ   g        ©r-   Úotherr,   )r/   r0   r2   ÚcdivÚminr1   Úzerosry   rö   ÚdotÚtoÚfloat16r3   )"rž   rŸ   r    rû   rü   rý   r   r   r   r   r9   ÚMÚ	num_pid_mÚ	num_pid_nÚnum_pid_in_groupÚgroup_idÚfirst_pid_mÚgroup_size_mÚpid_mÚpid_nÚoffs_amÚoffs_bnÚoffs_kÚa_ptrsÚb_ptrsÚaccumulatorÚkr€   r�   ÚcÚoffs_cmÚoffs_cnÚc_ptrsÚc_masks"                                     r$   Ústrange_config_matmul_kernelr  Y  sx  € ôN �mŠm Ñ#ˆÜ�GŠG�E‹NˆØ�‹6�l RÓ'à‰AØ�!‹VàÜ—G’G˜AÓ,ˆ	Ü—G’G˜AÓ,ˆ	Ø'Ñ3ÐØÑ*ˆØÑ-ˆÜ˜9 {Ñ2°LÓAˆØ Ñ 6¸,ÑFÑGˆØÑ'¨LÑ8ˆà˜<Ñ'¬"¯)ª)°A°|Ó*DÑDÈÑIˆØ˜<Ñ'¬"¯)ª)°A°|Ó*DÑDÈÑIˆÜ—’˜1˜lÓ+ˆØ˜'¢! T 'Ñ*¨V°Dº!°G©_Ñ<Ñ=ˆØ˜&¢ D ™/¨G°Dº!°GÑ,<Ñ<Ñ=ˆä—h’h Ð;Ä2Ç:Á:ÑNˆÜ”r—w’w˜qÓ/Ö0ˆAÜ—’˜ V¨D²!¨G¡_°q¸1¸|Ñ;KÑ7KÑ%KÐSVÑWˆAÜ—’˜ VªA¨t¨G¡_°q¸1¸|Ñ;KÑ7KÑ%KÐSVÑWˆAÜŸ&š&  A {Ó3ˆKØ�lÑ"ˆFØ�lÑ"ŠFñ 1ð �N‰Nœ2Ÿ:™:Ó&ˆà˜,Ñ&¬¯ª°1°lÓ)CÑCˆØ˜,Ñ&¬¯ª°1°lÓ)CÑCˆØ˜¢ D Ñ)Ñ)¨G°Dº!°GÑ,<Ñ<ˆØš!˜T˜'Ñ" QÑ&¨7°4º°7Ñ+;¸aÑ+?Ñ@ˆÜ
�Š�˜ Ó(r&   c                 óâ   • [         R                  " SS9n[         R                  " SU5      X2-  -   n[         R                  " U/S[         R                  S9n[         R
                  " X-   XTU:  S9  g)zq
This kernel contains a triple-quote docstring w/ double quotes.
Make sure that codegen sanitizes the docstring.
r   r*   ç      ð?rÿ   r,   N©r/   r0   r1   rÕ   ry   r3   ©r7   Únumelr'   r9   r;   Úoness         r$   Ú#kernel_with_docstring_double_quotesr$  ¦  sW   € ô �mŠm Ñ#ˆÜ—)’)˜A˜zÓ*¨SÑ-=Ñ=ˆÜ�wŠw˜
�| S´·
±
Ñ;ˆÜ
�Š�Ñ" D¸©Ó?r&   c                 óâ   • [         R                  " SS9n[         R                  " SU5      X2-  -   n[         R                  " U/S[         R                  S9n[         R
                  " X-   XTU:  S9  g)z¬
This kernel contains a triple-quote docstring w/ single quotes
Make sure that codegen sanitizes the docstring.
To prevent it from being linted to double quotes: """!!!"""
r   r*   r  rÿ   r,   Nr   r!  s         r$   Ú#kernel_with_docstring_single_quotesr&  ±  sW   € ô �mŠm Ñ#ˆÜ—)’)˜A˜zÓ*¨SÑ-=Ñ=ˆÜ�wŠw˜
�| S´·
±
Ñ;ˆÜ
�Š�Ñ" D¸©Ó?r&   c                 óâ   • [         R                  " SS9n[         R                  " SU5      X2-  -   n[         R                  " U/S[         R                  S9n[         R
                  " X-   XTU:  S9  g)z,Docstring with literal backslash escapes: 
\r   r*   r  rÿ   r,   Nr   r!  s         r$   Ú"kernel_with_backslash_in_docstringr(  ½  sW   € ô �mŠm Ñ#ˆÜ—)’)˜A˜zÓ*¨SÑ-=Ñ=ˆÜ�wŠw˜
�| S´·
±
Ñ;ˆÜ
�Š�Ñ" D¸©Ó?r&   c           	      ó  • [         R                  " SS9n[         R                  " SU5      XC-  -   n[         R                  " X-   XR:  S9n[         R                  " SSU/[         R
                  SSS9n[         R                  " X-   XuU:  S9  g )	Nr   r*   r,   z{
            {
                cos.approx.f32 $0, $1;
                ex2.approx.f32 $0, $0;
            }
                ú=r, rTrV   ©ÚasmÚconstraintsÚargsrœ   rÒ   rÓ   ©r/   r0   r1   r2   r×   ry   r3   ©r¿   r7   r"  r'   r9   r;   ÚdataÚcos_pows           r$   Úkernel_inline_asm_double_quotesr3  Å  s‚   € ô �mŠm Ñ#ˆÜ—)’)˜A˜zÓ*¨SÑ-=Ñ=ˆÜ�wŠw�vÑ'¨g©oÑ>ˆÜ×+Ò+ðð !Ø�Ü—*‘*ØØñ
ˆô 	�Š�Ñ" G¸E±/ÓBr&   c           	      ó  • [         R                  " SS9n[         R                  " SU5      XC-  -   n[         R                  " X-   XR:  S9n[         R                  " SSU/[         R
                  SSS9n[         R                  " X-   XuU:  S9  g )	Nr   r*   r,   z»
            {
                // double quotes to pacify the linter """!!!"""
                cos.approx.f32 $0, $1;
                ex2.approx.f32 $0, $0;
            }
                r*  TrV   r+  r/  r0  s           r$   Úkernel_inline_asm_single_quotesr5  Û  s‚   € ô �mŠm Ñ#ˆÜ—)’)˜A˜zÓ*¨SÑ-=Ñ=ˆÜ�wŠw�vÑ'¨g©oÑ>ˆÜ×+Ò+ðð !Ø�Ü—*‘*ØØñ
ˆô 	�Š�Ñ" G¸E±/ÓBr&   c           	      ó  • [         R                  " SS9n[         R                  " SU5      XC-  -   n[         R                  " X-   XR:  S9n[         R                  " SSU/[         R
                  SSS9n[         R                  " X-   XuU:  S9  g ©	Nr   r*   r,   zK
            v_sin_f32 $0, $1
            v_exp_f32 $0, $0
                z=v, vTrV   r+  r/  r0  s           r$   Ú$kernel_inline_asm_rocm_double_quotesr8  ò  ó‚   € ô �mŠm Ñ#ˆÜ—)’)˜A˜zÓ*¨SÑ-=Ñ=ˆÜ�wŠw�vÑ'¨g©oÑ>ˆÜ×+Ò+ðð !Ø�Ü—*‘*ØØñ

ˆô 	�Š�Ñ" G¸E±/ÓBr&   c           	      ó  • [         R                  " SS9n[         R                  " SU5      XC-  -   n[         R                  " X-   XR:  S9n[         R                  " SSU/[         R
                  SSS9n[         R                  " X-   XuU:  S9  g r7  r/  r0  s           r$   Ú$kernel_inline_asm_rocm_single_quotesr;    r9  r&   c                 ó  • [         R                  " SS9nXe-  nU[         R                  " SU5      -   nXƒ:  n	[         R                  " X-   U	S9n
U(       a  [         R                  " X-   U	S9nX«-   nOU
n[         R                  " X(-   XÉS9  g r)   r.   )r5   r6   r7   r8   Úadd_xyr'   r9   r:   r;   r-   r<   r=   r>   s                r$   Úadd_kernel_with_boolean_paramr>    s}   € ô �mŠm Ñ#ˆØÑ&ˆØ¤§	¢	¨!¨ZÓ 8Ñ8ˆØÑ#ˆÜ�GŠG�GÑ%¨DÑ1ˆÞÜ—’˜Ñ)°Ñ5ˆAØ‘U‰FàˆFÜ
�Š�Ñ" FÓ6r&   c                 ód  • [         R                  " SS9nXe-  nU[         R                  " SU5      -   nX„:  n	[         R                  " X-   U	S9n
[         R                  " X-   U	S9n[         R                  " X(-   U	SS9S:g  n[         R                  " XÊU-   U
5      n[         R
                  " X8-   XÙS9  g)z6Kernel that loads a bool tensor and uses it as a mask.r   r*   r,   r   N)r/   r0   r1   r2   r¹   r3   )r5   r6   Úmask_ptrr7   r8   r'   r9   r:   r;   Úvalidr<   r=   Úkeepr>   s                 r$   Ú"masked_add_kernel_with_bool_tensorrC  /  sœ   € ô �mŠm Ñ#ˆØÑ&ˆØ¤§	¢	¨!¨ZÓ 8Ñ8ˆØÑ$ˆÜ�GŠG�GÑ%¨EÑ2ˆÜ�GŠG�GÑ%¨EÑ2ˆÜ�wŠw�xÑ)°¸QÑ?À1ÑDˆÜ—’˜$ A¡ qÓ)ˆÜ
�Š�Ñ" FÓ7r&   Úblock_sizesÚnew_apic           	      óp  • U(       a3  [         R                  R                  R                  R	                  X5      $ [        U5      S:X  aZ  [         R                  R                  R                  U R                  5       U R                  S5      US   U R                  5       5      $ [        U5      S:w  a  [        S[        U5       35      e[         R                  R                  R                  U R                  5       U R                  S5      U R                  S5      US   US   U R                  5       5      $ )NrV   r   r   z$Expected len(block_sizes) == 2, got )r!   ÚtoolsÚtensor_descriptorÚTensorDescriptorÚfrom_tensorÚlenÚexperimental_descriptorÚcreate_1d_tma_descriptorÚdata_ptrÚsizeÚelement_sizeÚAssertionErrorÚcreate_2d_tma_descriptor)ÚtensorrD  rE  s      r$   Úcreate_tensor_descriptor_shimrT  D  s  € ö Ü—<‘<×1Ñ1×BÑB×NÑNØóð ô �;Ó 1Ó$Ü—|‘|×;Ñ;×TÑTØ—O‘OÓ%Ø—K‘K “NØ ‘NØ×'Ñ'Ó)ó	ð ô �{Ó# qÓ(Ü(Ø>¼sÀ;Ó?OÐ>PÐQóð ô —|‘|×;Ñ;×TÑTØ—O‘OÓ%Ø—K‘K “NØ—K‘K “NØ ‘NØ ‘NØ×'Ñ'Ó)óð r&   )r'   útl.constexpr)rF   rU  r'   rU  )rS   rU  rT   rU  )r'   rU  r¼   rU  )rÂ   rU  rÃ   rU  )r¢   rU  rÜ   rU  )T)FÚunittestÚ&torch.testing._internal.inductor_utilsr   r   r   r   Útorch.utils._tritonr   Ú
skipUnlessÚrequires_cuda_and_tritonÚrequires_xpu_and_tritonÚrequires_gpu_and_tritonÚrequires_gpur!   r   r/   r   Úlistr"   r%   Újitr?   rB   rG   rJ   ÚautotunerM   rP   rR   rd   ri   rq   rt   r‚   rŠ   r�   r‘   r§   r±   r´   r·   rº   r½   rÍ   rß   rá   Ú	constexprrë   rí   Útriton.languager2   r3   rï   rñ   rô   rø   rú   r  r$  r&  r(  r3  r5  r8  r;  r>  rC  ÚintÚboolrT  rf   r&   r$   Ú<module>re     s/	  ðó ÷ó õ +ð $×.Ò.ØÐ3óÐ ð #×-Ò-ØÐ1óÐ ð #×-Ò-Ø×DÐ-×DÐ1DØóÐ ð ×"Ò" 7¨NÓ;€á‡<‚<ÛÝ%ãð5 $ v§}¡}Ñ"5ô 5ðp ‡Z�Zð7ð
 #ó7ó ð7ð  ‡Z�Zð7ð
 #ó7ó ð7ð  ‡Z�Zð7ð
 $ð7ð #ó7ó ð7ð( ‡Z�Zð@ð $ð@ð #ó@ó ð@ð* ‡_‚_à�MŠM˜<¨Ð-¸!ÀqÑIØ�MŠM˜<¨Ð-¸!ÀqÑIØ�MŠM˜<¨Ð,¸ÀaÑHØ�MŠM˜<¨Ð,¸ÀaÑHð	
ð ñð ‡Z�Zð7ð
 #ó7ó óð7ð  ‡_‚_à�MŠM˜<¨Ð-¸!ÀqÑIØ�MŠM˜<¨Ð-¸!ÀqÑIØ�MŠM˜<¨Ð,¸ÀaÑHØ�MŠM˜<¨Ð,¸ÀaÑHð	
ð ñð ‡Z�Zð7ð
 #ó7ó óð7ð  ‡_‚_à�MŠM˜<¨Ð,¸ÀaÑHð
ð ñ	ð ‡Z�Zð7ð #ó	7ó óð7ð$ ‡_‚_à�MŠMØ!$°cÑ:ÀqÐTUñð �MŠMØ!$°cÑ:ÀqÐTUñð �MŠMØ!#°RÑ8ÀQÐRSñð �MŠMØ!#°RÑ8ÀQÐRSñð
ð ñð" ‡Z�ZðJð %ðJð %óJó ó#ð$Jò,ð ‡_‚_à�MŠM˜<¨Ð-¸!ÀqÑIØ�MŠM˜<¨Ð,¸ÀaÑHð
ð ØØØ.Ð0IÐJñ	ð ‡Z�Zð7ð
 #ó7ó ó	ð7ð  ‡Z�Zð7ð #ó7ó ð7ð" ‡Z�Zð
ð #ó	
ó ð
ð< ‡Z�Zð
ð %ð	
ð
 %ó
ó ð
ðB ‡Z�Zð
ð #ó	
ó ð
ð4 ‡Z�Zð
ð %ð	
ð
 %ó
ó ð
ð: ‡Z�ZðA
ð #óA
ó ðA
ðF ‡Z�Zð4
ð #ó4
ó ð4
ðl ‡Z�Zð7ð #ó	7ó ð7ð ‡Z�Zð3ð #ó3ó ð3ð ‡Z�Zñ&ó ð&ð ‡Z�Zð2ð #ð	2ð
 #ó2ó ð2ð$ ‡Z�Zð3ð
 %ð3ð %ó3ó ð3ð& ‡Z�Zð-Ø"ð-Ø+9ó-ó ð-ð  ‡Z�Zð-Ø"ð-Ø+9ó-ó ð-ð  ‡Z�Zð+
ð
 —L‘Ló+
ó ð+
ðZ ‡Z�Zð
ð —L‘Ló	
ó ð
÷B ,à‡Z�Zð4ð
 #ó4ó ð4ð  ‡Z�Zð7ð
 #ó7ó ð7ð& ‡Z�Zð<ð
 #ó<ó ð<ð  ‡Z�Zð;ð
 #ó;ó ð;ð, ‡Z�Zð7ð
 #ó7ó ð7ð  ‡_‚_à�MŠMà$&Ø$&Ø$&Ø$%ñ	ð Øñ	ð �MŠMà$'Ø$&Ø$&Ø$%ñ	ð Øñ	ð
ò,  ñ/ð2 ‡Z�Zð1)ð —l‘lð1)ð —l‘lð1)ð —l‘lð1)ð —l‘ló1)ó ó3ð41)ðf ‡Z�Zð@ÈÏÉó @ó ð@ð ‡Z�Zð	@ÈÏÉó 	@ó ð	@ð ‡Z�Zð@ÀrÇ|Á|ó @ó ð@ð ‡Z�ZðCØ,.¯L©LóCó ðCð* ‡Z�ZðCØ,.¯L©LóCó ðCð, ‡Z�ZðCØ,.¯L©LóCó ðCð& ‡Z�ZðCØ,.¯L©LóCó ðCð& ‡Z�Zð7ð #ó7ó ð7ð( ‡Z�Zð8ð #ó8ó ð8ð* 9=ñØ! #™YðØ15öðU! r&   