+
    –°³jhÁ  ã                  ó0  € ^ RI Ht ^ RIt^ RIt^ RIt^ RIt^ RIHt ^ RIH	t	H
t
 ^ RIt^ RIHt ^ RIt^ RIHt ^ RIHt ^ RIHt ^ R	IHtHt ^ R
IHt ^RIHtHtHt ^RIH t H!t!H"t" ^RI#H$t$H%t%H&t&H't'H(t(H)t) ^RI*H+t+H,t,H-t- ]
'       d   ^RI.H/t/H0t0 ^RI1H2t2H3t3 ^RI#H4t4 ]Pj                  ! ]64      t7]Pp                  R]Pr                  R]Pt                  R]Pv                  R]Px                  R]Pz                  R]P|                  R]P~                  R]P€                  R]P‚                  R]P„                  R]P†                  R/tDR R ltE ! R  R!]4      tF ! R" R#](4      tG]GP‘                  R$4       ]GP“                  4         ! R% R&],4      tJ ! R' R(]-4      tKR# ))é    )Ú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ÚushortÚuintÚulongÚfloatÚhalfÚbfloatc               ó    € V ^8„  d   QhRRRR/# )é   Úvalz&float | int | bool | str | CSEVariableÚreturnÚstr© )Úformats   "Ú\/home/luca/screen-checking/.venv/lib/python3.14/site-packages/torch/_inductor/codegen/mps.pyÚ__annotate__r6   :   s   € ÷ ñ Ð>ð À3ñ ó    c                ó  € \        V \        4      '       dC   V \        P                  8X  d   R # V \        P                  ) 8X  d   R# W 8w  d   R# \	        V 4      # \        V \
        4      '       d   V '       d   R# R# \	        V 4      # )Ú	HUGE_VALFz
-HUGE_VALFÚNANÚtrueÚfalse)Ú
isinstancer+   ÚtorchÚinfr2   r"   )r0   s   &r5   Úvalue_to_metalr@   :   sc   € Ü�#”u×ÒØ”%—)‘)ÔÙØ”U—Y‘Y�JÔÙØŒZÙÜ�3‹xˆÜ	�Cœ×	Ò	ßˆvÐ) 'Ð)Üˆs‹8€Or7   c                  óô   € ] tR t^HtRtR R ltR R ltR R lt]R R	 l4       t	R
 R lt
R R ltR R ltR R ltR R ltR R ltR R ltR R ltR R ltR R lt]tR R ltR  R! ltR" R# ltR$tR%# )&ÚMetalExprPrinterz/Converts sympy expression to Metal code snippetc               ó    € V ^8„  d   QhRRRR/# ©r/   Úexprú
sympy.Exprr1   r2   r3   )r4   s   "r5   r6   ÚMetalExprPrinter.__annotate__K   s   € ÷ .ñ . Jð .°3ñ .r7   c                	óª   € VP                   w  r#V P                  V4      pV P                  V4      pVP                  '       d
   R V RV R2# RV RV R2# )úc10::metal::floor_divide(ú, Ú)úmetal::floor(ú) / (©ÚargsÚdoprintÚ
