Ë
    ÌÓ^j4‡  ã            
       ó0  — d dl Z d dlmZmZmZ d dlmZ  e j                  ed«      Z e j                  ed«      Z	 e j                  exs ed«      Z
 e j                  ed«      Z e«       �r*d dlZd dlmZ d dlZd	eej"                     fd
„Zej&                  	 	 dNd„«       Zej&                  	 	 dNd„«       Zej&                  	 	 	 	 dOd„«       Zej&                  	 	 	 	 dOd„«       Z ej0                   ej"                  ddidd¬«       ej"                  ddidd¬«       ej"                  ddidd¬«       ej"                  ddidd¬«      gg ¬«      ej&                  	 	 dNd„«       «       Z ej0                   ej"                  ddidd¬«       ej"                  ddidd¬«       ej"                  ddidd¬«       ej"                  ddidd¬«      gg ¬«      ej&                  	 	 dNd„«       «       Z ej0                   ej"                  ddidd¬«      gg ¬«      ej&                  	 	 dNd„«       «       Z ej0                   ej"                  dddœdd¬«       ej"                  dddœdd¬«       ej"                  dddœdd¬«       ej"                  dddœdd¬«      gg ¬«      ej&                  	 	 	 	 dPd„«       «       Zd„ Z ej0                   ej"                  ddidd¬«       ej"                  ddidd¬«      gg dd d!ei¬"«      ej&                  	 	 dNd#„«       «       Zej&                  	 	 dNd$„«       Zej&                  	 	 dNd%„«       Z ej&                  	 	 	 	 dPd&„«       Z!ej&                  	 	 dNd'„«       Z"ej&                  	 	 	 	 dPd(„«       Z#ej&                  	 	 dNd)„«       Z$ej&                  	 	 dNd*„«       Z%ej&                  	 	 dNd+„«       Z&ej&                  	 	 dNd,„«       Z'ej&                  d-„ «       Z(ej&                  	 	 	 	 dQd.„«       Z)ej&                  	 	 	 	 dRd/„«       Z*ej&                  	 	 	 	 dSd0„«       Z+ej&                  	 	 	 	 dSd1„«       Z,ej&                  dejZ                  fd2„«       Z.ej&                  dejZ                  fd3„«       Z/d d4l0m1Z1m2Z2 ej&                  	 	 dNd5„«       Z3ej&                  	 	 dNd6„«       Z4ej&                  	 	 dNd7„«       Z5ej&                  	 	 dNd8„«       Z6ej&                  	 	 dNd9„«       Z7 ej0                   ej"                  ddddd:œdd¬«       ej"                  ddd;dd:œdd¬«      gg d<¢¬«      ej&                  d=ejZ                  d>ejZ                  d?ejZ                  d@ejZ                  fdA„«       «       Z8ej&                  dejZ                  fdB„«       Z9ej&                  dejZ                  fdC„«       Z:ej&                  dejZ                  fdD„«       Z;ej&                  dejZ                  fdE„«       Z<ej&                  dejZ                  fdF„«       Z=ej&                  dejZ                  fdG„«       Z>ej&                  dejZ                  fdH„«       Z?ej&                  	 	 dNdI„«       Z@ej&                  	 	 dNdJ„«       ZA	 dTdKeeB   dLeCfdM„ZDyy)Ué    N)ÚHAS_CUDA_AND_TRITONÚHAS_GPUÚHAS_XPU_AND_TRITON)Ú
has_tritonzrequires cuda and tritonzrequires xpu and tritonzrequires gpu and tritonzrequires gpu)ÚlanguageÚreturnc                  ó2  — t         j                  j                  rBt        j                  ddddddddœdd¬«      t        j                  ddddddddœdd¬«      g} | S t        j                  ddddd	œdd¬«      t        j                  ddd
dd	œdd¬«      g} | S )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    úo/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/torch/testing/_internal/triton_utils.pyÚ_get_strange_configsr$      sØ   € Ü�=‰=×Òä—‘à(*Ø(*Ø(*Ø()Ø02Ø()Ø!"ñð  !Øôô —‘à(+Ø(*Ø(*Ø()Ø02Ø()Ø!"ñð  !ØôðˆGðf ˆô+ —‘à(*Ø(*Ø(*Ø()ñ	ð  !Øô	ô —‘à(+Ø(*Ø(*Ø()ñ	ð  !Øô	ðˆGð, ˆó    Ú
BLOCK_SIZEc                 ó  — t        j                  d¬«      }||z  }|t        j                  d|«      z   }||k  }t        j                  | |z   |¬«      }	t        j                  ||z   |¬«      }
|	|
z   }t        j                  ||z   ||¬«       y ©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>   V   óƒ   € ô �m‰m Ô#ˆØ˜JÑ&ˆØ¤§	¡	¨!¨ZÓ 8Ñ8ˆØ˜Ñ#ˆÜ�G‰G�G˜gÑ%¨DÔ1ˆÜ�G‰G�G˜gÑ%¨DÔ1ˆØ�Q‘ˆÜ
�‰�˜7Ñ" F°Ö6r%   c                 ó  — t        j                  d¬«      }||z  }|t        j                  d|«      z   }||k  }t        j                  | |z   |¬«      }	t        j                  ||z   |¬«      }
|	|
z
  }t        j                  ||z   ||¬«       y r(   r-   r3   s               r#   Ú
sub_kernelrA   g   r?   r%   c                 ó.  — t        j                  d¬«      }||z  }|t        j                  d|«      z   }||k  }	t        j                  | |z   |	¬«      }
|dk(  r t        j                  ||z   |	¬«      }|
|z   }n|
}t        j                  ||z   ||	¬«       y ©Nr   r)   r+   Útwor-   )r4   r5   r6   r7   ÚARGS_PASSEDr&   r8   r9   r:   r,   r;   r<   r=   s                r#   Úadd_kernel_with_optional_paramrF   x   s�   € ô �m‰m Ô#ˆØ˜JÑ&ˆØ¤§	¡	¨!¨ZÓ 8Ñ8ˆØ˜Ñ#ˆÜ�G‰G�G˜gÑ%¨DÔ1ˆØ˜%ÒÜ—‘˜ 'Ñ)°Ô5ˆAØ˜‘U‰FàˆFÜ
�‰�˜7Ñ" F°Ö6r%   c                 ó:  — t        j                  d¬«      }||z  }|t        j                  d|«      z   }	|	|k  }
t        j                  | |	|z  z   |
¬«      }|dk(  r t        j                  ||	z   |
¬«      }||z   }n|}t        j                  ||	|z  z   ||
¬«       y rC   r-   )r4   r5   r6   r7   ÚstriderE   r&   r8   r9   r:   r,   r;   r<   r=   s                 r#   Ú-add_kernel_with_none_param_and_equal_to_1_argrI   �   sš   € ô �m‰m Ô#ˆØ˜JÑ&ˆØ¤§	¡	¨!¨ZÓ 8Ñ8ˆØ˜Ñ#ˆÜ�G‰G�G˜g¨Ñ.Ñ.°TÔ:ˆØ˜%ÒÜ—‘˜ 'Ñ)°Ô5ˆAØ˜‘U‰FàˆFÜ
�‰�˜7 VÑ+Ñ+¨V¸$Ö?r%   r   r   r   r   r   r   )r"   Úkeyc                 ó  — t        j                  d¬«      }||z  }|t        j                  d|«      z   }||k  }t        j                  | |z   |¬«      }	t        j                  ||z   |¬«      }
|	|
z   }t        j                  ||z   ||¬«       y r(   r-   r3   s               r#   Úadd_kernel_autotunedrL   £   óƒ   € ô" �m‰m Ô#ˆØ˜JÑ&ˆØ¤§	¡	¨!¨ZÓ 8Ñ8ˆØ˜Ñ#ˆÜ�G‰G�G˜gÑ%¨DÔ1ˆÜ�G‰G�G˜gÑ%¨DÔ1ˆØ�Q‘ˆÜ
�‰�˜7Ñ" F°Ö6r%   c                 ó  — t        j                  d¬«      }||z  }|t        j                  d|«      z   }||k  }t        j                  | |z   |¬«      }	t        j                  ||z   |¬«      }
|	|
z
  }t        j                  ||z   ||¬«       y r(   r-   r3   s               r#   Úsub_kernel_autotunedrO   ½   rM   r%   r
   r   c                 ó  — t        j                  d¬«      }||z  }|t        j                  d|«      z   }||k  }t        j                  | |z   |¬«      }	t        j                  ||z   |¬«      }
