Ë
    ÄÓ^ja½  ã                  óÖ  — d dl mZ d dlZd dlZd dlZd dlZd dlmZ d dlm	Z	m
Z
 d dlZd dlmZ d dlZd dlmZ d dlmZ d dlmZ d d	lmZmZ d d
lmZ ddlmZmZmZ ddlm Z m!Z!m"Z" ddl#m$Z$m%Z%m&Z&m'Z'm(Z(m)Z) ddl*m+Z+m,Z,m-Z- e
rddl.m/Z/m0Z0 ddl1m2Z2m3Z3 ddl#m4Z4  ejj                  e6«      Z7ejp                  dejr                  dejt                  dejv                  dejx                  dejz                  dej|                  dej~                  dej€                  di	ZAd'd„ZB G d„ de«      ZC G d „ d!e(«      ZDeDj‹                  d"«       eDj�                  «         G d#„ d$e,«      ZG G d%„ d&e-«      ZHy)(é    )ÚannotationsN)ÚPath)ÚAnyÚTYPE_CHECKING)Ú
PRECEDENCE)Ú_embed_headers)Ú
OrderedSet)ÚMin)Ú
CppPrinterÚExprPrinter)ÚValueRangesé   )ÚceildivÚget_bounds_index_exprÚget_kernel_metadata)ÚopsÚ
OpsWrapperÚVé   )ÚCSEVariableÚDeferredLineÚDTYPE_TO_COMPUTATION_DTYPEÚIndentedBufferÚOpOverridesÚPythonPrinter)ÚIterationRangesEntryÚ
SIMDKernelÚSIMDScheduling)ÚReductionTypeÚ	StoreMode)Ú	SchedulerÚSchedulerNode)ÚOpVarTÚboolÚcharÚshortÚintÚlongÚucharÚfloatÚhalfÚbfloatc                óØ   — t        | t        «      r:| t        j                  k(  ry| t        j                   k(  ry| | k7  ryt	        | «      S t        | t
        «      r| rdS dS t	        | «      S )NÚ	HUGE_VALFz
-HUGE_VALFÚNANÚtrueÚfalse)Ú
isinstancer*   ÚtorchÚinfÚstrr$   )Úvals    úf/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/torch/_inductor/codegen/mps.pyÚvalue_to_metalr8   7   s_   € Ü�#”uÔØ”%—)‘)ÒØØ”U—Y‘Y�JÒØØ�CŠZØÜ�3‹xˆÜ	�CœÔ	ÙˆvÐ) 'Ð)Üˆs‹8€Oó    c                  ó¦   — e Zd ZdZdd„Zdd„Zdd„Zedd„«       Zdd„Z	dd„Z
dd„Zdd	„Zdd
„Zdd„Zdd„Zdd„Zdd„Zdd„ZeZdd„Zdd„Zdd„Zy)ÚMetalExprPrinterz/Converts sympy expression to Metal code snippetc                ó    — |j                   \  }}| j                  |«      }| j                  |«      }|j                  r	d|› d|› d�S d|› d|› d�S )Núc10::metal::floor_divide(ú, ú)úmetal::floor(ú) / (©ÚargsÚdoprintÚ
is_integer)ÚselfÚexprÚxÚdivs       r7   Ú_print_FloorDivz MetalExprPrinter._print_FloorDivH   s[   € Ø—‘‰ˆˆ3Ø�L‰L˜‹OˆØ�l‰l˜3ÓˆØ�?Š?Ø.¨q¨c°°C°5¸Ð:Ð:Ø˜q˜c  s e¨1Ð-Ð-r9   c                ó  — |j                   \  }}}|dk(  xr ||dz
  z  dk7  }| j                  |«      }|dk7  r0| j                  |«      }|j                  r
d|› d|› d�}n	d|› d|› d�}| j                  |«      }|r	d|› d	|› d�S d|› d
|› d�S )Ni   r   r   ú(rA   r?   r@   zc10::metal::safe_mod(r>   z) % (rB   )rF   rG   rH   rI   ÚmodÚuse_safe_mods         r7   Ú_print_ModularIndexingz'MetalExprPrinter._print_ModularIndexingP   s¸   € Ø—i‘i‰ˆˆ3�à˜e‘|Ò>¨°°a±©¸QÑ(>ˆà�L‰L˜‹OˆØ�!Š8Ø—,‘,˜sÓ#ˆCØ�ŠØ˜�s˜% ˜u AÐ&‘à# A 3 e¨C¨5°Ð2�Ø�l‰l˜3ÓˆÙØ*¨1¨#¨R°¨u°AÐ6Ð6Ø�1�#�U˜3˜%˜qÐ!Ð!r9   c                óœ   — t        t        | j                  |j                  «      «      }|d   }|dd  D ]  }| j	                  |||«      }Œ |S )Nr   r   )ÚlistÚmapÚ_printrC   Ú_print_binary_min_max)rF   rG   ÚfnrC   ÚresultÚargs         r7   Ú_print_min_maxzMetalExprPrinter._print_min_maxa   sT   € ä”C˜Ÿ™ T§Y¡YÓ/Ó0ˆØ�a‘ˆØ˜˜�8ò 	AˆCØ×/Ñ/°¸¸RÓ@‰Fð	Aàˆr9   c                óJ   — d| › d|› d| › d�}d| › d|› d|› d�}d|› d|› d|› d�S )Nústatic_cast<decltype(ú+ú)>(r?   zmetal::rL   r>   © )ÚaÚbrU   Ú
typecast_aÚ
typecast_bs        r7   rT   z&MetalExprPrinter._print_binary_min_maxi   sQ   € à,¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø,¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø˜˜˜A˜j˜\¨¨J¨<°qÐ9Ð9r9   c                ó&   — | j                  |d«      S )NÚmin©rX   ©rF   rG   s     r7   Ú
_print_MinzMetalExprPrinter._print_Mino   ó   € Ø×"Ñ" 4¨Ó/Ð/r9   c                ó&   — | j                  |d«      S )NÚmaxrd   re   s     r7   Ú
_print_MaxzMetalExprPrinter._print_Maxr   rg   r9   c                óz   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › d�S )Nr   úmetal::abs(r   r?   ©ÚlenrC   rS   re   s     r7   Ú
_print_AbszMetalExprPrinter._print_Absu   s9   € Ü�4—9‘9‹~ Ò"Ð"Ð"à˜TŸ[™[¨¯©°1©Ó6Ð7°qÐ9Ð9r9   c                óz   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › d�S )Nr   zstatic_cast<long>(metal::rint(r   ú))rm   re   s     r7   Ú_print_RoundToIntz"MetalExprPrinter._print_RoundToIntz   s9   € Ü�4—9‘9‹~ Ò"Ð"Ð"à/°·±¸D¿I¹IÀa¹LÓ0IÐ/JÈ"ÐMÐMr9   c                óä   — t        |j                  «      dk(  sJ ‚|j                  \  }}|j                  r|dk  sJ ‚t        d|› d�«      ‚| j	                  |t
        d   «      }d|› d|› d| › d	�S )