is_integer)ÚselfrE   ÚxÚdivs   &&  r5   Ú_print_FloorDivÚ MetalExprPrinter._print_FloorDivK   s[   € Ø—‘‰ˆØ�L‰L˜‹OˆØ�l‰l˜3ÓˆØ�?�?ˆ?Ø.¨q¨c°°C°5¸Ð:Ð:Ø˜q˜c  s e¨1Ð-Ð-r7   c               ó    € V ^8„  d   QhRRRR/# rD   r3   )r4   s   "r5   r6   rG   S   s   € ÷ "ñ "¨:ð "¸#ñ "r7   c                	óP  € VP                   w  r#pVR 8H  ;'       d    WD^,
          ,          ^ 8g  pV P                  V4      pV^8w  d7   V P                  V4      pVP                  '       d   RV RV R2pM	RV RV R2pV P                  V4      pV'       d
   RV RV R2# RV RV R2# )i   Ú(rM   rK   rL   zc10::metal::safe_mod(rJ   z) % (rN   )rR   rE   rS   rT   ÚmodÚuse_safe_mods   &&    r5   Ú_print_ModularIndexingÚ'MetalExprPrinter._print_ModularIndexingS   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Ð!Ð!r7   c               ó$   € V ^8„  d   QhRRRRRR/# )r/   rE   rF   Úfnr2   r1   r3   )r4   s   "r5   r6   rG   d   s!   € ÷ ñ  :ð °3ð ¸3ñ r7   c                	ó®   € \        \        V P                  VP                  4      4      pV^ ,          pVR,           F  pV P	                  WEV4      pK  	  V# )r   :é   NN)ÚlistÚmapÚ_printrO   Ú_print_binary_min_max)rR   rE   r_   rO   ÚresultÚargs   &&&   r5   Ú_print_min_maxÚMetalExprPrinter._print_min_maxd   sK   € ä”C˜Ÿ™ T§Y¡YÓ/Ó0ˆØ�a•ˆØ˜—8�8ˆCØ×/Ñ/°¸RÓ@ŠFñ àˆr7   c               ó(   € V ^8„  d   QhRRRRRRRR/# )r/   Úar2   Úbr_   r1   r3   )r4   s   "r5   r6   rG   m   s(   € ÷ :ñ : ð :¨ð :°#ð :¸#ñ :r7   c                	óJ   € R V  RV RV  R2pR V  RV RV R2pRV RV RV R2# )ústatic_cast<decltype(Ú+ú)>(rK   zmetal::rY   rJ   r3   )rk   rl   r_   Ú
typecast_aÚ
typecast_bs   &&&  r5   re   Ú&MetalExprPrinter._print_binary_min_maxl   sQ   € à,¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø,¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø˜˜˜A˜j˜\¨¨J¨<°qÐ9Ð9r7   c               ó    € V ^8„  d   QhRRRR/# rD   r3   )r4   s   "r5   r6   rG   r   ó   € ÷ 0ñ 0˜zð 0¨cñ 0r7   c                	ó&   € V P                  VR 4      # )Úmin©rh   ©rR   rE   s   &&r5   Ú
_print_MinÚMetalExprPrinter._print_Minr   ó   € Ø×"Ñ" 4¨Ó/Ð/r7   c               ó    € V ^8„  d   QhRRRR/# rD   r3   )r4   s   "r5   r6   rG   u   ru   r7   c                	ó&   € V P                  VR 4      # )Úmaxrx   ry   s   &&r5   Ú
_print_MaxÚMetalExprPrinter._print_Maxu   r|   r7   c               ó    € V ^8„  d   QhRRRR/# rD   r3   )r4   s   "r5   r6   rG   x   s   € ÷ :ñ :˜zð :¨cñ :r7   c                	óÄ   € \        VP                  4      ^8w  d"   \        R\        VP                  4       24      hRV P                  VP                  ^ ,          4       R2# )ra   úexpected 1 arg, got úmetal::abs(rK   ©ÚlenrO   ÚAssertionErrorrd   ry   s   &&r5   Ú
_print_AbsÚMetalExprPrinter._print_Absx   sO   € Üˆt�y‰y‹>˜QÔÜ Ð#7¼¸D¿I¹I»Ð7GÐ!HÓIÐIà˜TŸ[™[¨¯©°1­Ó6Ð7°qÐ9Ð9r7   c               ó    € V ^8„  d   QhRRRR/# rD   r3   )r4   s   "r5   r6   rG   ~   s   € ÷ Nñ N jð N°Sñ Nr7   c                	óÄ   € \        VP                  4      ^8w  d"   \        R\        VP                  4       24      hRV P                  VP                  ^ ,          4       R2# )ra   r„   zstatic_cast<long>(metal::rint(ú))r†   ry   s   &&r5   Ú_print_RoundToIntÚ"MetalExprPrinter._print_RoundToInt~   sO   € Üˆt�y‰y‹>˜QÔÜ Ð#7¼¸D¿I¹I»Ð7GÐ!HÓIÐIà/°·±¸D¿I¹IÀa½LÓ0IÐ/JÈ"ÐMÐMr7   c               ó    € V ^8„  d   QhRRRR/# rD   r3   )r4   s   "r5   r6   rG   „   s   € ÷ ]ñ ]¨
ð ]°sñ ]r7   c                	óT  € \        VP                  4      ^8w  d"   \        R\        VP                  4       24      hVP                  w  r#VP                  '       d%   V^ 8¼  d   \        RV 24      h\	        RV R24      hV P                  V\        R,          4      pRV RV RV)  R	2# )
r/   úexpected 2 args, got zexpected ndigits < 0, got zOFor integer inputs, only non-negative ndigits are currently supported, but got Ú.ÚMulz!static_cast<float>(metal::rint(1eú * z) * 1erK   )r‡   rO   rˆ   rQ   Ú
ValueErrorÚparenthesizer   )rR   rE   ÚnumberÚndigitsÚ
number_strs   &&   r5   Ú_print_RoundDecimalÚ$MetalExprPrinter._print_RoundDecimal„   sµ   € Üˆt�y‰y‹>˜QÔÜ Ð#8¼¸T¿Y¹Y»Ð8HÐ!IÓJÐJØŸ)™)‰ˆØ××Ðà˜!Œ|Ü$Ð'AÀ'ÀÐ%KÓLÐLÜØaÐbiÐajÐjkÐlóð ð ×&Ñ& v¬z¸%Õ/@ÓAˆ
Ø2°7°)¸3¸z¸lÈ&ÐRYÐQYÐPZÐZ[Ð\Ð\r7   c               ó    € V ^8„  d   QhRRRR/# rD   r3   )r4   s   "r5   r6   rG   ’   s   € ÷ bñ b jð b°Sñ br7   c                	ól   € VP                   w  r#R V P                  V4       RV P                  V4       R2# )ústatic_cast<float>(z) / static_cast<float>(rK   )rO   rd   )rR   rE   ÚlhsÚrhss   &&  r5   Ú_print_IntTrueDivÚ"MetalExprPrinter._print_IntTrueDiv’   s;   € Ø—9‘9‰ˆð % T§[¡[°Ó%5Ð$6Ð6MÈdÏkÉkÐZ]ÓN^ÐM_Ð_`ÐaÐar7   c               ó    € V ^8„  d   QhRRRR/# rD   r3   )r4   s   "r5   r6   rG   ˜   s   € ÷ Xñ X¨
ð X°sñ Xr7   c                	óÎ   € \        VP                  4      ^8w  d"   \        R\        VP                  4       24      h\        V P                  VP                  4      w  r#RV RV R2# )r/   r’   z'metal::precise::pow(static_cast<float>(z), static_cast<float>(r�   )r‡   rO   rˆ   rc   rP   )rR   rE   rS   Úys   &&  r5   Ú_print_PowByNaturalÚ$MetalExprPrinter._print_PowByNatural˜   s]   € Üˆt�y‰y‹>˜QÔÜ Ð#8¼¸T¿Y¹Y»Ð8HÐ!IÓJÐJÜ�4—<‘< §¡Ó+‰ˆØ8¸¸Ð;QÐRSÐQTÐTVÐWÐWr7   c               ó    € V ^8„  d   QhRRRR/# rD   r3   )r4   s   "r5   r6   rG   ž   s   € ÷ *ñ * :ð *°#ñ *r7   c                	óÈ   € \        VP                  4      ^8w  d"   \        R\        VP                  4       24      hV P                  VP                  ^ ,          4      pRV R2# )ra   r„   rŸ   rK   ©r‡   rO   rˆ   rP   ©rR   rE   rS   s   && r5   Ú_print_ToFloatÚMetalExprPrinter._print_ToFloatž   sS   € Üˆt�y‰y‹>˜QÔÜ Ð#7¼¸D¿I¹I»Ð7GÐ!HÓIÐIØ�L‰L˜Ÿ™ 1�Ó&ˆØ$ Q C qÐ)Ð)r7   c               ó    € V ^8„  d   QhRRRR/# rD   r3   )r4   s   "r5   r6   rG   ¤   s   € ÷ ñ  ð °ñ r7   c                	ód   € VP                   '       d   \        \        V4      4      # \        V4      # ©N)rQ   r2   r%   ry   s   &&r5   Ú_print_FloatÚMetalExprPrinter._print_Float¤   s%   € Ø�?�?ˆ?ô ”s˜4“y“>Ð!ä�t“9Ðr7   c               ó    € V ^8„  d   QhRRRR/# rD   r3   )r4   s   "r5   r6   rG   ­   s   € ÷ Jñ J jð J°Sñ Jr7   c                	óÈ   € \        VP                  4      ^8w  d"   \        R\        VP                  4       24      hV P                  VP                  ^ ,          4      pRV R2# )ra   r„   z1static_cast<int>(metal::floor(static_cast<float>(z)))r«   r¬   s   && r5   Ú_print_FloorToIntÚ"MetalExprPrinter._print_FloorToInt­   sS   € Üˆt�y‰y‹>˜QÔÜ Ð#7¼¸D¿I¹I»Ð7GÐ!HÓIÐIØ�L‰L˜Ÿ™ 1�Ó&ˆØBÀ1À#ÀSÐIÐIr7   c               ó    € V ^8„  d   QhRRRR/# rD   r3   )r4   s   "r5   r6   rG   µ   s   € ÷ 6ñ 6 jð 6°Sñ 6r7   c                	óÈ   € \        VP                  4      ^8w  d"   \        R\        VP                  4       24      hV P                  VP                  ^ ,          4      pRV R2# )ra   r„   zstatic_cast<int>(metal::trunc(r�   r«   r¬   s   && r5   Ú_print_TruncToIntÚ"MetalExprPrinter._print_TruncToIntµ   sS   € Üˆt�y‰y‹>˜QÔÜ Ð#7¼¸D¿I¹I»Ð7GÐ!HÓIÐIØ�L‰L˜Ÿ™ 1�Ó&ˆØ/°¨s°"Ð5Ð5r7   c               ó    € V ^8„  d   QhRRRR/# rD   r3   )r4   s   "r5   r6   rG   »   s   € ÷ ,ñ ,¨jð ,¸Sñ ,r7   c                	óÈ   € \        VP                  4      ^8w  d"   \        R\        VP                  4       24      hV P                  VP                  ^ ,          4      pRV R2# )ra   r„   zmetal::precise::log2(rK   r«   r¬   s   && r5   Ú_print_OpaqueUnaryFn_log2Ú*MetalExprPrinter._print_OpaqueUnaryFn_log2»   sS   € Üˆt�y‰y‹>˜QÔÜ Ð#7¼¸D¿I¹I»Ð7GÐ!HÓIÐIØ�L‰L˜Ÿ™ 1�Ó&ˆØ& q c¨Ð+Ð+r7   c               ó    € V ^8„  d   QhRRRR/# rD   r3   )r4   s   "r5   r6   rG   Á   s   € ÷ "ñ " ð "°ñ "r7   c                	óJ   a € V 3R  lVP                    4       w  r#pV RV RV 2# )c              3  ól   <"  € T F)  pSP                  V\        R ,          R,
          4      x € K+  	  R# 5i)ÚAtomg      à?N)r—   r   )Ú.0rg   rR   s   & €r5   Ú	<genexpr>Ú0MetalExprPrinter._print_Where.<locals>.<genexpr>Â   s.   øé € ð 
ÙHQÀˆD×Ñ˜c¤:¨fÕ#5¸Õ#;×<Ð<Ë	ùs   ƒ14ú ? z : )rO   )rR   rE   ÚcÚpÚqs   f&   r5   Ú_print_WhereÚMetalExprPrinter._print_WhereÁ   s3   ø€ ô
ØHLÏ	Ê	ó
‰ˆˆað ��C˜�s˜#˜a˜SÐ!Ð!r7   r3   N)Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__Ú__doc__rU   r\   rh   Ústaticmethodre   rz   r€   r‰   rŽ   r›   r¢   r§   r­   r²   r¶   Ú_print_floorrº   r¾   rË   Ú__static_attributes__r3   r7   r5   rB   rB   H   sz   † Ù9õ.õ"õ"ð ô:ó ð:õ
0õ0õ:õNõ]õbõXõ*õõJð %€Lõ6õ,÷"ñ "r7   rB   c                  óþ  € ] tR t^ÈtRt]R`R R ll4       t]R R l4       t]R R l4       t]R	 R
 l4       t	]R R l4       t
]R R l4       t]R R l4       t]R R l4       t]R R l4       t]R R l4       t]R R l4       t]R R l4       t]R R l4       t]R R l4       t]R R  l4       t]R! R" l4       t]R# R$ l4       t]R% R& l4       t]R' R( l4       t]R) R* l4       t]R+ R, l4       t]R- R. l4       t]R/ R0 l4       t]R1 R2 l4       t]R3 R4 l4       t]R5 R6 l4       t]R7 R8 l4       t ]R9 R: l4       t!]R; R< l4       t"]R= R> l4       t#]R? R@ l4       t$]RA RB l4       t%]RC RD l4       t&]RE RF l4       t']RG RH l4       t(]RI RJ l4       t)]RK RL l4       t*]RM RN l4       t+]RO RP l4       t,]RQ RR l4       t-]RS RT l4       t.]RU RV l4       t/]RW RX l4       t0RY RZ lt1R[ R\ lt2]3R] R^ l4       t4R_t5R# )aÚMetalOverrideszXImplements Metal-specific overrides for ops. Base class emits Python-friendly overrides.Nc          
     ó,   € V ^8„  d   QhRRRRRRRRR	R
/# )r/   rS   r   Údtypeútorch.dtypeÚ	src_dtypeztorch.dtype | NoneÚuse_compute_typesr"   r1   r2   r3   )r4   s   "r5   r6   ÚMetalOverrides.__annotate__Ì   s:   € ÷ <ñ <Øð<àð<ð &ð<ð  ð	<ð
 
ñ<r7   c                	óŠ   € V\         P                  8X  d   \        P                  R 4       RV  R2# R\        V,           RV  R2# )z>float64 cast requested, probably from tensorify_python_scalarsrŸ   rK   ústatic_cast<ú>()r>   ÚdoubleÚlogÚwarningÚDTYPE_TO_METAL)rS   rØ   rÚ   rÛ   s   &&&&r5   Úto_dtypeÚMetalOverrides.to_dtypeË   sK   € ð ”E—L‘LÔ Ü�K‰KØPôð )¨¨¨1Ð-Ð-Øœn¨UÕ3Ð4°B°q°c¸Ð;Ð;r7   c               ó(   € V ^8„  d   QhRRRRRRRR/# )r/   rS   r   rØ   rÙ   rÚ   r1   r2   r3   )r4   s   "r5   r6   rÜ   Ú   s2   € ÷ bñ bØðbØ*ðbØ7Bðbà	ñbr7   c                	óF   € R \         V,           R\         V,           RV  R2# )zas_type<z>(static_cast<rß   r�   ©rã   )rS   rØ   rÚ   s   &&&r5   Úto_dtype_bitcastÚMetalOverrides.to_dtype_bitcastÙ   s/   € ð œ.¨Õ/Ð0°¼~ÈiÕ?XÐ>YÐY[Ð\]Ð[^Ð^`ÐaÐar7   c               ó$   € V ^8„  d   QhRRRRRR/# )r/   r0   zbool | float | intrØ   rÙ   r1   r2   r3   )r4   s   "r5   r6   rÜ   à   s"   € ÷ #ñ #Ð(ð #°ð #Àñ #r7   c                	ó   € \        V 4      # r±   ©r@   )r0   rØ   s   &&r5   ÚconstantÚMetalOverrides.constantß   s   € ä˜cÓ"Ð"r7   c               ó$   € V ^8„  d   QhRRRRRR/# ©r/   rE   rF   rØ   rÙ   r1   r2   r3   )r4   s   "r5   r6   rÜ   ä   s!   € ÷ (ñ (˜ð (¨Kð (¸Cñ (r7   c                	ó@  € \         P                  P                  \         P                  P                  V 4      4      p\         P                  P                  P                  \         P                  P                  V\        V 4      R 7      p\        P                  ! W14      # ))Úbounds)
r   ÚkernelÚindex_to_strÚprepare_indexingÚcseÚgenerateÚcomputer   r   rä   )rE   rØ   Úidx_strÚvars   &&  r5   Ú
index_exprÚMetalOverrides.index_exprã   sj   € ä—(‘(×'Ñ'¬¯©×(AÑ(AÀ$Ó(GÓHˆÜ�h‰h�l‰l×#Ñ#Ü�H‰H×Ñ˜gÔ.CÀDÓ.Ið $ó 
ˆô �|Š|˜CÓ'Ð'r7   c               ó$   € V ^8„  d   QhRRRRRR/# rñ   r3   )r4   s   "r5   r6   rÜ   ì   s!   € ÷ 6ñ 6˜ð 6¨Kð 6¸Cñ 6r7   c                	ó,   € \         P                  W4      # r±   )rÖ   rü   )rE   rØ   s   &&r5   Ú
value_exprÚMetalOverrides.value_exprë   s   € ô ×(Ñ(¨Ó5Ð5r7   c               ó(   € V ^8„  d   QhRRRRRRRR/# )r/   Úmaskr   ÚbodyrF   Úotherr1   r2   r3   )r4   s   "r5   r6   rÜ   ò   s(   € ÷  ñ  �[ð  ¨
ð  ¸;ð  È3ñ  r7   c           
     	óŒ  € \        V4      p\        4       p\        P                  P	                  V4      ;_uu_ 4        VP                  4       ;_uu_ 4        \        P                  ! 4       \        P                  P                  n	        \        P                  P                  ;P                  R ,          un
        V! 4       pRRR4       RRR4       V  RVP                  4        RV 2p\        P                  P                  P                  V4      pV'       Ego   \        P                  P                  P                  XP                  R7      p\        P                  P                  P                  Wg4       \        P                  P                   P#                  \$        VP                  ,           RV R2RV  R2.4       \        P                  P                   P                  4       ;_uu_ 4        \        P                  P                   P'                  V4       \        P                  P                   P)                  V RV R	V R
24       RRR4       \        P                  P                   P)                  RV RV R	V R
24       V#   + '       g   i     ELÓ; i  + '       g   i     ELß; i  + '       g   i     Lm; i)Ú_scopedNÚ:©rØ   Ú Ú;úif (ú) {z = static_cast<decltype(rp   ú);z} else )r@   r   r   rô   Úswap_buffersÚindentÚ	itertoolsÚcountr÷   Úiter_buffer_idsÚname_prefixÚgetvalueÚtry_getÚnewvarrØ   Úputrù   Ú
writelinesrã   ÚspliceÚ	writeline)r  r  r  Ú	other_strÚscoped_bodyÚrcÚ	cache_keyrû   s   &&&     r5   ÚmaskedÚMetalOverrides.maskedñ   sð  € ô
 # 5Ó)ˆ	Ü$Ó&ˆÜ�X‰X×"Ñ" ;×/Õ/°×1CÑ1C×1EÕ1Eô ,5¯?ª?Ó+<ŒA�H‰H�L‰LÔ(ä�H‰H�L‰L×$Ò$¨	Õ1Õ$Ù“ˆB÷ 2F×/ð �f˜A˜k×2Ñ2Ó4Ð5°Q°y°kÐBˆ	Ü�h‰h�l‰l×"Ñ" 9Ó-ˆßˆsÜ—(‘(—,‘,×%Ñ%¨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 2F×1EÐ1Eú×/×/Ð/ú÷$ +×*ús1   ½JÁA+J	ÃJÇ2AJ3ÊJÊ	JÊJ0	Ê3K	c               ó(   € V ^8„  d   QhRRRRRRRR/# )r/   rk   r!   rl   rÈ   r1   r2   r3   )r4   s   "r5   r6   rÜ     s.   € ÷ Nñ N�ð N˜Fð N vð N°#ñ Nr7   c           	     	ó0   € V  R V RV R\        V4       R2# )rÇ   z : static_cast<decltype(rp   rK   rí   )rk   rl   rÈ   s   &&&r5   ÚwhereÚMetalOverrides.where  s*   € ð ��C˜�sÐ2°1°#°S¼ÈÓ9JÐ8KÈ1ÐMÐMr7   c               ó$   € V ^8„  d   QhRRRRRR/# )r/   rk   r!   rl   r1   r2   r3   )r4   s   "r5   r6   rÜ     s!   € ÷ 2ñ 2�Vð 2 ð 2¨3ñ 2r7   c                	ó   € R V  RV R2# )zc10::metal::remainder(rJ   rK   r3   ©rk   rl   s   &&r5   Ú	remainderÚMetalOverrides.remainder  s   € à'¨ s¨"¨Q¨C¨qÐ1Ð1r7   c               ó$   € V ^8„  d   QhRRRRRR/# ©r/   rk   r   rl   r1   r2   r3   )r4   s   "r5   r6   rÜ     ó!   € ÷ >ñ >�;ð > ;ð >°3ñ >r7   c                	óD   € R V  RV RV  R2pR V  RV RV R2pRV RV R2# )rn   ro   rp   rK   zc10::metal::max(rJ   r3   ©rk   rl   rq   rr   s   &&  r5   ÚmaximumÚMetalOverrides.maximum  óM   € ð -¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø,¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø! * ¨R°
¨|¸1Ð=Ð=r7   c               ó$   € V ^8„  d   QhRRRRRR/# r,  r3   )r4   s   "r5   r6   rÜ   &  r-  r7   c                	óD   € R V  RV RV  R2pR V  RV RV R2pRV RV R2# )rn   ro   rp   rK   zc10::metal::min(rJ   r3   r/  s   &&  r5   ÚminimumÚMetalOverrides.minimum$  r2  r7   c               ó$   € V ^8„  d   QhRRRRRR/# r,  r3   )r4   s   "r5   r6   rÜ   -  s!   € ÷ ñ �kð  kð °cñ r7   c                	ó   € V  R V 2# )z || r3   r(  s   &&r5   Ú
logical_orÚMetalOverrides.logical_or+  ó   € ð ��D˜˜ˆ}Ðr7   c               ó$   € V ^8„  d   QhRRRRRR/# r,  r3   )r4   s   "r5   r6   rÜ   2  s!   € ÷ ñ �{ð  {ð °sñ r7   c                	ó   € V  R V 2# )z && r3   r(  s   &&r5   Úlogical_andÚMetalOverrides.logical_and0  r;  r7   c               ó    € V ^8„  d   QhRRRR/# ©r/   rS   r   r1   r2   r3   )r4   s   "r5   r6   rÜ   7  ó   € ÷ $ñ $�ð $ ñ $r7   c                	ó   € R V  R2# )zmetal::isnan(rK   r3   ©rS   s   &r5   ÚisnanÚMetalOverrides.isnan5  ó   € ð ˜q˜c Ð#Ð#r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   <  rB  r7   c                	ó   € R V  R2# )zmetal::isinf(rK   r3   rD  s   &r5   ÚisinfÚMetalOverrides.isinf:  rG  r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   A  ó   € ÷ +ñ +ˆ{ð +˜sñ +r7   c                	ó   € R V  R2# )zmetal::precise::log(rK   r3   rD  s   &r5   rá   ÚMetalOverrides.log?  ó   € ð & a S¨Ð*Ð*r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   F  rM  r7   c                	ó   € R V  R2# )zmetal::precise::exp(rK   r3   rD  s   &r5   ÚexpÚMetalOverrides.expD  rP  r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   K  s   € ÷ "ñ "ˆ{ð "˜sñ "r7   c                	ó   € R V  R2# )r…   rK   r3   rD  s   &r5   ÚabsÚMetalOverrides.absI  s   € ð ˜Q˜C˜qÐ!Ð!r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   P  s   € ÷ &ñ &�;ð & 3ñ &r7   c                	ó   € R V  R2# )zmetal::signbit(rK   r3   rD  s   &r5   ÚsignbitÚMetalOverrides.signbitN  s   € ð !   1Ð%Ð%r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   U  rM  r7   c                	ó   € R V  R2# )zmetal::precise::sin(rK   r3   rD  s   &r5   ÚsinÚMetalOverrides.sinS  rP  r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   Y  s   € ÷ (ñ (�ð ( ñ (r7   c                	ó   € R V  R2# )zc10::metal::sinc(rK   r3   rD  s   &r5   ÚsincÚMetalOverrides.sincX  s   € à" 1 # QÐ'Ð'r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   ^  rM  r7   c                	ó   € R V  R2# )zmetal::precise::cos(rK   r3   rD  s   &r5   ÚcosÚMetalOverrides.cos\  rP  r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   c  rM  r7   c                	ó   € R V  R2# )zmetal::precise::tan(rK   r3   rD  s   &r5   ÚtanÚMetalOverrides.tana  rP  r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   h  ó   € ÷ ,ñ ,�ð , ñ ,r7   c                	ó   € R V  R2# )zmetal::precise::asin(rK   r3   rD  s   &r5   ÚasinÚMetalOverrides.asinf  ó   € ð ' q c¨Ð+Ð+r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   m  rn  r7   c                	ó   € R V  R2# )zmetal::precise::acos(rK   r3   rD  s   &r5   ÚacosÚMetalOverrides.acosk  rr  r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   r  rn  r7   c                	ó   € R V  R2# )zmetal::precise::atan(rK   r3   rD  s   &r5   ÚatanÚMetalOverrides.atanp  rr  r7   c               ó$   € V ^8„  d   QhRRRRRR/# )r/   rS   r   r¦   r1   r2   r3   )r4   s   "r5   r6   rÜ   w  s!   € ÷ 4ñ 4�ð 4 ð 4°ñ 4r7   c                	ó   € R V  RV R2# )z::metal::precise::atan2(rJ   rK   r3   )rS   r¦   s   &&r5   Úatan2ÚMetalOverrides.atan2u  s   € ð *¨!¨¨B¨q¨c°Ð3Ð3r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   |  rn  r7   c                	ó   € R V  R2# )zmetal::precise::sqrt(rK   r3   rD  s   &r5   ÚsqrtÚMetalOverrides.sqrtz  rr  r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   €  s   € ÷ 3ñ 3ˆ{ð 3˜sñ 3r7   c                	ó   € R V  RV  R2# )rn   z)>(-rK   r3   rD  s   &r5   ÚnegÚMetalOverrides.neg  s   € ð ' q c¨¨a¨S°Ð2Ð2r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   ‡  ó   € ÷ -ñ -�ð - ñ -r7   c                	ó   € R V  R2# )zmetal::precise::rsqrt(rK   r3   rD  s   &r5   ÚrsqrtÚMetalOverrides.rsqrt…  ó   € ð (¨ s¨!Ð,Ð,r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   Œ  rn  r7   c                	ó   € R V  R2# )zmetal::precise::tanh(rK   r3   rD  s   &r5   ÚtanhÚMetalOverrides.tanhŠ  rr  r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   ‘  rˆ  r7   c                	ó   € R V  R2# )zmetal::precise::atanh(rK   r3   rD  s   &r5   ÚatanhÚMetalOverrides.atanh�  rŒ  r7   c               ó$   € V ^8„  d   QhRRRRRR/# r,  r3   )r4   s   "r5   r6   rÜ   •  s!   € ÷ 5ñ 5�Kð 5 Kð 5°Cñ 5r7   c                	ó   € R V  RV R2# )rI   rJ   rK   r3   r(  s   &&r5   ÚfloordivÚMetalOverrides.floordiv”  s   € ð +¨1¨#¨R°¨s°!Ð4Ð4r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   ›  rB  r7   c                	ó   € R V  R2# )rL   rK   r3   rD  s   &r5   ÚfloorÚMetalOverrides.floor™  rG  r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ      ó   € ÷ #ñ #�ð # ñ #r7   c                	ó   € R V  R2# )zmetal::sign(rK   r3   rD  s   &r5   ÚsignÚMetalOverrides.signž  ó   € ð ˜a˜S Ð"Ð"r7   c               ó$   € V ^8„  d   QhRRRRRR/# r,  r3   )r4   s   "r5   r6   rÜ   ¥  s!   € ÷ :ñ :�ð : ð :°ñ :r7   c                	óD   € R V  RV RV  R2pR V  RV RV R2pRV RV R2# )rn   ro   rp   rK   zmetal::fmod(rJ   r3   r/  s   &&  r5   ÚfmodÚMetalOverrides.fmod£  sM   € ð -¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø,¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø˜j˜\¨¨J¨<°qÐ9Ð9r7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   ¬  rB  r7   c                	ó   € R V  R2# )úmetal::trunc(rK   r3   rD  s   &r5   ÚtruncÚMetalOverrides.truncª  rG  r7   c               ó$   € V ^8„  d   QhRRRRRR/# r,  r3   )r4   s   "r5   r6   rÜ   ±  s!   € ÷ ñ �Kð  Kð °Cñ r7   c                	óÈ   € V  R V 2pV P                   e   V P                   P                  '       g+   VP                   e#   VP                   P                  '       d   RV R2# V# )z / r©  rK   )rØ   Úis_floating_point)rk   rl   Úquots   && r5   ÚtruncdivÚMetalOverrides.truncdiv¯  sY   € ð ��C˜�sˆ|ˆØ�G‰GÒ A§G¡G×$=×$=Ð$=Ø�G‰GÒ A§G¡G×$=×$=Ð$=à" 4 &¨Ð*Ð*Øˆr7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   »  rž  r7   c                	ó   € R V  R2# )zmetal::ceil(rK   r3   rD  s   &r5   ÚceilÚMetalOverrides.ceil¹  r¢  r7   c               ó$   € V ^8„  d   QhRRRRRR/# ©r/   Úseedr   Úoffsetr1   r2   r3   )r4   s   "r5   r6   rÜ   ¿  s!   € ÷ 5ñ 5�;ð 5¨ð 5¸ñ 5r7   c                	óf   € \         P                  P                  P                  R 4       RV  RV R2# )Úrandomzc10::metal::rand(rJ   rK   ©r   rô   ÚheadersÚadd©r¸  r¹  s   &&r5   ÚrandÚMetalOverrides.rand¾  s/   € ä	�‰×Ñ×Ñ˜XÔ&Ø" 4 &¨¨6¨(°!Ð4Ð4r7   c               ó$   € V ^8„  d   QhRRRRRR/# r·  r3   )r4   s   "r5   r6   rÜ   Ä  s!   € ÷ 6ñ 6�Kð 6¨ð 6¸ñ 6r7   c                	óf   € \         P                  P                  P                  R 4       RV  RV R2# )r»  zc10::metal::randn(rJ   rK   r¼  r¿  s   &&r5   ÚrandnÚMetalOverrides.randnÃ  s/   € ä	�‰×Ñ×Ñ˜XÔ&Ø# D 6¨¨F¨8°1Ð5Ð5r7   c          
     ó,   € V ^8„  d   QhRRRRRRRRRR/# )r/   r¸  r   r¹  ÚlowÚhighr1   r2   r3   )r4   s   "r5   r6   rÜ   É  s;   € ÷ Iñ IØðIØ#.ðIØ5@ðIØHSðIà	ñIr7   c           	     	ór   € \         P                  P                  P                  R 4       RV  RV RV RV R2	# )r»  zc10::metal::randint64(rJ   rK   r¼  )r¸  r¹  rÇ  rÈ  s   &&&&r5   Ú	randint64ÚMetalOverrides.randint64È  s=   € ô 	
�‰×Ñ×Ñ˜XÔ&Ø'¨ v¨R°¨x°r¸#¸¸bÀÀÀaÐHÐHr7   c               ó    € V ^8„  d   QhRRRR/# rA  r3   )r4   s   "r5   r6   rÜ   Ñ  s   € ÷ #ñ #�ð # ñ #r7   c                	ó   € R V  R2# )zmetal::rint(rK   r3   rD  s   &r5   ÚroundÚMetalOverrides.roundÏ  r¢  r7   c               ó$   € V ^8„  d   QhRRRRRR/# r,  r3   )r4   s   "r5   r6   rÜ   Õ  s!   € ÷ 1ñ 1ˆ{ð 1˜{ð 1¨sñ 1r7   c                	óD   € R V  RV RV  R2pR V  RV RV R2pRV RV R2# )rn   ro   rp   rK   zmetal::pow(rJ   r3   )rk   rl   Úcast_aÚcast_bs   &&  r5   Ú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Ð0r7   c               ó$   € V ^8„  d   QhRRRRRR/# )r/   rk   r   Únamer2   r1   r3   )r4   s   "r5   r6   rÜ   Ú  s!   € ÷ *ñ * ð *°3ð *¸3ñ *r7   c                	óf   € \         P                  P                  P                  R 4       RV RV R2# )Úspecial_mathúc10::metal::rY   rK   r¼  )rR   rk   r×  s   &&&r5   Ú_special_unaryÚMetalOverrides._special_unaryÚ  s/   € Ü	�‰×Ñ×Ñ˜^Ô,Ø˜d˜V 1 Q C qÐ)Ð)r7   c               ó(   € V ^8„  d   QhRRRRRRRR/# )r/   rk   r   rl   r×  r2   r1   r3   )r4   s   "r5   r6   rÜ   Þ  s(   € ÷ /ñ / ð /°ð /ÀCð /ÈCñ /r7   c                	ól   € \         P                  P                  P                  R 4       RV RV RV R2# )rÙ  rÚ  rY   rJ   rK   r¼  )rR   rk   rl   r×  s   &&&&r5   Ú_special_binaryÚMetalOverrides._special_binaryÞ  s5   € Ü	�‰×Ñ×Ñ˜^Ô,Ø˜d˜V 1 Q C r¨!¨¨AÐ.Ð.r7   c               ó   € V ^8„  d   QhRR/# ©r/   r1   ÚNoner3   )r4   s   "r5   r6   rÜ   ã  s   € ÷ >ñ >¨ñ >r7   c           
     	ó  € R F/  p\        W\        P                  ! V P                  VR7      4       K1  	  \        P                  ! V P                  RR7      V n        R F7  p\        V V\        P                  ! V P                  VR,           R7      4       K9  	  R F/  p\        W\        P                  ! V P
                  VR7      4       K1  	  R F7  p\        V V\        P                  ! V P
                  VR,           R7      4       K9  	  R# )	Úerf)r×  Ú	log_gammaÚ_forwardN)	rå  ÚerfcÚerfinvÚi0Úi0eÚi1Úi1eÚdigammaÚspherical_bessel_j0)
Ú	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)Ú	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   & r5   Ú_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öó
r7   r3   ©NT)6rÍ   rÎ   rÏ   rÐ   rÑ   rÒ   rä   ré   rî   rü   r   r   r$  r)  r0  r5  r9  r>  rE  rJ  rá   rS  rW  r[  r_  rc  rg  rk  rp  ru  ry  r}  r�  r…  rŠ  r�  r“  r—  r›  r   r¥  rª  r°  r´  rÀ  rÄ  rÊ  rÎ  rÔ  rÛ  rß  Úclassmethodr  rÔ   r3   r7   r5   rÖ   rÖ   È   s„  † Ùbàö<ó ð<ð ôbó ðbð
 ô#ó ð#ð ô(ó ð(ð ô6ó ð6ð
 ô ó ð ðD ôNó ðNð ô2ó ð2ð ô>ó ð>ð
 ô>ó ð>ð
 ôó ðð ôó ðð ô$ó ð$ð ô$ó ð$ð ô+ó ð+ð ô+ó ð+ð ô"ó ð"ð ô&ó ð&ð ô+ó ð+ð ô(ó ð(ð ô+ó ð+ð ô+ó ð+ð ô,ó ð,ð ô,ó ð,ð ô,ó ð,ð ô4ó ð4ð ô,ó ð,ð ô3ó ð3ð
 ô-ó ð-ð ô,ó ð,ð ô-ó ð-ð ô5ó ð5ð ô$ó ð$ð ô#ó ð#ð ô:ó ð:ð
 ô$ó ð$ð ôó ðð ô#ó ð#ð ô5ó ð5ð ô6ó ð6ð ôIó ðIð ô#ó ð#ð ô1ó ð1õ
*õ/ð ô>ó ô>r7   rÖ   Úmpsc                  ó®  a € ] tR tRt$ Rt]tRtRtRt	^ t
]! 4       P                  t]! 4       P                  t]! 4       P                  t]t]! R.4      tR]R&   . tR	]R
&   R V 3R lltR R ltR R ltR*R R lltR R ltRRR]P:                  ! 4       3R R lltR R ltR R lt R R lt!R R  lt"R+R! R" llt#R,R# R$ llt$R% R& lt%R' R( lt&R)t'V ;t(# )-ÚMetalKerneli(  z;Implement Metal codegen based on the SIMDKernel abstractionr  úauto i   ÚutilszOrderedSet[str]r½  zlist[IterationRangesEntry]Úmultistage_reduction_entryc               ó$   € V ^8„  d   QhRRRRRR/# )r/   Útilingzdict[str, sympy.Expr]Úkwargsr   r1   rã  r3   )r4   s   "r5   r6   ÚMetalKernel.__annotate__7  s(   € ÷ -ñ -à%ð-ð ð-ð 
ñ	-r7   c                	ó^   <€ \         SV `  ! V3/ VB  \        P                  ! 4       V n        R # r±   )ÚsuperÚ__init__r  r  Úacc_var_ids)rR   r  r  Ú	__class__s   &&,€r5   r  ÚMetalKernel.__init__7  s&   ø€ ô
 	‰Ò˜Ñ* 6Ò*Ü$Ÿ?š?Ó,ˆÖr7   c               ó    € V ^8„  d   QhRRRR/# )r/   rØ   rÙ   r1   r2   r3   )r4   s   "r5   r6   r  ?  s   € ÷ %ñ % +ð %°#ñ %r7   c                	ó   € \         V,          # r±   rè   )rR   rØ   s   &&r5   Údtype_to_strÚMetalKernel.dtype_to_str?  s   € Ü˜eÕ$Ð$r7   c               ó$   € V ^8„  d   QhRRRRRR/# )r/   r×  r2   ÚindexrF   r1   r   r3   )r4   s   "r5   r6   r  B  s&   € ÷ @ñ @˜ð @ Zð @°Kñ @r7   c                óŠ  € V P                   P                  V4      pV P                  V4      p\        P                  P                  V4      pV RV P                  V4       R2pV\        P                  \        P                  39   d   RV R2p\        P                  pV P                  P                  V P                  WTR7      # )z"Codegen a load from an InputBufferÚ[Ú]rŸ   rK   r	  )rO   Úinputrö   r   ÚgraphÚ	get_dtyperõ   r>   Úfloat16Úbfloat16Úfloat32r÷   rø   Úloads)rR   r×  r&  rû   rØ   Úlines   &&&   r5   ÚloadÚMetalKernel.loadB  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Ð Ó?Ð?r7   Nc          
     ó,   € V ^8„  d   QhRRRRRRRRR	R
/# )r/   r×  r2   r&  rF   Úvaluer   Úmoder   r1   rã  r3   )r4   s   "r5   r6   r  P  s4   € ÷ <ñ <Øð<Ø *ð<Ø3>ð<ØFOð<à	ñ<r7   c                	ó|  € V P                   P                  V4      pV P                  V4      pV P                  \        P
                  P                  V4      4      pR V RV R2pVf   V RV P                  V4       RV R2pM]VR8X  dI   V P                  P                  R4       R	V R
2p	RV	 RV R2p
V	 RV
 RV P                  V4       RV R2pM\        RV 24      hV P                  '       d'   V P                  P                  \        W4      4       R# V P                  P                  \        W4      4       R# )rÞ   rß   rK   Nr(  z] = r  Ú
atomic_addÚatomiczc10::metal::AtomicType<Ú>zreinterpret_cast<device z
::type *>(z::atomic_add(rJ   r  zUnimplemented store mode )rO   Úoutputrö   r#  r   r+  r,  rõ   r½  r¾  ÚRuntimeErrorÚinside_reductionrù   r  r   Ústores)rR   r×  r&  r5  r6  rû   Ú	dtype_strÚcast_valr1  Úatomic_typeÚcast_vars   &&&&&      r5   ÚstoreÚMetalKernel.storeP  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×"Ñ"¤<°Ó#;Ö<à�K‰K×!Ñ!¤,¨tÓ":Ö;r7   c               ó(   € V ^8„  d   QhRRRRRRRR/# )	r/   r×  r2   r&  rF   r5  r   r1   rã  r3   )r4   s   "r5   r6   r  e  s)   € ÷ 	8ñ 	8 Cð 	8°
ð 	8À;ð 	8ÐSWñ 	8r7   c                	ó˜  € V P                   P                  V4      pV P                  V4      pV P                  \        P
                  P                  V4      4      p\        R  V P                   4       4      pV RV P                  V4       RV RV R2pRVP                   RV 2pV P                  P                  \        W4      4       R# )c              3  óL   "  € T F  qP                   '       g   K  Vx € K  	  R # 5ir±   ©Úis_reduction©rÄ   Úts   & r5   rÅ   Ú.MetalKernel.store_reduction.<locals>.<genexpr>j  s   é € ÐKÑ(8 1¿N½NŸQšQÓ(8ùó   ‚$š
$r(  z] = static_cast<rß   r  r  z == 0) N)rO   r;  rö   r#  r   r+  r,  ÚnextÚrange_treesrõ   r×  r>  r  r   )rR   r×  r&  r5  rû   r?  Úreduction_dimr1  s   &&&&    r5   Ústore_reductionÚMetalKernel.store_reductione  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Ö7r7   Tc               ó0   € V ^8„  d   QhRRRRRRRRR	R
RR/# )r/   rØ   zstr | torch.dtypeÚ
elem_countz
int | NoneÚdefault_valuez
Any | NoneÚis_threadgroupr"   ró   zValueRanges[Any]r1   r   r3   )r4   s   "r5   r6   r  p  sF   € ÷ ñ à ðð ðð "ð	ð
 ðð !ðð 
ñr7   c                	óî  € \        V\        P                  4      '       d   V P                  V4      pR \	        V P
                  4       2p\        P                  P                  WeV4      pV'       d   RMRpW� RV 2,          pV'       d   VRV P                  V4       R2,          pVe    V'       d   \        R4      hVRV 2,          pV P                  P                  W€P                  ,           4       V# )Útmp_acc_zthreadgroup Ú r
  r(  r)  z+Thread group var can not have default valueú = )r=   r>   rØ   r#  rN  r  r   rô   Úcreate_cse_varÚsexprrˆ   Úindexing_coder  Úsuffix)	rR   rØ   rT  rU  rV  ró   Úvar_namerû   Úvar_defs	   &&&&&&   r5   Ú_new_idxvarÚMetalKernel._new_idxvarp  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ØÒ$ßÜ$Ð%RÓSÐSØ˜˜]˜OÐ,Õ,ˆGØ×Ñ×$Ñ$ W¯{©{Õ%:Ô;Øˆ
r7   c          
     ó,   € V ^8„  d   QhRRRRRRRRRR/# ©	r/   rØ   rÙ   rÚ   Úreduction_typer   r5  z%CSEVariable | tuple[CSEVariable, ...]r1   r3   )r4   s   "r5   r6   r  ‡  s<   € ÷ ñ àðð ðð &ð	ð
 5ðð 
/ñr7   c                óÐ   € W#V3pWPP                   P                  9   d   V P                   P                  V,          # V P                  WW44      pW`P                   P                  V&   V# )z)Caching wrapper around _reduction_nocache)r÷   Úreduction_cacheÚ_reduction_nocache)rR   rØ   rÚ   re  r5  r  rf   s   &&&&&  r5   Ú	reductionÚMetalKernel.reduction‡  s\   € ð °Ð6ˆ	àŸ™×0Ñ0Ô0Ø—8‘8×+Ñ+¨IÕ6Ð6Ø×(Ñ(¨¸>ÓQˆØ.4�‰× Ñ  Ñ+Øˆr7   c          
     ó,   € V ^8„  d   QhRRRRRRRRRR/# rd  r3   )r4   s   "r5   r6   r  —  sC   € ÷ z2ñ z2àðz2ð ðz2ð &ð	z2ð
 5ðz2ð 