|	|
z   }t        j                  ||z   ||¬«       y r(   r-   )r4   r5   r7   r&   r6   r8   r9   r:   r,   r;   r<   r=   s               r#   Ú&add_kernel_autotuned_weird_param_orderrQ   ×   sƒ   € ô  �m‰m Ô#ˆØ˜JÑ&ˆØ¤§	¡	¨!¨ZÓ 8Ñ8ˆØ˜Ñ#ˆÜ�G‰G�G˜gÑ%¨DÔ1ˆÜ�G‰G�G˜gÑ%¨DÔ1ˆØ�Q‘ˆÜ
�‰�˜7Ñ" F°Ö6r%   )ÚBLOCK_SIZE_XÚBLOCK_SIZE_Yc                 óØ  — t        j                  d«      |z  }|t        j                  d|«      d d …d f   z   }||k  }	t        j                  d«      |z  }
|
t        j                  d|«      d d d …f   z   }||k  }|}|}t        j                  | |||z  z   z   |	|z  «      }t        j                  | |||z  z   z   |	|z  «      }||z   }t        j                  ||||z  z   z   ||	|z  «       y )Nr   é   r-   )r4   r5   r6   Ú
x_elementsÚ
y_elementsrR   rS   ÚxoffsetÚxindexÚxmaskÚyoffsetÚyindexÚymaskÚx1Úy0Útmp0Útmp1Útmp2s                     r#   Úadd_kernel_2d_autotunedrc   ð   só   € ô6 —-‘- Ó" \Ñ1ˆØœ2Ÿ9™9 Q¨Ó5²a¸°gÑ>Ñ>ˆØ˜Ñ#ˆÜ—-‘- Ó" \Ñ1ˆØœ2Ÿ9™9 Q¨Ó5°dºA°gÑ>Ñ>ˆØ˜Ñ#ˆØˆØˆÜ�w‰w�w "¨
°R©Ñ"8Ñ9¸5À5¹=ÓIˆÜ�w‰w�w "¨
°R©Ñ"8Ñ9¸5À5¹=ÓIˆØ�d‰{ˆÜ
�‰�˜B *¨r¡/Ñ2Ñ3°T¸5À5¹=ÕIr%   c                 ó   — | S )N© )r"   Ú_Ú__s      r#   Ú_dummy_early_config_prunerh     s   € Øˆr%   é
   é   Úearly_config_prune)r"   rJ   ÚwarmupÚrepÚprune_configs_byc                 ó  — t        j                  d¬«      }||z  }|t        j                  d|«      z   }||k  }t        j                  | |z   |¬«      }	t        j                  ||z   |¬«      }