Nr   r   zOFor integer inputs, only non-negative ndigits are currently supported, but got ú.ÚMulz!static_cast<float>(metal::rint(1eú * z) * 1er?   )rn   rC   rE   Ú
ValueErrorÚparenthesizer   )rF   rG   ÚnumberÚndigitsÚ
number_strs        r7   Ú_print_RoundDecimalz$MetalExprPrinter._print_RoundDecimal   s�   € Ü�4—9‘9‹~ Ò"Ð"Ð"ØŸ)™)‰ˆ�Ø×Òà˜Q’;Ð�;ÜØaÐbiÐajÐjkÐlóð ð ×&Ñ& v¬z¸%Ñ/@ÓAˆ
Ø2°7°)¸3¸z¸lÈ&ÐRYÐQYÐPZÐZ[Ð\Ð\r9   c                ón   — |j                   \  }}d| j                  |«      › d| j                  |«      › d�S )Nústatic_cast<float>(z) / static_cast<float>(r?   )rC   rS   )rF   rG   ÚlhsÚrhss       r7   Ú_print_IntTrueDivz"MetalExprPrinter._print_IntTrueDiv‹   s=   € Ø—9‘9‰ˆˆSð % T§[¡[°Ó%5Ð$6Ð6MÈdÏkÉkÐZ]ÓN^ÐM_Ð_`ÐaÐar9   c                óŽ   — t        |j                  «      dk(  sJ ‚t        | j                  |j                  «      \  }}d|› d|› d�S )Nr   z'metal::precise::pow(static_cast<float>(z), static_cast<float>(rq   )rn   rC   rR   rD   )rF   rG   rH   Úys       r7   Ú_print_PowByNaturalz$MetalExprPrinter._print_PowByNatural‘   sI   € Ü�4—9‘9‹~ Ò"Ð"Ð"Ü�4—<‘< §¡Ó+‰ˆˆ1Ø8¸¸Ð;QÐRSÐQTÐTVÐWÐWr9   c                ó~   — t        |j                  «      dk(  sJ ‚| j                  |j                  d   «      }d|› d�S )Nr   r   r~   r?   ©rn   rC   rD   ©rF   rG   rH   s      r7   Ú_print_ToFloatzMetalExprPrinter._print_ToFloat–   s=   € Ü�4—9‘9‹~ Ò"Ð"Ð"Ø�L‰L˜Ÿ™ 1™Ó&ˆØ$ Q C qÐ)Ð)r9   c                óX   — |j                   rt        t        |«      «      S t        |«      S ©N)rE   r5   r'   re   s     r7   Ú_print_FloatzMetalExprPrinter._print_Float›   s#   € Ø�?Š?ô ”s˜4“y“>Ð!ä�t“9Ðr9   c                ó~   — t        |j                  «      dk(  sJ ‚| j                  |j                  d   «      }d|› d�S )Nr   r   z1static_cast<int>(metal::floor(static_cast<float>(z)))r†   r‡   s      r7   Ú_print_FloorToIntz"MetalExprPrinter._print_FloorToInt¤   s=   € Ü�4—9‘9‹~ Ò"Ð"Ð"Ø�L‰L˜Ÿ™ 1™Ó&ˆØBÀ1À#ÀSÐIÐIr9   c                ó~   — t        |j                  «      dk(  sJ ‚| j                  |j                  d   «      }d|› d�S )Nr   r   zstatic_cast<int>(metal::trunc(rq   r†   r‡   s      r7   Ú_print_TruncToIntz"MetalExprPrinter._print_TruncToInt«   s=   € Ü�4—9‘9‹~ Ò"Ð"Ð"Ø�L‰L˜Ÿ™ 1™Ó&ˆØ/°¨s°"Ð5Ð5r9   c                ó~   — t        |j                  «      dk(  sJ ‚| j                  |j                  d   «      }d|› d�S )Nr   r   zmetal::precise::log2(r?   r†   r‡   s      r7   Ú_print_OpaqueUnaryFn_log2z*MetalExprPrinter._print_OpaqueUnaryFn_log2°   s=   € Ü�4—9‘9‹~ Ò"Ð"Ð"Ø�L‰L˜Ÿ™ 1™Ó&ˆØ& q c¨Ð+Ð+r9   c                óJ   ‡ — ˆ fd„|j                   D «       \  }}}|› d|› d|› �S )Nc              3  óV   •K  — | ]   }‰j                  |t        d    dz
  «      –— Œ" y­w)ÚAtomg      à?N)rx   r   )Ú.0rW   rF   s     €r7   ú	<genexpr>z0MetalExprPrinter._print_Where.<locals>.<genexpr>¶   s-   øè ø€ ò 
ØADˆD×Ñ˜c¤:¨fÑ#5¸Ñ#;×<ñ
ùs   ƒ&)ú ? z : )rC   )rF   rG   ÚcÚpÚqs   `    r7   Ú_print_WherezMetalExprPrinter._print_Whereµ   s5   ø€ ó
ØHLÏ	É	ô
‰ˆˆ1ˆað ��C˜�s˜#˜a˜SÐ!Ð!r9   N)rG   ú
sympy.ExprÚreturnr5   )rG   rœ   rU   r5   r�   r5   )r^   r5   r_   r5   rU   r5   r�   r5   )Ú__name__Ú
__module__Ú__qualname__Ú__doc__rJ   rO   rX   ÚstaticmethodrT   rf   rj   ro   rr   r|   r�   r„   rˆ   r‹   r�   Ú_print_floorr�   r‘   r›   r]   r9   r7   r;   r;   E   su   „ Ù9ó.ó"ó"ð ò:ó ð:ó
0ó0ó:ó
Nó

]óbóXó
*ó
óJð
 %€Ló6ó
,ô
"r9   r;   c                  ót  — e Zd ZdZe	 	 d1	 	 	 	 	 	 	 	 	 d2d„«       Ze	 	 	 	 	 	 	 	 d3d„«       Zed4d„«       Zed5d„«       Zed5d„«       Z	ed6d„«       Z
ed7d	„«       Zed8d
„«       Zed9d„«       Zed9d„«       Zed9d„«       Zed9d„«       Zed:d„«       Zed:d„«       Zed:d„«       Zed:d„«       Zed:d„«       Zed:d„«       Zed:d„«       Zed:d„«       Zed:d„«       Zed:d„«       Zed:d„«       Zed:d„«       Zed:d„«       Zed;d„«       Zed:d„«       Zed:d„«       Z ed:d„«       Z!ed:d „«       Z"ed:d!„«       Z#ed9d"„«       Z$ed:d#„«       Z%ed:d$„«       Z&ed9d%„«       Z'ed:d&„«       Z(ed9d'„«       Z)ed:d(„«       Z*ed<d)„«       Z+ed<d*„«       Z,e	 	 	 	 	 	 	 	 	 	 d=d+„«       Z-ed:d,„«       Z.ed9d-„«       Z/d>d.„Z0d?d/„Z1e2d@d0„«       Z3y)AÚMetalOverrideszXImplements Metal-specific overrides for ops. Base class emits Python-friendly overrides.Nc                ó~   — |t         j                  k(  rt        j                  d«       d| › d�S dt        |   › d| › d�S )Nz>float64 cast requested, probably from tensorify_python_scalarsr~   r?   ústatic_cast<ú>()r3   ÚdoubleÚlogÚwarningÚDTYPE_TO_METAL)rH   ÚdtypeÚ	src_dtypeÚuse_compute_typess       r7   Úto_dtypezMetalOverrides.to_dtype¿   sK   € ð ”E—L‘LÒ Ü�K‰KØPôð )¨¨¨1Ð-Ð-Øœn¨UÑ3Ð4°B°q°c¸Ð;Ð;r9   c                ó6   — dt         |   › dt         |   › d| › d�S )Nzas_type<z>(static_cast<r¨   rq   ©r¬   )rH   r­   r®   s      r7   Úto_dtype_bitcastzMetalOverrides.to_dtype_bitcastÍ   s/   € ð œ.¨Ñ/Ð0°¼~ÈiÑ?XÐ>YÐY[Ð\]Ð[^Ð^`ÐaÐar9   c                ó   — t        | «      S rŠ   ©r8   )r6   r­   s     r7   ÚconstantzMetalOverrides.constantÓ   s   € ä˜cÓ"Ð"r9   c                ó@  — t         j                  j                  t         j                  j                  | «      «      }t         j                  j                  j                  t         j                  j                  |t        | «      ¬«      }t        j                  ||«      S )N)Úbounds)
r   ÚkernelÚindex_to_strÚprepare_indexingÚcseÚgenerateÚcomputer   r   r°   )rG   r­   Úidx_strÚvars       r7   Ú
index_exprzMetalOverrides.index_expr×   sl   € ä—(‘(×'Ñ'¬¯©×(AÑ(AÀ$Ó(GÓHˆÜ�h‰h�l‰l×#Ñ#Ü�H‰H×Ñ˜gÔ.CÀDÓ.Ið $ó 
ˆô �|‰|˜C Ó'Ð'r9   c                ó.   — t         j                  | |«      S rŠ   )r¥   rÁ   )rG   r­   s     r7   Ú
value_exprzMetalOverrides.value_exprß   s   € ô ×(Ñ(¨¨uÓ5Ð5r9   c           	     ó  — t        |«      }t        «       }t        j                  j	                  |«      5  |j                  «       5  t        j                  «       t        j                  j                  _	        t        j                  j                  xj                  dz  c_
         |«       }d d d «       d d d «       | › d|j                  «       › d|› �}t        j                  j                  j                  |«      }|�sct        j                  j                  j                  j                  ¬«      }t        j                  j                  j                  ||«       t        j                  j                   j#                  t$        |j                     › d|› d�d| › d�g«       t        j                  j                   j                  «       5  t        j                  j                   j'                  |«       t        j                  j                   j)                  |› d|› d	|› d
