Ë
    g^(hÍi  ã                  ó€  — d dl mZ d dlZd dlZd dlmZmZmZ d dlZd dl	m
Z
 d dlZd dlmZ d dlmZ ddlmZmZ dd	lmZmZmZ d
dlmZmZmZmZmZmZ d
dlm Z m!Z!m"Z" erd dlm#Z# ddl$m%Z%m&Z& ddl'm(Z(m)Z) d
dlm*Z* ejV                  dejX                  dejZ                  dej\                  dej^                  dej`                  dejb                  dejd                  dejf                  di	Z4d%d„Z5 G d„ de«      Z6 G d„ de«      Z7e7jq                  d«        G d „ d!e!«      Z9ejt                  d&d"„«       Z; G d#„ d$e"«      Z<y)'é    )ÚannotationsN)ÚAnyÚOptionalÚTYPE_CHECKING)Ú
PRECEDENCE)ÚExprPrinter)ÚValueRangesé   )Ú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                óØ   — t        | t        «      r:| t        j                  k(  ry| t        j                   k(  ry| | k7  ryt	        | «      S t        | t
        «      r| rdS dS t	        | «      S )NÚ	HUGE_VALFz
-HUGE_VALFÚNANÚtrueÚfalse)Ú
isinstancer&   ÚtorchÚinfÚstrr    )Úvals    úY/var/www/skyplay_api_hub/venv/lib/python3.12/site-packages/torch/_inductor/codegen/mps.pyÚvalue_to_metalr4   2   s_   € Ü�#”uÔØ”%—)‘)ÒØØ”U—Y‘Y�JÒØØ�CŠZØÜ�3‹xˆÜ	�CœÔ	ÙˆvÐ) 'Ð)Üˆs‹8€Oó    c                  óL   — e Zd Zd
d„Zd
d„Zd
d„Zd
d„Zd
d„Zd
d„Zd
d„Z	d
d„Z
y	)ÚMetalExprPrinterc                ó    — |j                   \  }}| j                  |«      }| j                  |«      }|j                  r	d|› d|› d�S d|› d|› d�S )Nú(ú) / (ú)úmetal::floor(©ÚargsÚdoprintÚ
is_integer)ÚselfÚexprÚxÚdivs       r3   Ú_print_FloorDivz MetalExprPrinter._print_FloorDivA   s[   € Ø—‘‰ˆˆ3Ø�L‰L˜‹OˆØ�l‰l˜3ÓˆØ�?Š?Ø�q�c˜˜s˜e 1Ð%Ð%Ø˜q˜c  s e¨1Ð-Ð-r5   c                óâ   — |j                   \  }}}| j                  |«      }|dk7  r0| j                  |«      }|j                  r
d|› d|› d�}n	d|› d|› d�}| j                  |«      }d|› d|› d�S )Nr   r9   r:   r;   r<   z) % (r=   )rA   rB   rC   rD   Úmods        r3   Ú_print_ModularIndexingz'MetalExprPrinter._print_ModularIndexingI   s…   € Ø—i‘i‰ˆˆ3�Ø�L‰L˜‹OˆØ�!Š8Ø—,‘,˜sÓ#ˆCØ�ŠØ˜�s˜% ˜u AÐ&‘à# A 3 e¨C¨5°Ð2�Ø�l‰l˜3ÓˆØ�1�#�U˜3˜%˜qÐ!Ð!r5   c                ó®   — t        |j                  «      dk7  rt        d«      ‚ddj                  t	        | j
                  |j                  «      «      › d�S )Nr
   z$metal::min only supported for 2 argszmetal::min(ú, r;   ©Úlenr>   ÚRuntimeErrorÚjoinÚmapÚ_print©rA   rB   s     r3   Ú
_print_MinzMetalExprPrinter._print_MinU   óF   € Üˆt�y‰y‹>˜QÒÜÐEÓFÐFØ˜TŸY™Y¤s¨4¯;©;¸¿	¹	Ó'BÓCÐDÀAÐFÐFr5   c                ó®   — t        |j                  «      dk7  rt        d«      ‚ddj                  t	        | j
                  |j                  «      «      › d�S )Nr
   z$metal::max only supported for 2 argszmetal::max(rJ   r;   rK   rQ   s     r3   Ú
_print_MaxzMetalExprPrinter._print_MaxZ   rS   r5   c                óz   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › d�S )Nr   úmetal::abs(r   r;   ©rL   r>   rP   rQ   s     r3   Ú
_print_AbszMetalExprPrinter._print_Abs_   s9   € Ü�4—9‘9‹~ Ò"Ð"Ð"Ø˜TŸ[™[¨¯©°1©Ó6Ð7°qÐ9Ð9r5   c                óz   — t        |j                  «      dk(  sJ ‚d| j                  |j                  d   «      › d�S )Nr   zstatic_cast<long>(metal::rint(r   z))rX   rQ   s     r3   Ú_print_RoundToIntz"MetalExprPrinter._print_RoundToIntc   s9   € Ü�4—9‘9‹~ Ò"Ð"Ð"Ø/°·±¸D¿I¹IÀa¹LÓ0IÐ/JÈ"ÐMÐMr5   c                óä   — t        |j                  «      dk(  sJ ‚|j                  \  }}|j                  r|dk  sJ ‚t        d|› d�«      ‚| j	                  |t
        d   «      }d|› d|› d| › d	�S )
Nr
   r   zOFor integer inputs, only non-negative ndigits are currently supported, but got ú.ÚMulz!static_cast<float>(metal::rint(1eú * z) * 1er;   )rL   r>   r@   Ú
ValueErrorÚparenthesizer   )rA   rB   ÚnumberÚndigitsÚ
number_strs        r3   Ú_print_RoundDecimalz$MetalExprPrinter._print_RoundDecimalg   s�   € Ü�4—9‘9‹~ Ò"Ð"Ð"ØŸ)™)‰ˆ�Ø×Òà˜Q’;Ð�;ÜØaÐbiÐajÐjkÐlóð ð ×&Ñ& v¬z¸%Ñ/@ÓAˆ
Ø2°7°)¸3¸z¸lÈ&ÐRYÐQYÐPZÐZ[Ð\Ð\r5   c                ón   — |j                   \  }}d| j                  |«      › d| j                  |«      › d�S )Nústatic_cast<float>(z) / static_cast<float>(r;   )r>   rP   )rA   rB   ÚlhsÚrhss       r3   Ú_print_IntTrueDivz"MetalExprPrinter._print_IntTrueDivs   s;   € Ø—9‘9‰ˆˆSà$ T§[¡[°Ó%5Ð$6Ð6MÈdÏkÉkÐZ]ÓN^ÐM_Ð_`ÐaÐar5   N)rB   ú
sympy.ExprÚreturnr1   )Ú__name__Ú
__module__Ú__qualname__rE   rH   rR   rU   rY   r[   re   rj   © r5   r3   r7   r7   @   s/   „ ó.ó
"óGó
Gó
:óNó
]ôbr5   r7   c                  óÞ  — e Zd Ze	 	 d5	 	 	 	 	 	 	 	 	 d6d„«       Ze	 	 	 	 	 	 	 	 d7d„«       Zed8d„«       Zed9d„«       Zed:d„«       Zed;d„«       Z	ed<d„«       Z
ed=d	„«       Zed=d
„«       Zed=d„«       Zed=d„«       Zed>d„«       Zed>d„«       Zed>d„«       Zed>d„«       Zed>d„«       Zed>d„«       Zed>d„«       Zed>d„«       Zed>d„«       Zed>d„«       Zed>d„«       Zed>d„«       Zed>d„«       Zed>d„«       Zed?d„«       Zed>d„«       Zed>d„«       Zed>d„«       Z ed>d„«       Z!ed>d „«       Z"ed>d!„«       Z#ed>d"„«       Z$ed>d#„«       Z%ed>d$„«       Z&ed=d%„«       Z'ed>d&„«       Z(ed>d'„«       Z)ed=d(„«       Z*ed>d)„«       Z+ed=d*„«       Z,ed>d+„«       Z-ed@d,„«       Z.ed@d-„«       Z/e	 	 	 	 	 	 	 	 	 	 dAd.„«       Z0ed>d/„«       Z1ed=d0„«       Z2ed=d1„«       Z3ed>d2„«       Z4ed>d3„«       Z5ed>d4„«       Z6y)BÚMetalOverridesNc                ó"   — dt         |   › d| › d�S )Nzstatic_cast<ú>(r;   ©ÚDTYPE_TO_METAL)rC   ÚdtypeÚ	src_dtypeÚuse_compute_typess       r3   Úto_dtypezMetalOverrides.to_dtypez   s   € ð œn¨UÑ3Ð4°B°q°c¸Ð;Ð;r5   c                ó"   — dt         |   › d| › d�S )Nz*reinterpret_cast<thread z*>(&r;   ru   )rC   rw   rx   s      r3   Úto_dtype_bitcastzMetalOverrides.to_dtype_bitcastƒ   s   € ð +¬>¸%Ñ+@Ð*AÀÀaÀSÈÐJÐJr5   c                ó   — t        | «      S ©N©r4   )r2   rw   s     r3   ÚconstantzMetalOverrides.constant‰   s   € ä˜cÓ"Ð"r5   c                ó@  — t         j                  j                  t         j                  j                  | «      «      }t         j                  j                  j                  t         j                  j                  |t        | «      ¬«      }t        j                  ||«      S )N)Úbounds)
r   ÚkernelÚindex_to_strÚprepare_indexingÚcseÚgenerateÚcomputer   r   rz   )rB   rw   Úidx_strÚvars       r3   Ú
index_exprzMetalOverrides.index_expr�   sl   € ä—(‘(×'Ñ'¬¯©×(AÑ(AÀ$Ó(GÓHˆÜ�h‰h�l‰l×#Ñ#Ü�H‰H×Ñ˜gÔ.CÀDÓ.Ið $ó 
ˆô �|‰|˜C Ó'Ð'r5   c                óê   — t         j                  j                  | |«      5 } |«       }d d d «       j                  j                  rt        |«      }t        j                  ||«      S # 1 sw Y   ŒAxY wr~   )r   rƒ   Ú
mask_loadsr‚   Úis_boolr    r   Úwhere)ÚmaskÚbodyÚotherÚnew_maskÚresults        r3   ÚmaskedzMetalOverrides.masked•   sa   € ô �X‰X× Ñ   uÓ-ð 	°Ù“VˆF÷	ð �=‰=× Ò Ü˜“KˆEä�y‰y˜ 6¨5Ó1Ð1÷	ð 	ús   ¡A)Á)A2c                ó(   — | › d|› dt        |«      › �S )Nz ? z : r   )ÚaÚbÚcs      r3   r�   zMetalOverrides.where    s   € à��C˜�s˜#œn¨QÓ/Ð0Ð1Ð1r5   c                ód  — t        |t        «      r)|j                  �|j                  j                  s| › d|› �S t        | t        «      r#| j                  t        j
                  k7  rd| › d�n| }t        |t        «      r#|j                  t        j
                  k7  rd|› d�n|}|› d|› d|› d|› d�S )Nú % rg   r;   z - z * metal::floor(ú / )r.   r   rw   Úis_floating_pointr/   r&   ©r—   r˜   Úfloat_aÚfloat_bs       r3   Ú	remainderzMetalOverrides.remainder¤   s¹   € ô �qœ+Ô&Ø—‘Ð#Ø—G‘G×-Ò-à�S˜˜A˜3�<Ðô ˜!œ[Ô)¨a¯g©g¼¿¹Ò.Dð " !  AÑ&àð 	ô ˜!œ[Ô)¨a¯g©g¼¿¹Ò.Dð " !  AÑ&àð 	ð
 �˜#˜g˜YÐ&6°w°i¸sÀ7À)È1ÐMÐMr5   c                óD   — d| › d|› d| › d�}d| › d|› d|› d�}d|› d|› d�S )Nústatic_cast<decltype(ú+ú)>(r;   zc10::metal::max(rJ   rp   ©r—   r˜   Ú