|	|
z   }t        j                  ||z   ||¬«       y r(   r-   r3   s               r#   Ú*add_kernel_autotuned_with_unsupported_argsrp     sƒ   € ô$ �m‰m Ô#ˆØ˜JÑ&ˆØ¤§	¡	¨!¨ZÓ 8Ñ8ˆØ˜Ñ#ˆÜ�G‰G�G˜gÑ%¨DÔ1ˆÜ�G‰G�G˜gÑ%¨DÔ1ˆØ�Q‘ˆÜ
�‰�˜7Ñ" F°Ö6r%   c                 ó$  — t        j                  d¬«      }||z  }|t        j                  d|«      z   }||k  }	t        j                  | |z   |	¬«      }
t        j                  ||z   |	¬«      }|
|z   |z  }t        j                  ||z   ||	¬«       y r(   r-   )r4   r5   r6   r7   Úscaling_factorr&   r8   r9   r:   r,   r;   r<   r=   s                r#   Úadd_kernel_with_scalingrs   6  sˆ   € ô �m‰m Ô#ˆØ˜JÑ&ˆØ¤§	¡	¨!¨ZÓ 8Ñ8ˆØ˜Ñ#ˆÜ�G‰G�G˜gÑ%¨DÔ1ˆÜ�G‰G�G˜gÑ%¨DÔ1ˆØ�a‘%˜>Ñ)ˆÜ
�‰�˜7Ñ" F°Ö6r%   c                 ó  — t        j                  d¬«      }||z  }t        j                  | |g|gt         j                  «      }t        j                  ||g|gt         j                  «      }||z   }t        j                  |||g«       y ©Nr   r)   ©r.   r/   Ú_experimental_descriptor_loadÚfloat32Ú_experimental_descriptor_store©	Úin_desc_ptr0Úin_desc_ptr1Úout_desc_ptrr&   r8   ÚoffsetÚaÚbr=   s	            r#   Úadd_kernel_with_tma_1d_old_apir�   H  sŠ   € ô �m‰m Ô#ˆØ�zÑ!ˆä×,Ñ,ØØˆHØˆLÜ�J‰Jó	
ˆô ×,Ñ,ØØˆHØˆLÜ�J‰Jó	
ˆð �Q‘ˆä
×)Ñ)ØØØˆHõ	
r%   c                 óT  — t        j                  d¬«      }t        j                  d¬«      }||z  }||z  }t        j                  | ||g||gt         j                  «      }	t        j                  |||g||gt         j                  «      }
|	|
z   }t        j                  ||||g«       y ©Nr   r)   rU   rv   ©r{   r|   r}   rR   rS   Úpid_xÚpid_yÚoffset_xÚoffset_yr;   r<   r=   s               r#   Úadd_kernel_with_tma_2d_old_apir‰   g  s±   € ô —‘ 1Ô%ˆÜ—‘ 1Ô%ˆØ˜<Ñ'ˆØ˜<Ñ'ˆä×,Ñ,ØØ�xÐ Ø˜<Ð(Ü�J‰Jó	
ˆô ×,Ñ,ØØ�xÐ Ø˜<Ð(Ü�J‰Jó	
ˆð �Q‘ˆä
×)Ñ)ØØØ�xÐ õ	
r%   c                 óÐ   — t        j                  d¬«      }||z  }t        j                  | |g«      }t        j                  ||g«      }||z   }t        j                  ||g|«       y ru   ©r.   r/   Úload_tensor_descriptorÚstore_tensor_descriptorrz   s	            r#   Úadd_kernel_with_tma_1d_new_apirŽ   ‰  sr   € ô �m‰m Ô#ˆØ�zÑ!ˆä×%Ñ%ØØˆHó
ˆô ×%Ñ%ØØˆHó
ˆð
 �Q‘ˆä
×"Ñ"ØØˆHØõ	
r%   c                 ó  — t        j                  d¬«      }t        j                  d¬«      }||z  }||z  }t        j                  | ||g«      }	t        j                  |||g«      }
|	|
z   }t        j                  |||g|«       y rƒ   r‹   r„   s               r#   Úadd_kernel_with_tma_2d_new_apir�   ¤  s“   € ô —‘ 1Ô%ˆÜ—‘ 1Ô%ˆØ˜<Ñ'ˆØ˜<Ñ'ˆä×%Ñ%ØØ�xÐ ó
ˆô ×%Ñ%ØØ�xÐ ó
ˆð
 �Q‘ˆä
