ó
    Eñi²  ã                  óæ  • S SK Jr  S SKrS SKrS SKrS SKrS SKJr  S SKJ	r	J
r
Jr  S SKrS SKJr  S SKrS SKJr  S SKJr  S SKJrJr  S S	KJr  S
SKJrJrJr  S
SKJrJ r J!r!  SSK"J#r#J$r$J%r%J&r&J'r'J(r(  SSK)J*r*J+r+J,r,  \(       a  S SKJ-r-  S
SK.J/r/J0r0  S
SK1J2r2J3r3  SSK"J4r4  \Rj                  " \65      r7\Rp                  S\Rr                  S\Rt                  S\Rv                  S\Rx                  S\Rz                  S\R|                  S\R~                  S\R€                  S0	rAS'S jrB " S S\5      rC " S  S!\'5      rD\DR‹                  S"5        \DR�                  5          " S# S$\+5      rG " S% S&\,5      rHg)(é    )ÚannotationsN)ÚPath)ÚAnyÚOptionalÚTYPE_CHECKING)Ú
PRECEDENCE)Ú_embed_headers)Ú
OrderedSet)Ú
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)ÚUnion)ÚReductionTypeÚ	StoreMode)Ú	SchedulerÚSchedulerNode)ÚOpVarTÚboolÚcharÚshortÚintÚlongÚucharÚfloatÚhalfÚbfloatc                óú   • [        U [        5      (       a<  U [        R                  :X  a  gU [        R                  * :X  a  gX :w  a  g[	        U 5      $ [        U [
        5      (       a  U (       a  S$ S$ [	        U 5      $ )NÚ	HUGE_VALFz
-HUGE_VALFÚNANÚtrueÚfalse)Ú
isinstancer+   ÚtorchÚinfÚstrr%   )Úvals    ÚX/home/mande/repo/quber/.venv/lib/python3.13/site-packages/torch/_inductor/codegen/mps.pyÚvalue_to_metalr9   8   sc   € Ü�#”u×ÑØ”%—)‘)ÓØØ”U—Y‘Y�JÓØØ‹ZØÜ�3‹xˆÜ	�Cœ×	Ñ	ÞˆvÐ) 'Ð)Üˆs‹8€Oó    c                  ó²   • \ rS rSrSrSS jrSS jrSS jrSS jrSS jr	SS jr
SS	 jrSS
 jrSS jrSS jrSS jrSS jr\rSS jrSS jrSS jrSrg)ÚMetalExprPrinteréF   z/Converts sympy expression to Metal code snippetc                ó¨   • UR                   u  p#U R                  U5      nU R                  U5      nUR                  (       a	  SU SU S3$ SU SU S3$ )Núc10::metal::floor_divide(ú, Ú)úmetal::floor(ú) / (©ÚargsÚdoprintÚ
is_integer)ÚselfÚexprÚxÚdivs       r8   Ú_print_FloorDivÚ MetalExprPrinter._print_FloorDivI   sY   € Ø—‘‰ˆØ�L‰L˜‹OˆØ�l‰l˜3ÓˆØ�?�?Ø.¨q¨c°°C°5¸Ð:Ð:Ø˜q˜c  s e¨1Ð-Ð-r:   c                ó8  • UR                   u  p#nUS:H  =(       a
    XDS-
  -  S:g  nU R                  U5      nUS:w  a5  U R                  U5      nUR                  (       a
  SU SU S3nO	SU SU S3nU R                  U5      nU(       a	  SU S	U S3$ SU S
U S3$ )Ni   r   r   Ú(rC   rA   rB   zc10::metal::safe_mod(r@   z) % (rD   )rH   rI   rJ   rK   ÚmodÚuse_safe_mods         r8   Ú_print_ModularIndexingÚ'MetalExprPrinter._print_ModularIndexingQ   s´   € Ø—i‘i‰ˆ�à˜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Ð!Ð!r:   c                óÐ   • [        UR                  5      S:w  a  [        S5      e[        U R                  UR                  5      u  p#SU SU SU S3nSU SU SU S3nSU SU S3$ )	Nr   z$metal::min only supported for 2 argsústatic_cast<decltype(Ú+ú)>(rA   zmetal::min(r@   ©ÚlenrE   ÚRuntimeErrorÚmapÚ_print©rH   rI   ÚaÚbÚ
typecast_aÚ
typecast_bs         r8   Ú
_print_MinÚMetalExprPrinter._print_Minb   ó}   € Üˆt�y‰y‹>˜QÓÜÐEÓFÐFä�4—;‘; §	¡	Ó*‰ˆØ,¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø,¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø˜Z˜L¨¨:¨,°aÐ8Ð8r:   c                óÐ   • [        UR                  5      S:w  a  [        S5      e[        U R                  UR                  5      u  p#SU SU SU S3nSU SU SU S3nSU SU S3$ )	Nr   z$metal::max only supported for 2 argsrU   rV   rW   rA   zmetal::max(r@   rX   r]   s         r8   Ú
_print_MaxÚMetalExprPrinter._print_Maxk   rd   r:   c                ó|   • [        UR                  5      S:X  d   eSU R                  UR                  S   5       S3$ )Nr   úmetal::abs(r   rA   ©rY   rE   r\   ©rH   rI   s     r8   Ú
_print_AbsÚMetalExprPrinter._print_Abst   s9   € Ü�4—9‘9‹~ Ó"Ð"Ð"à˜TŸ[™[¨¯©°1©Ó6Ð7°qÐ9Ð9r:   c                ó|   • [        UR                  5      S:X  d   eSU R                  UR                  S   5       S3$ )Nr   zstatic_cast<long>(metal::rint(r   ú))rj   rk   s     r8   Ú_print_RoundToIntÚ"MetalExprPrinter._print_RoundToInty   s9   € Ü�4—9‘9‹~ Ó"Ð"Ð"à/°·±¸D¿I¹IÀa¹LÓ0IÐ/JÈ"ÐMÐMr:   c                óð   • [        UR                  5      S:X  d   eUR                  u  p#UR                  (       a  US:  d   e[        SU S35      eU R	                  U[
        S   5      nSU SU SU*  S	3$ )
Nr   r   zOFor integer inputs, only non-negative ndigits are currently supported, but got Ú.ÚMulz!static_cast<float>(metal::rint(1eú * z) * 1erA   )rY   rE   rG   Ú
ValueErrorÚparenthesizer   )rH   rI   ÚnumberÚndigitsÚ
number_strs        r8   Ú_print_RoundDecimalÚ$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[Ð\Ð\r:   c                ól   • UR                   u  p#SU R                  U5       SU R                  U5       S3$ )Nústatic_cast<float>(z) / static_cast<float>(rA   )rE   r\   )rH   rI   ÚlhsÚrhss       r8   Ú_print_IntTrueDivÚ"MetalExprPrinter._print_IntTrueDivŠ   s;   € Ø—9‘9‰ˆð % T§[¡[°Ó%5Ð$6Ð6MÈdÏkÉkÐZ]ÓN^ÐM_Ð_`ÐaÐar:   c                óŽ   • [        UR                  5      S:X  d   e[        U R                  UR                  5      u  p#SU SU S3$ )Nr   z'metal::precise::pow(static_cast<float>(z), static_cast<float>(ro   )rY   rE   r[   rF   )rH   rI   rJ   Úys       r8   Ú_print_PowByNaturalÚ$MetalExprPrinter._print_PowByNatural�   sG   € Ü�4—9‘9‹~ Ó"Ð"Ð"Ü�4—<‘< §¡Ó+‰ˆØ8¸¸Ð;QÐRSÐQTÐTVÐWÐWr:   c                ó€   • [        UR                  5      S:X  d   eU R                  UR                  S   5      nSU S3$ )Nr   r   r~   rA   ©rY   rE   rF   ©rH   rI   rJ   s      r8   Ú_print_ToFloatÚMetalExprPrinter._print_ToFloat•   s=   € Ü�4—9‘9‹~ Ó"Ð"Ð"Ø�L‰L˜Ÿ™ 1™Ó&ˆØ$ Q C qÐ)Ð)r:   c                ób   • UR                   (       a  [        [        U5      5      $ [        U5      $ ©N)rG   r6   r(   rk   s     r8   Ú_print_FloatÚMetalExprPrinter._print_Floatš   s#   € Ø�?�?ô ”s˜4“y“>Ð!ä�t“9Ðr:   c                ó€   • [        UR                  5      S:X  d   eU R                  UR                  S   5      nSU S3$ )Nr   r   z1static_cast<int>(metal::floor(static_cast<float>(z)))rˆ   r‰   s      r8   Ú_print_FloorToIntÚ"MetalExprPrinter._print_FloorToInt£   s=   € Ü�4—9‘9‹~ Ó"Ð"Ð"Ø�L‰L˜Ÿ™ 1™Ó&ˆØBÀ1À#ÀSÐIÐIr:   c                ó€   • [        UR                  5      S:X  d   eU R                  UR                  S   5      nSU S3$ )Nr   r   zstatic_cast<int>(metal::trunc(ro   rˆ   r‰   s      r8   Ú_print_TruncToIntÚ"MetalExprPrinter._print_TruncToIntª   s=   € Ü�4—9‘9‹~ Ó"Ð"Ð"Ø�L‰L˜Ÿ™ 1™Ó&ˆØ/°¨s°"Ð5Ð5r:   c                ó€   • [        UR                  5      S:X  d   eU R                  UR                  S   5      nSU S3$ )Nr   r   zmetal::precise::log2(rA   rˆ   r‰   s      r8   Ú_print_OpaqueUnaryFn_log2Ú*MetalExprPrinter._print_OpaqueUnaryFn_log2¯   s=   € Ü�4—9‘9‹~ Ó"Ð"Ð"Ø�L‰L˜Ÿ™ 1™Ó&ˆØ& q c¨Ð+Ð+r:   c                óJ   ^ • U 4S jUR                    5       u  p#nU SU SU 3$ )Nc              3  ó\   >#   • U  H!  nTR                  U[        S    S-
  5      v •  M#     g7f)ÚAtomg      à?N)rw   r   )Ú.0ÚargrH   s     €r8   Ú	<genexpr>Ú0MetalExprPrinter._print_Where.<locals>.<genexpr>µ   s.   øé € ð 
ÚHQÀˆD×Ñ˜c¤:¨fÑ#5¸Ñ#;×<Ð<Ê	ùs   ƒ),ú ? z : )rE   )rH   rI   ÚcÚpÚqs   `    r8   Ú_print_WhereÚMetalExprPrinter._print_Where´   s3   ø€ ô
ØHLÏ	Ê	ó
‰ˆˆað ��C˜�s˜#˜a˜SÐ!Ð!r:   © N)rI   ú
sympy.ExprÚreturnr6   )Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__Ú__doc__rL   rR   rb   rf   rl   rp   r{   r�   r…   rŠ   rŽ   r‘   Ú_print_floorr”   r—   r¤   Ú__static_attributes__r¦   r:   r8   r<   r<   F   s\   † Ù9ô.ô"ô"9ô9ô:ô
Nô

]ôbôXô
*ô
ôJð
 %€Lô6ô
,÷
"r:   r<   c                  óÆ  • \ rS rSrSr\  S2         S3S jj5       r\        S4S j5       r\S5S j5       r\S6S j5       r	\S7S j5       r
\S8S	 j5       r\S9S
 j5       r\S:S j5       r\S:S j5       r\S:S j5       r\S:S j5       r\S;S j5       r\S;S j5       r\S;S j5       r\S;S j5       r\S;S j5       r\S;S j5       r\S;S j5       r\S;S j5       r\S;S j5       r\S;S j5       r\S;S j5       r\S;S j5       r\S;S j5       r\S<S j5       r\S;S j5       r\S;S j5       r \S;S j5       r!\S;S  j5       r"\S;S! j5       r#\S:S" j5       r$\S;S# j5       r%\S;S$ j5       r&\S:S% j5       r'\S;S& j5       r(\S:S' j5       r)\S;S( j5       r*\S=S) j5       r+\S=S* j5       r,\          S>S+ j5       r-\S;S, j5       r.\S:S- j5       r/S?S. jr0S@S/ jr1\2SAS0 j5       r3S1r4g)BÚMetalOverridesé»   zXImplements Metal-specific overrides for ops. Base class emits Python-friendly overrides.Nc                ó€   • U[         R                  :X  a  [        R                  S5        SU  S3$ S[        U    SU  S3$ )Nz>float64 cast requested, probably from tensorify_python_scalarsr~   rA   ústatic_cast<ú>()r4   ÚdoubleÚlogÚwarningÚDTYPE_TO_METAL)rJ   ÚdtypeÚ	src_dtypeÚuse_compute_typess       r8   Úto_dtypeÚMetalOverrides.to_dtype¾   sK   € ð ”E—L‘LÓ Ü�K‰KØPôð )¨¨¨1Ð-Ð-Øœn¨UÑ3Ð4°B°q°c¸Ð;Ð;r:   c                ó6   • S[         U    S[         U    SU  S3$ )Nzas_type<z>(static_cast<rµ   ro   ©r¹   )rJ   rº   r»   s      r8   Úto_dtype_bitcastÚMetalOverrides.to_dtype_bitcastÌ   s/   € ð œ.¨Ñ/Ð0°¼~ÈiÑ?XÐ>YÐY[Ð\]Ð[^Ð^`ÐaÐar:   c                ó   • [        U 5      $ r�   ©r9   )r7   rº   s     r8   ÚconstantÚMetalOverrides.constantÒ   s   € ä˜cÓ"Ð"r:   c                ó:  • [         R                  R                  [         R                  R                  U 5      5      n[         R                  R                  R                  [         R                  R                  U[        U 5      S9n[        R                  " X15      $ )N)Úbounds)
r   ÚkernelÚindex_to_strÚprepare_indexingÚcseÚgenerateÚcomputer   r   r½   )rI   rº   Úidx_strÚvars       r8   Ú
index_exprÚMetalOverrides.index_exprÖ   sj   € ä—(‘(×'Ñ'¬¯©×(AÑ(AÀ$Ó(GÓHˆÜ�h‰h�l‰l×#Ñ#Ü�H‰H×Ñ˜gÔ.CÀDÓ.Ið $ð 
ˆô �|Š|˜CÓ'Ð'r:   c           	     ó  • [        U5      n[        5       n[        R                  R	                  U5         UR                  5          [        R                  " 5       [        R                  R                  l	        S[        R                  R                  l
        U" 5       nS S S 5        S S S 5        U  SUR                  5        SU 3n[        R                  R                  R                  U5      nU(       Gd_  [        R                  R                  R                  WR                  S9n[        R                  R                  R                  Xg5        [        R                  R                   R#                  [$        UR                      SU S3SU  S3/5        [        R                  R                   R                  5          [        R                  R                   R'                  U5        [        R                  R                   R)                  U SU S	U S
35        S S S 5        [        R                  R                   R)                  SU SU S	U S
35        U$ ! , (       d  f       GNÁ= f! , (       d  f       GNË= f! , (       d  f       Ng= f)NÚtmp_scoped_Ú:©rº   Ú Ú;úif (ú) {z = static_cast<decltype(rW   ú);z} else )r9   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           r8   ÚmaskedÚMetalOverrides.maskedÞ   sÖ  € ô
 # 5Ó)ˆ	Ü$Ó&ˆÜ�X‰X×"Ñ" ;Õ/°×1CÑ1CÕ1Eô ,5¯?ª?Ó+<ŒA�H‰H�L‰LÔ(Ø'4ŒA�H‰H�L‰LÔ$Ù“ˆB÷ 2F×/ð �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ôð ˆ
÷3 2FÖ1Eú×/Ö/ú÷" +Õ*ús1   µI+ÁAIÂI+Ç AI=É
I(	É#I+É+
I:É=
Jc           	     ó0   • U  SU SU S[        U5       S3$ )Nr    z : static_cast<decltype(rW   rA   rÄ   )r^   r_   r¡   s      r8   ÚwhereÚMetalOverrides.where   s(   € à��C˜�sÐ2°1°#°S¼ÈÓ9JÐ8KÈ1ÐMÐMr:   c                ó   • SU  SU S3$ )Nzc10::metal::remainder(r@   rA   r¦   ©r^   r_   s     r8   Ú	remainderÚMetalOverrides.remainder  s   € à'¨ s¨"¨Q¨C¨qÐ1Ð1r:   c                óD   • SU  SU SU  S3nSU  SU SU S3nSU SU S3$ )NrU   rV   rW   rA   zc10::metal::max(r@   r¦   ©r^   r_   r`   ra   s       r8   ÚmaximumÚMetalOverrides.maximum  óK   € à,¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø,¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø! * ¨R°
¨|¸1Ð=Ð=r:   c                óD   • SU  SU SU  S3nSU  SU SU S3nSU SU S3$ )NrU   rV   rW   rA   zc10::metal::min(r@   r¦   rú   s       r8   ÚminimumÚMetalOverrides.minimum  rý   r:   c                ó   • U  SU 3$ )Nz || r¦   rö   s     r8   Ú
logical_orÚMetalOverrides.logical_or  ó   € à��D˜˜ˆ}Ðr:   c                ó   • U  SU 3$ )Nz && r¦   rö   s     r8   Úlogical_andÚMetalOverrides.logical_and  r  r:   c                ó   • SU  S3$ )Nzmetal::isnan(rA   r¦   ©rJ   s    r8   ÚisnanÚMetalOverrides.isnan  ó   € à˜q˜c Ð#Ð#r:   c                ó   • SU  S3$ )Nzmetal::isinf(rA   r¦   r	  s    r8   ÚisinfÚMetalOverrides.isinf   r  r:   c                ó   • SU  S3$ )Nzmetal::precise::log(rA   r¦   r	  s    r8   r·   ÚMetalOverrides.log$  ó   € à% a S¨Ð*Ð*r:   c                ó   • SU  S3$ )Nzmetal::precise::exp(rA   r¦   r	  s    r8   ÚexpÚMetalOverrides.exp(  r  r:   c                ó   • SU  S3$ )Nri   rA   r¦   r	  s    r8   ÚabsÚMetalOverrides.abs,  s   € à˜Q˜C˜qÐ!Ð!r:   c                ó   • SU  S3$ )Nzmetal::signbit(rA   r¦   r	  s    r8   ÚsignbitÚMetalOverrides.signbit0  s   € à    1Ð%Ð%r:   c                ó   • SU  S3$ )Nzmetal::precise::sin(rA   r¦   r	  s    r8   ÚsinÚMetalOverrides.sin4  r  r:   c                ó   • SU  S3$ )Nzc10::metal::sinc(rA   r¦   r	  s    r8   ÚsincÚMetalOverrides.sinc8  s   € à" 1 # QÐ'Ð'r:   c                ó   • SU  S3$ )Nzmetal::precise::cos(rA   r¦   r	  s    r8   ÚcosÚMetalOverrides.cos<  r  r:   c                ó   • SU  S3$ )Nzmetal::precise::tan(rA   r¦   r	  s    r8   ÚtanÚMetalOverrides.tan@  r  r:   c                ó   • SU  S3$ )Nzmetal::precise::asin(rA   r¦   r	  s    r8   ÚasinÚMetalOverrides.asinD  ó   € à& q c¨Ð+Ð+r:   c                ó   • SU  S3$ )Nzmetal::precise::acos(rA   r¦   r	  s    r8   ÚacosÚMetalOverrides.acosH  r+  r:   c                ó   • SU  S3$ )Nzmetal::precise::atan(rA   r¦   r	  s    r8   ÚatanÚMetalOverrides.atanL  r+  r:   c                ó   • SU  SU S3$ )Nz::metal::precise::atan2(r@   rA   r¦   )rJ   r„   s     r8   Úatan2ÚMetalOverrides.atan2P  s   € à)¨!¨¨B¨q¨c°Ð3Ð3r:   c                ó   • SU  S3$ )Nzmetal::precise::sqrt(rA   r¦   r	  s    r8   ÚsqrtÚMetalOverrides.sqrtT  r+  r:   c                ó   • SU  SU  S3$ )NrU   z)>(-rA   r¦   r	  s    r8   ÚnegÚMetalOverrides.negX  s   € ð ' q c¨¨a¨S°Ð2Ð2r:   c                ó   • SU  S3$ )Nzmetal::precise::rsqrt(rA   r¦   r	  s    r8   ÚrsqrtÚMetalOverrides.rsqrt^  ó   € à'¨ s¨!Ð,Ð,r:   c                ó   • SU  S3$ )Nzmetal::precise::tanh(rA   r¦   r	  s    r8   ÚtanhÚMetalOverrides.tanhb  r+  r:   c                ó   • SU  S3$ )Nzmetal::precise::atanh(rA   r¦   r	  s    r8   ÚatanhÚMetalOverrides.atanhf  r>  r:   c                ó   • SU  SU S3$ )Nr?   r@   rA   r¦   rö   s     r8   ÚfloordivÚMetalOverrides.floordivj  s   € ð +¨1¨#¨R°¨s°!Ð4Ð4r:   c                ó   • SU  S3$ )NrB   rA   r¦   r	  s    r8   ÚfloorÚMetalOverrides.flooro  r  r:   c                ó   • SU  S3$ )Nzmetal::sign(rA   r¦   r	  s    r8   ÚsignÚMetalOverrides.signs  ó   € à˜a˜S Ð"Ð"r:   c                óD   • SU  SU SU  S3nSU  SU SU S3nSU SU S3$ )NrU   rV   rW   rA   zmetal::fmod(r@   r¦   rú   s       r8   ÚfmodÚMetalOverrides.fmodw  sK   € à,¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø,¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø˜j˜\¨¨J¨<°qÐ9Ð9r:   c                ó   • SU  S3$ )Númetal::trunc(rA   r¦   r	  s    r8   ÚtruncÚMetalOverrides.trunc}  r  r:   c                óÀ   • U  SU 3nU R                   b  U R                   R                  (       d(  UR                   b!  UR                   R                  (       a  SU S3$ U$ )Nz / rS  rA   )rº   Úis_floating_point)r^   r_   Úquots      r8   ÚtruncdivÚMetalOverrides.truncdiv�  sQ   € à��C˜�sˆ|ˆØ�G‰GÑ A§G¡G×$=×$=Ø�G‰GÑ A§G¡G×$=×$=à" 4 &¨Ð*Ð*Øˆr:   c                ó   • SU  S3$ )Nzmetal::ceil(rA   r¦   r	  s    r8   ÚceilÚMetalOverrides.ceilŠ  rN  r:   c                óf   • [         R                  R                  R                  S5        SU  SU S3$ )NÚrandomzc10::metal::rand(r@   rA   ©r   rÉ   ÚheadersÚadd©ÚseedÚoffsets     r8   ÚrandÚMetalOverrides.randŽ  s/   € ä	�‰×Ñ×Ñ˜XÔ&Ø" 4 &¨¨6¨(°!Ð4Ð4r:   c                óf   • [         R                  R                  R                  S5        SU  SU S3$ )Nr_  zc10::metal::randn(r@   rA   r`  rc  s     r8   ÚrandnÚMetalOverrides.randn“  s/   € ä	�‰×Ñ×Ñ˜XÔ&Ø# D 6¨¨F¨8°1Ð5Ð5r:   c           	     ór   • [         R                  R                  R                  S5        SU  SU SU SU S3	$ )Nr_  zc10::metal::randint64(r@   rA   r`  )rd  re  ÚlowÚhighs       r8   Ú	randint64ÚMetalOverrides.randint64˜  s=   € ô 	
�‰×Ñ×Ñ˜XÔ&Ø'¨ v¨R°¨x°r¸#¸¸bÀÀÀaÐHÐHr:   c                ó   • SU  S3$ )Nzmetal::rint(rA   r¦   r	  s    r8   ÚroundÚMetalOverrides.roundŸ  rN  r:   c                óD   • SU  SU SU  S3nSU  SU SU S3nSU SU S3$ )NrU   rV   rW   rA   zmetal::pow(r@   r¦   )r^   r_   Úcast_aÚcast_bs       r8   ÚpowÚMetalOverrides.pow£  sK   € à(¨¨¨1¨Q¨C¨s°1°#°QÐ7ˆØ(¨¨¨1¨Q¨C¨s°1°#°QÐ7ˆØ˜V˜H B v h¨aÐ0Ð0r:   c                óf   • [         R                  R                  R                  S5        SU SU S3$ )NÚspecial_mathúc10::metal::rO   rA   r`  )rH   r^   Únames      r8   Ú_special_unaryÚMetalOverrides._special_unary©  s/   € Ü	�‰×Ñ×Ñ˜^Ô,Ø˜d˜V 1 Q C qÐ)Ð)r:   c                ól   • [         R                  R                  R                  S5        SU SU SU S3$ )Nry  rz  rO   r@   rA   r`  )rH   r^   r_   r{  s       r8   Ú_special_binaryÚMetalOverrides._special_binary­  s5   € Ü	�‰×Ñ×Ñ˜^Ô,Ø˜d˜V 1 Q C r¨!¨¨AÐ.Ð.r:   c           
     óì  • S H,  n[        X[        R                  " U R                  US95        M.     [        R                  " U R                  SS9U l        S H0  n[        U U[        R                  " U R                  US-   S95        M2     S H,  n[        X[        R                  " U R
                  US95        M.     S H0  n[        U U[        R                  " U R
                  US-   S95        M2     g )N)ÚerfÚerfinvÚi0Úi0eÚi1Úi1eÚdigammaÚspherical_bessel_j0)r{  Ú	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Úpartialmethodr|  Úlgammar  )Úclsr{  s     r8   Ú_initialize_special_opsÚ&MetalOverrides._initialize_special_ops±  sä   € ó	
ˆDô �Cœy×6Ò6°s×7IÑ7IÐPTÑUÖVñ	
ô ×,Ò,¨S×-?Ñ-?ÀkÑRˆŒ
ó
ˆDô ØØÜ×'Ò'¨×(:Ñ(:ÀÈ
ÑARÑSöñ
ó&
ˆDô �Cœy×6Ò6°s×7JÑ7JÐQUÑVÖWñ
ó
ˆDô ØØÜ×'Ò'¨×(;Ñ(;À$ÈÑBSÑTöò
r:   r¦   ©NT)
rJ   r   rº   útorch.dtyper»   zOptional[torch.dtype]r¼   r%   r¨   r6   )rJ   r   rº   r¬  r»   r¬  r¨   r6   )r7   zUnion[bool, float, int]rº   r¬  r¨   r6   )rI   r§   rº   r¬  r¨   r6   )ré   r   rê   r§   rë   r   r¨   r6   )r^   r$   r_   r$   r¡   r$   r¨   r6   )r^   r$   r_   r$   r¨   r6   )r^   r   r_   r   r¨   r6   )rJ   r   r¨   r6   )rJ   r   r„   r   r¨   r6   )rd  r   re  r   r¨   r6   )
rd  r   re  r   rl  r   rm  r   r¨   r6   )r^   r   r{  r6   r¨   r6   )r^   r   r_   r   r{  r6   r¨   r6   ©r¨   ÚNone)5r©   rª   r«   r¬   r­   Ústaticmethodr½   rÁ   rÅ   rÑ   rð   ró   r÷   rû   rÿ   r  r  r
  r  r·   r  r  r  r  r   r#  r&  r)  r-  r0  r3  r6  r9  r<  r@  rC  rF  rI  rL  rP  rT  rY  r\  rf  ri  rn  rq  rv  r|  r  Úclassmethodr©  r¯   r¦   r:   r8   r±   r±   »   s÷  † Ùbàð ,0Ø"&ð	<Øð<àð<ð )ð<ð  ð	<ð
 
ô<ó ð<ð ðbØðbØ*ðbØ7Bðbà	óbó ðbð
 ó#ó ð#ð ó(ó ð(ð óó ððB ó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ô
*ô/ð ó=ó ó=r:   r±   Úmpsc                  ó
  ^ • \ rS rSr% Sr\rSrSrSr	Sr
\" 5       R                  r\" 5       R                  r\" 5       R                  r\r\" S/5      rS\S	'   / rS
\S'         SU 4S jjrSS jrSS jr S         S S jjrS!S jrSSS\R:                  " 5       4           S"S jjr          S#S jr          S#S jr S$S jr!S%S jr"SS&S jjr# S'       S(S jjr$          S)S jr%Sr&U =r'$ )*ÚMetalKerneliö  z;Implement Metal codegen based on the SIMDKernel abstractionrØ   úauto i   é    ÚutilszOrderedSet[str]ra  zlist[IterationRangesEntry]Úmultistage_reduction_entryc                ó\   >• [         TU ]  " U40 UD6  [        R                  " 5       U l        g r�   )ÚsuperÚ__init__rÞ   rß   Úacc_var_ids)rH   ÚtilingÚkwargsÚ	__class__s      €r8   rº  ÚMetalKernel.__init__  s&   ø€ ô
 	‰Ò˜Ñ* 6Ò*Ü$Ÿ?š?Ó,ˆÕr:   c                ó   • [         U   $ r�   rÀ   )rH   rº   s     r8   Údtype_to_strÚMetalKernel.dtype_to_str  s   € Ü˜eÑ$Ð$r:   c                ó‚  • U R                   R                  U5      nU R                  U5      n[        R                  R                  U5      nU SU R                  U5       S3nU[        R                  [        R                  4;   a  SU S3n[        R                  nU R                  R                  U R                  XTS9$ )z"Codegen a load from an InputBufferÚ[Ú]r~   rA   rÖ   )rE   ÚinputrË   r   ÚgraphÚ	get_dtyperÊ   r4   Úfloat16Úbfloat16Úfloat32rÌ   rÍ   Úloads)rH   r{  ÚindexrÐ   rº   Úlines         r8   ÚloadÚMetalKernel.load  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Ð Ð?Ð?r:   Nc                ór  • U R                   R                  U5      nU R                  U5      nU R                  [        R
                  R                  U5      5      nSU SU S3nUc  U SU R                  U5       SU S3nO\US:X  aH  U R                  R                  S5        S	U S
3n	SU	 SU S3n
U	 SU
 SU R                  U5       SU S3nO[        SU 35      eU R                  (       a%  U R                  R                  [        X5      5        g U R                  R                  [        X5      5        g )Nr´   rµ   rA   rÄ  ú] = rØ   Ú
atomic_addÚatomiczc10::metal::AtomicType<Ú>zreinterpret_cast<device z
::type *>(z::atomic_add(r@   rÛ   zUnimplemented store mode )rE   ÚoutputrË   rÁ  r   rÇ  rÈ  rÊ   ra  rb  rZ   Úinside_reductionrÎ   rè   r   Ústores)rH   r{  rÍ  ÚvalueÚmoderÐ   Ú	dtype_strÚcast_valrÎ  Úatomic_typeÚcast_vars              r8   ÚstoreÚMetalKernel.store  s8  € ð �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×"Ñ"¤<°Ó#;Õ<à�K‰K×!Ñ!¤,¨tÓ":Õ;r:   c                ó–  • U R                   R                  U5      nU R                  U5      nU R                  [        R
                  R                  U5      5      n[        S U R                   5       5      nU SU R                  U5       SU SU S3nSUR                   SU 3nU R                  R                  [        X5      5        g )Nc              3  óJ   #   • U  H  oR                   (       d  M  Uv •  M     g 7fr�   ©Úis_reduction©rœ   Úts     r8   rž   Ú.MetalKernel.store_reduction.<locals>.<genexpr>8  s   é € ÐKÒ(8 1¿N½NŸQ™QÒ(8ùó   ‚#š	#rÄ  z] = static_cast<rµ   rÛ   rÙ   z == 0) )rE   rÖ  rË   rÁ  r   rÇ  rÈ  ÚnextÚrange_treesrÊ   r{  rØ  rè   r   )rH   r{  rÍ  rÙ  rÐ   rÛ  Úreduction_dimrÎ  s           r8   Ústore_reductionÚMetalKernel.store_reduction3  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Õ7r:   Tc                ó¼  • [        U[        R                  5      (       a  U R                  U5      nS[	        U R
                  5       3n[        R                  R                  XeU5      nU(       a  SOSnX� SU 3-  nU(       a  USU R                  U5       S3-  nUb  U(       a   S5       eUSU 3-  nU R                  R                  X€R                  -   5        U$ )	NÚtmp_acc_zthreadgroup Ú r×   rÄ  rÅ  z+Thread group var can not have default valueú = )r3   r4   rº   rÁ  ré  r»  r   rÉ   Úcreate_cse_varÚsexprÚindexing_coderè   Úsuffix)	rH   rº   Ú
elem_countÚdefault_valueÚis_threadgrouprÈ   Úvar_namerÐ   Úvar_defs	            r8   Ú_new_idxvarÚMetalKernel._new_idxvar>  sÕ   € ô �eœUŸ[™[×)Ñ)Ø×%Ñ% eÓ,ˆEØœd 4×#3Ñ#3Ó4Ð5Ð6ˆÜ�h‰h×%Ñ% h¸Ó>ˆÞ$2‘.¸ˆØ�W˜A˜h˜ZÐ(Ñ(ˆÞØ˜˜4Ÿ:™: jÓ1Ð2°!Ð4Ñ4ˆGØÑ$Þ%ÐTÐ'TÓTÐ%Ø˜˜]˜OÐ,Ñ,ˆGØ×Ñ×$Ñ$ W¯{©{Ñ%:Ô;Øˆ
r:   c                óÆ   • X#U4nXPR                   R                  ;   a  U R                   R                  U   $ U R                  XX45      nX`R                   R                  U'   U$ )z)Caching wrapper around _reduction_nocache)rÌ   Úreduction_cacheÚ_reduction_nocache)rH   rº   r»   Úreduction_typerÙ  rï   Úresults          r8   Ú	reductionÚMetalKernel.reductionT  s\   € ð °Ð6ˆ	àŸ™×0Ñ0Ó0Ø—8‘8×+Ñ+¨IÑ6Ð6Ø×(Ñ(¨¸>ÓQˆØ.4�‰× Ñ  Ñ+Øˆr:   c                ó¾  • U R                   (       d   eU R                  (       a   eSAS jnSnSnU R                   H”  nUR                  (       d  M  U(       a  US-  nXhR                   SU 3-  n[        UR                  [        R                  5      (       a  XxR                  -  nMn  U[        R                  " UR                   S3SSS9-  nM–     [        R                  " XpR                  5      nU R                  U5      n	[        U[        R                  5      (       a  [        XpR                  5      OU R                  n
US	:X  a‰  U R!                  U5      nU R"                  R%                  U S
35        U R"                  R%                  S5        U R&                  R)                  SU SU S35        U R*                  R%                  S5        U$ U R,                  R/                  S5        US;   a«  [0        U   nU R!                  XÊ5      nU R2                  (       d  UnOAUS:X  a  SOSu  nnU R!                  XÏSS9nU R&                  R)                  U SU SU S35        U R4                  R7                  U R*                  SU SU SU SU SU	 S3[0        U   S9$ US;   aÎ  U R!                  X*5      n[8        U   nSU S U S3nU R2                  (       d  UnO[UR;                  S!5      (       a  S"OS!nS#U S$U S%3nU R!                  UUSS9nU R&                  R)                  U S&U SU SU S'35        U R4                  R7                  U R*                  SU SU SU SU SU	 S3[0        U   S9$ US(;   Gab  U R!                  X*5      nU R!                  X5      n[8        U   nSU S U S3nU R2                  (       d  UnS[8        U    S U S3nOÏUR;                  S!5      (       a  S"OS!nS#U S$U S%3nU R!                  UUSS9nU R!                  US)SS9n[=        S* U R>                  RA                  5        5       5      nUS+:X  a  S,OS-nURB                  (       a  S.U S/3OSnU R&                  R)                  SU SU SU U SU S0U S1U S0UR                   S235        U R4                  R7                  U R*                  SU SU SU SU SU SU SU	 S3US9$ US3:X  Ga)  U R2                  (       dz  U R!                  X'5      nU R&                  R)                  U S4U S5U S35        U R4                  R7                  U R&                  SU SU SU	 S3[D        RF                  S9nU" U5      $ U R!                  S6U5      nU S4U S73nU R"                  R)                  U S835        U R&                  R%                  U S9U S:U S;35        U R4                  R7                  U R*                  S<U SU S3[D        RF                  S9nU" U5      $ US=:X  Ga@  [        U[H        5      (       d   S>5       eU R!                  S6U5      nU S4U S73nS?US)    SUS    SUS@    S3nU R"                  R)                  U S835        U R2                  (       aC  U R"                  R)                  U S835        U R&                  R%                  U S9U SU S'35        O!U R&                  R%                  U S0U S35        U R4                  R7                  U R2                  (       a  U R*                  OU R&                  SU SU SU	 S3[D        RF                  S9nU" U5      $ [K        U5      e)Bz]Codegen a reduction operation.
Only sum and prod operations are somewhat reasonable optimizedc           
     ó¢   • [         R                  " S Vs/ s H)  n[        U  SU 3U R                  U R                  5      PM+     sn5      $ s  snf )NÚxyzrs   )r   Ú_unwrapr   rÈ   rº   )Úres3ræ  s     r8   Ú_unwrap_helperÚ6MetalKernel._reduction_nocache.<locals>._unwrap_helperp  sE   € ä×%Ò%ÙNSÓTÊeÈ” ˜v Q q c˜]¨D¯K©K¸¿¹ÖDÉeÑTóð ùÚTs   •0Arð  r   ú + ru   ÚnumelT©ÚintegerÚpositiveÚanyz	 = false;z7threadgroup_barrier(metal::mem_flags::mem_threadgroup);z
                if (z) {
                    z' = true;
                }
            Úreduction_utils)ÚprodÚsumr  )r   rV   )r   Ú*F)r÷  rø  r×   z= rØ   zc10::metal::threadgroup_rO   r@   rA   rÖ   )ÚmaxÚminr´   rµ   r  Úlowestz::metal::numeric_limits<z>::z()z = ::c10::metal::rÛ   )ÚargminÚargmaxr   c              3  óJ   #   • U  H  oR                   (       d  M  Uv •  M     g 7fr�   rã  rå  s     r8   rž   Ú1MetalKernel._reduction_nocache.<locals>.<genexpr>Ù  s   é € ð Ú=˜!ÇÅ—A‘AÒ=ùrè  r  rÕ  Ú<z || ::metal::isnan(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ä  r{  r3   r  ÚsympyÚIntegerÚSymbolÚprefixÚMinÚmax_threadgroup_sizeró  r   Úsimd_group_sizerû  rô  rè   rÎ   rç   rØ  ra  rb  r   r·  rÌ   rÍ   r¹   Úendswithré  Úrange_tree_nodesÚvaluesrW  r4   rË  ÚtupleÚNotImplementedError)rH   rº   r»   r   rÙ  r	  Úreduction_idxÚacc_buf_sizeÚrdÚacc_buf_size_strÚshmem_buf_sizeÚaccÚ	acc_dtypeÚacc_bufr7   Údefault_valÚreduction_opÚsrc_metal_typeÚ
cast_valueÚlim_fnÚ	limit_valÚdata_acc_bufÚidx_acc_bufÚidx_valÚidx_varÚcmp_opÚ
nan_suffixÚwf_resÚacc_thread_varÚ	inp_values                                 r8   rÿ  ÚMetalKernel._reduction_nocached  sT  € ð ×$×$Ð$Ð$Ø—?—?Ð"Ð"ô	ð ˆØˆØ×"Ô"ˆBØ—?—?ÙÞØ Ñ&�Ø§¡˜y¨¨L¨>Ð:Ñ:ˆMä˜"Ÿ(™(¤E§M¡M×2Ñ2Ø§¡Ñ(’à¤§¢Ø—y‘y�k Ð'°Àñ!ñ ’ñ #ô —y’y ×/HÑ/HÓIˆØŸ:™: 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Ø ‘à%3×%<Ñ%<¸U×%CÑ%C™È�Ø6°~Ð6FÀcÈ&ÈÐQSÐT�	Ø×&Ñ&Ø¨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‘à%3×%<Ñ%<¸U×%CÑ%C™È�Ø6°~Ð6FÀcÈ&ÈÐQSÐT�	Ø×&Ñ&Ø¨YÀuð 'ð �ð ×*Ñ*¨5ÀÐRWÐ*ÐX�Üñ Ø#×4Ñ4×;Ñ;Ô=óó �ð !/°(Ó :™À�ð !×2×2ð *¨%¨°Ñ3àð ð
 —‘×#Ñ#ð )Ø�G˜1˜V˜H A c U¨:¨,ð 7Ø�E˜˜U˜Gð $Ø�I˜S §¡ ð /ð%ô ð —8‘8×$Ñ$Ø—‘Ø*¨>Ð*:¸!¸L¸>ÈÈKÈ=ÐXZØ�%�r˜'˜ " ] O°2Ð6FÐ5GÀqðJàð	 %ð ð ð Ð-Ô-Ø×2×2Ø×*Ñ*¨9ÓC�Ø—‘×#Ñ# w i¨q°°¸tÀEÀ7È!Ð$LÔMØŸ™×*Ñ*Ø—L‘LØ.¨~Ð.>¸aÀ¸yÈÐK[ÐJ\Ð\]Ð^ÜŸ-™-ð +ð �ñ
 & fÓ-Ð-Ø×&Ñ& x°Ó>ˆGØ '˜y¨¨-¨¸Ð:ˆNØ×Ñ×%Ñ%¨Ð(8¸Ð&@ÔAØ�L‰L×"Ñ"Ø!Ð"Ð"CÀNÐCSÐS\Ð]bÐ\cÐcpÐqôð —X‘X×&Ñ&Ø—‘Ø:¸7¸)À2ÀlÀ^ÐSTÐUÜ—m‘mð 'ð ˆFñ
 " &Ó)Ð)ØÐ.Ô.Ü˜e¤U×+Ñ+ÐUÐ-UÓUÐ+Ø×&Ñ& x°Ó>ˆ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Ð1r:   c                ó	  ^• U R                  TR                  5      nU R                  U5      nTR                  (       aW  [	        TR
                  R                  [        R                  5      (       a]  TR
                  R                  U R                  ::  a9  U R                  R                  U R                   STR                   SU S35        g [	        TR
                  R                  [        R                  5      (       a  TR
                  R                  O,[        R                  " TR
                  R                   S3SSS9n[!        U4S jU R"                   5       5      nTR
                  R                   S3nU(       Gdf  U R"                  R%                  T5        U['        U R                  S	-
  5      -   ['        U R                  5      -  nU R                  U5      nTR
                  R                  n	U R(                  R                  S
TR
                  R                   STR
                  R                   SU STR
                  R                   S3	5        U R(                  R+                  5          [	        U[        R                  5      (       aS  U R(                  R                  U R                   SU SU R                   STR
                  R                   SU	 S3
5        OHU R(                  R                  U R                   SU SU SU	 STR
                  R                   S3
5        [	        U[        R                  5      (       d  XpR                  -  U:w  a"  U R(                  R                  SU SU S35        UR-                  TR
                  R                  U5      n
U R(                  R                  U R                   STR                   SU
 S35        S S S 5        g U R(                  R+                  5          UR-                  TR
                  R                  U5      n
U R(                  R                  U R                   STR                   SU
 S35        S S S 5        g ! , (       d  f       g = f! , (       d  f       g = f)Nr×   rñ  rØ   r  Tr  c              3  óR   >#   • U  H  oR                   TR                   L v •  M     g 7fr�   )Úroot)rœ   ÚeÚentrys     €r8   rž   Ú=MetalKernel.codegen_iteration_ranges_entry.<locals>.<genexpr>-  s    øé € ð %
Ú*I Q�F‰F�e—j‘jÕ Ò*Iùs   ƒ$'Ú_linear_idxr   z	for(auto z
_cnt = 0; z_cnt < z; ++z_cnt) {ru   z_cnt + r  z_cnt;rÙ   ú >= z) break;)Úrename_indexingrI   ró  rä  r3   rG  r  r!  r"  r&  rô  rè   Úindex_dtyper{  r#  r$  r  r·  Úappendr+   rê   rÝ   Úreplace)rH   rI  rÑ   Ú	index_strÚacc_sizeÚroot_already_processedÚlinear_idx_nameÚ	loop_sizeÚloop_size_strÚ	root_nameÚsub_index_strs    `         r8   Úcodegen_iteration_ranges_entryÚ*MetalKernel.codegen_iteration_ranges_entry  sÌ  ø€ Ø×)Ñ)¨%¯*©*Ó5ˆ
Ø—J‘J˜zÓ*ˆ	à×!×!Ü�u—z‘z×'Ñ'¬¯©×7Ñ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Ñ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Ñ6Ø ×#<Ñ#<Ñ<ÀÓHà—I‘I×'Ñ'¨$¨Ð.?¸tÀHÀ:ÈXÐ(VÔWð !*× 1Ñ 1°%·*±*·/±/À?Ó S�Ø—	‘	×#Ñ#Ø×'Ñ'Ð(¨¨%¯*©*¨°S¸¸ÀqÐIô÷) $Ð#ð2 —‘×!Ñ!Õ#Ø )× 1Ñ 1°%·*±*·/±/À?Ó S�Ø—	‘	×#Ñ#Ø×'Ñ'Ð(¨¨%¯*©*¨°S¸¸ÀqÐIô÷ $Ð#÷3 $Õ#ú÷2 $Õ#ús   É6E,Q-ÐAQ>Ñ-
Q;Ñ>
Rc                ó6  • U R                   (       Ga8  U R                  R                  5          U R                  R                  U R                  5        U R                  R                  U R
                  5        SSS5        U R                  R                  S[        U R                   5      -  5        U R                  R                  [        S U R                  R                  R                  5        5       5      5        U R                   (       a;  U R                   R                  5       R                  5         U R                   (       a  M;  OJU R                  R                  U R                  5        U R                  R                  U R
                  5        U R                  R                  U R                  5        U R                  R!                  5         U R
                  R!                  5         U R                  R!                  5         g! , (       d  f       GN˜= f)zØ
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  ój   #   • U  H)  n[        U[        5      (       a  UOU4  H  nUv •  M	     M+     g 7fr�   )r3   r+  )rœ   ÚitemÚvs      r8   rž   Ú+MetalKernel.codegen_body.<locals>.<genexpr>w  s7   é € ð â A˜Ü&0°´u×&=Ñ&=™dÀDÀ7ÓJ˜õ áJñ Ú Aùs   ‚13)r·  rê   rÝ   rç   rÌ  rÎ   rè   rY   rÌ   Ú
invalidater
   rþ  r*  ÚpopÚcache_clearrØ  Úclear)rH   s    r8   Úcodegen_bodyÚMetalKernel.codegen_bodyc  sc  € ð ×*×*Ð*Ø—‘×!Ñ!Õ#Ø—	‘	× Ñ  §¡Ô,Ø—	‘	× Ñ  §¡Ô.÷ $ð �I‰I×Ñ ¤c¨$×*IÑ*IÓ&JÑ JÔKð
 �H‰H×ÑÜñ à $§¡× 8Ñ 8× ?Ñ ?Ô Aóó ôð ×1×1Ø×/Ñ/×3Ñ3Ó5×AÑAÔCð ×1×1Ñ1øð �I‰I×Ñ˜TŸZ™ZÔ(Ø�I‰I×Ñ˜TŸ\™\Ô*Ø�	‰	×Ñ˜Ÿ™Ô%Ø�
‰
×ÑÔØ�‰×ÑÔØ�‰×ÑÕ÷1 $Ö#ús   ­AH	È	
Hc                ó  • U R                  5         [        5       n[        R                  R                  (       a  UR                  S5        OUR                  S5        U R                  5       nUR                  5          [        R                  R                  (       d)  U R                   H  nUR                  SU S35        M     OuU R                   Vs/ s H	  nSU S3PM     nn[        U[        [        5      R                  R                  R                  S-  /[        5       5      nUR                  U5        U R                  (       a|  [        R                   " S U R"                   5       5      n[%        U[&        R(                  5      (       a  [+        XpR,                  5      OU R,                  nUR                  SU S35        UR                  S	5        UR                  5          U R.                  R0                  R3                  5        H\  u  pšX�R4                  ;   a  M  U R7                  [        R                  R9                  U	5      5      nUR                  S
U SU
 S35        M^     U R.                  R:                  R3                  5        H¥  u  pš[        R                  R9                  U	5      nU[<        R>                  :X  aD  [        R                  RA                  U	5      nUb  URC                  5       / :w  a  [E        S5      eSnOU R7                  U5      nUR                  SU SU
 S35        M§     U R.                  RF                  RI                  5        H  n
UR                  SU
 S35        M     U HM  n[%        URJ                  [&        R(                  5      (       a  M.  UR                  SURL                   S35        MO     SU R                  ;   a  UR                  S5        [O        U5      S:  d   S5       e[O        U5      S:”  a  S[O        U5       3OSn[O        U5      S:X  a  US   RP                  OSnU R                  (       a  SOSnUR                  U SU SU 35        U R                  (       a  UR                  U S35        SSS5        UR                  S5        UR                  5          [O        U5      S:”  aC  [S        U5       H4  u  nnUR                  S URP                   S![U        S"U-   5       S#35        M6     URW                  U RX                  5        URW                  U RZ                  5        SSS5        UR                  S$5        SSS5        [        R                  R                  (       a!  UR                  S%5        UR]                  5       $ UR                  S&5        UR]                  5       $ s  snf ! , (       d  f       GN@= f! , (       d  f       NŸ= f! , (       d  f       N—= f)'z3Called at the end to generate a final kernel stringz(R"MTL(zcompile_mps_shader('''z#include <c10/metal/z.h>Úincludec              3  ó^   #   • U  H#  oR                   (       d  M  UR                  v •  M%     g 7fr�   )rä  r  rå  s     r8   rž   Ú-MetalKernel.codegen_kernel.<locals>.<genexpr>£  s   é € ð 1Ú%5 ¿½“G�A—G–GÒ%5ùs   ‚-š-z$[[max_total_threads_per_threadgroup(z)]]zkernel void generated_kernel(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");z'''))/re  r   r   rÇ  Úcpp_wrapperrè   Úactive_range_treesrÝ   ra  r	   r   Ú__file__Úparentr
   r×  Úmathr  rê  r3   r!  r"  r  r&  rE   Úoutput_buffersÚitemsÚremoved_buffersrÁ  rÈ  Úinput_buffersr4   Úfloat64Útry_get_bufferÚget_sizerZ   Úsizevarsr*  r  r$  rY   r{  Ú	enumerateÚchrrç   rô  rê   râ   )rH   r{  ÚcodeÚidx_varsÚheaderra  Úheader_contentsÚtotal_reduction_sizeÚthreadgroup_sizeÚouterÚinnerrÛ  rº   Ú	outer_bufr>  Úthread_pos_dtypeÚthread_pos_var_nameÚthread_pos_suffixÚidxrÐ   s                       r8   Úcodegen_kernelÚMetalKernel.codegen_kernelˆ  sâ  € à×ÑÔÜÓˆä�7‰7××Ø�N‰N˜9Õ%à�N‰NÐ3Ô4à×*Ñ*Ó,ˆØ�[‰[�]Ü—7‘7×&×&Ø"Ÿlœl�FØ—N‘NÐ%9¸&¸ÀÐ#EÖFò +ð FJÇ\Â\óÚEQ¸6Ð*¨6¨(°#Ó6Á\ð ð ô #1ØÜœ(“^×*Ñ*×1Ñ1×8Ñ8¸9ÑDÐEÜ“Ló#�ð
 —‘˜Ô/à×$×$Ü'+§y¢yñ 1Ø%)×%5Ò%5ó1ó (Ð$ô "Ð"6¼¿¹×FÑFô Ð,×.GÑ.GÔHà×2Ñ2ð !ð
 —‘Ø:Ð;KÐ:LÈCÐPôð �N‰NÐ:Ô;Ø—‘•Ø$(§I¡I×$<Ñ$<×$BÑ$BÖ$D‘L�EØ× 4Ñ 4Ó4Ù Ø $× 1Ñ 1´!·'±'×2CÑ2CÀEÓ2JÓ K�IØ—N‘N W¨Y¨K°r¸%¸ÀÐ#BÖCñ	 %Eð
 %)§I¡I×$;Ñ$;×$AÑ$AÖ$C‘L�EÜŸ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ñ %Dð "ŸY™Y×/Ñ/×6Ñ6Ö8�EØ—N‘N _°U°G¸1Ð#=Ö>ñ 9ó  (�GÜ! '§-¡-´·±×?Ñ?ÙàŸ™¨¸¿¹Ð8HÈÐ'OÖPñ	  (ð ˜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öñ %8ð —‘˜D×.Ñ.Ô/Ø—‘˜DŸI™IÔ&÷ ð �N‰N˜3Ô÷k ôn �7‰7××Ø�N‰N˜9Ô%ð �}‰}‹Ðð �N‰N˜6Ô"à�}‰}‹Ðùòo÷2 –ú÷^ •ú÷[ �]úsR   Á=AW3ÃWÃ$DW3Ç)J.WÒ)W3Ó B	W"Õ	W3×W3×
W	×W3×"
W0	×,W3×3
Xc           
     óÖ	  • [         R                  R                  nU R                  R                   H  nUR                  U5        M     U R                  R                  5       u  pgph[        Xx5       V	V
s0 s H  u  pš[        U	5      U
_M     nn	n
/ U R                  R                  R                  5       QU R                  R                  R                  5       QnU Vs/ s H  oÝU R                  ;  d  M  UPM     nnXÀR                  R                   Vs/ s H  n[        U5      PM     sn-  nU Vs/ s H  oÛU   PM	     nnU R                   Hð  n[        UR                  [         R"                  [$        45      (       a  M4  [        UR                  [         R&                  5      (       a  UR                  nO3[         R                  R                  R)                  X5      R*                  nUR,                  (       a  U R.                  (       d  MÁ  UR1                  [        U5      5        UR1                  [$        5        Mò     [         R                  R2                  (       a  U R4                  OU R6                  nSS jn[9        U R;                  5       5      S:”  aš  U R;                  5        Vs/ s HQ  nU" UR,                  (       a+  [         R<                  " UR                  U R>                  5      OUR                  5      PMS     nnUR1                  U" US5      5        UR1                  [@        5        O*[         R                  R2                  (       a  [C        S5      eU R.                  (       a�  U R;                  5        Vs/ s HG  nUR,                  (       a1  U" [         R<                  " UR                  U R>                  5      5      OSPMI     nnUR1                  U" US5      5        UR1                  [@        5        O6[         R                  R2                  (       a  US/-  nUR1                  S5        SU RD                  ;   aY  [         R                  R2                  (       d:  UR1                  S	[9        U Vs/ s H  oÝc  M  S
U;  d  M  UPM     sn5       35        URG                  UU[H        RJ                  " S5      SUS9  gs  sn
n	f s  snf s  snf s  snf s  snf s  snf s  snf )z 
Codegens a call to this kernel
Úthreadsc                óÎ   • [         R                  R                  (       a+  U  Vs/ s H	  nSU S3PM     n nSSR                  U 5       S3$ U SSR                  U 5       S3$ s  snf )Nzstatic_cast<uint64_t>(rA   Ú{r@   r\  z=[rÅ  )r   rÇ  rq  Újoin)r�  Úkwargræ  s      r8   Úformat_threadsÚ/MetalKernel.call_kernel.<locals>.format_threads  si   € Ü�w‰w×"×"ÙBIÓJÂ'¸QÐ3°A°3°aÓ8Á'�Ð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_sizeNrl  zerror_buf_idx=Ú=r±  F)ÚdeviceÚtritonÚ	arg_types)r�  z	list[str]r”  r6   r¨   r6   )&r   rÇ  Úwrapper_coderE   r}  Úensure_size_computedÚpython_argdefsÚzipr6   rv  Úkeysry  rx  rê  r3   r  r!  r"  r(   r#  Úgenerate_numel_exprr‡  rä  r×  rO  rq  ÚcexprÚpexprrY   rr  r%  r&  ÚlistrZ   ra  Úgenerate_kernel_callr4   rš  )rH   r{  ÚnodeÚdeallocate_wsÚwrapperr_  Ú_Ú	call_argsrœ  Úcall_argÚarg_typeÚarg_name_to_typerE   r�   ÚtreerI   Úexpr_printerr•  r�  s                      r8   Úcall_kernelÚMetalKernel.call_kernelñ  sÏ  € ô —'‘'×&Ñ&ˆà—‘×#Ô#ˆAØ×(Ñ(¨Ö+ñ $ð &*§Y¡Y×%=Ñ%=Ó%?Ñ"ˆ�aä>AÀ)Ô>Wô
Ú>WÑ(:¨ŒC�‹M˜8Ò#Ñ>Wð 	ñ 
ð S�—‘×)Ñ)×.Ñ.Ó0ÐR°4·9±9×3JÑ3J×3OÑ3OÓ3QÐRˆÙ#ÓGšt˜°$×2FÑ2FÑ'F—™tˆÐGØ§¡×!3Ò!3Ó4Ò!3˜A”�Q–Ñ!3Ñ4Ñ4ˆÙ6:Ó;²d¨s cÔ*±dˆ	Ð;ð ×$Ô$ˆDÜ˜$Ÿ*™*¤u§}¡}´cÐ&:×;Ñ;áÜ˜DŸJ™J¬¯©×5Ñ5Ø—z‘z‘ä—w‘w×+Ñ+×?Ñ?ÀÓK×QÑQ�à×$×$¨×(=×(=Ñ(=Ø—‘œC ›IÔ&Ø× Ñ ¤Ö%ñ %ô &'§W¡W×%8×%8�t—z’z¸d¿j¹jˆô	9ô ˆt×&Ñ&Ó(Ó)¨AÓ-ð ×0Ñ0Ô2óò 3�Añ à—~—~ô —I’I˜aŸg™g t×'@Ñ'@ÔAàŸ™öñ
 3ð ð ð �K‰K™ w°	Ó:Ô;Ø×ÑœTÕ"ä�w‰w×"×"Ü"Ð#CÓDÐDà× × ð
 ×0Ñ0Ô2ó	ò 3�Að —>—>ñ œUŸYšY q§w¡w°×0IÑ0IÓJÔKàòñ 3ð	 ð ð �K‰K™ w°Ó=Ô>Ø×ÑœTÕ"ä�w‰w×"×"ð ˜˜‘�Ø× Ñ  Ô&ð �d—l‘lÓ"¬1¯7©7×+>×+>Ø�K‰KØ ¤±TÓ%`²T¨c£cÐQTÐ\_ÑQ_§c±TÑ%`Ó!aÐ bÐcôð 	×$Ñ$ØØÜ—<’< Ó&ØØð 	%ò 	
ùóS
ùò
 HùÚ4ùÚ;ùò8ùò ùò& &as>   Á4SÃSÃ4SÄSÄ2SÊ-ASÎAS!Ñ:S&ÒS&ÒS&c                óà  • U(       d  U(       d  g U R                  U5      nU R                  U5      nU(       a  U(       a  SU SU SU S3nOU(       a  U S3nOU SU 3n[        R                  R                  (       a*  U R                  R                  U R                  SU S3SS	9  g U R                  R                  S
5        U R                  R                  SU S3SU SU S3SS/5        g )NrO   z < 0 || rL  rA   z < 0rÙ   z) returnF)Ú
assignmentrl  rÚ   z,    TORCH_REPORT_ERROR(error_buf, "Index ", z, " out of range [0, ", z, ")");z    return;r\  )
rÊ   r   rÇ  rq  rÌ   rÍ   rÎ   ra  rb  ræ   )rH   rI   ÚsizeÚlowerÚupperÚexpr_strÚsize_strÚ	conditions           r8   Úcheck_boundsÚMetalKernel.check_boundsN  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Ø!Øð	õr:   )r»  )r¼  zdict[str, sympy.Expr]r½  r   r¨   r®  )rº   r¬  r¨   r6   )r{  r6   rÍ  r§   r¨   r   r�   )
r{  r6   rÍ  r§   rÙ  r   rÚ  r!   r¨   r®  )r{  r6   rÍ  r§   rÙ  r   r¨   r®  )rº   zUnion[str | torch.dtype]rö  zOptional[int]r÷  zOptional[Any]rø  r%   rÈ   zValueRanges[Any]r¨   r   )
rº   r¬  r»   r¬  r   r    rÙ  ú+Union[CSEVariable, tuple[CSEVariable, ...]]r¨   r½  )rI  r   r¨   r®  r­  )r{  zOptional[str]r¨   r6   r«  )r{  r6   r§  r   r¨  r%   r¨   r®  )
rI   r§   rµ  r§   r¶  r%   r·  r%   r¨   r®  )(r©   rª   r«   r¬   r­   r±   Ú	overridesrõ  Únewvar_prefixr&  r'  r   rF   r¤  r   r£  r<   ró  Úkexprr
   ra  Ú__annotations__r·  rº  rÁ  rÏ  rß  rì  r   Úunknownrû  r  rÿ  rY  re  r�  r±  r»  r¯   Ú__classcell__©r¾  s   @r8   r³  r³  ö  s  ø‡ ÙEà€IØ€FØ€MØÐØ€OÙ‹O×#Ñ#€EÙ‹L× Ñ €EÙÓ×&Ñ&€EØ€EÙ)¨7¨)Ó4€Gˆ_Ó4Ø=?ÐÐ :Ó?ð-à%ð-ð ð-ð 
÷	-ô%ô@ð SWð<Øð<Ø *ð<Ø3>ð<ØFOð<à	õ<ô*	8ð %)Ø'+Ø#Ø#.×#6Ò#6Ó#8ðà'ðð "ðð %ð	ð
 ðð !ðð 
õð,àðð ðð &ð	ð
 ;ðð 
5ôð s2àðs2ð ðs2ð &ð	s2ð
 ;ðs2ð 
5ôs2ôjHôT#öJgðT BFð[
Øð[
Ø"ð[
Ø:>ð[
à	õ[
ðz"Øð"Ø&0ð"Ø9=ð"ØFJð"à	÷"ò "r:   r³  c                  óJ   ^ • \ rS rSr\rSU 4S jjr        SS jrSrU =r	$ )ÚMetalSchedulingis  c                óÖ   >• [         TU ]  U5        [        R                  R                  nUb<  [        R                  R
                  (       d  UR                  R                  S5        g g g )NzDfrom torch._inductor.runtime.runtime_utils import compile_mps_shader)r¹  rº  r   rÇ  r�  rq  r‚  rç   )rH   Ú	schedulerr©  r¾  s      €r8   rº  ÚMetalScheduling.__init__v  sQ   ø€ Ü‰Ñ˜Ô#Ü—'‘'×&Ñ&ˆØÑÜ—7‘7×&×&Ø—‘×%Ñ%ØZõð 'ð r:   c                ó\  • [         R                  R                  nXR                  ;   a  UR                  U   nU$ SUR	                  5        3nU nXTR                  U'   [         R                  R
                  (       a	  SU S3U-   n[        X$5      u  pxU SU 3n	UR                  XaU	SS9  U$ )NÚmps_lib_zconst char* z
_source = Ú
F)Úgpu)r   rÇ  r�  Úsrc_to_kernelÚnext_kernel_suffixrq  r   Údefine_kernel)
rH   Úsrc_codeÚnode_schedulerÉ   r©  Úkernel_nameÚmps_lib_nameÚoriginsÚdetailed_originsÚmetadata_comments
             r8   rÐ  ÚMetalScheduling.define_kernel  sÁ   € ô —'‘'×&Ñ&ˆØ×,Ñ,Ó,Ø!×/Ñ/°Ñ9ˆKð" Ðð & g×&@Ñ&@Ó&BÐ%CÐDˆLà)˜NˆKØ.9×!Ñ! (Ñ+ä�w‰w×"×"à)¨,¨°zÐBÀXÑM�ä(;¸MÓ(SÑ%ˆGØ") ¨"Ð-=Ð,>Ð?ÐØ×!Ñ! ,Ð:JÐPUÐ!ÑVàÐr:   r¦   )rÈ  zOptional[Scheduler]r¨   r®  )rÑ  r6   rÒ  zlist[SchedulerNode]rÉ   r³  r¨   r6   )
r©   rª   r«   r¬   r³  Úkernel_typerº  rÐ  r¯   rÃ  rÄ  s   @r8   rÆ  rÆ  s  s7   ø† Ø€K÷ðØðØ,?ðØITðà	÷ò r:   rÆ  )r7   z)Union[float, int, bool, str, CSEVariable]r¨   r6   )IÚ
__future__r   r¥  rÞ   Úloggingru  Úpathlibr   Útypingr   r   r   r!  Úsympy.printing.precedencer   r4   Útorch.utils._cpp_embed_headersr	   Útorch.utils._ordered_setr
   Ú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   r   Úops_handlerr    r!   rÈ  r"   r#   r$   Ú	getLoggerr©   r·   r%   Úint8Úint16Úint32Úint64Úuint8r+   r,   rÊ  r¹   r9   r<   r±   Ú_initialize_pointwise_overridesr©  r³  rÆ  r¦   r:   r8   Ú<module>rï     s(  ðõ #ã Û Û Û Ý ß /Ñ /ã Ý 0ã Ý 9Ý /ß OÝ 7ç GÑ Gß ,Ñ ,÷÷ ÷ CÑ Bö Ýç6ß4Ýà×Ò˜Ó!€ð 
‡J�J�Ø	‡J�J�Ø	‡K�K�Ø	‡K�K�Ø	‡K�K�Ø	‡K�K�Ø	‡K�K�Ø	‡J�J�Ø	‡N�N�Hð
€ôôr"�|ô r"ôjt�[ô tðn	 × .Ñ .¨uÔ 5Ø × &Ñ &Ô (ôz	�*ô z	ôz"�nõ "r:   