�«       d d d «       t        j                  j                   j)                  d|› d|› d	|› d
�«       |S # 1 sw Y   �Œ»xY w# 1 sw Y   �ŒÀxY w# 1 sw Y   ŒXxY w)NÚ_scopedú:©r­   ú ú;úif (ú) {z = static_cast<decltype(r\   ú);z} else )r8   r   r   r¹   Úswap_buffersÚindentÚ	itertoolsÚcountr¼   Úiter_buffer_idsÚname_prefixÚgetvalueÚtry_getÚnewvarr­   Úputr¾   Ú
writelinesr¬   ÚspliceÚ	writeline)ÚmaskÚbodyÚotherÚ	other_strÚscoped_bodyÚrcÚ	cache_keyrÀ   s           r7   ÚmaskedzMetalOverrides.maskedå   sø  € ô
 # 5Ó)ˆ	Ü$Ó&ˆÜ�X‰X×"Ñ" ;Ó/ñ 	°×1CÑ1CÓ1Eñ 	ô ,5¯?©?Ó+<ŒA�H‰H�L‰LÔ(ä�H‰H�L‰L×$Ò$¨	Ñ1Õ$Ù“ˆB÷	÷ 	ð �f˜A˜k×2Ñ2Ó4Ð5°Q°y°kÐBˆ	Ü�h‰h�l‰l×"Ñ" 9Ó-ˆÚÜ—(‘(—,‘,×%Ñ%¨B¯H©HÐ%Ó5ˆCÜ�H‰H�L‰L×Ñ˜Y¨Ô,Ü�H‰H×Ñ×'Ñ'Ü" 2§8¡8Ñ,Ð-¨Q¨s¨e°1Ð5¸¸d¸VÀ4Ð7HÐIôô —‘×!Ñ!×(Ñ(Ó*ñ Ü—‘× Ñ ×'Ñ'¨Ô4Ü—‘× Ñ ×*Ñ*Ø�eÐ3°C°5¸¸B¸4¸rÐBô÷ô
 �H‰H×Ñ×&Ñ&Ø˜3˜%Ð7¸°u¸CÀ	¸{È"ÐMôð ˆ
÷5	ñ 	ú÷ 	ñ 	ú÷$ð ús1   µI2ÁA&I%Â,I2ÇAI?É%I/	É*I2É2I<É?Jc           	     ó0   — | › d|› d|› dt        |«      › d�S )Nr—   z : static_cast<decltype(r\   r?   rµ   )r^   r_   r˜   s      r7   ÚwherezMetalOverrides.where  s*   € ð ��C˜�sÐ2°1°#°S¼ÈÓ9JÐ8KÈ1ÐMÐMr9   c                ó   — d| › d|› d�S )Nzc10::metal::remainder(r>   r?   r]   ©r^   r_   s     r7   Ú	remainderzMetalOverrides.remainder  s   € à'¨ s¨"¨Q¨C¨qÐ1Ð1r9   c                óD   — d| › d|› d| › d�}d| › d|› d|› d�}d|› d|› d�S )NrZ   r[   r\   r?   zc10::metal::max(r>   r]   ©r^   r_   r`   ra   s       r7   ÚmaximumzMetalOverrides.maximum  óM   € ð -¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø,¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø! * ¨R°
¨|¸1Ð=Ð=r9   c                óD   — d| › d|› d| › d�}d| › d|› d|› d�}d|› d|› d�S )NrZ   r[   r\   r?   zc10::metal::min(r>   r]   rè   s       r7   ÚminimumzMetalOverrides.minimum  rê   r9   c                ó   — | › d|› �S )Nz || r]   rå   s     r7   Ú
logical_orzMetalOverrides.logical_or  ó   € ð ��D˜˜ˆ}Ðr9   c                ó   — | › d|› �S )Nz && r]   rå   s     r7   Úlogical_andzMetalOverrides.logical_and$  rï   r9   c                ó   — d| › d�S )Nzmetal::isnan(r?   r]   ©rH   s    r7   ÚisnanzMetalOverrides.isnan)  ó   € ð ˜q˜c Ð#Ð#r9   c                ó   — d| › d�S )Nzmetal::isinf(r?   r]   ró   s    r7   ÚisinfzMetalOverrides.isinf.  rõ   r9   c                ó   — d| › d�S )Nzmetal::precise::log(r?   r]   ró   s    r7   rª   zMetalOverrides.log3  ó   € ð & a S¨Ð*Ð*r9   c                ó   — d| › d�S )Nzmetal::precise::exp(r?   r]   ró   s    r7   ÚexpzMetalOverrides.exp8  rù   r9   c                ó   — d| › d�S )Nrl   r?   r]   ró   s    r7   ÚabszMetalOverrides.abs=  s   € ð ˜Q˜C˜qÐ!Ð!r9   c                ó   — d| › d�S )Nzmetal::signbit(r?   r]   ró   s    r7   ÚsignbitzMetalOverrides.signbitB  s   € ð !   1Ð%Ð%r9   c                ó   — d| › d�S )Nzmetal::precise::sin(r?   r]   ró   s    r7   ÚsinzMetalOverrides.sinG  rù   r9   c                ó   — d| › d�S )Nzc10::metal::sinc(r?   r]   ró   s    r7   ÚsinczMetalOverrides.sincL  s   € à" 1 # QÐ'Ð'r9   c                ó   — d| › d�S )Nzmetal::precise::cos(r?   r]   ró   s    r7   ÚcoszMetalOverrides.cosP  rù   r9   c                ó   — d| › d�S )Nzmetal::precise::tan(r?   r]   ró   s    r7   ÚtanzMetalOverrides.tanU  rù   r9   c                ó   — d| › d�S )Nzmetal::precise::asin(r?   r]   ró   s    r7   ÚasinzMetalOverrides.asinZ  ó   € ð ' q c¨Ð+Ð+r9   c                ó   — d| › d�S )Nzmetal::precise::acos(r?   r]   ró   s    r7   ÚacoszMetalOverrides.acos_  r
  r9   c                ó   — d| › d�S )Nzmetal::precise::atan(r?   r]   ró   s    r7   ÚatanzMetalOverrides.atand  r
  r9   c                ó   — d| › d|› d�S )Nz::metal::precise::atan2(r>   r?   r]   )rH   rƒ   s     r7   Úatan2zMetalOverrides.atan2i  s   € ð *¨!¨¨B¨q¨c°Ð3Ð3r9   c                ó   — d| › d�S )Nzmetal::precise::sqrt(r?   r]   ró   s    r7   ÚsqrtzMetalOverrides.sqrtn  r
  r9   c                ó   — d| › d| › d�S )NrZ   z)>(-r?   r]   ró   s    r7   ÚnegzMetalOverrides.negs  s   € ð ' q c¨¨a¨S°Ð2Ð2r9   c                ó   — d| › d�S )Nzmetal::precise::rsqrt(r?   r]   ró   s    r7   ÚrsqrtzMetalOverrides.rsqrty  ó   € ð (¨ s¨!Ð,Ð,r9   c                ó   — d| › d�S )Nzmetal::precise::tanh(r?   r]   ró   s    r7   ÚtanhzMetalOverrides.tanh~  r
  r9   c                ó   — d| › d�S )Nzmetal::precise::atanh(r?   r]   ró   s    r7   ÚatanhzMetalOverrides.atanhƒ  r  r9   c                ó   — d| › d|› d�S )Nr=   r>   r?   r]   rå   s     r7   ÚfloordivzMetalOverrides.floordivˆ  s   € ð +¨1¨#¨R°¨s°!Ð4Ð4r9   c                ó   — d| › d�S )Nr@   r?   r]   ró   s    r7   ÚfloorzMetalOverrides.floor�  rõ   r9   c                ó   — d| › d�S )Nzmetal::sign(r?   r]   ró   s    r7   ÚsignzMetalOverrides.sign’  ó   € ð ˜a˜S Ð"Ð"r9   c                óD   — d| › d|› d| › d�}d| › d|› d|› d�}d|› d|› d�S )NrZ   r[   r\   r?   zmetal::fmod(r>   r]   rè   s       r7   ÚfmodzMetalOverrides.fmod—  sM   € ð -¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø,¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø˜j˜\¨¨J¨<°qÐ9Ð9r9   c                ó   — d| › d�S )Númetal::trunc(r?   r]   ró   s    r7   ÚtrunczMetalOverrides.truncž  rõ   r9   c                ó¨   — | › d|› �}| j                   �| j                   j                  s"|j                   �|j                   j                  rd|› d�S |S )Nz / r&  r?   )r­   Úis_floating_point)r^   r_   Úquots      r7   ÚtruncdivzMetalOverrides.truncdiv£  sS   € ð ��C˜�sˆ|ˆØ�G‰GÐ A§G¡G×$=Ò$=Ø�G‰GÐ A§G¡G×$=Ò$=à" 4 &¨Ð*Ð*Øˆr9   c                ó   — d| › d�S )Nzmetal::ceil(r?   r]   ró   s    r7   ÚceilzMetalOverrides.ceil­  r"  r9   c                óf   — t         j                  j                  j                  d«       d| › d|› d�S )NÚrandomzc10::metal::rand(r>   r?   ©r   r¹   ÚheadersÚadd©ÚseedÚoffsets     r7   ÚrandzMetalOverrides.rand²  s/   € ä	�‰×Ñ×Ñ˜XÔ&Ø" 4 &¨¨6¨(°!Ð4Ð4r9   c                óf   — t         j                  j                  j                  d«       d| › d|› d�S )Nr/  zc10::metal::randn(r>   r?   r0  r3  s     r7   ÚrandnzMetalOverrides.randn·  s/   € ä	�‰×Ñ×Ñ˜XÔ&Ø# D 6¨¨F¨8°1Ð5Ð5r9   c           	     ór   — t         j                  j                  j                  d«       d| › d|› d|› d|› d�	S )Nr/  zc10::metal::randint64(r>   r?   r0  )r4  r5  ÚlowÚhighs       r7   Ú	randint64zMetalOverrides.randint64¼  s=   € ô 	
�‰×Ñ×Ñ˜XÔ&Ø'¨ v¨R°¨x°r¸#¸¸bÀÀÀaÐHÐHr9   c                ó   — d| › d�S )Nzmetal::rint(r?   r]   ró   s    r7   ÚroundzMetalOverrides.roundÃ  r"  r9   c                óD   — d| › d|› d| › d�}d| › d|› d|› d�}d|› d|› d�S )NrZ   r[   r\   r?   zmetal::pow(r>   r]   )r^   r_   Úcast_aÚcast_bs       r7   ÚpowzMetalOverrides.powÈ  sK   € à(¨¨¨1¨Q¨C¨s°1°#°QÐ7ˆØ(¨¨¨1¨Q¨C¨s°1°#°QÐ7ˆØ˜V˜H B v h¨aÐ0Ð0r9   c                óf   — t         j                  j                  j                  d«       d|› d|› d�S )NÚspecial_mathúc10::metal::rL   r?   r0  )rF   r^   Únames      r7   Ú_special_unaryzMetalOverrides._special_unaryÎ  s/   € Ü	�‰×Ñ×Ñ˜^Ô,Ø˜d˜V 1 Q C qÐ)Ð)r9   c                ól   — t         j                  j                  j                  d«       d|› d|› d|› d�S )NrD  rE  rL   r>   r?   r0  )rF   r^   r_   rF  s       r7   Ú_special_binaryzMetalOverrides._special_binaryÒ  s5   € Ü	�‰×Ñ×Ñ˜^Ô,Ø˜d˜V 1 Q C r¨!¨¨AÐ.Ð.r9   c           
     óô  — dD ].  }t        | |t        j                  | j                  |¬«      «       Œ0 t        j                  | j                  d¬«      | _        dD ]1  }t        | |t        j                  | j                  |dz   ¬«      «       Œ3 dD ].  }t        | |t        j                  | j
                  |¬«      «       Œ0 dD ]1  }t        | |t        j                  | j
                  |dz   ¬«      «       Œ3 y )N)ÚerfÚerfinvÚi0Úi0eÚi1Úi1eÚdigammaÚspherical_bessel_j0)rF  Ú	log_gamma)
Ú	bessel_j0Ú	bessel_j1Ú	bessel_y0Ú	bessel_y1Úmodified_bessel_i0Úmodified_bessel_i1Úmodified_bessel_k0Úmodified_bessel_k1Úscaled_modified_bessel_k0Úscaled_modified_bessel_k1Ú_forward)Ú	polygammaÚigammaÚigammacÚzeta)
Úchebyshev_polynomial_tÚchebyshev_polynomial_uÚchebyshev_polynomial_vÚchebyshev_polynomial_wÚhermite_polynomial_hÚhermite_polynomial_heÚshifted_chebyshev_polynomial_tÚshifted_chebyshev_polynomial_uÚshifted_chebyshev_polynomial_vÚshifted_chebyshev_polynomial_w)ÚsetattrÚ	functoolsÚpartialmethodrG  ÚlgammarI  )ÚclsrF  s     r7   Ú_initialize_special_opsz&MetalOverrides._initialize_special_opsÖ  s   € ð	
ò 
	WˆDô �C˜œy×6Ñ6°s×7IÑ7IÐPTÔUÕVð
	Wô ×,Ñ,¨S×-?Ñ-?ÀkÔRˆŒ
ð
ò 	ˆDô ØØÜ×'Ñ'¨×(:Ñ(:ÀÈ
ÑARÔSõð	ð&
ò 	XˆDô �C˜œy×6Ñ6°s×7JÑ7JÐQUÔVÕWð	Xð
ò 	ˆDô ØØÜ×'Ñ'¨×(;Ñ(;À$ÈÑBSÔTõñ	r9   ©NT)
rH   r   r­   útorch.dtyper®   ztorch.dtype | Noner¯   r$   r�   r5   )rH   r   r­   rt  r®   rt  r�   r5   )r6   zbool | float | intr­   rt  r�   r5   )rG   rœ   r­   rt  r�   r5   )rÚ   r   rÛ   rœ   rÜ   r   r�   r5   )r^   r#   r_   r#   r˜   r#   r�   r5   )r^   r#   r_   r#   r�   r5   )r^   r   r_   r   r�   r5   )rH   r   r�   r5   )rH   r   rƒ   r   r�   r5   )r4  r   r5  r   r�   r5   )
r4  r   r5  r   r:  r   r;  r   r�   r5   )r^   r   rF  r5   r�   r5   )r^   r   r_   r   rF  r5   r�   r5   ©r�   ÚNone)4rž   rŸ   r    r¡   r¢   r°   r³   r¶   rÁ   rÃ   rá   rã   ræ   ré   rì   rî   rñ   rô   r÷   rª   rû   rý   rÿ   r  r  r  r  r	  r  r  r  r  r  r  r  r  r  r  r!  r$  r'  r+  r-  r6  r8  r<  r>  rB  rG  rI  Úclassmethodrr  r]   r9   r7   r¥   r¥   ¼   s  „ Ùbàð )-Ø"&ð	<Øð<àð<ð &ð<ð  ð	<ð
 