×"Ñ"ØØ�xÐ Øõ	
r%   c                 ó  — |}|dz   }|dz   }	t         j                  j                  j                  || ||g||g| j                  j
                  ¬«       t         j                  j                  j                  ||||g||g|j                  j
                  ¬«       t         j                  j                  j                  |	|||g||g|j                  j
                  ¬«       t         j                  j                  j                  |«       t         j                  j                  j                  |«       t         j                  j                  j                  |	«       t        j                  d¬«      }
t        j                  d¬«      }|
|z  }||z  }t        j                  |||g||gt         j                  «      }t        j                  |||g||gt         j                  «      }||z   }t        j                  |	|||g«       y )Nr   é   )Údesc_ptrÚglobal_addressÚ	load_sizeÚglobal_sizeÚ
element_tyr   r)   rU   )r.   ÚextraÚcudaÚ&experimental_device_tensormap_create2dÚdtyper—   Ú)experimental_tensormap_fenceproxy_acquirer/   rw   rx   ry   )Ú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Ë  € ð ˆ
Ø ‘_ˆ
Ø ‘_ˆ
Ü
�‰�‰×<Ñ<ØØ Ø! :Ð.Ø˜A˜Ø—{‘{×-Ñ-ð 	=ô 	
ô 	�‰�‰×<Ñ<ØØ Ø! :Ð.Ø˜A˜Ø—{‘{×-Ñ-ð 	=ô 	
ô 	�‰�‰×<Ñ<ØØ Ø! :Ð.Ø˜A˜Ø—{‘{×-Ñ-ð 	=ô 	
ô 	�‰�‰×?Ñ?À
ÔKÜ
�‰�‰×?Ñ?À
ÔKÜ
�‰�‰×?Ñ?À
ÔKä—‘ 1Ô%ˆÜ—‘ 1Ô%ˆØ˜:Ñ%ˆØ˜:Ñ%ˆô ×,Ñ,ØØ�xÐ Ø˜Ð$Ü�J‰Jó	
ˆô ×,Ñ,ØØ�xÐ Ø˜Ð$Ü�J‰Jó	
ˆð �Q‘ˆô 	×)Ñ)ØØØ�xÐ õ	
r%   c                 óÆ  — t        j                  | ||g|dg||g¬«      }t        j                  |||g|dg||g¬«      }t        j                  |||g|dg||g¬«      }	t        j                  d¬«      }
t        j                  d¬«      }|
|z  }||z  }t        j                  |||g«      }t        j                  |||g«      }||z   }t        j                  |	||g|«       y )NrU   )ÚbaseÚshapeÚstridesÚblock_shaper   r)   )r.   Úmake_tensor_descriptorr/   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  € ô ×*Ñ*ØØ�a�&Ø˜�FØ# ZÐ0ô	
ˆô ×*Ñ*ØØ�a�&Ø˜�FØ# ZÐ0ô	
ˆô ×*Ñ*ØØ�a�&Ø˜�FØ# ZÐ0ô	
ˆô —‘ 1Ô%ˆÜ—‘ 1Ô%ˆØ˜:Ñ%ˆØ˜:Ñ%ˆô ×%Ñ%ØØ�xÐ ó
ˆô ×%Ñ%ØØ�xÐ ó
ˆð �Q‘ˆô 	×"Ñ"ØØ�xÐ Øõ	
r%   c                 óê   — t        j                  d¬«      }||z  }|t        j                  d|«      z   }||k  }t        j                  | |z   |¬«      }d|z  }	t        j                  ||z   |	|¬«       y ©Nr   r)   r+   r   r-   )
r4   r6   r7   r&   r8   r9   r:   r,   r;   r=   s
             r#   Úmul2_kernelr³   =  sn   € ô �m‰m Ô#ˆØ˜JÑ&ˆØ¤§	¡	¨!¨ZÓ 8Ñ8ˆØ˜Ñ#ˆÜ�G‰G�G˜gÑ%¨DÔ1ˆØ�Q‘ˆÜ
�‰�˜7Ñ" F°Ö6r%   c                 óê   — t        j                  d¬«      }||z  }|t        j                  d|«      z   }||k  }t        j                  | |z   |¬«      }d|z  }t        j                  | |z   ||¬«       y r²   r-   )	Úptrr7   r&   r8   r9   r:   r,   r;   r=   s	            r#   Úmul2_inplace_kernelr¶   L  sl   € ô �m‰m Ô#ˆØ˜JÑ&ˆØ¤§	¡	¨!¨ZÓ 8Ñ8ˆØ˜Ñ#ˆÜ�G‰G�C˜'‘M¨Ô-ˆØ�Q‘ˆÜ
�‰��w‘ ¨TÖ2r%   c                 ó6   — t        j                  | dk\  | d«      S )Nr   )r.   Úwhere)r;   s    r#   Ú	zero_negsr¹   Z  s   € ä�x‰x˜˜Q™  1Ó%Ð%r%   c                 ó2  — t        j                  d¬«      }||z  }|t        j                  d|«      z   }||k  }|dk(  rt        | ||¬«       n|dk(  rt	        | | |||¬«       t        j
                  | |z   |¬«      }	t        j                  ||z   |	|¬«       y )Nr   r)   r¶   )r&   r>   r+   )r.   r/   r0   r¶   r>   r1   r2   )