typecast_aÚ
typecast_bs       r3   ÚmaximumzMetalOverrides.maximum¹   óK   € à,¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø,¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø! * ¨R°
¨|¸1Ð=Ð=r5   c                óD   — d| › d|› d| › d�}d| › d|› d|› d�}d|› d|› d�S )Nr£   r¤   r¥   r;   zc10::metal::min(rJ   rp   r¦   s       r3   ÚminimumzMetalOverrides.minimum¿   rª   r5   c                ó   — | › d|› �S )Nz || rp   ©r—   r˜   s     r3   Ú
logical_orzMetalOverrides.logical_orÅ   ó   € à��D˜˜ˆ}Ðr5   c                ó   — | › d|› �S )Nz && rp   r®   s     r3   Úlogical_andzMetalOverrides.logical_andÉ   r°   r5   c                ó   — d| › d�S )Nzmetal::isnan(r;   rp   ©rC   s    r3   ÚisnanzMetalOverrides.isnanÍ   ó   € à˜q˜c Ð#Ð#r5   c                ó   — d| › d�S )Nzmetal::isinf(r;   rp   r´   s    r3   ÚisinfzMetalOverrides.isinfÑ   r¶   r5   c                ó   — d| › d�S )Nzmetal::log(r;   rp   r´   s    r3   ÚlogzMetalOverrides.logÕ   ó   € à˜Q˜C˜qÐ!Ð!r5   c                ó   — d| › d�S )Nzmetal::exp(r;   rp   r´   s    r3   ÚexpzMetalOverrides.expÙ   r»   r5   c                ó   — d| › d�S )NrW   r;   rp   r´   s    r3   ÚabszMetalOverrides.absÝ   r»   r5   c                ó   — d| › d�S )Nzmetal::signbit(r;   rp   r´   s    r3   ÚsignbitzMetalOverrides.signbitá   ó   € à    1Ð%Ð%r5   c                ó   — d| › d�S )Nzmetal::precise::sin(r;   rp   r´   s    r3   ÚsinzMetalOverrides.sinå   ó   € à% a S¨Ð*Ð*r5   c                ó   — d| › d�S )Nzc10::metal::sinc(r;   rp   r´   s    r3   ÚsinczMetalOverrides.sincé   ó   € à" 1 # QÐ'Ð'r5   c                ó   — d| › d�S )Nzmetal::precise::cos(r;   rp   r´   s    r3   ÚcoszMetalOverrides.cosí   rÅ   r5   c                ó   — d| › d�S )Nzc10::metal::i0(r;   rp   r´   s    r3   Úi0zMetalOverrides.i0ñ   rÂ   r5   c                ó   — d| › d�S )Nzc10::metal::i1(r;   rp   r´   s    r3   Úi1zMetalOverrides.i1õ   rÂ   r5   c                ó   — d| › d�S )Nzc10::metal::erf(r;   rp   r´   s    r3   ÚerfzMetalOverrides.erfù   s   € à! !  AÐ&Ð&r5   c                ó   — d| › d�S )Nzc10::metal::erfinv(r;   rp   r´   s    r3   ÚerfinvzMetalOverrides.erfinvý   s   € à$ Q C qÐ)Ð)r5   c                ó   — d| › d�S )Nzc10::metal::log_gamma(r;   rp   r´   s    r3   ÚlgammazMetalOverrides.lgamma  s   € à'¨ s¨!Ð,Ð,r5   c                ó   — d| › d|› d�S )Nzc10::metal::polygamma(rJ   r;   rp   )rC   Úys     r3   Ú	polygammazMetalOverrides.polygamma  s   € à'¨ s¨"¨Q¨C¨qÐ1Ð1r5   c                ó   — d| › d�S )Nzc10::metal::digamma(r;   rp   r´   s    r3   ÚdigammazMetalOverrides.digamma	  rÅ   r5   c                ó   — d| › d�S )Nzmetal::tan(r;   rp   r´   s    r3   ÚtanzMetalOverrides.tan  r»   r5   c                ó   — d| › d�S )Nzmetal::asin(r;   rp   r´   s    r3   ÚasinzMetalOverrides.asin  ó   € à˜a˜S Ð"Ð"r5   c                ó   — d| › d�S )Nzmetal::acos(r;   rp   r´   s    r3   ÚacoszMetalOverrides.acos  rÞ   r5   c                ó   — d| › d�S )Nzmetal::atan(r;   rp   r´   s    r3   ÚatanzMetalOverrides.atan  rÞ   r5   c                ó   — d| › d�S )Nzmetal::sqrt(r;   rp   r´   s    r3   ÚsqrtzMetalOverrides.sqrt  rÞ   r5   c                ó   — d| › d�S )Nzmetal::rsqrt(r;   rp   r´   s    r3   ÚrsqrtzMetalOverrides.rsqrt!  r¶   r5   c                ó   — d| › d�S )Nzmetal::tanh(r;   rp   r´   s    r3   ÚtanhzMetalOverrides.tanh%  rÞ   r5   c                ó   — d| › d�S )Nzmetal::atanh(r;   rp   r´   s    r3   ÚatanhzMetalOverrides.atanh)  r¶   r5   c                óH   — | › d|› �}| › d|› �}d| › d|› d|› d|› d|› d|› d	�S )
Nrœ   r›   z((z
 < 0) != (z	 < 0) ? (z != 0 ? z - 1 : z) : r;   rp   )r—   r˜   ÚquotÚrems       r3   ÚfloordivzMetalOverrides.floordiv-  sR   € ð ��C˜�sˆ|ˆØ��3�q�cˆlˆØ�A�3�j   9¨S¨E°¸$¸¸wÀtÀfÈDÐQUÐPVÐVWÐXÐXr5   c                ó   — d| › d�S )Nr<   r;   rp   r´   s    r3   ÚfloorzMetalOverrides.floor4  r¶   r5   c                ó   — d| › d�S )Nzmetal::sign(r;   rp   r´   s    r3   ÚsignzMetalOverrides.sign8  rÞ   r5   c                óD   — d| › d|› d| › d�}d| › d|› d|› d�}d|› d|› d�S )Nr£   r¤   r¥   r;   zmetal::fmod(rJ   rp   r¦   s       r3   ÚfmodzMetalOverrides.fmod<  sK   € à,¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø,¨Q¨C¨q°°°3°q°c¸Ð;ˆ
Ø˜j˜\¨¨J¨<°qÐ9Ð9r5   c                ó   — d| › d�S )Númetal::trunc(r;   rp   r´   s    r3   ÚtrunczMetalOverrides.truncB  r¶   r5   c                ó¨   — | j                   t        j                  k7  rd| › d�n| }|j                   t        j                  k7  rd|› d�n|}d|› d|› d�S )Nrg   r;   rö   ú/)rw   r/   r&   rž   s       r3   ÚtruncdivzMetalOverrides.truncdivF  sY   € ð 12·±¼5¿;¹;Ò0FÐ'¨ s¨!Ñ,ÈAˆØ01·±¼5¿;¹;Ò0FÐ'¨ s¨!Ñ,ÈAˆØ˜w˜i q¨¨	°Ð3Ð3r5   c                ó   — d| › d�S )Nzmetal::ceil(r;   rp   r´   s    r3   ÚceilzMetalOverrides.ceilN  rÞ   r5   c                ó   — d| › d|› d�S )Nzc10::metal::rand(rJ   r;   rp   ©ÚseedÚoffsets     r3   ÚrandzMetalOverrides.randR  s   € à" 4 &¨¨6¨(°!Ð4Ð4r5   c                ó   — d| › d|› d�S )Nzc10::metal::randn(rJ   r;   rp   rþ   s     r3   ÚrandnzMetalOverrides.randnV  s   € à# D 6¨¨F¨8°1Ð5Ð5r5   c           	     ó    — d| › d|› d|› d|› d�	S )Nzc10::metal::randint64(rJ   r;   rp   )rÿ   r   ÚlowÚhighs       r3   Ú	randint64zMetalOverrides.randint64Z  s%   € ð (¨ v¨R°¨x°r¸#¸¸bÀÀÀaÐHÐHr5   c                ó   — d| › d�S )Nzmetal::round(r;   rp   r´   s    r3   ÚroundzMetalOverrides.round`  r¶   r5   c                óD   — d| › d|› d| › d�}d| › d|› d|› d�}d|› d|› d�S )Nr£   r¤   r¥   r;   zmetal::pow(rJ   rp   )r—   r˜   Úcast_aÚcast_bs       r3   ÚpowzMetalOverrides.powd  sK   € à(¨¨¨1¨Q¨C¨s°1°#°QÐ7ˆØ(¨¨¨1¨Q¨C¨s°1°#°QÐ7ˆØ˜V˜H B v h¨aÐ0Ð0r5   c                ó   — d| › d|› d�S )Nzc10::metal::zeta(rJ   r;   rp   r®   s     r3   ÚzetazMetalOverrides.zetaj  s   € à" 1 # R¨ s¨!Ð,Ð,r5   c                ó   — d| › d�S )Nz c10::metal::spherical_bessel_j0(r;   rp   r´   s    r3   Úspherical_bessel_j0z"MetalOverrides.spherical_bessel_j0n  s   € à1°!°°AÐ6Ð6r5   c                ó   — d| › d�S )Nzc10::metal::xlog1py(r;   rp   r´   s    r3   Úxlog1pyzMetalOverrides.xlog1pyr  rÅ   r5   c                ó   — d| › d�S )Nzc10::metal::entr(r;   rp   r´   s    r3   ÚentrzMetalOverrides.entrv  rÈ   r5   )NT)
rC   r   rw   útorch.dtyperx   zOptional[torch.dtype]ry   r    rl   r1   )rC   r   rw   r  rx   r  rl   r1   )r2   zUnion[bool, float, int]rw   r  rl   r1   )rB   rk   rw   r  rl   r1   )r�   r   r‘   rk   r’   r   rl   r1   )r—   r   r˜   r   r™   r   rl   r1   )r—   r   r˜   r   rl   r1   )r—   r   r˜   r   rl   r1   )rC   r   rl   r1   )rC   r   rÖ   r   rl   r1   )rÿ   r   r   r   rl   r1   )
rÿ   r   r   r   r  r   r  r   rl   r1   )7rm   rn   ro   Ústaticmethodrz   r|   r€   r‹   r•   r�   r¡   r©   r¬   r¯   r²   rµ   r¸   rº   r½   r¿   rÁ   rÄ   rÇ   rÊ   rÌ   rÎ   rÐ   rÒ   rÔ   r×   rÙ   rÛ   rÝ   rà   râ   rä   ræ   rè   rê   rî   rð   rò   rô   r÷   rú   rü   r  r  r  r	  r  r  r  r  r  rp   r5   r3   rr   rr   y   s‹  „ Øð ,0Ø"&ð	<Øð<àð<ð )ð<ð  ð	<ð
 