/ñz2r7   c                óŽ  € V P                   '       g   \        R4      hV P                  '       d   \        R4      hR R lpRp^pV P                   Fª  pVP                  '       g   K  V'       d
   VR,          pWhP
                   RV 2,          p\        VP                  \        P                  4      '       d   WxP                  ,          pK}  V\        P                  ! VP                   R2R	R	R
7      ,          pK¬  	  \        WpP                  4      pV P                  V4      p	\        V\        P                  4      '       d   TMV P                  p
\        V\        P                  4      '       d   \        WpP                   4      MV P                   pVR8X  dŠ   V P#                  V4      pV P$                  P'                  V R24       V P$                  P'                  R4       V P(                  P+                  RV RV R24       V P,                  P'                  R4       V# V P.                  P1                  R4       VR=9   d½   \2        V,          pV P#                  WÛ4      pV P4                  '       g   TpMFVR8X  d   R>MR?w  ppV P#                  VVRR7      pV P(                  P+                  V RV RV R24       V P6                  P9                  V P,                  RV RV RV RV RV	 R2\2        V,          R7      # VR@9   dÏ   V P#                  W+4      p\:        V,          pRV RV R2pV P4                  '       g   TpMOVR8X  d   R MR!pR"V R#V R$2pV P#                  VVRR7      pV P(                  P+                  V R%V RV RV R&24       V P6                  P9                  V P,                  RV RV RV RV RV	 R2\2        V,          R7      # VRA9   EdD   V P#                  W+4      pV P#                  W4      p\:        V,          pRV RV R2pV P4                  '       g   TpR\:        V,           RV R2pM¤VR'8X  d   R MR!pR"V R#V R$2pV P#                  VVRR7      pV P#                  V^ RR7      p\=        R( V P>                  PA                  4        4       4      pV P(                  P+                  R)V R#V R*V RV R+V R,V R-V R,VP
                   R.24       V P6                  P9                  V P,                  RV RV RV RV RV RV RV	 R2VR7      # VR/8X  Ed5   V P4                  '       g   V P#                  R0V
