+
    ˜°³ji‡  ã                   ó   € ^ RI t ^ RIHtHtHtHt ^ RIHt ] P                  ! ]R4      t	] P                  ! ]R4      t
] P                  ! ];'       g    ];'       g    ]R4      t] P                  ! ]R4      t]! 4       '       EdŠ   ^ RIt^ RIHt ^ RItR	 R
 lt]P$                  R R l4       t]P$                  R R l4       t]P$                  R R l4       t]P$                  R R l4       t]P.                  ! ]P0                  ! R^€/^^R7      ]P0                  ! R^€/^^R7      ]P0                  ! R^@/^^R7      ]P0                  ! R^@/^^R7      .. R7      ]P$                  R R l4       4       t]P.                  ! ]P0                  ! R^€/^^R7      ]P0                  ! R^€/^^R7      ]P0                  ! R^@/^^R7      ]P0                  ! R^@/^^R7      .. R7      ]P$                  R R l4       4       t]P.                  ! ]P0                  ! R^/^^R7      .. R7      ]P$                  R R l4       4       t]P.                  ! ]P0                  ! R^€R^€/^^R7      ]P0                  ! R^€R^€/^^R7      ]P0                  ! R^@R^@/^^R7      ]P0                  ! R^@R^@/^^R7      .. R7      ]P$                  R R l4       4       tR  t]P.                  ! ]P0                  ! R^€/^^R7      ]P0                  ! R^@/^^R7      .. ^
^R!]/R"7      ]P$                  R# R$ l4       4       t]P$                  R% R& l4       t]P$                  R' R( l4       t ]P$                  R) R* l4       t!]P$                  R+ R, l4       t"]P$                  R- R. l4       t#]P$                  R/ R0 l4       t$]P$                  R1 R2 l4       t%]P$                  R3 R4 l4       t&]P$                  R5 R6 l4       t']P$                  R7 4       t(]P$                  R8 R9 l4       t)]P$                  R: R; l4       t*]P$                  R< R= l4       t+]P$                  R> R? l4       t,]P$                  R@ RA l4       t-]P$                  RB RC l4       t.^ RDI/H0t0H1t1 ]P$                  RE RF l4       t2]P$                  RG RH l4       t3]P$                  RI RJ l4       t4]P$                  RK RL l4       t5]P$                  RM RN l4       t6]P.                  ! ]P0                  ! RO^RP^RQ^RR^/^^R7      ]P0                  ! RO^€RP^@RQ^ RR^/^^R7      .. RiOR7      ]P$                  RS RT l4       4       t7]P$                  RU RV l4       t8]P$                  RW RX l4       t9]P$                  RY RZ l4       t:]P$                  R[ R\ l4       t;]P$                  R] R^ l4       t<]P$                  R_ R` l4       t=]P$                  Ra Rb l4       t>]P$                  Rc Rd l4       t?]P$                  Re Rf l4       t@RjRg Rh lltAR# R# )ké    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)Úlanguagec                óN   € V ^8„  d   QhR\         \        P                  ,          /# )é   Úreturn)ÚlistÚtritonÚConfig)Úformats   "Úe/home/luca/screen-checking/.venv/lib/python3.14/site-packages/torch/testing/_internal/triton_utils.pyÚ__annotate__r       s   € ÷ 5ñ 5¤$¤v§}¡}Õ"5ñ 5ó    c                  ój  € \         P                  P                  '       dQ   \        P                  ! R ^R^R^R^R^R^R^/^^R7      \        P                  ! R ^€R^@R^R^R^R^R^/^^R7      .p V # \        P                  ! R ^R^R^R^/^^R7      \        P                  ! R ^€R^@R^ R^/^^R7      .p V # )ÚBLOCK_SIZE_MÚBLOCK_SIZE_NÚBLOCK_SIZE_KÚGROUP_SIZE_MÚmatrix_instr_nonkdimÚwaves_per_euÚkpack©Ú
num_stagesÚ	num_warps)ÚtorchÚversionÚhipr   r   )Úconfigss    r   Ú_get_strange_configsr"       s  € Ü�=‰=××Ðä—’à&¨Ø&¨Ø&¨Ø&¨Ø.°Ø&¨Ø ðð  !Øôô —’à&¨Ø&¨Ø&¨Ø&¨Ø.°Ø&¨Ø ðð  !ØôðˆGðf ˆô+ —’à&¨Ø&¨Ø&¨Ø&¨ð	ð  !Øô	ô —’à&¨Ø&¨Ø&¨Ø&¨ð	ð  !Øô	ðˆGð, ˆr   c                ó   € V ^8„  d   QhRR/# ©r
   Ú
BLOCK_SIZEútl.constexpr© )r   s   "r   r   r   Y   ó   € ÷ 7ñ 7ð
 #ñ7r   c                 óL  € \         P                  ! ^ R7      pWT,          pV\         P                  ! ^ V4      ,           pWs8  p\         P                  ! W,           VR7      p	\         P                  ! W,           VR7      p
Wš,           p\         P                  ! W',           W¸R7       R# ©r   ©Úaxis©ÚmaskN©Ú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                ó   € V ^8„  d   QhRR/# r$   r'   )r   s   "r   r   r   j   r(   r   c                 óL  € \         P                  ! ^ R7      pWT,          pV\         P                  ! ^ V4      ,           pWs8  p\         P                  ! W,           VR7      p	\         P                  ! W,           VR7      p
Wš,
          p\         P                  ! W',           W¸R7       R# r*   r/   r5   s   &&&&&       r   Ú