ò<ó ð<ð ðKØðKØ*ðKØ7BðKà	òKó ðKð
 ò#ó ð#ð ò(ó ð(ð ò2ó ð2ð ò2ó ð2ð òNó ðNð( ò>ó ð>ð
 ò>ó ð>ð
 òó ðð òó ðð ò$ó ð$ð ò$ó ð$ð ò"ó ð"ð ò"ó ð"ð ò"ó ð"ð ò&ó ð&ð ò+ó ð+ð ò(ó ð(ð ò+ó ð+ð ò&ó ð&ð ò&ó ð&ð ò'ó ð'ð ò*ó ð*ð ò-ó ð-ð ò2ó ð2ð ò+ó ð+ð ò"ó ð"ð ò#ó ð#ð ò#ó ð#ð ò#ó ð#ð ò#ó ð#ð ò$ó ð$ð ò#ó ð#ð ò$ó ð$ð òYó ðYð ò$ó ð$ð ò#ó ð#ð ò:ó ð:ð
 ò$ó ð$ð ò4ó ð4ð ò#ó ð#ð ò5ó ð5ð ò6ó ð6ð ðIØðIØ#.ðIØ5@ðIØHSðIà	òIó ðIð
 ò$ó ð$ð ò1ó ð1ð
 ò-ó ð-ð ò7ó ð7ð ò+ó ð+ð ò(ó ñ(r5   rr   Úmpsc                  óH  ‡ — e Zd ZeZdZdZdZ e«       j                  Z
 e«       j                  ZeZ	 	 	 	 	 	 dˆ fd„Zdd„Zdd„Z	 d	 	 	 	 	 	 	 	 	 dd„Zd ej&                  «       f	 	 	 	 	 	 	 dd	„Z	 	 	 	 	 	 	 	 	 	 dd
„Zdd„Zdd„Zddd„Zddd„Z	 	 	 	 	 	 	 	 	 	 dd„Zˆ xZS )ÚMetalKernelú;úauto i   c                óf   •— t        ‰| �  |fi |¤Ž t        j                  «       | _        d| _        y )NF)ÚsuperÚ__init__Ú	itertoolsÚcountÚacc_var_idsÚmultistage_reduction)rA   ÚtilingÚkwargsÚ	__class__s      €r3   r  zMetalKernel.__init__‡  s.   ø€ ô
 	‰Ñ˜Ñ* 6Ò*Ü$Ÿ?™?Ó,ˆÔØ$)ˆÕ!r5   c                ó   — t         |   S r~   ru   )rA   rw   s     r3   Údtype_to_strzMetalKernel.dtype_to_str�  s   € Ü˜eÑ$Ð$r5   c                ó  — | j                   j                  |«      }| j                  |«      }|› d| j                  |«      › d�}| j                  j                  | j                  |t        j                  j                  |«      ¬«      S )z"Codegen a load from an InputBufferú[ú]©rw   )