ò<ó ð<ð ðbØðbØ*ðbØ7Bðbà	òbó ðbð
 ò#ó ð#ð ò(ó ð(ð ò6ó ð6ð
 ò ó ð ðD òNó ðNð ò2ó ð2ð ò>ó ð>ð
 ò>ó ð>ð
 òó ðð òó ðð ò$ó ð$ð ò$ó ð$ð ò+ó ð+ð ò+ó ð+ð ò"ó ð"ð ò&ó ð&ð ò+ó ð+ð ò(ó ð(ð ò+ó ð+ð ò+ó ð+ð ò,ó ð,ð ò,ó ð,ð ò,ó ð,ð ò4ó ð4ð ò,ó ð,ð ò3ó ð3ð
 ò-ó ð-ð ò,ó ð,ð ò-ó ð-ð ò5ó ð5ð ò$ó ð$ð ò#ó ð#ð ò:ó ð:ð
 ò$ó ð$ð òó ðð ò#ó ð#ð ò5ó ð5ð ò6ó ð6ð ðIØðIØ#.ðIØ5@ðIØHSðIà	òIó ðIð ò#ó ð#ð ò1ó ð1ó
*ó/ð ò=ó ñ=r9   r¥   Úmpsc                  óæ  ‡ — e Zd ZU dZeZdZdZdZdZ	 e
«       j                  Z e«       j                  Z e«       j                  ZeZ edg«      Zded<   g Zd	ed
<   	 	 	 	 	 	 dˆ fd„Zdd„Zdd„Z	 d	 	 	 	 	 	 	 	 	 dd„Zd d„Zddd ej8                  «       f	 	 	 	 	 	 	 	 	 	 	 d!d„Z	 	 	 	 	 	 	 	 	 	 d"d„Z	 	 	 	 	 	 	 	 	 	 d"d„Zd#d„Z d$d„Z!d%d&d„Z"	 d'	 	 	 	 	 	 	 d(d„Z#d)d„Z$	 	 	 	 	 	 	 	 	 	 d*d„Z%ˆ xZ&S )+ÚMetalKernelz;Implement Metal codegen based on the SIMDKernel abstractionrÉ   úauto i   é    ÚutilszOrderedSet[str]r1  zlist[IterationRangesEntry]Úmultistage_reduction_entryc                óX   •— t        ‰| �  |fi |¤Ž t        j                  «       | _        y rŠ   )ÚsuperÚ__init__rÏ   rÐ   Úacc_var_ids)rF   ÚtilingÚkwargsÚ	__class__s      €r7   r�  zMetalKernel.__init__*  s&   ø€ ô
 	‰Ñ˜Ñ* 6Ò*Ü$Ÿ?™?Ó,ˆÕr9   c                ó   — t         |   S rŠ   r²   )rF   r­   s     r7   Údtype_to_strzMetalKernel.dtype_to_str2  s   € Ü˜eÑ$Ð$r9   c                ó†  — | j                   j                  |«      }| j                  |«      }t        j                  j                  |«      }|› d| j                  |«      › d�}|t        j                  t        j                  fv rd|› d�}t        j                  }| j                  j                  | j                  ||¬«      S )z"Codegen a load from an InputBufferú[ú]r~   r?   rÇ   )rC   Úinputr»   r   ÚgraphÚ	get_dtyperº   r3   Úfloat16Úbfloat16Úfloat32r¼   r½   Úloads)rF   rF  ÚindexrÀ   r­   Úlines         r7   ÚloadzMetalKernel.load5  s¤   € à�i‰i�o‰o˜dÓ#ˆØ×%Ñ% eÓ,ˆÜ—‘×!Ñ! $Ó'ˆØ��a˜×)Ñ)¨%Ó0Ð1°Ð3ˆØ”U—]‘]¤E§N¡NÐ3Ñ3ð )¨¨¨aÐ0ˆDÜ—M‘MˆEØ�x‰x× Ñ  §¡¨T¸Ð Ó?Ð?r9   Nc                óh  — | j                   j                  |«      }| j                  |«      }| j                  t        j
                  j                  |«      «      }d|› d|› d�}|€|› d| j                  |«      › d|› d�}n[|dk(  rH| j                  j                  d«       d	|› d
�}	d|	› d|› d�}
|	› d|
› d| j                  |«      › d|› d�}nt        d|› �«      ‚| j                  r&| j                  j                  t        ||«      «       y | j                  j                  t        ||«      «       y )Nr§   r¨   r?   r‰  ú] = rÉ   Ú
atomic_addÚatomiczc10::metal::AtomicType<ú>zreinterpret_cast<device z
::type *>(z::atomic_add(r>   rÌ   zUnimplemented store mode )rC   Úoutputr»   r‡  r   rŒ  r�  rº   r1  r2  ÚRuntimeErrorÚinside_reductionr¾   rÙ   r   Ústores)rF   rF  r’  ÚvalueÚmoderÀ   Ú	dtype_strÚcast_valr“  Úatomic_typeÚcast_vars              r7   ÚstorezMetalKernel.storeC  s<  € ð �i‰i×Ñ˜tÓ$ˆØ×%Ñ% eÓ,ˆØ×%Ñ%¤a§g¡g×&7Ñ&7¸Ó&=Ó>ˆ	Ø! ) ¨B¨u¨g°QÐ7ˆØˆ<Ø�U˜!˜D×-Ñ-¨eÓ4Ð5°T¸(¸À1ÐE‰DØ�\Ò!Ø�L‰L×Ñ˜XÔ&Ø3°I°;¸aÐ@ˆKØ1°+°¸jÈÈÈQÐOˆHØ!�] -°¨z¸¸D×<MÑ<MÈeÓ<TÐ;UÐUWÐX`ÐWaÐacÐd‰DäÐ!:¸4¸&ÐAÓBÐBØ× Ò Ø�L‰L×"Ñ"¤<°°dÓ#;Õ<à�K‰K×!Ñ!¤,¨t°TÓ":Õ;r9   c                ó˜  — | j                   j                  |«      }| j                  |«      }| j                  t        j
                  j                  |«      «      }t        d„ | j                  D «       «      }|› d| j                  |«      › d|› d|› d�}d|j                  › d|› �}| j                  j                  t        ||«      «       y )Nc              3  ó:   K  — | ]  }|j                   sŒ|–— Œ y ­wrŠ   ©Úis_reduction©r•   Úts     r7   r–   z.MetalKernel.store_reduction.<locals>.<genexpr>]  s   è ø€ ÒK 1¸A¿N»NœQÑKùó   ‚”r‰  z] = static_cast<r¨   rÌ   rÊ   z == 0) )rC   rš  r»   r‡  r   rŒ  r�  ÚnextÚrange_treesrº   rF  r�  rÙ   r   )rF   rF  r’  rž  rÀ   r   Úreduction_dimr“  s           r7   Ústore_reductionzMetalKernel.store_reductionX  s¹   € Ø�i‰i×Ñ˜tÓ$ˆØ×%Ñ% eÓ,ˆØ×%Ñ%¤a§g¡g×&7Ñ&7¸Ó&=Ó>ˆ	äÑK¨×(8Ñ(8ÔKÓKˆà��a˜×)Ñ)¨%Ó0Ð1Ð1AÀ)ÀÈBÈuÈgÐUWÐXˆØ�m×(Ñ(Ð)¨°°Ð7ˆØ�‰×Ñœl¨4°Ó6Õ7r9   Tc                ó˜  — t        |t        j                  «      r| j                  |«      }dt	        | j
                  «      › �}t        j                  j                  |||«      }|rdnd}||› d|› �z  }|r|d| j                  |«      › d�z  }|�|rJ d«       ‚|d|› �z  }| j                  j                  || j                  z   «       |S )	NÚtmp_acc_zthreadgroup Ú rÈ   r‰  rŠ  z+Thread group var can not have default valueú = )r2   r3   r­   r‡  r¬  r‚  r   r¹   Úcreate_cse_varÚsexprÚindexing_coderÙ   Úsuffix)	rF   r­   Ú
elem_countÚdefault_valueÚis_threadgroupr¸   Úvar_namerÀ   Úvar_defs	            r7   Ú_new_idxvarzMetalKernel._new_idxvarc  sØ   € ô �eœUŸ[™[Ô)Ø×%Ñ% eÓ,ˆEØœd 4×#3Ñ#3Ó4Ð5Ð6ˆÜ�h‰h×%Ñ% h°¸Ó>ˆÙ$2‘.¸ˆØ�e�W˜A˜h˜ZÐ(Ñ(ˆÙØ˜˜4Ÿ:™: jÓ1Ð2°!Ð4Ñ4ˆGØÐ$Ù%ÐTÐ'TÓTÐ%Ø˜˜]˜OÐ,Ñ,ˆGØ×Ñ×$Ñ$ W¨t¯{©{Ñ%:Ô;Øˆ
r9   c                óÌ   — |||f}|| j                   j                  v r| j                   j                  |   S | j                  ||||«      }|| j                   j                  |<   |S )z)Caching wrapper around _reduction_nocache)r¼   Úreduction_cacheÚ_reduction_nocache)rF   r­   r®   Úreduction_typerž  rà   rV   s          r7   Ú	reductionzMetalKernel.reductiony  sf   € ð  °Ð6ˆ	à˜Ÿ™×0Ñ0Ñ0Ø—8‘8×+Ñ+¨IÑ6Ð6Ø×(Ñ(¨°	¸>È5ÓQˆØ.4ˆ�‰× Ñ  Ñ+Øˆr9   c                ó$  — | j                   sJ ‚| j                  rJ ‚dAd„}d}d}| j                  D ]†  }|j                  sŒ|r|dz  }||j                  › d|› �z  }t        |j                  t        j                  «      r||j                  z  }Œ_|t        j                  |j                  › d�dd¬«      z  }Œˆ t        || j                  «      }| j                  |«      }	t        |t        j                  «      r|n| j                  }
t        |t        j                  «      rt        || j                  «      n| j                  }|d	k(  r‰| j!                  |«      }| j"                  j%                  |› d
�«       | j"                  j%                  d«       | j&                  j)                  d|› d|› d�«       | j*                  j%                  d«       |S | j,                  j/                  d«       |dv r­t0        |   }| j!                  ||«      }| j2                  s|}nD|dk(  rdnd\  }}| j!                  ||d¬«      }| j&                  j)                  |› d|› d|› d�«       | j4                  j7                  | j*                  d|› d|› d|› d|› d|	› d�t0        |   ¬«      S |dv r¿| j!                  ||«      }t8        |   }d|› d |› d�}| j2                  s|}nM|d!k(  rd"nd#}d$|› d%|› d&�}| j!                  ||d¬«      }| j&                  j)                  |› d'|› d|› d|› d(�«       | j4                  j7                  | j*                  d|› d|› d|› d|› d|	› d�t0        |   ¬«      S |d)v �r5| j!                  ||«      }| j!                  ||«      }t8        |   }d|› d |› d�}| j2                  s|}dt8        |   › d |› d�}n¢|d*k(  rd"nd#}d$|› d%|› d&�}| j!                  ||d¬«      }| j!                  |d+d¬«      }t;        d,„ | j<                  j?                  «       D «       «      }| j&                  j)                  d-|› d%|› d.|› d|› d/|› d0|› d1|› d0|j                  › d2�«       | j4                  j7                  | j*                  d|› d|› d|› d|› d|› d|› d|	› d�|¬«      S |d3k(  �r+| j2                  s~| j!                  ||
«      }| j&                  j)                  |› d4|› d5|› d�«       | j4                  j7                  | j&                  d|› d|› d|	› d�t@        jB                  ¬«      } ||«      S | j!                  d6|
«      }|› d4|› d7�}| j"                  j)                  |› d8�«       | j&                  j%                  |› d9|› d:|› d;�«       | j4                  j7                  | j*                  d<|› d|	› d�t@        jB                  ¬«      } ||«      S |d=k(  �r4t        |tD        «      sJ d>«       ‚| j!                  d6|
«      }|› d4|› d7�}d?|d+   › d|d   › d|d@   › d�}| j"                  j)                  |› d8�«       | j2                  rC| j"                  j)                  |› d8�«       | j&                  j%                  |› d9|› d|› d(�«       n!| j&                  j%                  |› d0|› d�«       | j4                  j7                  | j2                  r| j*                  n| j&                  d|› d|› d|	› d�t@        jB                  ¬«      } ||«      S tG        |«      ‚)BzeCodegen a reduction operation.
        Only sum and prod operations are somewhat reasonable optimizedc           
     óš   — t        j                  dD �cg c](  }t        | › d|› �| j                  | j                  «      ‘Œ* c}«      S c c}w )NÚxyzrt   )r   Ú_unwrapr   r¸   r­   )Úres3rª  s     r7   Ú_unwrap_helperz6MetalKernel._reduction_nocache.<locals>._unwrap_helper•  sA   € ä×%Ñ%ØNSÖTÈ” ˜v Q q c˜]¨D¯K©K¸¿¹ÕDÒTóð ùÚTs   ”-Ar²  r   ú + rv   ÚnumelT©ÚintegerÚpositiveÚanyz	 = false;z7threadgroup_barrier(metal::mem_flags::mem_threadgroup);z
                if (z) {
                    z' = true;
                }
            Úreduction_utils)ÚprodÚsumrÑ  )r   r[   )r   Ú*F)r¹  rº  rÈ   z= rÉ   zc10::metal::threadgroup_rL   r>   r?   rÇ   )ri   rc   r§   r¨   ri   ÚMaxOpÚMinOpz::c10::metal::ú<z>::identity()z = ::c10::metal::rÌ   )ÚargminÚargmaxr×  r   c              3  ó:   K  — | ]  }|j                   sŒ|–— Œ y ­wrŠ   r§  r©  s     r7   r–   z1MetalKernel._reduction_nocache.<locals>.<genexpr>  s   è ø€ ò ØÀÇÃ”Añùr«  z#
                if (::c10::metal::z>::replace(z)) {
                    r³  z;
                    z$;
                }
                Úwelford_reducer‰  r–  Úfloat3rŠ  z = 0.0;z! = ::c10::metal::welford_combine(z	, float3(z, 0.0, 1.0));z(c10::metal::threadgroup_welford_combine(Úwelford_combinez&Input to welford combine must be tuplezfloat3(r   )rÇ  r   r�   ztuple[CSEVariable, ...])$rœ  Ú
_load_maskr­  r¨  rF  r2   rÊ  ÚsympyÚIntegerÚSymbolÚprefixr
   Úmax_threadgroup_sizerµ  r   Úsimd_group_sizer½  r¶  rÙ   r¾   rØ   r�  r1  r2  r   r~  r¼   r½   r¬   r¬  Úrange_tree_nodesÚvaluesr3   r�  ÚtupleÚNotImplementedError)rF   r­   r®   rÁ  rž  rÈ  Úreduction_idxÚacc_buf_sizeÚrdÚacc_buf_size_strÚacc_buf_alloc_sizeÚshmem_buf_sizeÚaccÚ	acc_dtypeÚacc_bufr6   Údefault_valÚreduction_opÚsrc_metal_typeÚ
cast_valueÚ	op_structÚ	limit_valÚdata_acc_bufÚidx_acc_bufÚidx_valÚidx_varÚwf_resÚacc_thread_varÚ	inp_values                                r7   rÀ  zMetalKernel._reduction_nocache‰  sJ  € ð ×$Ò$Ð$Ð$Ø—?’?Ð"Ð"ó	ð ˆØˆØ×"Ñ"ò 	ˆBØ—?’?ØÙØ Ñ&�Ø §¡˜y¨¨L¨>Ð:Ñ:ˆMä˜"Ÿ(™(¤E§M¡MÔ2Ø §¡Ñ(‘à¤§¡Ø—y‘y�k Ð'°Àô!ñ ‘ð	ô ˜<¨×)BÑ)BÓCˆØŸ:™: lÓ3Ðô ˜,¬¯©Ô6ñ à×*Ñ*ð 	ô ˜,¬¯©Ô6ô �L $×"6Ñ"6Ô7à×%Ñ%ð 	ð ˜UÒ"Ø×"Ñ" 5Ó)ˆCØ×Ñ×(Ñ(¨C¨5°	Ð):Ô;Ø×Ñ×(Ñ(ØIôð �L‰L×ÑðØ�Gð Ø�Eð ðôð �K‰K×!Ñ!ØIôð ˆJà�‰×ÑÐ*Ô+à˜_Ñ,Ü2°9Ñ=ˆIØ×&Ñ& y°.ÓAˆGØ×2Ò2Ø‘ð !/°%Ò 7‘H¸Xñ *�˜\ð ×&Ñ&Ø¨[Èð 'ó �ð —‘×#Ñ# s e¨1¨\¨N¸"¸U¸GÀ1Ð$EÔFà—8‘8×$Ñ$Ø—‘Ø*¨>Ð*:¸!¸G¸9ÀBÀsÀeÈ2ÈmÈ_Ð\^Ð_oÐ^pÐpqÐrÜ0°Ñ7ð %ó ð ð
 ˜^Ñ+Ø×&Ñ& y°.ÓAˆGÜ+¨IÑ6ˆNØ'¨Ð'7°r¸%¸ÀÐBˆJØ×2Ò2Ø ‘à'5¸Ò'>™GÀG�	Ø,¨Y¨K°q¸Ð8HÈÐV�	Ø×&Ñ&Ø¨YÀuð 'ó �ð —‘×#Ñ#Ø�eÐ,¨^Ð,<¸A¸c¸UÀ"ÀZÀLÐPRÐSôð —8‘8×$Ñ$Ø—‘Ø*¨>Ð*:¸!¸G¸9ÀBÀsÀeÈ2ÈmÈ_Ð\^Ð_oÐ^pÐpqÐrÜ0°Ñ7ð %ó ð ð
 Ð1Ò1Ø×+Ñ+¨I°~ÓFˆLØ×*Ñ*¨5°.ÓAˆKÜ+¨IÑ6ˆNØ'¨Ð'7°r¸%¸ÀÐBˆJØ×2Ò2Ø �Ø(¬¸Ñ)>Ð(?¸rÀ-ÀÐPQÐR‘à'5¸Ò'A™GÀw�	Ø,¨Y¨K°q¸Ð8HÈÐV�	Ø×&Ñ&Ø¨YÀuð 'ó �ð ×*Ñ*¨5ÀÐRWÐ*ÓX�Üñ Ø#×4Ñ4×;Ñ;Ó=ôó �ð —‘×#Ñ#ð )#Ø#, +¨Q¨~Ð.>¸kÈ*ÈÐUWÐX[ÐW\ð ]Ø�E˜˜Z˜Lð )Ø�I˜S §¡ ð /ð%ô ð —8‘8×$Ñ$Ø—‘Ø*¨>Ð*:¸!¸L¸>ÈÈKÈ=ÐXZØ�%�r˜'˜ " ] O°2Ð6FÐ5GÀqðJàð	 %ó ð ð Ð-Ó-Ø×2Ò2Ø×*Ñ*¨9Ð6HÓI�Ø—‘×#Ñ# w i¨q°°¸tÀEÀ7È!Ð$LÔMØŸ™×*Ñ*Ø—L‘LØ.¨~Ð.>¸aÀ¸yÈÐK[ÐJ\Ð\]Ð^ÜŸ-™-ð +ó �ñ
 & fÓ-Ð-Ø×&Ñ& xÐ1CÓDˆGØ '˜y¨¨-¨¸Ð:ˆNØ×Ñ×%Ñ%¨Ð(8¸Ð&@ÔAØ�L‰L×"Ñ"Ø!Ð"Ð"CÀNÐCSÐS\Ð]bÐ\cÐcpÐqôð —X‘X×&Ñ&Ø—‘Ø:¸7¸)À2ÐFVÐEWÐWXÐYÜ—m‘mð 'ó ˆFñ
 " &Ó)Ð)ØÐ.Ó.Ü˜e¤UÔ+ÐUÐ-UÓUÐ+Ø×&Ñ& xÐ1CÓDˆGØ '˜y¨¨-¨¸Ð:ˆNØ! %¨¡( ¨2¨e°A©h¨Z°r¸%À¹(¸À1ÐEˆIØ×Ñ×%Ñ%¨Ð(8¸Ð&@ÔAØ×.Ò.Ø×"Ñ"×)Ñ)¨^Ð,<¸GÐ*DÔEØ—‘×&Ñ&Ø%Ð&Ð&GÈÐGWÐWYÐZcÐYdÐdfÐgõð —‘×&Ñ&¨.Ð)9¸¸Y¸KÀqÐ'IÔJØ—X‘X×&Ñ&Ø#×>Ò>�—’ÀDÇLÁLØ*¨>Ð*:¸!¸G¸9ÀBÐGWÐFXÐXYÐZÜ—m‘mð 'ó ˆFñ
 " &Ó)Ð)Ü! .Ó1Ð1r9   c                óÎ  ‡— | j                  ‰j                  «      }| j                  |«      }‰j                  rQt	        ‰j
                  j                  t        j                  «      r\‰j
                  j                  | j                  k  r9| j                  j                  | j                  › d‰j                  › d|› d�«       y t	        ‰j
                  j                  t        j                  «      r‰j
                  j                  n.t        j                  ‰j
                  j                  › d�dd¬«      }t!        ˆfd„| j"                  D «       «      }‰j
                  j                  › d�}|�s\| j"                  j%                  ‰«       |t'        | j                  d	z
  «      z   t'        | j                  «      z  }| j                  |«      }‰j
                  j                  }	| j(                  j                  d
‰j
                  j                  › d‰j
                  j                  › d|› d‰j
                  j                  › d�	«       | j(                  j+                  «       5  t	        |t        j                  «      rS| j(                  j                  | j                  › d|› d| j                  › d‰j
                  j                  › d|	› d�
«       nH| j(                  j                  | j                  › d|› d|› d|	› d‰j
                  j                  › d�
«       t	        |t        j                  «      s|| j                  z  |k7  r"| j(                  j                  d|› d|› d�«       |j-                  ‰j
                  j                  |«      }
| j(                  j                  | j                  › d‰j                  › d|
› d�«       d d d «       y | j(                  j+                  «       5  |j-                  ‰j
                  j                  |«      }
| j(                  j                  | j                  › d‰j                  › d|
› d�«       d d d «       y # 1 sw Y   y xY w# 1 sw Y   y xY w)NrÈ   r³  rÉ   rÊ  TrË  c              3  óN   •K  — | ]  }|j                   ‰j                   u –— Œ y ­wrŠ   ©Úroot)r•   ÚeÚentrys     €r7   r–   z=MetalKernel.codegen_iteration_ranges_entry.<locals>.<genexpr>T  s#   øè ø€ ò %
Ø%&ˆA�F‰F�e—j‘jÔ ñ%
ùs   ƒ"%Ú_linear_idxr   z	for(auto z
_cnt = 0; z_cnt < z; ++z_cnt) {rv   z_cnt + rÉ  z_cnt;rÊ   ú >= z) break;)Úrename_indexingrG   rµ  r¨  r2   r   rÊ  rÝ  rÞ  rá  r¶  rÙ   Úindex_dtyperF  rß  rà  rÎ  r~  Úappendr*   rÛ   rÎ   Úreplace)rF   r  rÁ   Ú	index_strÚacc_sizeÚroot_already_processedÚlinear_idx_nameÚ	loop_sizeÚloop_size_strÚ	root_nameÚsub_index_strs    `         r7   Úcodegen_iteration_ranges_entryz*MetalKernel.codegen_iteration_ranges_entry@  sÔ  ø€ Ø×)Ñ)¨%¯*©*Ó5ˆ
Ø—J‘J˜zÓ*ˆ	à×!Ò!Ü�u—z‘z×'Ñ'¬¯©Ô7Ø—
‘
× Ñ  D×$=Ñ$=Ò=à×Ñ×(Ñ(Ø×#Ñ#Ð$ A e§j¡j \°°Y°K¸qÐAôð ô ˜%Ÿ*™*×*Ñ*¬E¯M©MÔ:ð �J‰J×Òä—‘ §¡×!2Ñ!2Ð 3°5Ð9À4ÐRVÔWð 	ô "%ó %
Ø*.×*IÑ*Iô%
ó "
Ðð #ŸZ™Z×.Ñ.Ð/¨{Ð;ˆâ%Ø×+Ñ+×2Ñ2°5Ô9ð "¤E¨$×*CÑ*CÀaÑ*GÓ$HÑHÌUØ×)Ñ)óNñ ˆIð !ŸJ™J yÓ1ˆMàŸ
™
Ÿ™ˆIà�I‰I×ÑØ˜EŸJ™J×-Ñ-Ð.¨j¸¿¹×9JÑ9JÐ8KÈ7ÐS`ÐRaÐaeÐfk×fpÑfp×fwÑfwÐexð  yAð  Bôð —‘×!Ñ!Ó#ñ Ü˜h¬¯©Ô5Ø—I‘I×'Ñ'Ø×+Ñ+Ð,¨A¨oÐ->¸cØ×4Ñ4Ð5°S¸¿¹×9JÑ9JÐ8KÈ7ÐS\ÐR]Ð]^ð`õð
 —I‘I×'Ñ'Ø×+Ñ+Ð,¨A¨oÐ->¸cÀ-ÀÐPSÐT]ÐS^Ð^aÐbg×blÑbl×bsÑbsÐatÐtyÐzôô ˜x¬¯©Ô6Ø  4×#<Ñ#<Ñ<ÀÒHà—I‘I×'Ñ'¨$¨Ð.?¸tÀHÀ:ÈXÐ(VÔWð !*× 1Ñ 1°%·*±*·/±/À?Ó S�Ø—	‘	×#Ñ#Ø×'Ñ'Ð(¨¨%¯*©*¨°S¸¸ÀqÐIô÷)ð ð2 —‘×!Ñ!Ó#ñ Ø )× 1Ñ 1°%·*±*·/±/À?Ó S�Ø—	‘	×#Ñ#Ø×'Ñ'Ð(¨¨%¯*©*¨°S¸¸ÀqÐIô÷ð ÷3ð ú÷2ð ús   É"E"QÏ'AQÑQÑQ$c                ó˜  ‡— | j                   �rm| j                  j                  «       5  | j                  j                  | j                  «       | j                  j                  | j
                  «       ddd«       | j                  j                  dt        | j                   «      z  «       | j                  j                  t        d„ | j                  j                  j                  «       D «       «      «       | j                   D �cg c]  }|j                  ‘Œ }}| j                   j                  «        | j                  j                  «       D ]'  Št!        ˆfd„|D «       «      sŒ‰j#                  «        Œ) nJ| j                  j                  | j                  «       | j                  j                  | j
                  «       | j                  j                  | j$                  «       | j                  j                  «        | j
                  j                  «        | j$                  j                  «        y# 1 sw Y   �ŒÈxY wc c}w )a  
        Concat output code from index_code, loads, compute, stores,
        suffix into self.body.

        For pointwise kernels, this is called just once at the end.

        For reduction kernels, this generates a loop over the reduction
        axis.
        Nú}c              3  óT   K  — | ]   }t        |t        «      r|n|fD ]  }|–— Œ Œ" y ­wrŠ   )r2   rå  )r•   ÚitemÚvs      r7   r–   z+MetalKernel.codegen_body.<locals>.<genexpr>ž  s:   è ø€ ò àÜ&0°´uÔ&=™dÀDÀ7òð ô ðØñùs   ‚&(c              3  ó:   •K  — | ]  }‰j                   |u –— Œ y ­wrŠ   rÿ  )r•   ÚrÚnodes     €r7   r–   z+MetalKernel.codegen_body.<locals>.<genexpr>¨  s   øè ø€ Ò5¨!�t—y‘y A”~Ñ5ùs   ƒ)r~  rÛ   rÎ   rØ   r‘  r¾   rÙ   rn   r¼   Ú