sub_kernelrD   i   rA   r   c                ó    € V ^8„  d   QhRRRR/# ©r
   ÚARGS_PASSEDr&   r%   r'   )r   s   "r   r   r   {   s    € ÷ 7ñ 7ð
 $ð7ð #ñ7r   c                 ó`  € \         P                  ! ^ R7      pWe,          pV\         P                  ! ^ V4      ,           pWƒ8  p	\         P                  ! W,           V	R7      p
VR8X  d(   \         P                  ! W,           V	R7      pW«,           pMT
p\         P                  ! W(,           WÉR7       R# ©r   r+   r-   ÚtwoNr/   )r6   r7   r8   r9   rG   r%   r:   r;   r<   r.   r=   r>   r?   s   &&&&&&       r   Úadd_kernel_with_optional_paramrK   z   s‚   € ô �mŠm Ô#ˆØÕ&ˆØ¤§	¢	¨!¨ZÓ 8Õ8ˆØÑ#ˆÜ�GŠG�GÕ%¨DÔ1ˆØ˜%ÔÜ—’˜Õ)°Ô5ˆAØ•U‰FàˆFÜ
�Š�Õ" F×6r   c                ó    € V ^8„  d   QhRRRR/# rF   r'   )r   s   "r   r   r   �   s$   € ÷ @ñ @ð $ð@ð #ñ@r   c                 ó|  € \         P                  ! ^ R7      pWv,          pV\         P                  ! ^ V4      ,           p	W“8  p
\         P                  ! W	V,          ,           V
R7      pVR8X  d(   \         P                  ! W,           V
R7      pW¼,           pMTp\         P                  ! W)V,          ,           WÚR7       R# rI   r/   )r6   r7   r8   r9   ÚstriderG   r%   r:   r;   r<   r.   r=   r>   r?   s   &&&&&&&       r   Ú-add_kernel_with_none_param_and_equal_to_1_argrO   �   sŒ   € ô �mŠm Ô#ˆØÕ&ˆØ¤§	¢	¨!¨ZÓ 8Õ8ˆØÑ#ˆÜ�GŠG�G¨Õ.Õ.°TÔ:ˆØ˜%ÔÜ—’˜Õ)°Ô5ˆAØ•U‰FàˆFÜ
�Š� VÕ+Õ+¨V×?r   r%   r   )r!   Úkeyc                ó   € V ^8„  d   QhRR/# r$   r'   )r   s   "r   r   r   ¯   r(   r   c                 óL  € \         P                  ! ^ R7      pWT,          pV\         P                  ! ^ V4      ,           pWs8  p\         P                  ! W,           VR7      p	\         P                  ! W,           VR7      p
Wš,           p\         P                  ! W',           W¸R7       R# r*   r/   r5   s   &&&&&       r   Úadd_kernel_autotunedrS   ¥   óu   € ô" �mŠm Ô#ˆØÕ&ˆØ¤§	¢	¨!¨ZÓ 8Õ8ˆØÑ#ˆÜ�GŠG�GÕ%¨DÔ1ˆÜ�GŠG�GÕ%¨DÔ1ˆØ•ˆÜ
�Š�Õ" F×6r   c                ó   € V ^8„  d   QhRR/# r$   r'   )r   s   "r   r   r   É   r(   r   c                 óL  € \         P                  ! ^ R7      pWT,          pV\         P                  ! ^ V4      ,           pWs8  p\         P                  ! W,           VR7      p	\         P                  ! W,           VR7      p
Wš,
          p\         P                  ! W',           W¸R7       R# r*   r/   r5   s   &&&&&       r   Úsub_kernel_autotunedrW   ¿   rT   r   c                ó   € V ^8„  d   QhRR/# r$   r'   )r   s   "r   r   r   à   s   € ÷ 7ñ 7ð #ñ	7r   c                 óL  € \         P                  ! ^ R7      pWS,          pV\         P                  ! ^ V4      ,           pWr8  p\         P                  ! W,           VR7      p	\         P                  ! W,           VR7      p
Wš,           p\         P                  ! WG,           W¸R7       R# r*   r/   )r6   r7   r9   r%   r8   r:   r;   r<   r.   r=   r>   r?   s   &&&&&       r   Ú&add_kernel_autotuned_weird_param_orderrZ   Ù   su   € ô  �mŠm Ô#ˆØÕ&ˆØ¤§	¢	¨!¨ZÓ 8Õ8ˆØÑ#ˆÜ�GŠG�GÕ%¨DÔ1ˆÜ�GŠG�GÕ%¨DÔ1ˆØ•ˆÜ
�Š�Õ" F×6r   ÚBLOCK_SIZE_XÚBLOCK_SIZE_Yc                ó    € V ^8„  d   QhRRRR/# ©r
   r[   r&   r\   r'   )r   s   "r   r   r     s$   € ÷ Jñ Jð %ðJð %ñJr   c                 óZ  € \         P                  ! ^ 4      V,          pV\         P                  ! ^ V4      R,          ,           pWƒ8  p	\         P                  ! ^4      V,          p
V
\         P                  ! ^ V4      R,          ,           pW´8  pTpTp\         P                  ! WW>,          ,           ,           Wœ,          4      p\         P                  ! WWM,          ,           ,           Wœ,          4      pVV,           p\         P                  ! W-W>,          ,           ,           VWœ,          4       R# )r   N©ºNNNN©Nra   r/   )r6   r7   r8   Ú
x_elementsÚ
y_elementsr[   r\   ÚxoffsetÚxindexÚxmaskÚyoffsetÚyindexÚymaskÚx1Úy0Útmp0Útmp1Útmp2s   &&&&&&&           r   Úadd_kernel_2d_autotunedrp   ò   sÕ   € ô6 —-’- Ó" \Õ1ˆØœ2Ÿ9š9 Q¨Ó5°gÕ>Õ>ˆØÑ#ˆÜ—-’- Ó" \Õ1ˆØœ2Ÿ9š9 Q¨Ó5°gÕ>Õ>ˆØÑ#ˆØˆØˆÜ�wŠw�w¨
­Õ"8Õ9¸5½=ÓIˆÜ�wŠw�w¨
­Õ"8Õ9¸5½=ÓIˆØ�d�{ˆÜ
�Š� *¥/Õ2Õ3°T¸5½=ÖIr   c                 ó   € V # )Nr'   )r!   Ú_Ú__s   &*,r   Ú_dummy_early_config_prunert     s   € Øˆr   Úearly_config_prune)r!   rP   ÚwarmupÚrepÚprune_configs_byc                ó   € V ^8„  d   QhRR/# r$   r'   )r   s   "r   r   r   (  r(   r   c                 óL  € \         P                  ! ^ R7      pWT,          pV\         P                  ! ^ V4      ,           pWs8  p\         P                  ! W,           VR7      p	\         P                  ! W,           VR7      p
Wš,           p\         P                  ! W',           W¸R7       R# r*   r/   r5   s   &&&&&       r   Ú*add_kernel_autotuned_with_unsupported_argsr{     su   € ô$ �mŠm Ô#ˆØÕ&ˆØ¤§	¢	¨!¨ZÓ 8Õ8ˆØÑ#ˆÜ�GŠG�GÕ%¨DÔ1ˆÜ�GŠG�GÕ%¨DÔ1ˆØ•ˆÜ
�Š�Õ" F×6r   c                ó   € V ^8„  d   QhRR/# r$   r'   )r   s   "r   r   r   9  s   € ÷ 7ñ 7ð #ñ7r   c                 óZ  € \         P                  ! ^ R7      pWe,          pV\         P                  ! ^ V4      ,           pWƒ8  p	\         P                  ! W,           V	R7      p
\         P                  ! W,           V	R7      pW«,           V,          p\         P                  ! W(,           WÉR7       R# r*   r/   )r6   r7   r8   r9   Úscaling_factorr%   r:   r;   r<   r.   r=   r>   r?   s   &&&&&&       r   Úadd_kernel_with_scalingr   8  sz   € ô �mŠm Ô#ˆØÕ&ˆØ¤§	¢	¨!¨ZÓ 8Õ8ˆØÑ#ˆÜ�GŠG�GÕ%¨DÔ1ˆÜ�GŠG�GÕ%¨DÔ1ˆØ•%˜>Õ)ˆÜ
�Š�Õ" F×6r   c                ó   € V ^8„  d   QhRR/# r$   r'   )r   s   "r   r   r   K  s   € ÷ 
ñ 
ð #ñ	
r   c                 ó*  € \         P                  ! ^ R7      pWC,          p\         P                  ! V V.V.\         P                  4      p\         P                  ! VV.V.\         P                  4      pWg,           p\         P                  ! VVV.4       R# ©r   r+   N©r0   r1   Ú_experimental_descriptor_loadÚfloat32Ú_experimental_descriptor_store©	Úin_desc_ptr0Úin_desc_ptr1Úout_desc_ptrr%   r:   Ú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                ó    € V ^8„  d   QhRRRR/# r^   r'   )r   s   "r   r   r   j  s    € ÷ 
ñ 
ð %ð	
ð
 %ñ
r   c                 óh  € \         P                  ! ^ R7      p\         P                  ! ^R7      pWS,          pWd,          p\         P                  ! V Wx.W4.\         P                  4      p	\         P                  ! VWx.W4.\         P                  4      p
Wš,           p\         P                  ! VVWx.4       R# r‚   rƒ   ©rˆ   r‰   rŠ   r[   r\   Ú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                ó   € V ^8„  d   QhRR/# r$   r'   )r   s   "r   r   r   Œ  s   € ÷ 
ñ 
ð #ñ	
r   c                 óæ   € \         P                  ! ^ R7      pWC,          p\         P                  ! V V.4      p\         P                  ! VV.4      pWg,           p\         P                  ! VV.V4       R# r‚   ©r0   r1   Úload_tensor_descriptorÚstore_tensor_descriptorr‡   s	   &&&&     r   Úadd_kernel_with_tma_1d_new_apirœ   ‹  sn   € ô �mŠm Ô#ˆØÕ!ˆä×%Ò%ØØˆHó
ˆô ×%Ò%ØØˆHó
ˆð
 •ˆä
×"Ò"ØØˆHØö	
r   c                ó    € V ^8„  d   QhRRRR/# r^   r'   )r   s   "r   r   r   §  s    € ÷ 
ñ 
ð %ð	
ð
 %ñ
r   c                 ó$  € \         P                  ! ^ R7      p\         P                  ! ^R7      pWS,          pWd,          p\         P                  ! V Wx.4      p	\         P                  ! VWx.4      p
Wš,           p\         P                  ! VWx.V4       R# r‚   r™   r‘   s   &&&&&       r   Úadd_kernel_with_tma_2d_new_apirŸ   ¦  s‡   € ô —’ 1Ô%ˆÜ—’ 1Ô%ˆØÕ'ˆØÕ'ˆä×%Ò%ØØÐ ó
ˆô ×%Ò%ØØÐ ó
ˆð
 •ˆä
×"Ò"ØØÐ Øö	
r   c                ó   € V ^8„  d   QhRR/# r$   r'   )r   s   "r   r   r   Å  s   € ÷ A
ñ A
ð #ñA
r   c                 ó  € TpV^€,           pVR,           p	\         P                  P                  P                  VV Wf.W4.V P                  P
                  R7       \         P                  P                  P                  VVWf.W4.VP                  P
                  R7       \         P                  P                  P                  V	VWf.W4.VP                  P
                  R7       \         P                  P                  P                  V4       \         P                  P                  P                  V4       \         P                  P                  P                  V	4       \         P                  ! ^ R7      p
\         P                  ! ^R7      pW¦,          pW¶,          p\         P                  ! VWÍ.Wf.\         P                  4      p\         P                  ! VWÍ.Wf.\         P                  4      pWï,           p\         P                  ! V	VWÍ.4       R# )é€   é   )Údesc_ptrÚglobal_addressÚ	load_sizeÚglobal_sizeÚ
element_tyr+   N)r0   ÚextraÚcudaÚ&experimental_device_tensormap_create2dÚdtyper¨   Ú)experimental_tensormap_fenceproxy_acquirer1   r„   r…   r†   )Ú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                ó   € V ^8„  d   QhRR/# r$   r'   )r   s   "r   r   r   	  s   € ÷ 4
ñ 4
ð #ñ4
r   c                 óØ  € \         P                  ! V W4.V^.Wf.R7      p\         P                  ! VW4.V^.Wf.R7      p\         P                  ! VW4.V^.Wf.R7      p	\         P                  ! ^ R7      p
\         P                  ! ^R7      pW¦,          pW¶,          p\         P                  ! VWÍ.4      p\         P                  ! VWÍ.4      pWï,           p\         P                  ! V	WÍ.V4       R# )é   )ÚbaseÚshapeÚstridesÚblock_shaper+   N)r0   Úmake_tensor_descriptorr1   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                ó   € V ^8„  d   QhRR/# r$   r'   )r   s   "r   r   r   @  s   € ÷ 7ñ 7ð #ñ	7r   c                 ó  € \         P                  ! ^ R7      pWC,          pV\         P                  ! ^ V4      ,           pWb8  p\         P                  ! W,           VR7      p^V,          p	\         P                  ! W,           W—R7       R# r*   r/   )
r6   r8   r9   r%   r:   r;   r<   r.   r=   r?   s
   &&&&      r   Úmul2_kernelrÆ   ?  sd   € ô �mŠm Ô#ˆØÕ&ˆØ¤§	¢	¨!¨ZÓ 8Õ8ˆØÑ#ˆÜ�GŠG�GÕ%¨DÔ1ˆØ�Q•ˆÜ
�Š�Õ" F×6r   c                ó   € V ^8„  d   QhRR/# r$   r'   )r   s   "r   r   r   O  s   € ÷ 3ñ 3ð #ñ3r   c                 ó  € \         P                  ! ^ R7      pW2,          pV\         P                  ! ^ V4      ,           pWQ8  p\         P                  ! W,           VR7      p^V,          p\         P                  ! W,           W†R7       R# r*   r/   )	Úptrr9   r%   r:   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   € \         P                  ! V ^ 8¬  V ^ 4      # )r   )r0   Úwhere)r=   s   &r   Ú	zero_negsrÍ   \  s   € ä�xŠx˜˜Q™  1Ó%Ð%r   c                ó    € V ^8„  d   QhRRRR/# )r
   r%   r&   Ú
ACTIVATIONr'   )r   s   "r   r   r   a  s    € ÷ 2ñ 2ð #ð	2ð
 #ñ2r   c                 óT  € \         P                  ! ^ R7      pWS,          pV\         P                  ! ^ V4      ,           pWr8  pVR8X  d   \        WVR7       MVR8X  d   \	        W WVR7       \         P
                  ! W,           VR7      p	\         P                  ! W,           W˜R7       R# )r   r+   rÊ   )r%   r@   r-   N)r0   r1   r2   rÊ   r@   r3   r4   )
r6   r8   r9   r%   rÏ   r:   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                ó    € V ^8„  d   QhRRRR/# )r
   ÚX_BLOCK_SIZEr&   ÚY_BLOCK_SIZEr'   )r   s   "r   r   r   t  s    € ÷ 3ñ 3ð
 %ð3ð %ñ3r   c                 óò  € \         P                  ! ^ R7      p\         P                  ! ^R7      pWd,          pWu,          p	V\         P                  ! ^ V4      ,           p
V	\         P                  ! ^ V4      ,           pVR,          V,          V
R,          ,           pVR,          V,          V
R,          ,           p\         P                  ! W,           4      p\         P                  ! W,           VR,          4       R# )r   r+   Ng       @r`   rb   r/   )Úin_ptrr8   Úin_y_strideÚout_y_striderÓ   rÔ   Ú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ˆ	Ø Õ(¨;Õ6¸À7Õ9KÕKˆØ Õ(¨<Õ7¸)ÀGÕ:LÕLˆÜ�gŠg�fÕ*Ó+ˆÜ
�Š�Õ&¨¨c­	Ö2r   c                ó    € V ^8„  d   QhRRRR/# ©r
   r²   r&   ÚBLOCKr'   )r   s   "r   r   r   ˆ  ó   € ÷ -ñ -Ø"ð-Ø+9ñ-r   c           	      óÚ  € \         P                  ! V \         P                  ! ^ V4      ,           4      p\         P                  ! V\         P                  ! ^ V4      ,           4      p\         P                  ! V.V\         P                  4      p\         P
                  ! RRWVV.\         P                  R^R7      p\         P                  ! V\         P                  ! ^ V4      ,           V4       R# )r   úshf.l.wrap.b32 $0, $1, $2, $3;ú
=r,r, r, rT©r¬   Úis_pureÚpackN©r0   r3   r2   ÚfullÚint32Úinline_asm_elementwiser4   ©	ÚXÚYÚZr²   rå   r=   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                ó    € V ^8„  d   QhRRRR/# rä   r'   )r   s   "r   r   r   ™  ræ   r   c           	      óÚ  € \         P                  ! V \         P                  ! ^ V4      ,           4      p\         P                  ! V\         P                  ! ^ V4      ,           4      p\         P                  ! V.V\         P                  4      p\         P
                  ! RRWVV.\         P                  R^R7      p\         P                  ! V\         P                  ! ^ V4      ,           V4       R# )r   rè   ré   Frê   Nrí   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                ó8   € V ^8„  d   QhR\         P                  /# ©r
   r%   ©r0   Ú	constexpr)r   s   "r   r   r   ª  s   € ÷ +