r>   Úinputr…   r„   r†   r‡   Úloadsr   ÚgraphÚ	get_dtype)rA   ÚnameÚindexrŠ   Úlines        r3   ÚloadzMetalKernel.load“  sq   € à�i‰i�o‰o˜dÓ#ˆØ×%Ñ% eÓ,ˆØ��a˜×)Ñ)¨%Ó0Ð1°Ð3ˆØ�x‰x× Ñ  §¡¨T¼¿¹×9JÑ9JÈ4Ó9PÐ ÓQÐQr5   Nc                ó   — | j                   j                  |«      }| j                  |«      }| j                  t        j
                  j                  |«      «      }|› d| j                  |«      › d|› d|› d�}| j                  r&| j                  j                  t        ||«      «       y | j                  j                  t        ||«      «       y )Nr*  z] = static_cast<rt   ú);)r>   Úoutputr…   r(  r   r/  r0  r„   Úinside_reductionrˆ   Ú	writeliner   Ústores)rA   r1  r2  ÚvalueÚmoderŠ   Ú	dtype_strr3  s           r3   ÚstorezMetalKernel.storeš  s®   € ð �i‰i×Ñ˜tÓ$ˆØ×%Ñ% eÓ,ˆØ×%Ñ%¤a§g¡g×&7Ñ&7¸Ó&=Ó>ˆ	Ø��a˜×)Ñ)¨%Ó0Ð1Ð1AÀ)ÀÈBÈuÈgÐUWÐXˆØ× Ò Ø�L‰L×"Ñ"¤<°°dÓ#;Õ<à�K‰K×!Ñ!¤,¨t°TÓ":Õ;r5   c           	     óJ  — dt        | j                  «      › �}t        j                  j	                  |||«      }|r6| j
                  j                  d| j                  |«      › d|› d|› d�«       |S | j
                  j                  d| j                  |«      › d|› d�«       |S )NÚtmp_acc_zthreadgroup ú r*  z];r  )Únextr"  r   rƒ   Úcreate_cse_varÚindexing_coder9  r(  )rA   rw   Ú