r4   r6   r7   r&   Ú
ACTIVATIONr8   r9   r:   r,   r;   s
             r#   Úindirection_kernelr¼   ^  s’   € ô �m‰m Ô#ˆØ˜JÑ&ˆØ¤§	¡	¨!¨ZÓ 8Ñ8ˆØ˜Ñ#ˆØÐ.Ò.Ü ¨À
ÖKØ˜<Ò'Ü�w ¨°*ÈÕTÜ�G‰G�G˜gÑ%¨DÔ1ˆÜ
�‰�˜7Ñ" A¨DÖ1r%   c                 ó”  — t        j                  d¬«      }t        j                  d¬«      }||z  }||z  }	|t        j                  d|«      z   }
|	t        j                  d|«      z   }|d d …d f   |z  |
d d d …f   z   }|d d …d f   |z  |
d d d …f   z   }t        j                  | |z   «      }t        j                  ||z   |dz  «       y )Nr   r)   rU   g       @r-   )Úin_ptrr6   Ú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Ì   q  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                 ó²  — t        j                  | t        j                  d|«      z   «      }t        j                  |t        j                  d|«      z   «      }t        j                  |g|t         j                  «      }t        j
                  dd|||gt         j                  dd¬«      }t        j                  |t        j                  d|«      z   |«       y )Nr   úshf.l.wrap.b32 $0, $1, $2, $3;ú
=r,r, r, rTrU   ©r›   Úis_pureÚpack©r.   r1   r0   ÚfullÚint32Úinline_asm_elementwiser2   ©	Ú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                 ó²  — t        j                  | t        j                  d|«      z   «      }t        j                  |t        j                  d|«      z   «      }t        j                  |g|t         j                  «      }t        j
                  dd|||gt         j                  dd¬«      }t        j                  |t        j                  d|«      z   |«       y )Nr   rÎ   rÏ   FrU   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           
      óŠ  — t        j                  d¬«      }||z  }t        j                  t        j                  | |gdg|g|gdg¬«      dg¬«      }t        j                  t        j                  ||gdg|g|gdg¬«      dg¬«      }||z   }	t        j                  t        j                  ||gdg|g|gdg¬«      |	dg¬«       y ©Nr   r)   rU   )r¨   r©   rª   r:   r«   Úorder)Úboundary_check©r.   r/   r1   Úmake_block_ptrr2   )
Úx_ptrÚy_ptrÚ
output_ptrr7   r&   r8   r9   r;   r<   r=   s
             r#   Úadd_kernel_with_block_ptrrê   §  sç   € ô �m‰m Ô#ˆØ˜JÑ&ˆÜ�G‰GÜ×ÑØØ!�lØ˜Ø$˜Ø'˜LØ�côð ˜3ô

ˆô �G‰GÜ×ÑØØ!�lØ˜Ø$˜Ø'˜LØ�côð ˜3ô

ˆð �Q‘ˆÜ
�‰Ü×ÑØØ!�lØ˜Ø$˜Ø'˜LØ�côð Ø˜3ö	
r%   c                 ó,  — t        j                  d¬«      }||z  }t        j                  t        j                  | |dgddg|dg|dgddg¬«      dg¬«      }|}t        j                  t        j                  ||dgddg|dg|dgddg¬«      |dg¬«       y râ   rå   )rç   ré   r7   r&   r8   r9   r;   r=   s           r#   Úkernel_with_block_ptr_2drì   Õ  s¸   € ô �m‰m Ô#ˆØ˜JÑ&ˆÜ�G‰GÜ×ÑØØ! 1�oØ˜A˜Ø$ aÐ(Ø'¨˜OØ˜!�fôð ˜3ô