ñ +
ô
 —L‘Lñ+
r   c                 ó¦  € \         P                  ! ^ R7      pWT,          p\         P                  ! \         P                  ! V V.^.V.V.^ .R7      ^ .R7      p\         P                  ! \         P                  ! VV.^.V.V.^ .R7      ^ .R7      pWx,           p	\         P                  ! \         P                  ! VV.^.V.V.^ .R7      V	^ .R7       R# ©r   r+   )r»   r¼   r½   r<   r¾   Úorder)Úboundary_checkN©r0   r1   r3   Úmake_block_ptrr4   )
Úx_ptrÚy_ptrÚ
output_ptrr9   r%   r:   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                ó8   € V ^8„  d   QhR\         P                  /# rü   rý   )r   s   "r   r   r   Ø  s   € ÷ 
ñ 
ô —L‘Lñ	
r   c                 ó>  € \         P                  ! ^ R7      pWC,          p\         P                  ! \         P                  ! V V^.^^.V^ .V^.^^ .R7      ^ .R7      pTp\         P                  ! \         P                  ! VV^.^^.V^ .V^.^^ .R7      V^ .R7       R# r   r  )r  r  r9   r%   r:   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   )r3   r4   c                ó   € V ^8„  d   QhRR/# r$   r'   )r   s   "r   r   r   ü  s   € ÷ 4ñ 4ð
 #ñ4r   c                 ó
  € \         P                  ! ^ R7      pWT,          pV\         P                  ! ^ V4      ,           pWs8  p\        W,           VR7      p	\        W,           VR7      p