invalidater	   r¿  rä  r   Úclearrã  rÎ  Úcache_clearr�  )rF   r  Úrootsr  s      @r7   Úcodegen_bodyzMetalKernel.codegen_bodyŠ  s¡  ø€ ð ×*Ó*Ø—‘×!Ñ!Ó#ñ /Ø—	‘	× Ñ  §¡Ô,Ø—	‘	× Ñ  §¡Ô.÷/ð �I‰I×Ñ ¤c¨$×*IÑ*IÓ&JÑ JÔKð
 �H‰H×ÑÜñ à $§¡× 8Ñ 8× ?Ñ ?Ó Aôó ôð &*×%DÑ%DÖE �Q—V“VÐEˆEÐEØ×+Ñ+×1Ñ1Ô3Ø×-Ñ-×4Ñ4Ó6ò '�ÜÓ5¨uÔ5Õ5Ø×$Ñ$Õ&ñ'ð �I‰I×Ñ˜TŸZ™ZÔ(Ø�I‰I×Ñ˜TŸ\™\Ô*Ø�	‰	×Ñ˜Ÿ™Ô%Ø�
‰
×ÑÔØ�‰×ÑÔØ�‰×ÑÕ÷7/ñ /üò  Fs   ©AH:ÄIÈ:Ic                ót  — | j                  «        t        «       }|}t        j                  j                  r|j                  d«       n|j                  d«       | j                  «       }|j                  «       5  t        j                  j                  s'| j                  D ]  }|j                  d|› d�«       Œ nr| j                  D �cg c]  }d|› d�‘Œ
 }}t        |t        t        «      j                  j                  j                  dz  gt        «       «      }|j                  |«       | j                  rwt        j                   d„ | j"                  D «       «      }t%        |t&        j(                  «      rt+        || j,                  «      n| j,                  }	|j                  d|	› d�«       |j                  d	|› d