4      pV P(                  P+                  V R1V R2V R324       V P6                  P9                  V P(                  R4V RV RV	 R2\B        PD                  R7      pV! V4      # V P#                  R0V
4      pV R1V R52pV P$                  P+                  V R624       V P(                  P'                  V R7V R8V R924       V P6                  P9                  V P,                  R4V RV RV	 R2\B        PD                  R7      pV! V4      # VR:8X  EdW   \        V\F        4      '       g   \        R;4      hV P#                  R0V
4      pV R1V R52pR<V^ ,           RV^,           RV^,           R2pV P$                  P+                  V R624       V P4                  '       dD   V P$                  P+                  V R624       V P(                  P'                  V R7V RV R&24       M!V P(                  P'                  V R,V R24       V P6                  P9                  V P4                  '       d   V P,                  MV P(                  R4V RV RV	 R2\B        PD                  R7      pV! V4      # \I        V4      h)Bz]Codegen a reduction operation.
Only sum and prod operations are somewhat reasonable optimizedzexpected to be inside reductionz&expected no load mask during reductionc               ó    € V ^8„  d   QhRRRR/# )r/   Úres3r   r1   ztuple[CSEVariable, ...]r3   )r4   s   "r5   r6   Ú4MetalKernel._reduction_nocache.<locals>.__annotate__¥  s   € ÷ 	ñ 	 ð 	Ð1Hñ 	r7   c           
     ó    € \         P                  ! R  Uu. uF)  p\        V  RV 2V P                  V P                  4      NK+  	  up4      # u upi )Úxyzr“   )r   Ú_unwrapr   ró   rØ   )rn  rK  s   & r5   Ú_unwrap_helperÚ6MetalKernel._reduction_nocache.<locals>._unwrap_helper¥  sE   € ä×%Ò%ÙNSÓTÉeÈ” ˜v Q q c˜]¨D¯K©K¸¿¹ÖDÉeÑTóð ùÚTs   •/ArY  ú + r•   ÚnumelT©ÚintegerÚpositiveÚanyz	 = false;z7threadgroup_barrier(metal::mem_flags::mem_threadgroup);z
                if (z) {
                    z' = true;
                }
            Úreduction_utilsÚsumF)rU  rV  r
  z= r  zc10::metal::threadgroup_rY   rJ   rK   r	  r   rÞ   rß   ÚMaxOpÚMinOpz::c10::metal::Ú<z>::identity()z = ::c10::metal::r  Úargmaxc              3  óL   "  € T F  qP                   '       g   K  Vx € K  	  R # 5ir±   rH  rJ  s   & r5   rÅ   Ú1MetalKernel._reduction_nocache.<locals>.<genexpr>  s   é € ð Ù=˜!ÇÅ—A’AÓ=ùrM  z#
                if (::c10::metal::z>::replace(z)) {
                    rZ  z;
                    z$;
                }
                Úwelford_reduceÚfloat3r(  z] = float3(z, 0.0, 1.0);z(c10::metal::threadgroup_welford_combine(r)  z = 0.0;z! = ::c10::metal::welford_combine(z	, float3(z, 0.0, 1.0));Úwelford_combinez&Input to welford combine must be tuplezfloat3()Úprodr|  )r   ro   )ra   Ú*)r   rw   )Úargminr€  )%r=  rˆ   Ú