Wš,           p\	        W',           W¸R7       R# r*   r/   r5   s   &&&&&       r   Úadd_kernel_with_importr  û  si   € ô �mŠm Ô#ˆØÕ&ˆØ¤§	¢	¨!¨ZÓ 8Õ8ˆØÑ#ˆÜ�Õ"¨Ô.ˆÜ�Õ"¨Ô.ˆØ•ˆÜˆgÕ ×3r   c                ó   € V ^8„  d   QhRR/# r$   r'   )r   s   "r   r   r     s   € ÷ 7ñ 7ð
 #ñ7r   c                 ó”  € \         P                  ! ^ R7      pWT,          pV\         P                  ! ^ V4      ,           pWs8  p\         P                  ! W,           VR7      p	\         P                  ! W,           VR7      p
\         P                  ! ^ 4      ^ 8X  d
   Wš,           pMWš,          p\         P                  ! W',           W¸R7       R# r*   r/   r5   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                ó   € V ^8„  d   QhRR/# r$   r'   )r   s   "r   r   r   !  s   € ÷ <ñ <ð
 #ñ<r   c                 óL  € \         P                  ! ^ R7      pWT,          pV\         P                  ! ^ V4      ,           pWs8  p\         P                  ! W,           VR7      p	\         P                  ! W,           VR7      p