�«       |j                  «       5  | j.                  j0                  j3                  «       D ]Z  \  }
}|
| j4                  v rŒ| j7                  t        j                  j9                  |
«      «      }|j                  d|› d|› d�«       Œ\ | j.                  j:                  j3                  «       D ]¢  \  }
}t        j                  j9                  |
«      }|t<        j>                  k(  rBt        j                  jA                  |
«      }|�|jC                  «       g k7  rtE        d«      ‚d}n| j7                  |«      }|j                  d|› d|› d�«       Œ¤ | j.                  jF                  jI                  «       D ]  }|j                  d|› d�«       Œ |D ]F  }t%        |jJ                  t&        j(                  «      rŒ(|j                  d|jL                  › d�«       ŒH d| j                  v r|j                  d«       tO        |«      dk  sJ d«       ‚tO        |«      dkD  rdtO        |«      › �nd}tO        |«      dk(  r|d   jP                  nd}| j                  rdnd}|j                  |› d|› d|› �«       | j                  r|j                  |› d�«       ddd«       |j                  d «       |j                  «       5  tO        |«      dkD  rAtS        |«      D ]3  \  }}|j                  d!|jP                  › d"tU        d#|z   «      › d$�«       Œ5 |jW                  | jX                  «       |jW                  | jZ                  «       ddd«       |j                  d%«       ddd«       t        j                  j                  r!|j                  d&«       |j]                  «       S |j                  d'«       |j]                  «       S c c}w # 1 sw Y   �Œ3xY w# 1 sw Y   Œ�xY w# 1 sw Y   ŒƒxY w)(z3Called at the end to generate a final kernel stringz(R"MTL(úcompile_mps_shader('''z#include <c10/metal/z.h>Úincludec              3  óN   K  — | ]  }|j                   sŒ|j                  –— Œ y ­wrŠ   )r¨  rÊ  r©  s     r7   r–   z-MetalKernel.codegen_kernel.<locals>.<genexpr>Î  s   è ø€ ò 1Ø !¸¿»�A—G•Gñ1ùs   ‚%”%z$[[max_total_threads_per_threadgroup(z)]]úkernel void rL   zdevice z* ú,Nzfloat64 is not supported by MPSr*   z	constant zconstant long& znumel,Úerrorz,device c10::metal::ErrorMessages* error_buf,é   z%Up to 3 index variables are supportedr   Úuintr   Ú
thread_posr²  rÈ   z [[thread_position_in_grid]]z- group_pos [[thread_position_in_threadgroup]]rË   r{  z = thread_pos.éx   rÉ   r  z)MTL");ú'''))/r  r   r   rŒ  Úcpp_wrapperrÙ   Úactive_range_treesrÎ   r1  r   r   Ú__file__Úparentr	   rœ  ÚmathrÐ  r­  r2   rÝ  rÞ  rc   rá  rC   Úoutput_buffersÚitemsÚremoved_buffersr‡  r�  Úinput_buffersr3   Úfloat64Útry_get_bufferÚget_sizer›  Úsizevarsrä  rÊ  rà  rn   rF  Ú	enumerateÚchrrØ   r¶  rÛ   rÓ   )rF   rF  ÚcodeÚfn_nameÚidx_varsÚheaderr1  Úheader_contentsÚtotal_reduction_sizeÚthreadgroup_sizeÚouterÚinnerr   r­   Ú	outer_bufrù  Úthread_pos_dtypeÚthread_pos_var_nameÚthread_pos_suffixÚidxrÀ   s                        r7   Úcodegen_kernelzMetalKernel.codegen_kernel²  s/  € à×ÑÔÜÓˆØˆä�7‰7×ÒØ�N‰N˜9Õ%à�N‰NÐ3Ô4à×*Ñ*Ó,ˆØ�[‰[‹]ñ U	 Ü—7‘7×&Ò&Ø"Ÿl™lò G�FØ—N‘NÐ%9¸&¸ÀÐ#EÕFñGð FJÇ\Á\öØ;AÐ*¨6¨(°#Ò6ð�ð ô #1ØÜœ(“^×*Ñ*×1Ñ1×8Ñ8¸9ÑDÐEÜ“Ló#�ð
 —‘˜Ô/à×$Ò$Ü'+§y¡yñ 1Ø%)×%5Ñ%5ô1ó (Ð$ô "Ð"6¼¿¹ÔFô Ð,¨d×.GÑ.GÔHà×2Ñ2ð !ð
 —‘Ø:Ð;KÐ:LÈCÐPôð �N‰N˜\¨'¨°!Ð4Ô5Ø—‘“ñ -Ø$(§I¡I×$<Ñ$<×$BÑ$BÓ$Dò D‘L�E˜5Ø × 4Ñ 4Ñ4Ø Ø $× 1Ñ 1´!·'±'×2CÑ2CÀEÓ2JÓ K�IØ—N‘N W¨Y¨K°r¸%¸ÀÐ#BÕCð	Dð
 %)§I¡I×$;Ñ$;×$AÑ$AÓ$Cò 