ˆð ˆÜ
�‰Ü×ÑØØ! 1�oØ˜A˜Ø$ aÐ(Ø'¨˜OØ˜!�fôð Ø˜3ö	
r%   )r1   r2   c                 óâ   — t        j                  d¬«      }||z  }|t        j                  d|«      z   }||k  }t        | |z   |¬«      }	t        ||z   |¬«      }
|	|
z   }t	        ||z   ||¬«       y r(   r-   r3   s               r#   Úadd_kernel_with_importrî   ù  sw   € ô �m‰m Ô#ˆØ˜JÑ&ˆØ¤§	¡	¨!¨ZÓ 8Ñ8ˆØ˜Ñ#ˆÜ�˜7Ñ"¨Ô.ˆÜ�˜7Ñ"¨Ô.ˆØ�Q‘ˆÜˆg˜Ñ ¨dÖ3r%   c                 óZ  — t        j                  d¬«      }||z  }|t        j                  d|«      z   }||k  }t        j                  | |z   |¬«      }	t        j                  ||z   |¬«      }
t        j                  d«      dk(  r|	|
z   }n|	|
z  }t        j                  ||z   ||¬«       y r(   r-   r3   s               r#   Úcond_op_kernelrð   
  s�   € ô �m‰m Ô#ˆØ˜JÑ&ˆØ¤§	¡	¨!¨ZÓ 8Ñ8ˆØ˜Ñ#ˆÜ�G‰G�G˜gÑ%¨DÔ1ˆÜ�G‰G�G˜gÑ%¨DÔ1ˆÜ�=‰=˜Ó˜qÒ Ø˜‘U‰Fà˜‘UˆFÜ
�‰�˜7Ñ" F°Ö6r%   c                 ó  — t        j                  d¬«      }||z  }|t        j                  d|«      z   }||k  }t        j                  | |z   |¬«      }	t        j                  ||z   |¬«      }
|	|
z   }t        j                  ||z   ||¬«       y r(   )r.   r/   r0   r1   Ú
atomic_addr3   s               r#   Úatomic_add_kernelró     sƒ   € ô �m‰m Ô#ˆØ˜JÑ&ˆØ¤§	¡	¨!¨ZÓ 8Ñ8ˆØ˜Ñ#ˆÜ�G‰G�G˜gÑ%¨DÔ1ˆÜ�G‰G�G˜gÑ%¨DÔ1ˆØ�Q‘ˆÜ
�‰�g Ñ'¨°dÖ;r%   c                 ó¤  — t        j                  d¬«      }||z  }|t        j                  d|«      z   }||k  }t        j                  | |z   |¬«      }	t        j                  ||z   |¬«      }
t	        d«      D ]"  }|	|
z   }t        j
                  ||z   ||¬«       Œ$ d}|dkD  r,|dz  }|	|
z   }t        j
                  ||z   ||¬«       |dkD  rŒ+y y )Nr   r)   r+   r   rU   )r.   r/   r0   r1   Úranger2   )r4   r5   r6   r7   r&   r8   r9   r:   r,   r;   r<   rf   r=   Úis                 r#   Úadd_4_times_kernelr÷   /  sÔ   € ô �m‰m Ô#ˆØ˜JÑ&ˆØ¤§	¡	¨!¨ZÓ 8Ñ8ˆØ˜Ñ#ˆÜ�G‰G�G˜gÑ%¨DÔ1ˆÜ�G‰G�G˜gÑ%¨DÔ1ˆÜ�q“ò 	;ˆAØ˜‘UˆFÜ�H‰H�W˜wÑ&¨°TÖ:ð	;ð ˆØ�!ŠeØ�‰FˆAØ˜‘UˆFÜ�H‰H�W˜wÑ&¨°TÕ:ð �!�er%   c                 ó  — t        j                  d¬«      }||z  }|t        j                  d|«      z   }||k  }t        j                  | |z   |¬«      }	t        j                  ||z   |¬«      }
|	|
z   }t        j                  ||z   ||¬«       y r(   r-   )r4   r5   r7   r6   r&   r8   r9   r:   r,   r;   r<   r=   s               r#   Úadd_kernel_out_of_order_fn2rù   F  r?   r%   r   r   )ÚM_ptrÚNÚKr   r   r   r   c
                 óî  — t        j                  d¬«      }
t        j                  |«      }|dk(  r|dkD  rd}n|dk(  ry t        j                  ||«      }t        j                  ||«      }|	|z  }|
|z  }||	z  }t	        ||z
  |	«      }||
|z  |z  z   }|
|z  |z  }||z  t        j
                  d|«      z   |z  }||z  t        j
                  d|«      z   |z  }t        j
                  d|«      }| |d d …d f   |d d d …f   z   z   }||d d …d f   |d d d …f   z   z   }t        j                  ||ft         j                  ¬«      }t        t        j                  ||«      «      D ]s  }t        j                  ||d d d …f   |||z  z
  k  d¬«      }t        j                  ||d d …d f   |||z  z
  k  d¬«      }t        j                  |||«      }||z  }||z  }Œu |j                  t         j                  «      }||z  t        j
                  d|«      z   }||z  t        j
                  d|«      z   }||d d …d f   z   |d d d …f   z   } |d d …d f   |k  |d d d …f   |k  z  }!t        j                  | ||!¬«       y )	Nr   r)   r   i   ©r›   g        ©r,   Úotherr+   )r.   r/   r1   ÚcdivÚminr0   Úzerosrx   rõ   ÚdotÚtoÚfloat16r2   )"r�   rž   rŸ   rú   rû   rü   r   r   r   r   r8   Ú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  W  s’  € ôN �m‰m Ô#ˆÜ�G‰G�E‹NˆØ�Š6�l RÒ'à‰AØ�!ŠVàÜ—G‘G˜A˜|Ó,ˆ	Ü—G‘G˜A˜|Ó,ˆ	Ø'¨)Ñ3ÐØÐ*Ñ*ˆØ Ñ-ˆÜ˜9 {Ñ2°LÓAˆØ Ð&6Ñ 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 ¨lÐ;Ä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ð	#ð �N‰Nœ2Ÿ:™:Ó&ˆà˜,Ñ&¬¯©°1°lÓ)CÑCˆØ˜,Ñ&¬¯©°1°lÓ)CÑCˆØ˜¢ D Ñ)Ñ)¨G°Dº!°GÑ,<Ñ<ˆØš!˜T˜'Ñ" QÑ&¨7°4º°7Ñ+;¸aÑ+?Ñ@ˆÜ