Wš,           p\         P                  ! W',           W¸R7       R# r*   )r0   r1   r2   r3   Ú
atomic_addr5   s   &&&&&       r   Úatomic_add_kernelr     su   € ô �mŠm Ô#ˆØÕ&ˆØ¤§	¢	¨!¨ZÓ 8Õ8ˆØÑ#ˆÜ�GŠG�GÕ%¨DÔ1ˆÜ�GŠG�GÕ%¨DÔ1ˆØ•ˆÜ
�Š�gÕ'¨×;r   c                ó   € V ^8„  d   QhRR/# r$   r'   )r   s   "r   r   r   2  s   € ÷ ;ñ ;ð
 #ñ;r   c                 óä  € \         P                  ! ^ R7      pWT,          pV\         P                  ! ^ V4      ,           pWs8  p\         P                  ! W,           VR7      p	\         P                  ! W,           VR7      p
\	        ^4       F)  pWš,           p\         P
                  ! W',           WÈR7       K+  	  ^pV^ 8”  d2   V^,          pWš,           p\         P
                  ! W',           WÈR7       K8  R# r*   )r0   r1   r2   r3   Úranger4   )r6   r7   r8   r9   r%   r:   r;   r<   r.   r=   r>   rr   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Õ&¨×:ñ r   c                ó   € V ^8„  d   QhRR/# r$   r'   )r   s   "r   r   r   I  r(   r   c                 óL  € \         P                  ! ^ R7      pWT,          pV\         P                  ! ^ V4      ,           pWr8  p\         P                  ! W,           VR7      p	\         P                  ! W,           VR7      p