elem_countr‚   Úvar_namerŠ   s         r3   Ú_new_accvarzMetalKernel._new_accvar¦  s²   € ð œd 4×#3Ñ#3Ó4Ð5Ð6ˆÜ�h‰h×%Ñ% h°¸Ó>ˆÙØ×Ñ×(Ñ(Ø˜t×0Ñ0°Ó7Ð8¸¸(¸À1ÀZÀLÐPRÐSôð ˆ
ð ×Ñ×(Ñ(Ø˜t×0Ñ0°Ó7Ð8¸¸(¸À1ÐEôð ˆ
r5   c                ób	  — t        d„ | j                  D «       «      }t        |j                  | j                  «      }|dk(  r‰| j                  |«      }| j                  j                  |› d�«       | j                  j                  d«       | j                  j                  d|› d|› d�«       | j                  j                  d«       |S |dv rñ| j                  ||«      }| j                  rl|d	k(  rd
nd\  }	}
| j                  j                  |› d|j                  › d|	› d�«       | j                  j                  |› d|j                  › d|
› d|› d�«       n.| j                  j                  |› d|j                  › d|› d�«       | j                  j                  | j                  d|› d|› d|› d�t        |   ¬«      S |dv �r$| j                  ||«      }|› d|j                  › d�}t         |   }| j                  sV| j                  j                  |› d|› d|› d�«       | j                  j                  | j                  d|› d|› d|› d�|¬«      S |j#                  d«      rdnd}| j                  j                  |› d|› d|› d�«       |j%                  d «      ròt        d!„ | j&                  j)                  «       D «       «      }| j                  t*        j,                  |«      }|d"k(  rd#nd$}|› d|j                  › d�}| j                  j                  |› d%�«       | j                  j                  d|› d&|› d&|› d|› d'|› d(|› d'|j                  › d)�«       | j                  j                  | j                  |› d*|› d|› d|› d+�|¬«      S | j                  j                  |› d,|› d|› d|› d�«       | j                  j                  | j                  d|› d|› d|› d�|¬«      S |d-k(  r¸| j                  r
J d.|› �«       ‚| j                  ||«      }| j                  j                  |› d|j                  › d|› d�«       | j                  j                  | j                  d|› d|› d|› d�«      }t/        j0                  |› d/�|› d0�| j2                  j4                  f«      S t7        |«      ‚)1zCodegen a reduction operationc              3  ó:   K  — | ]  }|j                   sŒ|–— Œ y ­wr~   ©Úis_reduction©Ú.0Úts     r3   ú	<genexpr>z(MetalKernel.reduction.<locals>.<genexpr>À  s   è ø€ ÒK 1¸A¿N»NœQÑKùó   ‚”Úanyz	 = false;z7threadgroup_barrier(metal::mem_flags::mem_threadgroup);z
                if (z) {
                    z' = true;
                }
            )ÚprodÚsumrS  )r   r¤   )r   Ú*r*  z] = r  z] z= zc10::metal::threadgroup_r9   rJ   r;   r,  )ÚmaxÚminÚargminÚargmaxr+  z = static_cast<rt   r6  rU  Úlowestz = ::metal::numeric_limits<z>::z();Úargc              3  ó:   K  — | ]  }|j                   sŒ|–— Œ y ­wr~   rJ  rL  s     r3   rO  z(MetalKernel.reduction.<locals>.<genexpr>ø  s   è ø€ ò ØÀÇÃ”AñùrP  rX  ú>ú<z = -1;rA  ú = z;
                    z$;
                }
                z[c10::metal::threadgroup_z)]z = ::c10::metal::Úwelford_reducez+Multistage reduction not yet supported for z.xz.y)rB  Úrange_treesrV  ÚnumelÚmax_threadgroup_sizerG  rD  r9  rˆ   Úsplicer:  r#  r1  r†   r‡   r   rv   ÚendswithÚ