�‰�˜ Ö(r%   c                 óò   — t        j                  d¬«      }t        j                  d|«      ||z  z   }t        j                  |gdt         j                  ¬«      }t        j
                  | |z   |||k  ¬«       y)z‰
        This kernel contains a triple-quote docstring w/ double quotes.
        Make sure that codegen sanitizes the docstring.
        r   r)   ç      ð?rþ   r+   N©r.   r/   r0   rÔ   rx   r2   ©r6   Únumelr&   r8   r:   Úoness         r#   Ú#kernel_with_docstring_double_quotesr#  ¤  s]   € ô �m‰m Ô#ˆÜ—)‘)˜A˜zÓ*¨S°:Ñ-=Ñ=ˆÜ�w‰w˜
�| S´·
±
Ô;ˆÜ
�‰�˜7Ñ" D¨w¸©Ö?r%   c                 óò   — t        j                  d¬«      }t        j                  d|«      ||z  z   }t        j                  |gdt         j                  ¬«      }t        j
                  | |z   |||k  ¬«       y)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%  ¯  s]   € ô �m‰m Ô#ˆÜ—)‘)˜A˜zÓ*¨S°:Ñ-=Ñ=ˆÜ�w‰w˜
�| S´·
±
Ô;ˆÜ
�‰�˜7Ñ" D¨w¸©Ö?r%   c                 óò   — t        j                  d¬«      }t        j                  d|«      ||z  z   }t        j                  |gdt         j                  ¬«      }t        j
                  | |z   |||k  ¬«       y)z/Docstring with literal backslash escapes: 
 	 \r   r)   r  rþ   r+   Nr  r   s         r#   Ú"kernel_with_backslash_in_docstringr'  »  s]   € ô �m‰m Ô#ˆÜ—)‘)˜A˜zÓ*¨S°:Ñ-=Ñ=ˆÜ�w‰w˜
�| S´·
±
Ô;ˆÜ
�‰�˜7Ñ" D¨w¸©Ö?r%   c                 ó2  — t        j                  d¬«      }t        j                  d|«      ||z  z   }t        j                  | |z   ||k  ¬«      }t        j                  dd|gt         j
                  dd¬«      }t        j                  ||z   |||k  ¬«       y )	Nr   r)   r+   z{
            {
                cos.approx.f32 $0, $1;
                ex2.approx.f32 $0, $0;
            }
                ú=r, rTrU   ©ÚasmÚconstraintsÚargsr›   rÑ   rÒ   ©r.   r/   r0   r1   rÖ   rx   r2   ©r¾   r6   r!  r&   r8   r:   ÚdataÚcos_pows           r#   Úkernel_inline_asm_double_quotesr2  Ã  sŒ   € ô �m‰m Ô#ˆÜ—)‘)˜A˜zÓ*¨S°:Ñ-=Ñ=ˆÜ�w‰w�v Ñ'¨g¸©oÔ>ˆÜ×+Ñ+ðð !Ø�Ü—*‘*ØØô
ˆô 	�‰�˜7Ñ" G°'¸E±/ÖBr%   c                 ó2  — t        j                  d¬«      }t        j                  d|«      ||z  z   }t        j                  | |z   ||k  ¬«      }t        j                  dd|gt         j
                  dd¬«      }t        j                  ||z   |||k  ¬«       y )	Nr   r)   r+   z»
            {
                // double quotes to pacify the linter """!!!"""
                cos.approx.f32 $0, $1;
                ex2.approx.f32 $0, $0;
            }
                r)  TrU   r*  r.  r/  s           r#   Úkernel_inline_asm_single_quotesr4  Ù  sŒ   € ô �m‰m Ô#ˆÜ—)‘)˜A˜zÓ*¨S°:Ñ-=Ñ=ˆÜ�w‰w�v Ñ'¨g¸©oÔ>ˆÜ×+Ñ+ðð !Ø�Ü—*‘*ØØô
ˆô 	�‰�˜7Ñ" G°'¸E±/ÖBr%   c                 ó2  — t        j                  d¬«      }t        j                  d|«      ||z  z   }t        j                  | |z   ||k  ¬«      }t        j                  dd|gt         j
                  dd¬«      }t        j                  ||z   |||k  ¬«       y ©	Nr   r)   r+   zK
            v_sin_f32 $0, $1
            v_exp_f32 $0, $0
                z=v, vTrU   r*  r.  r/  s           r#   Ú$kernel_inline_asm_rocm_double_quotesr7  ð  óŒ   € ô �m‰m Ô#ˆÜ—)‘)˜A˜zÓ*¨S°:Ñ-=Ñ=ˆÜ�w‰w�v Ñ'¨g¸©oÔ>ˆÜ×+Ñ+ðð !Ø�Ü—*‘*ØØô

ˆô 	�‰�˜7Ñ" G°'¸E±/ÖBr%   c                 ó2  — t        j                  d¬«      }t        j                  d|«      ||z  z   }t        j                  | |z   ||k  ¬«      }t        j                  dd|gt         j
                  dd¬«      }t        j                  ||z   |||k  ¬«       y r6  r.  r/  s           r#   Ú$kernel_inline_asm_rocm_single_quotesr:    r8  r%   c                 ó(  — t        j                  d¬«      }||z  }|t        j                  d|«      z   }||k  }	t        j                  | |z   |	¬«      }