Wš,           p\         P                  ! W7,           W¸R7       R# r*   r/   )r6   r7   r9   r8   r%   r:   r;   r<   r.   r=   r>   r?   s   &&&&&       r   Úadd_kernel_out_of_order_fn2r  H  rA   r   r   r   r   r   c                ó˜   € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\         P                  /# )r
   r   r   r   r   rý   )r   s   "r   r   r   s  sD   € ÷ 1)ñ 1)ô —l‘lð1)ô —l‘lð1)ô —l‘lð1)ô —l‘lñ1)r   c
                 ó  € \         P                  ! ^ R7      p
\         P                  ! V4      pV^ 8X  d   V^ 8”  d   RpM	V^ 8X  d   R# \         P                  ! W¶4      p\         P                  ! WG4      pW�,          pW®,          pWù,          p\	        VV,
          V	4      pVW®,          V,          ,           pW®,          V,          pVV,          \         P
                  ! ^ V4      ,           V,          pVV,          \         P
                  ! ^ V4      ,           V,          p\         P
                  ! ^ V4      pV VR,          VR	,          ,           ,           pVVR,          VR	,          ,           ,           p\         P                  ! Wg3\         P                  R7      p\        \         P                  ! WX4      4       F�  p\         P                  ! VVR	,          VVV,          ,
          8  RR7      p\         P                  ! VVR,          VVV,          ,
          8  RR7      p\         P                  ! VVV4      pVV,          pVV,          pK‘  	  VP                  \         P                  4      pVV,          \         P
                  ! ^ V4      ,           pVV,          \         P
                  ! ^ V4      ,           pVVR,          ,           VR	,          ,           p VR,          V8  VR	,          V8  ,          p!\         P                  ! V VV!R7       R# )