startswithÚrange_tree_nodesÚvaluesr/   r$   r   Ú_unwrapÚfeaturesÚreduction_numelÚNotImplementedError)rA   rw   rx   Úreduction_typer;  Úreduction_dimÚacc_buf_sizeÚaccÚacc_bufÚdefault_valÚreduction_opÚacc_thread_varÚsrc_metal_typeÚlim_fnÚidx_varÚidx_acc_bufÚcmp_opÚidx_thread_varÚwf_ress                      r3   Ú	reductionzMetalKernel.reduction¸  sm  € ô ÑK¨×(8Ñ(8ÔKÓKˆÜ˜=×.Ñ.°×0IÑ0IÓJˆØ˜UÒ"Ø×"Ñ" 5Ó)ˆCØ×Ñ×(Ñ(¨C¨5°	Ð):Ô;Ø×Ñ×(Ñ(ØIôð �L‰L×ÑðØ�Gð Ø�Eð ðôð �K‰K×!Ñ!ØIôð ˆJØ˜_Ñ,Ø×&Ñ& y°,Ó?ˆGØ×(Ò(à .°%Ò 7‘H¸Xñ *�˜\ð ×"Ñ"×,Ñ,Ø�i˜q ×!3Ñ!3Ð 4°D¸¸ÀQÐGôð —‘×#Ñ#Ø�i˜q ×!3Ñ!3Ð 4°B°|°nÀBÀuÀgÈQÐOõð —‘×#Ñ# w i¨q°×1CÑ1CÐ0DÀDÈÈÈqÐ$QÔRØ—8‘8×$Ñ$Ø—‘Ø*¨>Ð*:¸!¸G¸9ÀBÀ|ÀnÐTUÐVÜ0°Ñ7ð %ó ð ð
 Ð?Ò?Ø×&Ñ& y°,Ó?ˆGØ '˜y¨¨-×*<Ñ*<Ð)=¸QÐ?ˆNÜ+¨IÑ6ˆNØ×,Ò,Ø—‘×#Ñ#Ø%Ð& o°nÐ5EÀRÈÀwÈbÐQôð —x‘x×(Ñ(Ø—K‘KØ.¨~Ð.>¸aÀ¸yÈÈ<È.ÐXYÐZØð )ó ð ð
 "0×!8Ñ!8¸Ô!?‘XÀUˆFØ×Ñ×(Ñ(Ø!Ð"Ð"=¸nÐ=MÈSÐQWÐPXÐX[Ð\ôð ×(Ñ(¨Ô/Üñ Ø#×4Ñ4×;Ñ;Ó=ôó �ð #×.Ñ.¬u¯z©z¸<ÓH�Ø .°(Ò :™À�Ø$/ =°°-×2DÑ2DÐ1EÀQÐ!G�Ø×"Ñ"×)Ñ)¨^Ð,<¸FÐ*CÔDØ—‘×#Ñ#ð )Ø�G˜1˜V˜H A nÐ%5ð 6Ø#Ð$ C¨ wð /Ø#Ð$ C¨¯© ~ð 6ð%ô ð —x‘x×(Ñ(Ø—K‘KØ"�mÐ#<¸^Ð<LÈAÈgÈYÐVXÐYeÐXfÐfhÐiØð )ó ð ð
 �L‰L×"Ñ"Ø!Ð"Ð"3°NÐ3CÀ1À^ÐDTÐTVÐW\ÐV]Ð]_Ð`ôð —8‘8×$Ñ$Ø—‘Ø*¨>Ð*:¸!¸G¸9ÀBÀ|ÀnÐTUÐVØð %ó ð ð
 Ð-Ò-Ø×0Ò0ð Ø=¸nÐ=MÐNóÐ0ð ×&Ñ& y°,Ó?ˆGØ�L‰L×Ñ 7 )¨1¨]×-?Ñ-?Ð,@ÀÀUÀGÈ1Ð MÔNØ—X‘X×&Ñ&Ø—‘Ø*¨>Ð*:¸!¸G¸9ÀBÀ|ÀnÐTUÐVóˆFô ×%Ñ%Ø�8˜2� 6 (¨" ¨t¯}©}×/LÑ/LÐMóð ô " .Ó1Ð1r5   c                ó  — | j                  |j                  «      }| j                  |«      }|j                  r(|j                  j
                  | j                  kD  | _        |j                  r| j                  s9| j                  j                  | j                  › d|j                  › d|› d�«       y |j                  j
                  | j                  z   dz
  | j                  z  }| j                  j                  d|j                  › d|j                  › d|› d|j                  › d	�	«       | j                  j                  «       5  | j                  j                  | j                  › d|j                  › d|› d
|› d|j                  › d�
«       || j                  z  |j                  j
                  k7  r@| j                  j                  d|j                  › d|j                  j
                  › d�«       d d d «       y # 1 sw Y   y xY w)NrA  r^  r  r   z	for(auto z
_cnt = 0; z_cnt < z; ++z_cnt) {r_   z + z_cnt;úif (ú >= z) break;)Úrename_indexingrB   ÚsexprrK  Úrootra  rb  r#  rD  r9  Úindex_dtyper1  r‘   Úindent)rA   Úentryr‹   Ú	index_strÚ	loop_sizes        r3   Úcodegen_iteration_ranges_entryz*MetalKernel.codegen_iteration_ranges_entry!  sÈ  € Ø×)Ñ)¨%¯*©*Ó5ˆ
Ø—J‘J˜zÓ*ˆ	Ø×ÒØ(-¯
©
×(8Ñ(8¸4×;TÑ;TÑ(TˆDÔ%Ø×!Ò!¨×)BÒ)BØ×Ñ×(Ñ(Ø×#Ñ#Ð$ A e§j¡j \°°Y°K¸qÐAôð ð
 �J‰J×Ñ˜t×8Ñ8Ñ8¸1Ñ<Ø×&Ñ&ñ'ˆ	ð 	�	‰	×ÑØ˜Ÿ
™
�| :¨e¯j©j¨\¸ÀÀÈ4ÐPU×PZÑPZÈ|Ð[cÐdô	
ð �Y‰Y×ÑÓñ 	WØ�I‰I×ÑØ×#Ñ#Ð$ A e§j¡j \°°Y°K¸sÀ9À+ÈSÐQV×Q[ÑQ[ÐP\Ð\aÐbôð ˜4×4Ñ4Ñ4¸¿
¹
×8HÑ8HÒHØ—	‘	×#Ñ# d¨5¯:©:¨,°d¸5¿:¹:×;KÑ;KÐ:LÈHÐ$UÔV÷	W÷ 	Wñ 	Wús   ÅB/G=Ç=Hc                óÎ  — | j                   r�| j                  j                  «       5  | j                  j                  | j                  «       | j                  j                  | j
                  «       ddd«       | j                  j                  d«       d| _         nJ| j                  j                  | j                  «       | j                  j                  | j
                  «       | j                  j                  | j                  «       | j                  j                  «        | j
                  j                  «        | j                  j                  «        y# 1 sw Y   ŒêxY w)a  
        Concat output code from index_code, loads, compute, stores,
        suffix into self.body.

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

        For reduction kernels, this generates a loop over the reduction
        axis.
        Nú}F)	r#  r‘   rƒ  rc  r.  rˆ   r9  r:  Úclear)rA   s    r3   Úcodegen_bodyzMetalKernel.codegen_body<  sì   € ð ×$Ò$Ø—‘×!Ñ!Ó#ñ /Ø—	‘	× Ñ  §¡Ô,Ø—	‘	× Ñ  §¡Ô.÷/ð �I‰I×Ñ Ô$Ø(-ˆDÕ%à�I‰I×Ñ˜TŸZ™ZÔ(Ø�I‰I×Ñ˜TŸ\™\Ô*Ø�	‰	×Ñ˜Ÿ™Ô%Ø�
‰
×ÑÔØ�‰×ÑÔØ�‰×ÑÕ÷/ð /ús   §AEÅE$c                ó  — | j                  «        t        «       }|j                  d«       | j                  «       }|j	                  «       5  |j                  dd¬«       | j                  r|j                  d«       |j                  d«       |j	                  «       5  | j                  j                  j                  «       D ]Z  \  }}|| j                  v rŒ| j                  t        j                  j                  |«      «      }|j                  d|› d|› d	�«       Œ\ | j                  j                  j                  «       D ]K  \  }}| j                  t        j                  j                  |«      «      }|j                  d
|› d|› d	�«       ŒM | j                  j                   j                  «       D ]  \  }}|j                  d|› d	�«       Œ t#        |«      dk  sJ d«       ‚t#        |«      dkD  rdt#        |«      › �nd}t#        |«      dk(  r|d   j$                  nd}| j                  rd	nd}	|j                  |› d|› d|	› �«       | j                  r|j                  |› d�«       ddd«       |j                  d«       |j	                  «       5  t#        |«      dkD  rAt'        |«      D ]3  \  }
}|j                  d|j$                  › dt)        d|
z   «      › d�«       Œ5 |j                  | j*                  «       |j                  | j,                  «       ddd«       |j                  d«       ddd«       |j                  d«       |j/                  «       S # 1 sw Y   ŒòxY w# 1 sw Y   ŒOxY w# 1 sw Y   ŒBxY w)z3Called at the end to generate a final kernel stringzcompile_mps_shader("""z�
            #include <c10/metal/random.h>
            #include <c10/metal/special_math.h>
            #include <c10/metal/utils.h>
            T)Ústripz&#include <c10/metal/reduction_utils.h>zkernel void generated_kernel(zdevice z* ú,z	constant zconstant long& é   z%Up to 3 index variables are supportedr   Úuintr   Ú
thread_posÚ rA  z [[thread_position_in_grid]]z- group_pos [[thread_position_in_threadgroup]]Nz) {r  z = thread_pos.éx   r  r‰  z"""))r‹  r   r9  Úactive_range_treesrƒ  rc  r8  r>   Úoutput_buffersÚitemsÚremoved_buffersr(  r   r/  r0  Úinput_buffersÚsizevarsrL   r1  Ú	enumerateÚchrrD  r‘   Úgetvalue)rA   r1  ÚcodeÚidx_varsÚouterÚinnerr=  Úthread_pos_dtypeÚthread_pos_var_nameÚthread_pos_suffixÚidxrŠ   s               r3   Úcodegen_kernelzMetalKernel.codegen_kernelT  s@  € à×ÑÔÜÓˆØ�‰Ð/Ô0Ø×*Ñ*Ó,ˆØ�[‰[‹]ñ /	 Ø�K‰Kðð
 ð ô ð ×$Ò$Ø—‘ÐGÔHØ�N‰NÐ:Ô;Ø—‘“ñ Ø$(§I¡I×$<Ñ$<×$BÑ$BÓ$Dò D‘L�E˜5Ø × 4Ñ 4Ñ4Ø Ø $× 1Ñ 1´!·'±'×2CÑ2CÀEÓ2JÓ K�IØ—N‘N W¨Y¨K°r¸%¸ÀÐ#BÕCð	Dð
 %)§I¡I×$;Ñ$;×$AÑ$AÓ$Cò F‘L�E˜5Ø $× 1Ñ 1´!·'±'×2CÑ2CÀEÓ2JÓ K�IØ—N‘N Y¨y¨k¸¸E¸7À!Ð#DÕEðFð %)§I¡I×$6Ñ$6×$<Ñ$<Ó$>ò ?‘L�E˜5Ø—N‘N _°U°G¸1Ð#=Õ>ð?ä˜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ô÷/ð4 �N‰N˜5Ô!Ø—‘“ñ 'Ü�x“= 1Ò$Ü$-¨hÓ$7ò ™˜˜SØŸ™Ø# C§H¡H :¨^¼CÀÀcÁ	»NÐ;KÈ1ÐMõðð —‘˜D×.Ñ.Ô/Ø—‘˜DŸI™IÔ&÷'ð �N‰N˜3Ô÷_/	 ð` 	�‰�vÔà�}‰}‹Ð÷Mð ú÷6'ð 'ú÷O/	 ð /	 ús?   ÁAM7ÂGMÉ.)M7ÊBM+ÌM7ÍM(	Í$M7Í+M4	Í0M7Í7N c           	     óL  — t         j                  j                  }g | j                  j                  j                  «       ¢| j                  j                  j                  «       ¢}|D �cg c]  }|| j                  vsŒ|‘Œ }}|| j                  j                  j                  «       D �cg c]  }t        |«      ‘Œ c}z  }t        | j                  «       «      dkD  r…| j                  «       D �cg c]S  }| j                  |j                  r*t        j                  |j                   | j"                  «      n|j                   «      ‘ŒU }}|ddj%                  |«      › d�gz  }| j&                  r{| j                  «       D �cg c]I  }|j                  r9| j                  t        j                  |j                   | j"                  «      «      nd‘ŒK }}|ddj%                  |«      › d�gz  }|j)                  ||t+        j,                  d«      d¬	«       y