F‘L�E˜5ÜŸG™G×-Ñ-¨eÓ4�Eà¤§¡Ò-Ü$%§G¡G×$:Ñ$:¸5Ó$A˜	Ø$Ð,°	×0BÑ0BÓ0DÈÒ0JÜ".Ð/PÓ"QÐQØ$+™	à$(×$5Ñ$5°eÓ$<˜	Ø—N‘N Y¨y¨k¸¸E¸7À!Ð#DÕEð
Fð "ŸY™Y×/Ñ/×6Ñ6Ó8ò ?�EØ—N‘N _°U°G¸1Ð#=Õ>ð?ð  (ò Q�GÜ! '§-¡-´·±Ô?ØàŸ™¨¸¿¹Ð8HÈÐ'OÕPð	Qð ˜dŸl™lÑ*Ø—N‘NÐ#QÔRä˜8“} qÒ(ÐQÐ*QÓQÐ(ä.1°(«m¸aÒ.?�dœ3˜x›=˜/Ñ*ÀVð !ô ),¨H«¸Ò(:�H˜Q‘K×$Ò$Àð $ð ,0×+@Ò+@¡CÀbÐ!Ø—‘Ø'Ð(¨Ð*=Ð)>Ð>ZÐ[lÐZmÐnôð ×(Ò(Ø—N‘NØ+Ð,Ð,YÐZô÷W-ð\ �N‰N˜5Ô!Ø—‘“ñ 'Ü�x“= 1Ò$Ü$-¨hÓ$7ò ™˜˜SØŸ™Ø# C§H¡H :¨^¼CÀÀcÁ	»NÐ;KÈ1ÐMõðð —‘˜D×.Ñ.Ô/Ø—‘˜DŸI™IÔ&÷'ð �N‰N˜3Ô÷kU	 ôn �7‰7×ÒØ�N‰N˜9Ô%ð �}‰}‹Ðð �N‰N˜6Ô"à�}‰}‹Ðùòo÷2-ñ -ú÷^'ð 'ú÷[U	 ð U	 úsR   Á:AV.Ã
VÃC?V.ÇJVÑ$)V.ÒBV"ÔV.ÖV.ÖV	ÖV.Ö"V+	Ö'V.Ö.V7c           
     ó	  — t         j                  j                  }| j                  j                  D ]  }|j                  |«       Œ | j                  j                  «       \  }}}}t        ||«      D �	�
ci c]  \  }	}
t        |	«      |
“Œ }}	}
g | j                  j                  j                  «       ¢| j                  j                  j                  «       ¢}|D �cg c]  }|| j                  vsŒ|‘Œ }}|| j                  j                  D �cg c]  }t        |«      ‘Œ c}z  }|D �cg c]  }||   ‘Œ	 }}| j                  D ]Ú  }t        |j                  t         j"                  t$        f«      rŒ.t        |j                  t         j&                  «      r|j                  }n4t         j                  j                  j)                  ||«      j*                  }|j,                  r| j.                  sŒ¬|j1                  t        |«      «       |j1                  t$        «       ŒÜ t         j                  j2                  r| j4                  n| j6                  }dd„}t9        | j;                  «       «      dkD  r‡| j;                  «       D �cg c]@  } ||j,                  r t=        |j                  | j>                  «      n|j                  «      ‘ŒB }}|j1                   ||d«      «       |j1                  t@        «       n%t         j                  j2                  rtC        d«      ‚| j.                  r}| j;                  «       D �cg c]6  }|j,                  r& |t=        |j                  | j>                  «      «      nd‘Œ8 }}|j1                   ||d«      «       |j1                  t@        «       n1t         j                  j2                  r|dgz  }|j1                  d«       d| jD                  v rMt         j                  j2                  s3|j1                  d	t9        |D �cg c]  }|€Œd
|vsŒ|‘Œ c}«      › �«       |jG                  ||tI        jJ                  d«      d|¬«       yc c}
}	w c c}w c c}w c c}w c c}w c c}w c c}w )z0
        Codegens a call to this kernel
        Úthreadsc                ó¾   — t         j                  j                  r(| D �cg c]  }d|› d�‘Œ
 } }ddj                  | «      › d�S |› ddj                  | «      › d�S c c}w )Nzstatic_cast<uint64_t>(r?   ú{r>   r  z=[rŠ  )r   rŒ  r+  Újoin)rJ  Úkwargrª  s      r7   Úformat_threadsz/MetalKernel.call_kernel.<locals>.format_threadsA  sh   € Ü�w‰w×"Ò"ØBIÖJ¸QÐ3°A°3°aÒ8ÐJ�ÐJØ˜DŸI™I gÓ.Ð/¨rÐ2Ð2à˜  4§9¡9¨WÓ#5Ð"6°aÐ8Ð8ùò Ks   ŸAr   zWe should always have threads?Ú1Ú
group_sizeNr%  zerror_buf_idx=ú=rx  F)ÚdeviceÚtritonÚ	arg_types)rJ  z	list[str]rN  r5   r�   r5   )&r   rŒ  Úwrapper_coderC   r7  Úensure_size_computedÚpython_argdefsÚzipr5   r0  Úkeysr3  r2  r­  r2   rÊ  rÝ  rÞ  r'   rß  Úgenerate_numel_exprrB  r¨  rœ  r  r+  ÚcexprÚpexprrn   r,  r
   rá  rQ   r›  r1  Úgenerate_kernel_callr3   rS  )rF   rF  r  Údeallocate_wsÚwrapperr  Ú_Ú	call_argsrU  Úcall_argÚarg_typeÚarg_name_to_typerC   rW   ÚtreerG   Úexpr_printerrO  rJ  s                      r7   Úcall_kernelzMetalKernel.call_kernel  sÂ  € ô —'‘'×&Ñ&ˆà—‘×#Ñ#ò 	,ˆAØ×(Ñ(¨Õ+ð	,ð &*§Y¡Y×%=Ñ%=Ó%?Ñ"ˆˆ9�a˜ä>AÀ)ÈYÓ>W÷
Ù(:¨°(ŒC�‹M˜8Ñ#ð
Ðñ 
ð S�—‘×)Ñ)×.Ñ.Ó0ÐR°4·9±9×3JÑ3J×3OÑ3OÓ3QÐRˆØ#ÖG˜ s°$×2FÑ2FÒ'F’ÐGˆÐGØ §¡×!3Ñ!3Ö4˜A”�Q•Ò4Ñ4ˆØ6:Ö;¨sÐ% cÓ*Ð;ˆ	Ð;ð ×$Ñ$ò 	&ˆDÜ˜$Ÿ*™*¤u§}¡}´cÐ&:Ô;àÜ˜DŸJ™J¬¯©Ô5Ø—z‘z‘ä—w‘w×+Ñ+×?Ñ?ÀÀdÓK×QÑQ�à×$Ò$¨×(=Ó(=Ø—‘œC ›IÔ&Ø× Ñ ¤Õ%ð	&ô &'§W¡W×%8Ò%8�t—z’z¸d¿j¹jˆó	9ô ˆt×&Ñ&Ó(Ó)¨AÒ-ð ×0Ñ0Ó2öð ñ à—~’~ô ˜Ÿ™ ×!:Ñ!:Ô;àŸ™õðˆGð ð �K‰K™ w°	Ó:Ô;Ø×ÑœTÕ"ä�w‰w×"Ò"Ü"Ð#CÓDÐDà× Ò ð
 ×0Ñ0Ó2ö	ð ð —>’>ñ œS §¡¨$×*CÑ*CÓDÔEàñðˆGð ð �K‰K™ w°Ó=Ô>Ø×ÑœTÕ"ä�w‰w×"Ò"ð ˜˜‘�Ø× Ñ  Ô&ð �d—l‘lÑ"¬1¯7©7×+>Ò+>Ø�K‰KØ ¤°TÖ%`¨c¸S¹_ÐQTÐ\_ÒQ_¢cÒ%`Ó!aÐ bÐcôð 	×$Ñ$ØØÜ—<‘< Ó&ØØð 	%õ 	