r   r+   i   N©r¬   g        ©r.   Úotherr-   r`   rb   )r0   r1   r3   ÚcdivÚminr2   Úzerosr…   r  ÚdotÚtoÚfloat16r4   )"r®   r¯   r°   ÚM_ptrÚNÚKr   r   r   r   r:   Ú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_kernelrA  Y  sP  € ô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Ó+ˆØ˜' 'Õ*¨V°G­_Õ<Õ=ˆØ˜& �/¨G°GÕ,<Õ<Õ=ˆä—h’h Ð;Ä2Ç:Á:ÔNˆÜ”r—w’w˜qÓ/Ö0ˆAÜ—’˜ V¨G¥_°q¸1¸|Õ;KÕ7KÑ%KÐSVÔWˆAÜ—’˜ V¨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ˆØ˜ Õ)Õ)¨G°GÕ,<Õ<ˆØ˜'Õ" QÑ&¨7°7Õ+;¸aÑ+?Õ@ˆÜ
�Š�˜ ×(r   c                ó8   € V ^8„  d   QhR\         P                  /# rü   rý   )r   s   "r   r   r   §  s   € ÷ @ñ @ÌÏÉñ @r   c                ó  € \         P                  ! ^ R7      p\         P                  ! ^ V4      W2,          ,           p\         P                  ! V.R\         P                  R7      p\         P
                  ! W,           WTV8  R7       R# )zq
This kernel contains a triple-quote docstring w/ double quotes.
Make sure that codegen sanitizes the docstring.
r+   ç      ð?r   r-   N©r0   r1   r2   rî   r…   r4   ©r8   Únumelr%   r:   r<   Úoness   &&&   r   Ú#kernel_with_docstring_double_quotesrI  ¦  sW   € ô �mŠm Ô#ˆÜ—)’)˜A˜zÓ*¨SÕ-=Õ=ˆÜ�wŠw˜
�| S´·
±
Ô;ˆÜ
�Š�Õ" D¸©×?r   c                ó8   € V ^8„  d   QhR\         P                  /# rü   rý   )r   s   "r   r   r   ²  s   € ÷ 	@ñ 	@ÌÏÉñ 	@r   c                ó  € \         P                  ! ^ R7      p\         P                  ! ^ V4      W2,          ,           p\         P                  ! V.R\         P                  R7      p\         P
                  ! W,           WTV8  R7       R# )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+   rD  r   r-   NrE  rF  s   &&&   r   Ú#kernel_with_docstring_single_quotesrL  ±  sW   € ô �mŠm Ô#ˆÜ—)’)˜A˜zÓ*¨SÕ-=Õ=ˆÜ�wŠw˜
�| S´·
±
Ô;ˆÜ
�Š�Õ" D¸©×?r   c                ó8   € V ^8„  d   QhR\         P                  /# rü   rý   )r   s   "r   r   r   ¾  s   € ÷ @ñ @ÄrÇ|Á|ñ @r   c                ó  € \         P                  ! ^ R7      p\         P                  ! ^ V4      W2,          ,           p\         P                  ! V.R\         P                  R7      p\         P
                  ! W,           WTV8  R7       R# )z,Docstring with literal backslash escapes: 
\r+   rD  r   r-   NrE  rF  s   &&&   r   Ú"kernel_with_backslash_in_docstringrO  ½  sW   € ô �mŠm Ô#ˆÜ—)’)˜A˜zÓ*¨SÕ-=Õ=ˆÜ�wŠw˜
�| S´·
±
Ô;ˆÜ
�Š�Õ" D¸©×?r   c                ó8   € V ^8„  d   QhR\         P                  /# rü   rý   )r   s   "r   r   r   Æ  s   € ÷ Cñ CÜ,.¯L©LñCr   c           	      óT  € \         P                  ! ^ R7      p\         P                  ! ^ V4      WC,          ,           p\         P                  ! W,           WR8  R7      p\         P                  ! RRV.\         P
                  R^R7      p\         P                  ! W,           WuV8  R7       R# )r   r+   r-   z{
            {
                cos.approx.f32 $0, $1;
                ex2.approx.f32 $0, $0;
            }
                ú=r, rT©ÚasmÚconstraintsÚargsr¬   rë   rì   N©r0   r1   r2   r3   rð   r…   r4   ©rÖ   r8   rG  r%   r:   r<   ÚdataÚcos_pows   &&&&    r   Úkernel_inline_asm_double_quotesr[  Å  s‚   € ô �mŠm Ô#ˆÜ—)’)˜A˜zÓ*¨SÕ-=Õ=ˆÜ�wŠw�vÕ'¨g©oÔ>ˆÜ×+Ò+ðð !Ø�Ü—*‘*ØØô
ˆô 	�Š�Õ" G¸E±/×Br   c                ó8   € V ^8„  d   QhR\         P                  /# rü   rý   )r   s   "r   r   r   Ü  s   € ÷ Cñ CÜ,.¯L©LñCr   c           	      óT  € \         P                  ! ^ R7      p\         P                  ! ^ V4      WC,          ,           p\         P                  ! W,           WR8  R7      p\         P                  ! RRV.\         P
                  R^R7      p\         P                  ! W,           WuV8  R7       R# )r   r+   r-   z»
            {
                // double quotes to pacify the linter """!!!"""
                cos.approx.f32 $0, $1;
                ex2.approx.f32 $0, $0;
            }
                rR  TrS  NrW  rX  s   &&&&    r   Úkernel_inline_asm_single_quotesr^  Û  s‚   € ô �mŠm Ô#ˆÜ—)’)˜A˜zÓ*¨SÕ-=Õ=ˆÜ�wŠw�vÕ'¨g©oÔ>ˆÜ×+Ò+ðð !Ø�Ü—*‘*ØØô
ˆô 	�Š�Õ" G¸E±/×Br   c                ó8   € V ^8„  d   QhR\         P                  /# rü   rý   )r   s   "r   r   r   ó  ó   € ÷ Cñ CÜ,.¯L©LñCr   c           	      óT  € \         P                  ! ^ R7      p\         P                  ! ^ V4      WC,          ,           p\         P                  ! W,           WR8  R7      p\         P                  ! RRV.\         P
                  R^R7      p\         P                  ! W,           WuV8  R7       R# ©r   r+   r-   zK
            v_sin_f32 $0, $1
            v_exp_f32 $0, $0
                z=v, vTrS  NrW  rX  s   &&&&    r   Ú$kernel_inline_asm_rocm_double_quotesrc  ò  ó‚   € ô �mŠm Ô#ˆÜ—)’)˜A˜zÓ*¨SÕ-=Õ=ˆÜ�wŠw�vÕ'¨g©oÔ>ˆÜ×+Ò+ðð !Ø�Ü—*‘*ØØô

ˆô 	�Š�Õ" G¸E±/×Br   c                ó8   € V ^8„  d   QhR\         P                  /# rü   rý   )r   s   "r   r   r     r`  r   c           	      óT  € \         P                  ! ^ R7      p\         P                  ! ^ V4      WC,          ,           p\         P                  ! W,           WR8  R7      p\         P                  ! RRV.\         P
                  R^R7      p\         P                  ! W,           WuV8  R7       R# rb  rW  rX  s   &&&&    r   Ú$kernel_inline_asm_rocm_single_quotesrg    rd  r   c                ó   € V ^8„  d   QhRR/# r$   r'   )r   s   "r   r   r     s   € ÷ 7ñ 7ð #ñ7r   c                 ób  € \         P                  ! ^ R7      pWe,          pV\         P                  ! ^ V4      ,           pWƒ8  p	\         P                  ! W,           V	R7      p