c c}w c c}w c c}w c c}w )zCodegen a call to this kernelr   z	threads=[rJ   r+  Ú1zgroup_size=[ÚcpuF)ÚdeviceÚtritonN)r   r/  Úwrapper_coder>   r•  Úkeysr˜  r—  r™  r1   rL   r”  ÚpexprrK  ÚsympyÚMinra  rb  rN   r8  Úgenerate_kernel_callr/   r©  )rA   r1  ÚnodeÚwrapperr>   rZ  ÚvÚthreadss           r3   Úcall_kernelzMetalKernel.call_kernelŽ  sæ  € ä—'‘'×&Ñ&ˆØR�—‘×)Ñ)×.Ñ.Ó0ÐR°4·9±9×3JÑ3J×3OÑ3OÓ3QÐRˆØ#ÖG˜ s°$×2FÑ2FÒ'F’ÐGˆÐGØ §¡×!3Ñ!3×!8Ñ!8Ó!:Ö;˜A”�Q•Ò;Ñ;ˆô ˆt×&Ñ&Ó(Ó)¨AÒ-ð ×0Ñ0Ó2öð ð —
‘
à—~’~ô —I‘I˜aŸg™g t×'@Ñ'@ÔAàŸ™õðˆGð ð �y §¡¨7Ó!3Ð 4°AÐ6Ð7Ñ7ˆDØ× Ò ð
 ×0Ñ0Ó2ö	ð ð —>’>ð —