_load_maskrO  rI  r×  r=   rv  ÚsympyÚIntegerÚSymbolÚprefixr
   Úmax_threadgroup_sizer\  r   Úsimd_group_sizera  r]  r  rù   r  r>  r½  r¾  r   r  r÷   rø   rã   rN  Úrange_tree_nodesÚvaluesr>   r/  ÚtupleÚNotImplementedError)rR   rØ   rÚ   re  r5  rs  Úreduction_idxÚacc_buf_sizeÚrdÚacc_buf_size_strÚacc_buf_alloc_sizeÚshmem_buf_sizeÚaccÚ	acc_dtypeÚacc_bufr0   Ú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   &&&&&                        r5   rh  ÚMetalKernel._reduction_nocache—  sr  € ð ×$×$Ð$Ü Ð!BÓCÐCØ�?�?ˆ?Ü Ð!IÓJÐJõ	ð ˆØˆØ×"Ô"ˆBØ—?—?�?ÙßØ Õ&�Ø§¡˜y¨¨L¨>Ð:Õ:ˆMä˜"Ÿ(™(¤E§M¡M×2Ò2Ø§¡Õ(’à¤§¢Ø—y‘y�k Ð'°Àô!õ ’ñ #ô ˜<×)BÑ)BÓCˆØŸ:™: lÓ3Ðô ˜,¬¯©×6Ò6ñ à×*Ñ*ð 	ô ˜,¬¯©×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Ð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Ð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Ð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Ð2Ø×*Ñ*¨8Ð5GÓH�Ø—‘×#Ñ#Ø�i˜q  ¨{¸5¸'ÀÐNôð Ÿ™×*Ñ*Ø—L‘LØ>¸w¸iÀrÈ-ÈÐXZÐ[kÐZlÐlmÐnÜŸ-™-ð +ó �ñ
 & fÓ-Ð-Ø×&Ñ& xÐ1CÓDˆGØ '˜y¨¨-¨¸Ð:ˆNØ×Ñ×%Ñ%¨Ð(8¸Ð&@ÔAØ�L‰L×"Ñ"Ø!Ð"Ð"CÀNÐCSÐS\Ð]bÐ\cÐcpÐqôð —X‘X×&Ñ&Ø—‘Ø:¸7¸)À2ÀmÀ_ÐTVÐWgÐVhÐhiÐjÜ—m‘mð 'ó ˆFñ
 " &Ó)Ð)ØÐ.Õ.Ü˜e¤U×+Ò+Ü$Ð%MÓNÐNØ×&Ñ& 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Ø:¸7¸)À2ÀmÀ_ÐTVÐWgÐVhÐhiÐjÜ—m‘mð 'ó ˆFñ
 " &Ó)Ð)Ü! .Ó1Ð1r7   c               ó    € V ^8„  d   QhRRRR/# )r/   Úentryr   r1   rã  r3   )r4   s   "r5   r6   r  S  s   € ÷ Hñ HÐ4Hð HÈTñ Hr7   c                	óä	  a€ V P                  SP                  4      pV P                  V4      pSP                  '       dZ   \	        SP
                  P                  \        P                  4      '       d`   SP
                  P                  V P                  8:  d;   V P                  P                  V P                   R SP                   RV R24       R# \	        SP
                  P                  \        P                  4      '       d   SP
                  P                  M/\        P                  ! SP
                  P                   R2RRR7      p\         ;QJ d)    V3R lV P"                   4       F  '       g   K   RM	  RM! V3R lV P"                   4       4      pSP
                  P                   R	2pV'       Egƒ   V P"                  P%                  S4       V\'        V P                  ^,
          4      ,           \'        V P                  4      ,          pV P                  V4      pSP
                  P                  p	V P(                  P                  R
SP
                  P                   RSP
                  P                   RV RSP
                  P                   R2	4       V P(                  P+                  4       ;_uu_ 4        \	        V\        P                  4      '       dT   V P(                  P                  V P                   R V RV P                   RSP
                  P                   RV	 R2
4       MHV P(                  P                  V P                   R V RV RV	 RSP
                  P                   R2
4       \	        V\        P                  4      '       g   WpP                  ,          V8w  d#   V P(                  P                  RV RV R24       VP-                  SP
                  P                  V4      p
V P(                  P                  V P                   R SP                   RV
 R24       RRR4       R# V P(                  P+                  4       ;_uu_ 4        VP-                  SP
                  P                  V4      p
V P(                  P                  V P                   R SP                   RV
 R24       RRR4       R#   + '       g   i     R# ; i  + '       g   i     R# ; i)r
  rZ  r  Nrv  Trw  c              3  óR   <"  € T F  qP                   SP                   J x € K  	  R # 5ir±   ©Úroot)rÄ   Úer¬  s   & €r5   rÅ   Ú=MetalKernel.codegen_iteration_ranges_entry.<locals>.<genexpr>g  s    øé € ð %
Ù*I Q�F‰F�e—j‘jÕ Ó*Iùs   ƒ$'FÚ_linear_idxz	for(auto z
_cnt = 0; z_cnt < z; ++z_cnt) {r•   z_cnt + ru  z_cnt;r  ú >= z) break;)Úrename_indexingrE   r\  rI  r=   r°  rv  rŠ  r‹  rŽ  r]  r  Úindex_dtyper×  rŒ  r�  rz  r  Úappendr+   r  r  Úreplace)rR   r¬  rü   Ú	index_strÚacc_sizeÚroot_already_processedÚlinear_idx_nameÚ	loop_sizeÚloop_size_strÚ	root_nameÚsub_index_strs   &f         r5   Úcodegen_iteration_ranges_entryÚ*MetalKernel.codegen_iteration_ranges_entryS  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ó%
§§¢ô %
Ø*.×*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   ËE3S
Ñ!ASÓ
S	ÓS/	c               ó   € V ^8„  d   QhRR/# râ  r3   )r4   s   "r5   r6   r  �  s   € ÷ &ñ &˜dñ &r7   c                ó,  a€ V P                   '       Edª   V P                  P                  4       ;_uu_ 4        V P                  P                  V P                  4       V P                  P                  V P
                  4       RRR4       V P                  P                  R\        V P                   4      ,          4       V P                  P                  \        R V P                  P                  P                  4        4       4      4       V P                    Uu. uF  qP                  NK  	  ppV P                   P                  4        V P                  P                  4        FU  o\         ;QJ d    V3R lV 4       F  '       g   K   RM	  RM! V3R lV 4       4      '       g   KE  SP#                  4        KW  	  MJV P                  P                  V P                  4       V P                  P                  V P
                  4       V P                  P                  V P$                  4       V P                  P                  4        V P
                  P                  4        V P$                  P                  4        R#   + '       g   i     EL; iu upi )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   "  € T F)  p\        V\        4      '       d   TMV3 F  pVx € K	  	  K+  	  R # 5ir±   )r=   r’  )rÄ   ÚitemÚvs   &  r5   rÅ   Ú+MetalKernel.codegen_body.<locals>.<genexpr>±  s7   é € ð á A˜Ü&0°´u×&=Ò&=™dÀDÀ7ÒJ˜õ áJñ Û Aùs   ‚13c              3  ó@   <"  € T F  pSP                   VJ x € K  	  R # 5ir±   r¯  )rÄ   ÚrÚnodes   & €r5   rÅ   rÉ  »  s   øé € Ð5©u¨!�t—y‘y A•~«uùs   ƒTF)r  r  r  r  r0  rù   r  r‡   r÷   Ú
invalidater	   rg  r‘  r°  Úclearr�  rz  Úcache_clearr>  )rR   r±  ÚrootsrÌ  s   &  @r5   Úcodegen_bodyÚMetalKernel.codegen_body�  s²  ø€ ð ×*×*Ñ*Ø—‘×!Ñ!×#Õ#Ø—	‘	× Ñ  §¡Ô,Ø—	‘	× Ñ  §¡Ô.÷ $ð �I‰I×Ñ ¤c¨$×*IÑ*IÓ&JÕ JÔKð
 �H‰H×ÑÜñ à $§¡× 8Ñ 8× ?Ñ ?Ô Aóó ôð &*×%DÒ%DÓEÑ%D —V”VÑ%DˆEÐEØ×+Ñ+×1Ñ1Ô3Ø×-Ñ-×4Ñ4Ö6�ß“3Ô5©uÓ5—3—3’3Ô5©uÓ5×5Ô5Ø×$Ñ$Ö&ò 7ð �I‰I×Ñ˜TŸZ™ZÔ(Ø�I‰I×Ñ˜TŸ\™\Ô*Ø�	‰	×Ñ˜Ÿ™Ô%Ø�
‰
×ÑÔØ�‰×ÑÔØ�‰×ÑÖ÷7 $×#Ð#üò  Fs   ·AI=ÄJÉ=J	c               ó    € V ^8„  d   QhRRRR/# )r/   r×  r2   r1   r3   )r4   s   "r5   r6   r  Å  s   € ÷ iñ i 3ð iÀñ ir7   c                ó�  € V P                  4        \        4       pTp\        P                  P                  '       d   VP                  R4       MVP                  R4       V P                  4       pVP                  4       ;_uu_ 4        \        P                  P                  '       g*   V P                   F  pVP                  RV R24       K  	  MxV P                   Uu. uF	  pRV R2NK  	  pp\        V\        \        4      P                  P                  P                  R,          .\        4       4      pVP                  V4       V P                  '       d~   \        P                   ! R V P"                   4       4      p\%        V\&        P(                  4      '       d   \+        W€P,                  4      MV P,                  p	VP                  RV	 R24       VP                  R	V R
24       VP                  4       ;_uu_ 4        V P.                  P0                  P3                  4        F]  w  r«W P4                  9   d   K  V P7                  \        P                  P9                  V
4      4      pVP                  RV RV R24       K_  	  V P.                  P:                  P3                  4        F¨  w  r«\        P                  P9                  V
4      pV\<        P>                  8X  dG   \        P                  PA                  V
4      pVe   VPC                  4       . 8w  d   \E        R4      hRpMV P7                  V4      pVP                  RV RV R24       Kª  	  V P.                  PF                  PI                  4        F  pVP                  RV R24       K  	  V FN  p\%        VPJ                  \&        P(                  4      '       d   K/  VP                  RVPL                   R24       KP  	  RV P                  9   d   VP                  R4       \O        V4      ^8¼  d   \Q        R4      h\O        V4      ^8”  d   R\O        V4       2MRp\O        V4      ^8X  d   V^ ,          PR                  MRpV P                  '       d   RMRpVP                  V RV RV 24       V P                  '       d   VP                  V R24       RRR4       VP                  R4       VP                  4       ;_uu_ 4        \O        V4      ^8”  dH   \U        V4       F8  w  ppVP                  RVPR                   R\W        ^xV,           4       R 24       K:  	  VPY                  V PZ                  4       VPY                  V P\                  4       RRR4       VP                  R!4       RRR4       \        P                  P                  '       d"   VP                  R"4       VP_                  4       # VP                  R#4       VP_                  4       # u upi   + '       g   i     ELP; i  + '       g   i     L¤; i  + '       g   i     Lž; i)$z3Called at the end to generate a final kernel stringz(R"MTL(úcompile_mps_shader('''z#include <c10/metal/z.h>Úincludec              3  ó`   "  € T F$  qP                   '       g   K  VP                  x € K&  	  R # 5ir±   )rI  rv  rJ  s   & r5   rÅ   Ú-MetalKernel.codegen_kernel.<locals>.<genexpr>á  s   é € ð 1Ù%5 ¿½”G�A—G–GÓ%5ùs   ‚.š.z$[[max_total_threads_per_threadgroup(z)]]úkernel void rY   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)   Ú
thread_posrY  r
  z [[thread_position_in_grid]]z- group_pos [[thread_position_in_threadgroup]]r  r  z = thread_pos.r  rÅ  z)MTL");ú'''))0rÑ  r   r   r+  Úcpp_wrapperr  Úactive_range_treesr  r½  r   r   Ú__file__Úparentr	   r=  Úmathr†  rO  r=   rŠ  r‹  rw   rŽ  rO   Úoutput_buffersÚitemsÚremoved_buffersr#  r,  Úinput_buffersr>   Úfloat64Útry_get_bufferÚget_sizer<  Úsizevarsr‘  rv  r�  r‡   rˆ   r×  Ú	enumerateÚchrr  r]  r  r  )rR   r×  ÚcodeÚfn_nameÚidx_varsÚheaderr½  Úheader_contentsÚtotal_reduction_sizeÚthreadgroup_sizeÚouterÚinnerr?  rØ   Ú	outer_bufr¦  Úthread_pos_dtypeÚthread_pos_var_nameÚthread_pos_suffixÚidxrû   s   &&                   r5   Ú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˜\¨'¨°!Ð4Ô5Ø—‘—•Ø$(§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ä�x“= AÔ%Ü(Ð)PÓQÐQä.1°(«m¸aÔ.?�dœ3˜x›=˜/Ñ*ÀVð !ô ),¨H«¸Ô(:�H˜Q•K×$Ò$Àð $ð ,0×+@×+@Ð+@¡CÀbÐ!Ø—‘Ø'Ð(¨Ð*=Ð)>Ð>ZÐ[lÐZmÐnôð ×(×(Ð(Ø—N‘NØ+Ð,Ð,YÐZô÷Y ð^ �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Ô÷m ôp �7‰7××ÐØ�N‰N˜9Ô%ð �}‰}‹Ðð �N‰N˜6Ô"à�}‰}‹Ðùòq÷2 —�ú÷` —ú÷] �]úse   ÂAX5Ã X	Ã/A,X5ÅB*X5ÈI=X	Ò.X	Ò3X	Ó1X5Ó8BX"	ÖX5Ø	X5ØXØ	X5Ø"X2Ø-X5Ø5Y	c               ó(   € V ^8„  d   QhRRRRRRRR/# )	r/   r×  r2   rÌ  r   Údeallocate_wsr"   r1   rã  r3   )r4   s   "r5   r6   r  0  s2   € ÷ [
ñ [
Øð[
Ø"ð[
Ø:>ð[
à	ñ[
r7   c           
     óÞ	  € \         P                  P                  pV P                  P                   F  pVP                  V4       K  	  V P                  P                  4       w  rgrh\        Wx4       U	U
u/ uF  w  rš\        V	4      V
bK  	  pp	p
. V P                  P                  P                  4       OV P                  P                  P                  4       OpV Uu. uF  qÝV P                  9  g   K  VNK  	  ppYÀP                  P                   Uu. uF  p\        V4      NK  	  up,          pV Uu. uF  qÛV,          NK  	  ppV P                   Fô  p\        VP                  \         P"                  \$        34      '       d   K5  \        VP                  \         P&                  4      '       d   VP                  pM3\         P                  P                  P)                  W4      P*                  pVP,                  '       d   V P.                  '       g   KÅ  VP1                  \        V4      4       VP1                  \$        4       Kö  	  \         P                  P2                  '       d   V P4                  MV P6                  pR R lp\9        V P;                  4       4      ^ 8”  d�   V P;                  4        Uu. uFG  pT! VP,                  '       d!   \=        VP                  V P>                  4      MVP                  4      NKI  	  ppVP1                  V! VR4      4       VP1                  \@        4       M+\         P                  P2                  '       d   \C        R4      hV P.                  '       d†   V P;                  4        Uu. uF=  pVP,                  '       d'   V! \=        VP                  V P>                  4      4      MRNK?  	  ppVP1                  V! VR4      4       VP1                  \@        4       M;\         P                  P2                  '       d   VR.,          pVP1                  R4       RV PD                  9   d\   \         P                  P2                  '       g<   TP1                  R	\9        V Uu. uF  qÝf   K  R
V9  g   K  VNK  	  up4       24       VPG                  VV\H        PJ                  ! R4      RVR7       R# u up
p	i u upi u upi u upi u upi u upi u upi )z 
Codegens a call to this kernel
c               ó$   € V ^8„  d   QhRRRRRR/# )r/   Úthreadsz	list[str]Úkwargr2   r1   r3   )r4   s   "r5   r6   Ú-MetalKernel.call_kernel.<locals>.__annotate__U  s!   € ÷ 	9ñ 	9 Ið 	9°cð 	9¸cñ 	9r7   c                óÎ   € \         P                  P                  '       d+   V  Uu. uF	  pR V R2NK  	  p pRRP                  V 4       R2# V RRP                  V 4       R2# u upi )zstatic_cast<uint64_t>(rK   Ú{rJ   rÅ  z=[r)  )r   r+  rÞ  Újoin)r  r  rK  s   && r5   Úformat_threadsÚ/MetalKernel.call_kernel.<locals>.format_threadsU  sl   € Ü�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_sizeNrÛ  zerror_buf_idx=Ú=r  F)ÚdeviceÚtritonÚ	arg_types)&r   r+  Úwrapper_coderO   rê  Úensure_size_computedÚpython_argdefsÚzipr2   rã  Úkeysræ  rå  rO  r=   rv  rŠ  r‹  r%   rŒ  Úgenerate_numel_exprrõ  rI  r=  r·  rÞ  ÚcexprÚpexprr‡   rß  r
   rŽ  rb   r<  r½  Úgenerate_kernel_callr>   r  )rR   r×  rÌ  rþ  ÚwrapperrÈ  Ú_Ú	call_argsr  Úcall_argÚarg_typeÚarg_name_to_typerO   rg   ÚtreerE   Úexpr_printerr  r  s   &&&&               r5   Úcall_kernelÚMetalKernel.call_kernel0  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Ð%8�t—z’z¸d¿j¹jˆõ	9ô ˆt×&Ñ&Ó(Ó)¨AÔ-ð ×0Ñ0Ô2óñ 3�Añ à—~—~�~ô ˜Ÿ™ ×!:Ñ!:Ô;àŸ™öñ
 3ð ð ð �K‰K™ w°	Ó:Ô;Ø×ÑœTÕ"ä�w‰w×"×"Ð"Ü"Ð#CÓDÐDà× × Ð ð
 ×0Ñ0Ô2ó	ñ 3�Að —>—>�>ñ œS §¡¨$×*CÑ*CÓDÔEàòñ 3ð	 ð ð �K‰K™ w°Ó=Ô>Ø×ÑœTÕ"ä�w‰w×"×"Ð"ð ˜˜•�Ø× Ñ  Ô&ð �d—l‘lÔ"¬1¯7©7×+>×+>Ð+>Ø�K‰KØ ¤±TÓ%`±T¨c¤cÐQTÐ\_ÑQ_§c c±TÑ%`Ó!aÐ bÐcôð 	×$Ñ$ØØÜ—<’< Ó&ØØð 	%ö 	
ùóS
ùò
 HùÚ4ùÚ;ùò8ùò ùò& &asH   Á4SÃSÃ3SÄSÄ4SÊ9S Ë4S ÎS%Î,S%Ñ9S*ÒS*ÒS*c               ó$   € V ^8„  d   QhRRRRRR/# )r/   Úcondr   Úmsgr2   r1   rã  r3   )r4   s   "r5   r6   r  �  s!   € ÷ ñ ¨ð ¸#ð À$ñ r7   c                	ó8  € \         P                  P                  '       d/   V P                  P	                  V P
                  R V R2RR7       R
# V P                  P                  R4       V P
                  P                  R V R2R\        V4       R2RR	.4       R
# )zif (!ú) returnF©Ú
assignmentrÛ  r  z"    TORCH_REPORT_ERROR(error_buf, r  ú    return;rÅ  N)
r   r+  rÞ  r÷   rø   rù   r½  r¾  r  Úrepr)rR   r#  r$  s   &&&r5   Údevice_assert_asyncÚMetalKernel.device_assert_async�  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Ø!Øð	ör7   c          
     ó,   € V ^8„  d   QhRRRRRRRRRR/# )	r/   rE   rF   ÚsizeÚlowerr"   Úupperr1   rã  r3   )r4   s   "r5   r6   r  ›  s4   € ÷ "ñ "Øð"Ø&0ð"Ø9=ð"ØFJð"à	ñ"r7   c                	óø  € V'       g   V'       g   R # V P                  V4      pV P                  V4      pV'       d   V'       d   RV RV RV R2pMV'       d   V R2pMV RV 2p\        P                  P                  '       d/   V P                  P                  V P                  RV R2RR	7       R # V P                  P                  R
4       V P                  P                  RV R2RV RV R2RR.4       R # )NrY   z < 0 || r´  rK   z < 0r  r&  Fr'  rÛ  r  z,    TORCH_REPORT_ERROR(error_buf, "Index ", z, " out of range [0, ", z, ")");r)  rÅ  )
rõ   r   r+  rÞ  r÷   rø   rù   r½  r¾  r  )rR   rE   r.  r/  r0  Úexpr_strÚsize_strÚ	conditions   &&&&&   r5   Úcheck_boundsÚ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Ø!Øð	ör7   )r  r±   )Úgenerated_kernelr  ))rÍ   rÎ   rÏ   rÐ   rÑ   rÖ   Ú	overridesr^  Únewvar_prefixrŽ  r�  r   rP   r  r   r  rB   r\  Úkexprr	   r½  Ú__annotations__r  r  r#  r2  rC  rQ  r   Úunknownra  ri  rh  rÁ  rÑ  rû  r   r+  r5  rÔ   Ú__classcell__©r  s   @r5   r  r  (  sÛ   ø‡ ÙEà€IØ€FØ€MØÐØ€OÙ‹O×#Ñ#€EÙ‹L× Ñ €EÙÓ×&Ñ&€EØ€EÙ)¨7¨)Ó4€Gˆ_Ó4Ø=?ÐÐ :Ó?÷-ð -õ%õ@÷<õ*	8ð "&Ø$(Ø#Ø#.×#6Ò#6Ó#8÷õ.õ z2õxHõT&÷Pi÷V[
õz÷"ó "r7   r  c                  ó>   a € ] tR tRt]tR V 3R lltR R ltRtV ;t	# )ÚMetalSchedulingiÀ  c               ó    € V ^8„  d   QhRRRR/# )r/   Ú	schedulerzScheduler | Noner1   rã  r3   )r4   s   "r5   r6   ÚMetalScheduling.__annotate__Ã  s   € ÷ $ñ $Ð"2ð $°tñ $r7   c                	ó&   <€ \         SV `  V4       R # r±   )r  r  )rR   rB  r  s   &&€r5   r  ÚMetalScheduling.__init__Ã  s   ø€ Ü‰Ñ˜Ö#r7   c               ó(   € V ^8„  d   QhRRRRRRRR/# )r/   Úsrc_coder2   Únode_schedulezlist[SchedulerNode]rô   r  r1   r3   )r4   s   "r5   r6   rC  Æ  s,   € ÷ -ñ -Øð-Ø,?ð-ØITð-à	ñ-r7   c           	     	óô  € \         P                  P                  pWP                  9   d   VP                  V,          # \         P                  P                  '       dY   R VP                  4        2pTpWdP                  V&   RV R2V,           p\        W$4      w  rxVP                  WQV RV 2RR7       V# RVP                  4        2p	W”P                  V&   Rp
VP                  V
4      \        V
4      ,           pVP                  R4      pWV p. pVP                  R4       FO  pVP                  4       P                  R	4      '       d   K*  VP                  VP                  R
RV	 R24      4       KQ  	  \!        \#        VP$                  4      4      pVP&                  P)                  V	 RV	: R24       V F  pVP&                  P)                  V4       K   	  VP&                  P)                  RV R24       V	# )Úmps_lib_zconst char* z
_source = Ú
F)ÚgpuÚgenerated_kernel_rÕ  rÝ  z#includezkernel void generated_kernel(rÙ  rY   z = async_compile.metal(z, '''z''', rK   )r   r+  r  Úsrc_to_kernelrÞ  Únext_kernel_suffixr   Údefine_kernelr&  r‡   ÚrindexÚsplitÚstripÚ
startswithr·  r¸  r*  Úsortedr½  rð  r  )rR   rG  rH  rô   r  Úmps_lib_nameÚkernel_nameÚoriginsÚdetailed_originsrî  Úmetal_src_startÚstartÚendÚ	metal_srcÚ
body_linesr1  Úheaders_reprs   &&&&             r5   rP  ÚMetalScheduling.define_kernelÆ  sá  € ô —'‘'×&Ñ&ˆØ×,Ñ,Ô,Ø×(Ñ(¨Õ2Ð2ä�7‰7××Ðà% g×&@Ñ&@Ó&BÐ%CÐDˆLØ&ˆKØ.9×!Ñ! (Ñ+Ø% l ^°:Ð>ÀÕIˆHÜ(;¸MÓ(SÑ%ˆGØ×!Ñ!Ø¨7¨)°2Ð6FÐ5GÐ(HÈeð "ô ð Ðð & g×&@Ñ&@Ó&BÐ%CÐDˆØ*1×Ñ˜hÑ'ð 3ˆØ—‘˜Ó/´#°oÓ2FÕFˆØ�o‰o˜fÓ%ˆØ 3Ð'ˆ	ð ˆ
Ø—O‘O DÖ)ˆDØ�z‰z‹|×&Ñ& z×2Ò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àˆr7   r3   )
rÍ   rÎ   rÏ   rÐ   r  Úkernel_typer  rP  rÔ   r=  r>  s   @r5   r@  r@  À  s   ø† Ø€K÷$ð $÷-ó -r7   r@  )LÚ
__future__r   r	  r  Úloggingrâ  Úpathlibr   Útypingr   r   rŠ  Úsympy.printing.precedencer   r>   Ú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   rB  r   r    r!   Ú	getLoggerrÍ   rá   r"   Úint8Úint16Úint32Úint64Úuint8Úuint16Úuint32Úuint64r+   r,   r.  rã   r@   rB   rÖ   Ú_initialize_pointwise_overridesr  r  r@  r3   r7   r5   Ú<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�Ø	‡L�L�(Ø	‡L�L�&Ø	‡L�L�'Ø	‡K�K�Ø	‡J�J�Ø	‡N�N�Hð€õ ô}"�|ô }"ô@Y�[ô Yðx
 × .Ñ .¨uÔ 5Ø × &Ñ &Ô (ôU
�*ô U
ôp3�nö 3r7   