V'       d(   \         P                  ! W,           V	R7      pW«,           pMT
p\         P                  ! W(,           WÉR7       R# r*   r/   )r6   r7   r8   r9   Úadd_xyr%   r:   r;   r<   r.   r=   r>   r?   s   &&&&&&       r   Úadd_kernel_with_boolean_paramrk    s}   € ô �mŠm Ô#ˆØÕ&ˆØ¤§	¢	¨!¨ZÓ 8Õ8ˆØÑ#ˆÜ�GŠG�GÕ%¨DÔ1ˆßÜ—’˜Õ)°Ô5ˆAØ•U‰FàˆFÜ
�Š�Õ" F×6r   c                ó   € V ^8„  d   QhRR/# r$   r'   )r   s   "r   r   r   0  s   € ÷ 8ñ 8ð #ñ8r   c                ó¼  € \         P                  ! ^ R7      pWe,          pV\         P                  ! ^ V4      ,           pW„8  p	\         P                  ! W,           V	R7      p
\         P                  ! W,           V	R7      p\         P                  ! W(,           V	^ R7      ^ 8g  p\         P                  ! WÊV,           V
4      p\         P
                  ! W8,           WÙR7       R# )z6Kernel that loads a bool tensor and uses it as a mask.r+   r-   r!  N)r0   r1   r2   r3   rÌ   r4   )r6   r7   Úmask_ptrr8   r9   r%   r:   r;   r<   Úvalidr=   r>   Úkeepr?   s   &&&&&&        r   Ú"masked_add_kernel_with_bool_tensorrq  /  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   c                óF   € V ^8„  d   QhR\         \        ,          R\        /# )r
   Úblock_sizesÚnew_api)r   ÚintÚbool)r   s   "r   r   r   D  s    € ÷ ñ Ü!¤#�YðÜ15ñr   c           	      óŽ  € V'       d4   \         P                  P                  P                  P	                  W4      # \        V4      ^8X  d_   \         P                  P                  P                  V P                  4       V P                  ^ 4      V^ ,          V P                  4       4      # \        V4      ^8w  d   \        R\        V4       24      h\         P                  P                  P                  V P                  4       V P                  ^ 4      V P                  ^4      V^ ,          V^,          V P                  4       4      # )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)Útensorrs  rt  s   &&&r   Úcreate_tensor_descriptor_shimr…  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)  r*  r+  )T)BÚ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   r0   r   r"   Újitr@   rD   rK   rO   Úautotuner   rS   rW   rZ   rp   rt   r{   r   rŽ   r–   rœ   rŸ   r·   rÃ   rÆ   rÊ   rÍ   rÑ   râ   r÷   rú   r  r  Útriton.languager3   r4   r  r  r  r  r  rA  rI  rL  rO  r[  r^  rc  rg  rk  rq  r…  r'   r   r   Ú<module>r‘     s  ðó ÷ó õ +ð $×.Ò.ØÐ3óÐ ð #×-Ò-ØÐ1óÐ ð #×-Ò-Ø×DÐDÐ-×DÐDÐ1DØóÐ ð ×"Ò" 7¨NÓ;€á‡<ƒ<ÛÝ%ãõ5ðp ‡Z�Zô7ó ð7ð  ‡Z�Zô7ó ð7ð  ‡Z�Zô7ó ð7ð( ‡Z�Zô@ó ð@ð* ‡_‚_à�MŠM˜<¨Ð-¸!ÀqÔIØ�MŠM˜<¨Ð-¸!ÀqÔIØ�MŠM˜<¨Ð,¸ÀaÔHØ�MŠM˜<¨Ð,¸ÀaÔHð	
ð ôð ‡Z�Zô7ó óð7ð  ‡_‚_à�MŠM˜<¨Ð-¸!ÀqÔIØ�MŠM˜<¨Ð-¸!ÀqÔIØ�MŠM˜<¨Ð,¸ÀaÔHØ�MŠM˜<¨Ð,¸ÀaÔHð	
ð ôð ‡Z�Zô7ó óð7ð  ‡_‚_à�MŠM˜<¨Ð,¸ÀaÔHð
ð ô	ð ‡Z�Zô7ó óð7ð$ ‡_‚_à�MŠMØ  n°cÐ:ÀqÐTUôð �MŠMØ  n°cÐ:ÀqÐTUôð �MŠMØ  ^°RÐ8ÀQÐRSôð �MŠMØ  ^°RÐ8ÀQÐRSôð
ð ôð" ‡Z�ZôJó ó#ð$Jò,ð ‡_‚_à�MŠM˜<¨Ð-¸!ÀqÔIØ�MŠM˜<¨Ð,¸ÀaÔHð
ð ØØØ.Ð0IÐJô	ð ‡Z�Zô7ó ó	ð7ð  ‡Z�Zô7ó ð7ð" ‡Z�Zô
ó ð
ð< ‡Z�Zô
ó ð
ðB ‡Z�Zô
ó ð
ð4 ‡Z�Zô
ó ð
ð: ‡Z�ZôA
ó ðA
ðF ‡Z�Zô4
ó ð4
ðl ‡Z�Zô7ó ð7ð ‡Z�Zô3ó ð3ð ‡Z�Zñ&ó ð&ð ‡Z�Zô2ó ð2ð$ ‡Z�Zô3ó ð3ð& ‡Z�Zô-ó ð-ð  ‡Z�Zô-ó ð-ð  ‡Z�Zô+
ó ð+
ðZ ‡Z�Zô
ó ð
÷B ,à‡Z�Zô4ó ð4ð  ‡Z�Zô7ó ð7ð& ‡Z�Zô<ó ð<ð  ‡Z�Zô;ó ð;ð, ‡Z�Zô7ó ð7ð  ‡_‚_à�MŠMà" BØ" BØ" BØ" Að	ð Øô	ð �MŠMà" CØ" BØ" BØ" Að	ð Øô	ð
ò,  ô/ð2 ‡Z�Zô1)ó ó3ð41)ðf ‡Z�Zô@ó ð@ð ‡Z�Zô	@ó ð	@ð ‡Z�Zô@ó ð@ð ‡Z�ZôCó ðCð* ‡Z�ZôCó ðCð, ‡Z�ZôCó ðCð& ‡Z�ZôCó ðCð& ‡Z�Zô7ó ð7ð( ‡Z�Zô8ó ð8÷(ñ ñU! r   