‘
œ5Ÿ9™9 Q§W¡W¨d×.GÑ.GÓHÔIàñðˆGð ð �| D§I¡I¨gÓ$6Ð#7°qÐ9Ð:Ñ:ˆDà×$Ñ$ØØÜ—<‘< Ó&Øð	 	%õ 	
ùò/ HùÚ;ùòùòs    Á)HÁ=HÂ*HÃ/AHÆ AH!c                óò   — |s|sy | j                  |«      }|r|› d�nd}|r|› d| j                  |«      › �nd}|r|r
d|› d|› d�}nd|› |› d�}| j                  j                  | j                  |d	¬
«       y )Nz < 0r’  r~  zif ((z) && (z	)) returnr}  z) returnF)Ú
assignment)r„   r†   r‡   rˆ   )	rA   rB   ÚsizeÚlowerÚupperÚexpr_strÚ
lower_exprÚ
upper_exprr3  s	            r3   Úcheck_boundszMetalKernel.check_bounds°  s–   € ñ ™Øð ×$Ñ$ TÓ*ˆÙ*/˜�z Ñ&°Rˆ
ÙCH˜�z  d×&7Ñ&7¸Ó&=Ð%>Ñ?Èbˆ
Ù‘UØ˜:˜, f¨Z¨L¸	ÐB‰Dà˜*˜ j \°Ð:ˆDØ�‰×Ñ˜$Ÿ,™,¨¸ÐÕ?r5   )r$  zdict[str, sympy.Expr]r%  r   rl   ÚNone)rw   r  rl   r1   )r1  r1   r2  rk   rl   r   r~   )
r1  r1   r2  rk   r;  r   r<  r   rl   r¿  )rw   r  rE  zOptional[int]r‚   zValueRanges[Any]rl   r   )
rw   r  rx   r  rl  r   r;  ú+Union[CSEVariable, tuple[CSEVariable, ...]]rl   rÀ  )r„  r   rl   r¿  ©rl   r¿  )r1  zOptional[str]rl   r1   )r1  r1   r±  r   rl   r¿  )
rB   rk   r¸  rk   r¹  r    rº  r    rl   r¿  )rm   rn   ro   rr   Ú	overridesÚsuffixÚnewvar_prefixrb  r   r?   r­  r7   r€  Úkexprr  r(  r4  r>  r	   ÚunknownrG  r{  r‡  r‹  r¥  rµ  r¾  Ú__classcell__©r&  s   @r3   r  r  ~  s\  ø„ Ø€IØ€FØ€MØÐÙ‹O×#Ñ#€EÙÓ×&Ñ&€EØ€Eð*à%ð*ð ð*ð 
õ	*ó%óRð SWð
<Øð
<Ø *ð
<Ø3>ð
<ØFOð
<à	ó
<ð %)Ø#6 ;×#6Ñ#6Ó#8ð	àðð "ðð !ð	ð
 
óð$g2àðg2ð ðg2ð &ð	g2ð
 ;ðg2ð 
5óg2óRWó6ô08ôt 
ðD@Øð@Ø&0ð@Ø9=ð@ØFJð@à	÷@r5   r  c                 ó2   — dd l } | j                  dd¬«       y )Nr   z�torch.compile for Metal is an early protoype and might not work as expected. For details see https://github.com/pytorch/pytorch/issues/150121r
   )Ú
stacklevel)ÚwarningsÚwarn)rË  s    r3   Ú_warn_prototyperÍ  Á  s   € ãà‡M�Mð	Làð õ r5   c                  ó<   ‡ — e Zd ZeZdˆ fd„Z	 	 	 	 	 	 	 	 dd„Zˆ xZS )ÚMetalSchedulingc                ó¨   •— t         ‰| �  |«       t        «        t        j                  j
                  }|�|j                  j                  d«       y y )NzDfrom torch._inductor.runtime.runtime_utils import compile_mps_shader)r  r  rÍ  r   r/  r«  Úheaderrc  )rA   Ú	schedulerr²  r&  s      €r3   r  zMetalScheduling.__init__Ï  sE   ø€ Ü‰Ñ˜Ô#ÜÔÜ—'‘'×&Ñ&ˆØÐØ�N‰N×!Ñ!ØVõð r5   c                ó  — t         j                  j                  }||j                  v r|j                  |   }|S d|j	                  «       › �}|› d�}||j                  |<   t        ||«      \  }}|› d|› �}	|j                  |||	«       |S )NÚmps_lib_z.generated_kernelú
)r   r/  r«  Úsrc_to_kernelÚnext_kernel_suffixr   Údefine_kernel)
rA   Úsrc_codeÚnode_schedulerƒ   r²  Úkernel_nameÚmps_lib_nameÚoriginsÚdetailed_originsÚmetadata_comments
             r3   rØ  zMetalScheduling.define_kernelØ  s¬   € ô —'‘'×&Ñ&ˆØ�w×,Ñ,Ñ,Ø!×/Ñ/°Ñ9ˆKð Ðð & g×&@Ñ&@Ó&BÐ%CÐDˆLØ)˜NÐ*;Ð<ˆKØ.9ˆG×!Ñ! (Ñ+Ü(;¸MÈ7Ó(SÑ%ˆGÐ%Ø") ¨"Ð-=Ð,>Ð?ÐØ×!Ñ! ,°Ð:JÔKàÐr5   )rÒ  zOptional[Scheduler]rl   r¿  )rÙ  r1   rÚ  zlist[SchedulerNode]rƒ   r  rl   r1   )rm   rn   ro   r  Úkernel_typer  rØ  rÇ  rÈ  s   @r3   rÏ  rÏ  Ì  s2   ø„ Ø€KõðØðØ,?ðØITðà	÷r5   rÏ  )r2   z)Union[float, int, bool, str, CSEVariable]rl   r1   rÁ  )=Ú
__future__r   Ú	functoolsr   Útypingr   r   r   r®  Úsympy.printing.precedencer   r/   Útorch.utils._sympy.printersr   ÚExprPrinter_Útorch.utils._sympy.value_rangesr	   Úutilsr   r   Úvirtualizedr   r   r   Úcommonr   r   r   r   r   r   Úsimdr   r   r   r   Úops_handlerr   r   rÒ  r   r   r   r    Úint8Úint16Úint32Úint64Úuint8r&   r'   Úbfloat16rv   r4   r7   rr   Ú_initialize_pointwise_overridesr  ÚcacherÍ  rÏ  rp   r5   r3   ú<module>rõ     s  ðõ #ã Û ß /Ñ /ã Ý 0ã Ý CÝ 7ç >ß ,Ñ ,÷÷ ÷ CÑ Bñ Ýç6ß4Ýð 
‡J�J�Ø	‡J�J�Ø	‡K�K�Ø	‡K�K�Ø	‡K�K�Ø	‡K�K�Ø	‡K�K�Ø	‡J�J�Ø	‡N�N�Hð
€óô6b�|ô 6bôr(�[ô (ðD × .Ñ .¨uÔ 5ô@@�*ô @@ðF
 ‡�òó ðô�nõ r5   