|r t        j                  ||z   |	¬«      }|
|z   }n|
}t        j                  ||z   ||	¬«       y r(   r-   )r4   r5   r6   r7   Úadd_xyr&   r8   r9   r:   r,   r;   r<   r=   s                r#   Úadd_kernel_with_boolean_paramr=    s‹   € ô �m‰m Ô#ˆØ˜JÑ&ˆØ¤§	¡	¨!¨ZÓ 8Ñ8ˆØ˜Ñ#ˆÜ�G‰G�G˜gÑ%¨DÔ1ˆÙÜ—‘˜ 'Ñ)°Ô5ˆAØ˜‘U‰FàˆFÜ
�‰�˜7Ñ" F°Ö6r%   c                 ó„  — t        j                  d¬«      }||z  }|t        j                  d|«      z   }||k  }	t        j                  | |z   |	¬«      }
t        j                  ||z   |	¬«      }t        j                  ||z   |	d¬«      dk7  }t        j                  ||
|z   |
«      }t        j
                  ||z   ||	¬«       y)z6Kernel that loads a bool tensor and uses it as a mask.r   r)   r+   rÿ   N)r.   r/   r0   r1   r¸   r2   )r4   r5   Úmask_ptrr6   r7   r&   r8   r9   r:   Úvalidr;   r<   Úkeepr=   s                 r#   Ú"masked_add_kernel_with_bool_tensorrB  -  s¬   € ô �m‰m Ô#ˆØ˜JÑ&ˆØ¤§	¡	¨!¨ZÓ 8Ñ8ˆØ˜*Ñ$ˆÜ�G‰G�G˜gÑ%¨EÔ2ˆÜ�G‰G�G˜gÑ%¨EÔ2ˆÜ�w‰w�x 'Ñ)°¸QÔ?À1ÑDˆÜ—‘˜$  A¡ qÓ)ˆÜ
�‰�˜7Ñ" F°Ö7r%   Úblock_sizesÚnew_apic           	      ód  — |r4t         j                  j                  j                  j	                  | |«      S t        |«      dk(  rZt         j                  j                  j                  | j                  «       | j                  d«      |d   | j                  «       «      S t        |«      dk7  rt        dt        |«      › �«      ‚t         j                  j                  j                  | j                  «       | j                  d«      | j                  d«      |d   |d   | j                  «       «      S )NrU   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)ÚtensorrC  rD  s      r#   Úcreate_tensor_descriptor_shimrS  B  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)rE   rT  r&   rT  )rR   rT  rS   rT  )r&   rT  r»   rT  )rÁ   rT  rÂ   rT  )r¡   rT  rÛ   rT  )T)EÚunittestÚ&torch.testing._internal.inductor_utilsr   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>   rA   rF   rI   ÚautotunerL   rO   rQ   rc   rh   rp   rs   r�   r‰   rŽ   r�   r¦   r°   r³   r¶   r¹   r¼   rÌ   rÞ   rà   Ú	constexprrê   rì   Útriton.languager1   r2   rî   rð   ró   r÷   rù   r  r#  r%  r'  r2  r4  r7  r:  r=  rB  ÚintÚboolrS  re   r%   r#   ú<module>rd     s]	  ðó ÷ñ õ
 +ð /˜8×.Ñ.ØÐ3óÐ ð .˜(×-Ñ-ØÐ1óÐ ð .˜(×-Ñ-ØÒ-Ð-Ð/HóÐ ð #ˆx×"Ñ" 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ð@ð $ð@ð #ò@ó ð@ð* €V‡_�_àˆF�M‰M˜<¨Ð-¸!ÀqÔIØˆF�M‰M˜<¨Ð-¸!ÀqÔIØˆF�M‰M˜<¨Ð,¸ÀaÔHØˆF�M‰M˜<¨Ð,¸ÀaÔHð	
ð ôð ‡Z�Zð7ð
 #ò7ó óð7ð  €V‡_�_àˆF�M‰M˜<¨Ð-¸!ÀqÔIØˆF�M‰M˜<¨Ð-¸!ÀqÔIØˆF�M‰M˜<¨Ð,¸ÀaÔHØˆF�M‰M˜<¨Ð,¸ÀaÔHð	
ð ôð ‡Z�Zð7ð
 #ò7ó óð7ð  €V‡_�_àˆF�M‰M˜<¨Ð,¸ÀaÔHð
ð ô	ð ‡Z�Zð7ð #ò	7ó óð7ð$ €V‡_�_àˆF�M‰MØ!$°cÑ:ÀqÐTUôð ˆF�M‰MØ!$°cÑ:ÀqÐTUôð ˆF�M‰MØ!#°RÑ8ÀQÐRSôð ˆF�M‰MØ!#°RÑ8ÀQÐRSôð
ð ôð" ‡Z�ZðJð %ðJð %òJó ó#ð$Jò,ð €V‡_�_àˆF�M‰M˜<¨Ð-¸!ÀqÔIØˆF�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ð  €V‡_�_àˆF�M‰Mà$&Ø$&Ø$&Ø$%ñ	ð Øô	ð ˆF�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%   