ùóS
ùò
 HùÚ4ùÚ;ùò8ùò ùò& &as=   Á5QÃQ$Ã0Q$ÄQ)Ä*Q.ÊAQ3Ì=;Q8ÐQ=ÐQ=Ð#Q=c                ó(  — t         j                  j                  r-| j                  j	                  | j
                  d|› d�d¬«       y | j                  j                  d«       | j
                  j                  d|› d�dt        |«      › d�d	d
g«       y )Nzif (!ú) returnF©Ú
assignmentr%  rË   z"    TORCH_REPORT_ERROR(error_buf, rÌ   ú    return;r  )
r   rŒ  r+  r¼   r½   r¾   r1  r2  r×   Úrepr)rF   ÚcondÚmsgs      r7   Údevice_assert_asynczMetalKernel.device_assert_asyncy  s€   € Ü�7‰7×ÒØ�H‰H×Ñ˜dŸl™l¨e°D°6¸Ð,BÈuÐÕUà�L‰L×Ñ˜WÔ%Ø�L‰L×#Ñ#à˜D˜6 Ð&Ø8¼¸c»¸À2ÐFØ!Øð	õr9   c                óª  — |s|sy | j                  |«      }| j                  |«      }|r|rd|› d|› d|› d�}n|r|› d�}n|› d|› �}t        j                  j                  r-| j                  j                  | j                  d|› d�d¬	«       y | j                  j                  d
«       | j                  j                  d|› d�d|› d|› d�ddg«       y )NrL   z < 0 || r  r?   z < 0rÊ   rj  Frk  r%  rË   z,    TORCH_REPORT_ERROR(error_buf, "Index ", z, " out of range [0, ", z, ")");rm  r  )
rº   r   rŒ  r+  r¼   r½   r¾   r1  r2  r×   )rF   rG   ÚsizeÚlowerÚupperÚexpr_strÚsize_strÚ	conditions           r7   Úcheck_boundszMetalKernel.check_bounds‡  sù   € ñ ™Øà×$Ñ$ TÓ*ˆØ×$Ñ$ TÓ*ˆñ ‘Uà˜H˜: X¨h¨Z°t¸H¸:ÀQÐG‰IÙØ#˜* DÐ)‰Ià#˜* D¨¨
Ð3ˆIô �7‰7×ÒØ�H‰H×ÑØ—‘  Y K¨xÐ8ÀUð õ ð
 �L‰L×Ñ˜WÔ%Ø�L‰L×#Ñ#à˜9˜+ TÐ*ØBÀ8À*ÐLdÐemÐdnÐnuÐvØ!Øð	õr9   )rƒ  zdict[str, sympy.Expr]r„  r   r�   rv  )r­   rt  r�   r5   )rF  r5   r’  rœ   r�   r   rŠ   )
rF  r5   r’  rœ   rž  r   rŸ  r    r�   rv  )rF  r5   r’  rœ   rž  r   r�   rv  )r­   zstr | torch.dtyper¸  z
int | Noner¹  z
Any | Nonerº  r$   r¸   zValueRanges[Any]r�   r   )
r­   rt  r®   rt  rÁ  r   rž  ú%CSEVariable | tuple[CSEVariable, ...]r�   rz  )r  r   r�   rv  ru  )Úgenerated_kernel)rF  r5   r�   r5   rs  )rF  r5   r  r   r_  r$   r�   rv  )ro  r   rp  r5   r�   rv  )
rG   rœ   rs  rœ   rt  r$   ru  r$   r�   rv  )'rž   rŸ   r    r¡   r¥   Ú	overridesr·  Únewvar_prefixrá  râ  r   rD   r]  r   r\  r;   rµ  Úkexprr	   r1  Ú__annotations__r~  r�  r‡  r”  r¤  r¯  r   Úunknownr½  rÂ  rÀ  r  r  rH  rh  rq  ry  Ú__classcell__©r…  s   @r7   rz  rz    s  ø… ÙEà€IØ€FØ€MØÐØ€OÙ‹O×#Ñ#€EÙ‹L× Ñ €EÙÓ×&Ñ&€EØ€EÙ)¨7¨)Ó4€Gˆ_Ó4Ø=?ÐÐ :Ó?ð-à%ð-ð ð-ð 
õ	-ó%ó@ð SWð<Øð<Ø *ð<Ø3>ð<ØFOð<à	ó<ó*	8ð "&Ø$(Ø#Ø#6 ;×#6Ñ#6Ó#8ðà ðð ðð "ð	ð
 ðð !ðð 
óð,àðð ðð &ð	ð
 5ðð 
/óð u2àðu2ð ðu2ð &ð	u2ð
 5ðu2ð 
/óu2ónHóT&ôPhðV BFð[
Øð[
Ø"ð[
Ø:>ð[
à	ó[
ózð"Øð"Ø&0ð"Ø9=ð"ØFJð"à	÷"r9   rz  c                  óL   ‡ — e Zd ZU eZdZded<   dˆ fd„Z	 	 	 	 	 	 	 	 dd„Zˆ xZ	S )ÚMetalSchedulingr   r'   Ú_kernel_fn_counterc                ó$   •— t         ‰| �  |«       y rŠ   )r€  r�  )rF   Ú	schedulerr…  s     €r7   r�  zMetalScheduling.__init__°  s   ø€ Ü‰Ñ˜Õ#r9   c           	     óâ  — t         j                  j                  }||j                  v r|j                  |   S t         j                  j                  rXd|j                  «       › �}|}||j                  |<   d|› d�|z   }t        ||«      \  }}|j                  |||› d|› �d¬«       |S d| j                  › �}	| xj                  dz  c_        |	|j                  |<   d	}
|j                  |
«      t        |
«      z   }|j                  d
«      }||| }g }|j                  d«      D ]G  }|j                  «       j                  d«      rŒ#|j                  |j!                  dd|	› d�«      «       ŒI t#        t%        |j&                  «      «      }|j(                  j+                  |	› d|	›d�«       |D ]  }|j(                  j+                  |«       Œ |j(                  j+                  d|› d�«       |	S )NÚmps_lib_zconst char* z
_source = ú
F)ÚgpuÚgenerated_kernel_r   r   r*  z#includezkernel void generated_kernel(r#  rL   z = async_compile.metal(z, '''z''', r?   )r   rŒ  rV  Úsrc_to_kernelr+  Únext_kernel_suffixr   Údefine_kernelr…  r’  rn   ÚrindexÚsplitÚstripÚ
startswithr  r  rn  Úsortedr1  r=  rÙ   )rF   Úsrc_codeÚnode_scheduler¹   r`  Úmps_lib_nameÚkernel_nameÚoriginsÚdetailed_originsr;  Úmetal_src_startÚstartÚendÚ	metal_srcÚ
body_linesr“  Úheaders_reprs                    r7   r�  zMetalScheduling.define_kernel³  s  € ô —'‘'×&Ñ&ˆØ�w×,Ñ,Ñ,Ø×(Ñ(¨Ñ2Ð2ä�7‰7×Òà% g×&@Ñ&@Ó&BÐ%CÐDˆLØ&ˆKØ.9ˆG×!Ñ! (Ñ+Ø% l ^°:Ð>ÀÑIˆHÜ(;¸MÈ7Ó(SÑ%ˆGÐ%Ø×!Ñ!Ø˜h¨7¨)°2Ð6FÐ5GÐ(HÈeð "ô ð Ðð & d×&=Ñ&=Ð%>Ð?ˆØ×Ò 1Ñ$ÕØ*1ˆ×Ñ˜hÑ'ð 3ˆØ—‘˜Ó/´#°oÓ2FÑFˆØ�o‰o˜fÓ%ˆØ˜U 3Ð'ˆ	ð ˆ
Ø—O‘O DÓ)ò 	ˆDØ�z‰z‹|×&Ñ& zÔ2ØØ×ÑØ—‘Ð<ÀÈWÈIÐUVÐ>WÓXõð	ô œF 6§>¡>Ó2Ó3ˆØ�‰× Ñ  G 9Ð,CÀGÀ;ÈeÐ!TÔUØò 	+ˆDØ�N‰N×$Ñ$ TÕ*ð	+à�‰× Ñ  5¨¨°aÐ!8Ô9àˆr9   )r‡  zScheduler | Noner�   rv  )r•  r5   r–  zlist[SchedulerNode]r¹   rz  r�   r5   )
rž   rŸ   r    rz  Úkernel_typer…  r  r�  r�  r�  r‚  s   @r7   r„  r„  ¬  s=   ø… Ø€KØÐ˜Óõ$ð.Øð.Ø,?ð.ØITð.à	÷.r9   r„  )r6   z&float | int | bool | str | CSEVariabler�   r5   )IÚ
__future__r   rn  rÏ   Úloggingr/  Úpathlibr   Útypingr   r   rÝ  Úsympy.printing.precedencer   r3   Útorch.utils._cpp_embed_headersr   Útorch.utils._ordered_setr	   Útorch.utils._sympy.functionsr
   Útorch.utils._sympy.printersr   r   ÚExprPrinter_Útorch.utils._sympy.value_rangesr   r}  r   r   r   Úvirtualizedr   r   r   Úcommonr   r   r   r   r   r   Úsimdr   r   r   Úops_handlerr   r    r‡  r!   r"   r#   Ú	getLoggerrž   rª   r$   Úint8Úint16Úint32Úint64Úuint8r*   r+   r�  r¬   r8   r;   r¥   Ú_initialize_pointwise_overridesrr  rz  r„  r]   r9   r7   ú<module>r¸     s'  ðõ #ã Û Û Û Ý ß %ã Ý 0ã Ý 9Ý /Ý ,ß OÝ 7ç GÑ Gß ,Ñ ,÷÷ ÷ CÑ Bñ ß6ß4Ýà€g×Ñ˜Ó!€ð 
‡J�J�Ø	‡J�J�Ø	‡K�K�Ø	‡K�K�Ø	‡K�K�Ø	‡K�K�Ø	‡K�K�Ø	‡J�J�Ø	‡N�N�Hð
€óôt"�|ô t"ônX�[ô Xðv
 × .Ñ .¨uÔ 5Ø × &Ñ &Ô (ôN
�*ô N
ôb5�nõ 5r9   