Ë
    7^(hÃç  ã                   óä  — d dl Z d dlZd dlZd dlmZmZ d dlZd dlZd dl	Z	d dl
mZ d dlmZ ddlmZ d dlmZ ddlmZ dd	lmZ  G d
„ d«      Z G d„ d«      Z G d„ d«      Z ed¬«       G d„ d«      «       Zd„ Zd„ Zd„ Zd„ Zd„ Z  ejB                  eejD                  g¬«      Z# ejB                  eejH                  g¬«      Z% ejB                  e ejL                  g¬«      Z' G d„ d«      Z( G d„ d«      Z)d„ Z*d„ Z+d „ Z, G d!„ d"«      Z- G d#„ d$e-«      Z. G d%„ d&e-«      Z/d'„ Z0d(„ Z1d)„ Z2d*„ Z3d+„ Z4 e)«       Z5d,„ Z6d-„ Z7 G d.„ d/«      Z8 G d0„ d1e jr                  «      Z: G d2„ d3«      Z; G d4„ d5«      Z<y)6é    N)ÚTupleÚList)Ú	dataclassé   )ÚInterpreterError)Úpartialé   )Úinterpreter)Úirc                   ó*   — e Zd Zd„ Zd„ Zd„ Zd„ Zd„ Zy)ÚTensorHandlec                 ó.   — || _         || _        i | _        y)a  
            data: numpy array
            dtype: triton type, either pointer_type or scalar_type.
            we don't store block_type here because the shape information is already available in the data field
            attr: a dictionary of attributes
        N)ÚdataÚdtypeÚattr)Úselfr   r   s      úX/var/www/skyplay_api_hub/venv/lib/python3.12/site-packages/triton/runtime/interpreter.pyÚ__init__zTensorHandle.__init__   s   € ð ˆŒ	ØˆŒ
Øˆ�	ó    c                 óH   — t        | j                  j                  «       «      S ©N)Úboolr   Úall©r   s    r   Ú__bool__zTensorHandle.__bool__   s   € Ü�D—I‘I—M‘M“OÓ$Ð$r   c                 óh   — | j                   }t        |d«      r|j                  }t        |d«      rŒ|S )NÚ
element_ty)r   Úhasattrr   )r   r   s     r   Úget_element_tyzTensorHandle.get_element_ty"   s1   € Ø—
‘
ˆÜ�e˜\Ô*Ø×$Ñ$ˆEô �e˜\Õ*àˆr   c                 ó^   — t        | j                  j                  «       | j                  «      S r   )r   r   Úcopyr   r   s    r   ÚclonezTensorHandle.clone(   s   € Ü˜DŸI™IŸN™NÓ,¨d¯j©jÓ9Ð9r   c                 ó"   — || j                   |<   y r   )r   )r   ÚkeyÚvalues      r   Úset_attrzTensorHandle.set_attr+   s   € Øˆ�	‰	�#Šr   N)Ú__name__Ú
__module__Ú__qualname__r   r   r   r"   r&   © r   r   r   r      s   „ ò	ò%òò:ór   r   c                   ó   — e Zd Zd„ Zd„ Zy)ÚBlockPointerHandlec                 óX   — || _         || _        || _        || _        || _        || _        y r   )ÚbaseÚshapeÚstridesÚoffsetsÚblock_shapeÚorder)r   r.   r/   r0   r1   r2   r3   s          r   r   zBlockPointerHandle.__init__1   s-   € ØˆŒ	ØˆŒ
ØˆŒØˆŒØ&ˆÔØˆ�
r   c                 óf  — | j                   j                  «       }|j                  dz  }t        j                  | j                   j
                  | j                  «      }t        j                  | j                  t        ¬«      }t        t        | j                  «      «      D ]á  }dgt        | j                  «      z  }| j                  |   ||<   | j                  |   j
                  t        j                  | j                  |   «      z   j                  |«      }|||z  | j                  |   j
                  z  j                  t        j                   «      z   }||v sŒ½||| j"                  |   j
                  k  z  |dk\  z  }Œã t%        || j                   j&                  j(                  «      }||fS )Né   ©r   r   r   )r.   r   Úprimitive_bitwidthÚnpÚbroadcast_tor   r2   Úonesr   ÚrangeÚlenr1   ÚarangeÚreshaper0   ÚastypeÚuint64r/   r   r   Úscalar)	r   Úboundary_checkÚdtype_ttÚn_bytesÚptrsÚmasksÚdimÚ
bcast_dimsÚoffs	            r   Úmaterialize_pointersz'BlockPointerHandle.materialize_pointers9   sb  € Ø—9‘9×+Ñ+Ó-ˆØ×-Ñ-°Ñ2ˆÜ�‰˜tŸy™yŸ~™~¨t×/?Ñ/?Ó@ˆÜ—‘˜×(Ñ(´Ô5ˆÜœ˜T×-Ñ-Ó.Ó/ò 	JˆCØ˜œs 4×#3Ñ#3Ó4Ñ4ˆJØ"×.Ñ.¨sÑ3ˆJ�s‰OØ—<‘< Ñ$×)Ñ)¬B¯I©I°d×6FÑ6FÀsÑ6KÓ,LÑL×UÑUÐV`ÓaˆCØ˜7 S™=¨4¯<©<¸Ñ+<×+AÑ+AÑA×IÑIÌ"Ï)É)ÓTÑTˆDØ�nÒ$Ø  t§z¡z°#¡×';Ñ';Ñ!;Ñ<ÀÀqÁÑI‘ð	Jô ˜D $§)¡)§/¡/×"8Ñ"8Ó9ˆØ�Uˆ{Ðr   N)r'   r(   r)   r   rJ   r*   r   r   r,   r,   /   s   „ òór   r,   c            	       óN   — e Zd Zdedee   dee   dee   fd„Zd„ Zdee   fd„Zy	)
ÚTensorDescHandler.   r/   r0   r2   c                 ó\   — || _         t        |«      | _        || _        || _        || _        y r   )r.   r<   Úndimr/   r0   r2   )r   r.   r/   r0   r2   s        r   r   zTensorDescHandle.__init__K   s*   € àˆŒ	Ü˜“JˆŒ	ØˆŒ
ØˆŒØ&ˆÕr   c                 óÎ  — | j                   j                  j                  «       dz  dk(  sJ d«       ‚t        | j                  «      | j
                  k(  sJ ‚t        | j                  «      | j
                  k(  sJ ‚| j                  d d D ])  }|j                  j                  «       dz  dk(  rŒ$J d«       ‚ | j                  d   j                  j                  «       dk(  sJ d«       ‚y )Né   r   zbase must be 16-byte alignedéÿÿÿÿzstride must be 16-byte alignedr   zlast dim must be contiguous)r.   r   Úitemr<   r0   rN   r2   )r   Ústrides     r   ÚvalidatezTensorDescHandle.validateS   sÒ   € Ø�y‰y�~‰~×"Ñ"Ó$ rÑ)¨QÒ.ÐNÐ0NÓNÐ.Ü�4—<‘<Ó  D§I¡IÒ-Ð-Ð-Ü�4×#Ñ#Ó$¨¯	©	Ò1Ð1Ð1à—l‘l 3 BÐ'ò 	RˆFØ—;‘;×#Ñ#Ó%¨Ñ*¨aÓ/ÐQÐ1QÓQÐ/ð	Rà�|‰|˜BÑ×$Ñ$×)Ñ)Ó+¨qÒ0ÐOÐ2OÓOÑ0r   r1   c                 óÆ  — t        |«      | j                  k(  sJ ‚| j                  j                  j                  }|j
                  dz  }|d   j                  |z  dz  dk(  sJ d«       ‚t        j                  | j                  j                  | j                  «      }t        j                  | j                  t        ¬«      }t        t        | j                  «      «      D ]Ò  }dgt        | j                  «      z  }| j                  |   ||<   ||   j                  t        j                  | j                  |   «      z   j                  |«      }|||z  | j                  |   j                  z  j!                  t        j"                  «      z   }|d|k  z  || j$                  |   j                  k  z  }ŒÔ t'        || j                  j                  j(                  «      }||fS )Nr5   rQ   rP   r   z*block offset start must be 16-byte alignedr6   r   )r<   rN   r.   r   r   r7   r   r8   r9   r2   r:   r   r;   r=   r>   r0   r?   r@   r/   r   rA   )	r   r1   Ú	scalar_tyÚitemsizerE   rF   rG   rH   rI   s	            r   rJ   z%TensorDescHandle.materialize_pointers\   s”  € Ü�7‹|˜tŸy™yÒ(Ð(Ð(Ø—I‘I—O‘O×.Ñ.ˆ	Ø×/Ñ/°1Ñ4ˆØ˜‘× Ñ  8Ñ+¨rÑ1°QÒ6ÐdÐ8dÓdÐ6ä�‰˜tŸy™yŸ~™~¨t×/?Ñ/?Ó@ˆÜ—‘˜×(Ñ(´Ô5ˆÜœ˜T×-Ñ-Ó.Ó/ò 	FˆCØ˜œs 4×#3Ñ#3Ó4Ñ4ˆJØ"×.Ñ.¨sÑ3ˆJ�s‰OØ˜3‘<×$Ñ$¤r§y¡y°×1AÑ1AÀ#Ñ1FÓ'GÑG×PÑPÐQ[Ó\ˆCØ˜8 c™>¨D¯L©L¸Ñ,=×,BÑ,BÑB×JÑJÌ2Ï9É9ÓUÑUˆDØ˜Q #™XÑ&¨#°·
±
¸3±×0DÑ0DÑ*DÑE‰Eð	Fô ˜D $§)¡)§/¡/×"8Ñ"8Ó9ˆØ�Uˆ{Ðr   N)	r'   r(   r)   r   r   Úintr   rT   rJ   r*   r   r   rL   rL   I   sJ   „ ð'˜\ð '°$°|Ñ2Dð 'ÈtÐT`ÑOað 'Ø" 3™ió'òPð¨D°Ñ,>ô r   rL   T)Úfrozenc                   ó¬   — e Zd ZU dZeed<   dZeed<   dZeed<   dZ	e
ed<   dZee
   ed	<   d
Zee
   ed<   dZe
ed<   dZee
   ed<   dZeed<   dZe
ed<   y)ÚInterpreterOptionsNÚextern_libsFÚdebugTÚsanitize_overflowÚarch)Úfp8e5Úfp8e5b16Úfp8e4nvÚfp8e4b8Úfp8e4b15Úsupported_fp8_dtypesr*   Údeprecated_fp8_dtypesÚtf32Údefault_dot_input_precision)rg   Útf32x3ÚieeeÚallowed_dot_input_precisionsr   Úmax_num_imprecise_acc_defaultr
   Úbackend_name)r'   r(   r)   r\   ÚdictÚ__annotations__r]   r   r^   r_   Ústrre   r   rf   rh   rk   rl   rX   rm   r*   r   r   r[   r[   n   sx   … à€K�ÓØ€Eˆ4ÓØ"Ð�tÓ"Ø€Dˆ#ÓØ'^Ð˜% ™*Ó^Ø(*Ð˜5 ™:Ó*Ø'-Ð Ó-Ø/IÐ  %¨¡*ÓIØ)*Ð! 3Ó*Ø%€L�#Ô%r   r[   c                 ó  — | t         j                  k(  rt         j                  S | t         j                  k(  rt         j                  S | t         j
                  k(  rt         j                  S | t         j                  k(  rt         j                  S | S r   )	r8   Úuint8Úint8Úuint16Úint16Úuint32Úint32r@   Úint64r6   s    r   Ú_get_signed_np_dtypery   |   s[   € Ø”—‘ÒÜ�w‰wˆØ”—	‘	ÒÜ�x‰xˆØ”—	‘	ÒÜ�x‰xˆØ”—	‘	ÒÜ�x‰xˆØ€Lr   c                 ó`  — t        | t        j                  «      r#t        j                  t        j
                  «      S i t        j                  t        j                  t        «      “t        j                  t        j                  t        j                  «      “t        j                  t        j                  t        j                  «      “t        j                  t        j                  t        j                  «      “t        j                  t        j                  t        j                  «      “t        j                  t        j                  t        j                  «      “t        j                  t        j                  t        j                  «      “t        j                  t        j                  t        j                  «      “t        j                  t        j                  t        j                  «      “t        j                   t        j                  t        j                   «      “t        j"                  t        j                  t        j"                  «      “t        j
                  t        j                  t        j
                  «      “t        j$                  t        j                  t        j                  «      “t        j&                  t        j                  t        j                  «      “t        j(                  t        j                  t        j                  «      “t        j*                  t        j                  t        j                  «      “t        j,                  t        j                  t        j                  «      “t        j.                  t        j                  t        j                  «      i¥}t        | t        j0                  «      rVt        | j2                  t        j                  «      r#t        j                  t        j
                  «      S || j2                     S ||    S r   )Ú
isinstanceÚtlÚpointer_typer8   r   r@   Úint1r   Úfloat16Úfloat32Úfloat64rs   rr   ru   rt   rw   rv   rx   Úbfloat16Úfloat8e5Úfloat8e5b16Ú
float8e4nvÚ
float8e4b8Úfloat8e4b15Ú
block_typer   )Útt_dtypeÚnp_typess     r   Ú_get_np_dtyper‹   ˆ   s~  € Ü�(œBŸO™OÔ,Ü�x‰xœŸ	™	Ó"Ð"ðÜ
�‰”—‘œ$“ðä
�
‰
”B—H‘HœRŸZ™ZÓ(ðô 	�
‰
”B—H‘HœRŸZ™ZÓ(ðô 	�
‰
”B—H‘HœRŸZ™ZÓ(ð	ô
 	�‰”—‘œ"Ÿ'™'Ó"ðô 	�‰”"—(‘(œ2Ÿ8™8Ó$ðô 	�‰”"—(‘(œ2Ÿ8™8Ó$ðô 	�	‰	”2—8‘8œBŸI™IÓ&ðô 	�‰”"—(‘(œ2Ÿ8™8Ó$ðô 	�	‰	”2—8‘8œBŸI™IÓ&ðô 	�‰”"—(‘(œ2Ÿ8™8Ó$ðô 	�	‰	”2—8‘8œBŸI™IÓ&ðô 	�‰”R—X‘XœbŸi™iÓ(ðô  	�‰”R—X‘XœbŸh™hÓ'ð!ô" 	�‰œŸ™¤§¡Ó*ð#ô$ 	�‰”r—x‘x¤§¡Ó)ð%ô& 	�‰”r—x‘x¤§¡Ó)ð'ô( 	�‰œŸ™¤§¡Ó*ñ)€Hô, �(œBŸM™MÔ*Ü�h×)Ñ)¬2¯?©?Ô;Ü—8‘8œBŸI™IÓ&Ð&Ø˜×+Ñ+Ñ,Ð,Ø�HÑÐr   c                 óÖ  — t        t        d|j                  › �«      }t        t        d|j                  › �«      }t        j                  | j	                  «       |¬«      }||j                  dz
  z	  dz  }|j                  |j
                  z
  dz
  }|j                  |j
                  z
  dz
  }	|d|j
                  z  dz
  z  }
|j                  }|j                  }||j
                  z	  d|z  dz
  z  j                  t        j                  «      }|dk(  }t        j                  |«      r˜t        j                  |t        j                  ¬«      }t        |j
                  «      D ]  }|
|z	  dz  }|j
                  |z
  ||dk(  <   Œ! |
dk(  }d||   z
  ||<   ||z
  |||z  <   |
|   ||   z  d|j
                  z  dz
  z  |
|<   t        j                  dt        j                  ||z
  |z   d|	z  dz
  «      «      }|j                  |«      }|j                  |«      }|j                  |j                  kD  rˆ|
|j
                  |j
                  z
  z	  d|j
                  z  dz
  z  }|t        j                  j                   k(  r*|
d|j
                  |j
                  z
  dz
  z  z  }||dkD  z   }|j                  |«      }n>|
j                  |«      |j
                  |j
                  z
  z  d|j
                  z  dz
  z  }|dk(  }t        j                  |«      r™||j
                  z	  d|z  dz
  z  j                  t        j                  «      }|dk7  }||z  }t        j                  |t        j                  ¬«      }d|z
  ||   |z
  z
  ||<   ||   ||   z	  d|j
                  ||   z
  z  z  ||<   ||j                  dz
  z  ||j
                  z  z  |z  }|j#                  | j$                  «      S )NÚuintr6   r   r   )Úgetattrr8   r7   Ú
frombufferÚtobytesÚfp_mantissa_widthÚexponent_biasr?   rw   ÚanyÚ
zeros_liker;   ÚmaximumÚminimumÚ_irÚROUNDING_MODEÚRTNEr>   r/   )ÚinputÚinput_dtypeÚoutput_dtypeÚrounding_modeÚinput_uint_dtypeÚoutput_unint_dtypeÚ	input_binÚsignÚinput_exponent_widthÚoutput_exponent_widthÚsignificandÚ
bias_inputÚbias_outputÚexponentÚsubnormal_indexÚbit_posÚiÚ	bit_indexÚzero_significand_indexÚexponent_outputÚsign_outputÚsignificand_outputÚcut_offÚnon_zero_exponent_indexÚshiftÚoutputs                             r   Ú_convert_floatr´   ¨   s+  € Üœr T¨+×*HÑ*HÐ)IÐ#JÓKÐÜ ¤ t¨L×,KÑ,KÐ+LÐ%MÓNÐÜ—‘˜eŸm™m›oÐ5EÔF€IØ˜+×8Ñ8¸1Ñ<Ñ=ÀÑE€DØ&×9Ñ9¸K×<YÑ<YÑYÐ\]Ñ]ÐØ(×;Ñ;¸l×>\Ñ>\Ñ\Ð_`Ñ`ÐØ  [×%BÑ%BÑ BÀaÑGÑH€KØ×*Ñ*€JØ×,Ñ,€KØ˜k×;Ñ;Ñ;ÀÐFZÑAZÐ^_Ñ@_Ñ`×hÑhÔik×iqÑiqÓr€HØ !‘m€OÜ	‡v�vˆoÔô
 —-‘- 	´·±Ô:ˆä�{×4Ñ4Ó5ò 	HˆAØ%¨Ñ*¨dÑ2ˆIà&1×&CÑ&CÀaÑ&GˆG�I ‘NÒ#ð	Hð "-°Ñ!1ÐØ$%¨°Ñ(@Ñ$@ˆ�Ñ!à=GÈ+Ñ=UˆÐ'¨/Ñ9Ñ:Ø(3°OÑ(DÈÐP_ÑH`Ñ(`Ø�+×/Ñ/Ñ/°1Ñ4ñ(6ˆ�OÑ$ô —j‘j ¤B§J¡J°¸:Ñ0EÈÑ0SÐWXÐ\qÑWqÐuvÑVvÓ$wÓx€OØ%×,Ñ,Ð-?Ó@€OØ—+‘+Ð0Ó1€KØ×%Ñ%¨×(GÑ(GÒGØ)¨k×.KÑ.KÈl×NlÑNlÑ.lÑmØ�,×0Ñ0Ñ0°AÑ5ñ7ÐàœC×-Ñ-×2Ñ2Ò2à! Q¨;×+HÑ+HÈ<×KiÑKiÑ+iÐlmÑ+mÑ%nÑoˆGØ!3°wÀ±{Ñ!CÐØ/×6Ñ6Ð7IÓJÑà)×0Ñ0Ð1CÓDØ+×=Ñ=À×@]Ñ@]Ñ]ñ_à#$¨×(FÑ(FÑ#FÈ!Ñ"KñMÐð &¨Ñ*€OÜ	‡v�vˆoÔð
  +×"?Ñ"?Ñ?ÀQÐJ^ÑE^ÐbcÑDcÑd×lÑlÔmo×muÑmuÓvˆØ"*¨a¡-Ðà)Ð,CÑCˆÜ—‘˜i¬r¯x©xÔ8ˆØ"# k¡/°h¸Ñ6OÐR\Ñ6\Ñ!]ˆˆoÑØ/AÀ/Ñ/RÐV[Ð\kÑVlÑ/lØ�,×0Ñ0°5¸Ñ3IÑIÑJñ/LÐ˜?Ñ+à˜l×=Ñ=ÀÑAÑBØ˜<×9Ñ9Ñ9ñ;Ø=OñP€Fà�>‰>˜%Ÿ+™+Ó&Ð&r   c                 ó,   — t        j                  | «      S r   )ÚmathÚerf©Úxs    r   Ú_erfrº   è   s   € ä�8‰8�A‹;Ðr   c                 ó6   — t        | «      t        |«      z  dz	  S )Né@   )rX   )ÚaÚbs     r   Ú
_umulhi_64r¿   í   s   € ô �‹F”S˜“V‰O Ñ"Ð"r   )Úotypesc                   ó   — e Zd Zed„ «       Zy)ÚExtraFunctionsc                 ód   — t        j                  |j                  | j                  ||«      |«      S r   )r|   ÚtensorÚcreate_fp_to_fpÚhandle)rš   Údst_tyÚfp_downcast_roundingÚ_builders       r   Ú_convert_custom_typesz$ExtraFunctions._convert_custom_typesú   s(   € ä�y‰y˜×1Ñ1°%·,±,ÀÐH\Ó]Ð_eÓfÐfr   N)r'   r(   r)   ÚstaticmethodrÊ   r*   r   r   rÂ   rÂ   ø   s   „ àñgó ñgr   rÂ   c                   ó´  — e Zd Zej                  j
                  ej                  j
                  ej                  j                  ej                  j                  ej                  j                  ej                  j                  ej                  j                  ej                  j                  iZ
ej                  j                  ej                  j                  ej                  j                  ej                  j                  ej                  j                  ej                  j                  ej                  j                   ej                  j                   ej                  j"                  ej                  j"                  ej                  j$                  ej                  j$                  ej                  j&                  ej                  j&                  ej                  j(                  ej                  j(                  ej                  j*                  ej                  j*                  ej                  j,                  ej                  j,                  i
Zd¤d„Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd	„ Zd
„ Z d„ Z!d„ Z"d„ Z#d„ Z$d„ Z%d„ Z&d„ Z'd„ Z(d„ Z)d„ Z*d„ Z+d„ Z,d„ Z-d„ Z.d„ Z/d„ Z0d„ Z1d„ Z2d„ Z3d„ Z4d„ Z5d „ Z6d!„ Z7d"„ Z8d#„ Z9d$„ Z:d%„ Z;d&„ Z<d'„ Z=d(„ Z>d)„ Z?d*„ Z@d+„ ZAd,„ ZBd-„ ZCd.„ ZDd/„ ZEd0„ ZFd1„ ZGd2„ ZHd3„ ZId4„ ZJd5„ ZKd6„ ZLd7„ ZMd8„ ZNd9„ ZOd:„ ZPd;„ ZQd<„ ZRd=„ ZSd>„ ZTd?„ ZUd@„ ZVdA„ ZWdB„ ZXdC„ ZYdD„ ZZdE„ Z[dF„ Z\dG„ Z]dH„ Z^dI„ Z_dJ„ Z`dK„ ZadL„ ZbdM„ ZcdN„ ZddO„ ZedP„ ZfdQ„ ZgdR„ ZhdS„ ZidT„ ZjdU„ ZkdV„ ZldW„ ZmdX„ ZndY„ ZodZ„ Zpd[„ Zqd\„ Zrd]„ Zsd^„ Ztd_„ Zud`„ Zvda„ Zwdb„ Zxdc„ Zydd„ Zzde„ Z{eJZ|eJZ}df„ Z~dg„ Zdh„ Z€di„ Z�dj„ Z‚dk„ Zƒdl„ Z„dm„ Z…dn„ Z†do„ Z‡dp„ Zˆdq„ Z‰dr„ ZŠds„ Z‹dt„ ZŒdu„ Z�dv„ ZŽdw„ Z�dx„ Z�dy„ Z‘dz„ Z’d{„ Z“d|„ Z”d}„ Z•d~„ Z–d„ Z—d€„ Z˜d�„ Z™d‚„ Zšdƒ„ Z›d„„ Zœd…„ Z�d†„ Zžd‡„ ZŸdˆ„ Z d‰„ Z¡dŠ„ Z¢d‹„ Z£dŒ„ Z¤d�„ Z¥dŽ„ Z¦d�„ Z§d�„ Z¨d‘„ Z©d’„ Zªd“„ Z«d”„ Z¬d•e­d–e®e­   d—e®e­   d˜e®e¯   fd™„Z°dše±d›e®e­   fdœ„Z²dše±d�e­d›e®e­   fdž„Z³dše±dŸe­d e­fd¡„Z´dše±d�e­dŸe­d e­fd¢„Zµd£„ Z¶y)¥ÚInterpreterBuilderNc                 ó˜   — d | _         t        «       | _        i | _        t        j
                  | j                  d<   d„ | j                  d<   y )NÚconvert_custom_typesc                  ó   — y)N)r   r   r   r*   )ÚlhsTypeÚrhsTypes     r   ú<lambda>z-InterpreterBuilder.__init__.<locals>.<lambda>  s   � r   Úmin_dot_size)r_   r[   ÚoptionsÚcodegen_fnsrÂ   rÊ   r   s    r   r   zInterpreterBuilder.__init__  sB   € ØˆŒ	Ü)Ó+ˆŒØˆÔÜ3A×3WÑ3Wˆ×ÑÐ/Ñ0Ù+Mˆ×Ñ˜Ò(r   c                 óÆ   — || j                   d   k  st        d«      ‚|| j                   d   k  st        d«      ‚|| j                   d   k  st        d«      ‚|||f| _        y )Nr   zx >= grid_dim[0]r   zy >= grid_dim[1]r	   zz >= grid_dim[2])Úgrid_dimÚ
ValueErrorÚgrid_idx©r   r¹   ÚyÚzs       r   Úset_grid_idxzInterpreterBuilder.set_grid_idx  sf   € Ø�4—=‘= Ñ#Ò#ÜÐ/Ó0Ð0Ø�4—=‘= Ñ#Ò#ÜÐ/Ó0Ð0Ø�4—=‘= Ñ#Ò#ÜÐ/Ó0Ð0Ø˜A˜q˜	ˆ�r   c                 ó   — |||f| _         y r   )rØ   )r   ÚnxÚnyÚnzs       r   Úset_grid_dimzInterpreterBuilder.set_grid_dim$  s   € Ø˜R ˜ˆ�r   c                 ó"   — t         j                  S r   )r|   r   r   s    r   Úget_half_tyzInterpreterBuilder.get_half_ty)  ó   € Ü�z‰zÐr   c                 ó"   — t         j                  S r   )r|   r‚   r   s    r   Úget_bf16_tyzInterpreterBuilder.get_bf16_ty,  ó   € Ü�{‰{Ðr   c                 ó"   — t         j                  S r   )r|   r€   r   s    r   Úget_float_tyzInterpreterBuilder.get_float_ty/  ræ   r   c                 ó"   — t         j                  S r   )r|   r�   r   s    r   Úget_double_tyz InterpreterBuilder.get_double_ty2  ræ   r   c                 ó"   — t         j                  S r   )r|   rs   r   s    r   Úget_int8_tyzInterpreterBuilder.get_int8_ty5  s   € Ü�w‰wˆr   c                 ó"   — t         j                  S r   )r|   rr   r   s    r   Úget_uint8_tyzInterpreterBuilder.get_uint8_ty8  ó   € Ü�x‰xˆr   c                 ó"   — t         j                  S r   )r|   ru   r   s    r   Úget_int16_tyzInterpreterBuilder.get_int16_ty;  rò   r   c                 ó"   — t         j                  S r   )r|   rt   r   s    r   Úget_uint16_tyz InterpreterBuilder.get_uint16_ty>  ó   € Ü�y‰yÐr   c                 ó"   — t         j                  S r   )r|   rw   r   s    r   Úget_int32_tyzInterpreterBuilder.get_int32_tyA  rò   r   c                 ó"   — t         j                  S r   )r|   rv   r   s    r   Úget_uint32_tyz InterpreterBuilder.get_uint32_tyD  r÷   r   c                 ó"   — t         j                  S r   )r|   rx   r   s    r   Úget_int64_tyzInterpreterBuilder.get_int64_tyG  rò   r   c                 ó"   — t         j                  S r   )r|   r@   r   s    r   Úget_uint64_tyz InterpreterBuilder.get_uint64_tyJ  r÷   r   c                 ó"   — t         j                  S r   )r|   r…   r   s    r   Úget_fp8e4nv_tyz!InterpreterBuilder.get_fp8e4nv_tyM  ó   € Ü�}‰}Ðr   c                 ó"   — t         j                  S r   )r|   r‡   r   s    r   Úget_fp8e4b15_tyz"InterpreterBuilder.get_fp8e4b15_tyP  ó   € Ü�~‰~Ðr   c                 ó"   — t         j                  S r   )r|   r†   r   s    r   Úget_fp8e4b8_tyz!InterpreterBuilder.get_fp8e4b8_tyS  r  r   c                 ó"   — t         j                  S r   )r|   rƒ   r   s    r   Úget_fp8e5_tyzInterpreterBuilder.get_fp8e5_tyV  ré   r   c                 ó"   — t         j                  S r   )r|   r„   r   s    r   Úget_fp8e5b16_tyz"InterpreterBuilder.get_fp8e5b16_tyY  r  r   c                 ó.   — t        j                  ||«      S r   )r|   r}   )r   Úelt_tyÚ
addr_spaces      r   Ú
get_ptr_tyzInterpreterBuilder.get_ptr_ty\  s   € Ü�‰˜v zÓ2Ð2r   c                 ó.   — t        j                  ||«      S r   )r|   rˆ   )r   r   r/   s      r   Úget_block_tyzInterpreterBuilder.get_block_ty_  s   € Ü�}‰}˜U EÓ*Ð*r   c                 ó~   — t        t        j                  |gt        j                  ¬«      t        j
                  «      S ©Nr6   )r   r8   ÚarrayÚbool_r|   r~   ©r   r%   s     r   Úget_int1zInterpreterBuilder.get_int1b  s$   € ÜœBŸH™H e W´B·H±HÔ=¼r¿w¹wÓGÐGr   c                 ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r  )r   r8   r  rr   r|   r  s     r   Ú	get_uint8zInterpreterBuilder.get_uint8e  ó$   € ÜœBŸH™H e W´B·H±HÔ=¼r¿x¹xÓHÐHr   c                 ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r  )r   r8   r  rs   r|   r  s     r   Úget_int8zInterpreterBuilder.get_int8h  s$   € ÜœBŸH™H e W´B·G±GÔ<¼b¿g¹gÓFÐFr   c                 ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r  )r   r8   r  rt   r|   r  s     r   Ú
get_uint16zInterpreterBuilder.get_uint16k  ó$   € ÜœBŸH™H e W´B·I±IÔ>ÄÇ	Á	ÓJÐJr   c                 ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r  )r   r8   r  ru   r|   r  s     r   Ú	get_int16zInterpreterBuilder.get_int16n  r  r   c                 ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r  )r   r8   r  rv   r|   r  s     r   Ú
get_uint32zInterpreterBuilder.get_uint32q  r  r   c                 ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r  )r   r8   r  rw   r|   r  s     r   Ú	get_int32zInterpreterBuilder.get_int32t  r  r   c                 ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r  )r   r8   r  r@   r|   r  s     r   Ú
get_uint64zInterpreterBuilder.get_uint64w  r  r   c                 ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r  )r   r8   r  rx   r|   r  s     r   Ú	get_int64zInterpreterBuilder.get_int64z  r  r   c                 ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r  )r   r8   r  r   r|   r  s     r   Úget_fp16zInterpreterBuilder.get_fp16}  ó$   € ÜœBŸH™H e W´B·J±JÔ?ÄÇÁÓLÐLr   c                 ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r  )r   r8   r  r€   r|   r  s     r   Úget_fp32zInterpreterBuilder.get_fp32€  r,  r   c                 ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r  )r   r8   r  r�   r|   r  s     r   Úget_fp64zInterpreterBuilder.get_fp64ƒ  r,  r   c                 óX   — t        t        j                  dgt        |«      ¬«      |«      S ©Nr   r6   )r   r8   r  r‹   )r   Útypes     r   Úget_null_valuez!InterpreterBuilder.get_null_value†  s!   € ÜœBŸH™H a S´¸dÓ0CÔDÀdÓKÐKr   c                 óÆ   — | j                   €t        d«      ‚t        t        j                  | j                   |   gt        j
                  ¬«      t        j
                  «      S )Nzgrid_idx is Noner6   )rÚ   rÙ   r   r8   r  rw   r|   ©r   Úaxiss     r   Úcreate_get_program_idz(InterpreterBuilder.create_get_program_idŠ  sD   € Ø�=‰=Ð ÜÐ/Ó0Ð0ÜœBŸH™H d§m¡m°DÑ&9Ð%:Ä"Ç(Á(ÔKÌRÏXÉXÓVÐVr   c                 ó˜   — t        t        j                  | j                  |   gt        j                  ¬«      t
        j                  «      S r  )r   r8   r  rØ   rw   r|   r6  s     r   Úcreate_get_num_programsz*InterpreterBuilder.create_get_num_programs�  s.   € ÜœBŸH™H d§m¡m°DÑ&9Ð%:Ä"Ç(Á(ÔKÌRÏXÉXÓVÐVr   c                 ó¬   — t        t        j                  |j                  t        ¬«      t
        j                  «      }d }| j                  ||||||«      S r  )r   r8   Ú	ones_liker   r   r|   r~   Úcreate_masked_load)r   ÚptrÚ_0Ú_1Úis_volatileÚmaskÚothers          r   Úcreate_loadzInterpreterBuilder.create_load“  sA   € ÜœBŸL™L¨¯©¼Ô>ÄÇÁÓHˆØˆØ×&Ñ& s¨D°%¸¸RÀÓMÐMr   c                 ó¦   — t        t        j                  |j                  t        ¬«      t
        j                  «      }| j                  |||d d «      S r  )r   r8   r<  r   r   r|   r~   Úcreate_masked_store)r   r>  Úvalr?  r@  rB  s         r   Úcreate_storezInterpreterBuilder.create_store˜  s:   € ÜœBŸL™L¨¯©¼Ô>ÄÇÁÓHˆØ×'Ñ'¨¨S°$¸¸dÓCÐCr   c                 ó  — |j                  «       }t        |«      }|€+t        t        j                  |j
                  |¬«      |«      }t        j                  |j
                  |j
                  |j
                  |«      }	t        |	|«      S r  )r   r‹   r   r8   r”   r   Ú_interpreterÚload)
r   rE   rB  rC  Úcache_modifierÚeviction_policyrA  rC   Údtype_npÚrets
             r   r=  z%InterpreterBuilder.create_masked_loadœ  si   € Ø×&Ñ&Ó(ˆÜ  Ó*ˆØˆ=Ü ¤§¡¨t¯y©yÀÔ!IÈ8ÓTˆEÜ×Ñ §	¡	¨4¯9©9°e·j±jÀ(ÓKˆÜ˜C Ó*Ð*r   c                 ól   — t        j                  |j                  |j                  |j                  «      S r   )rJ  Ústorer   )r   rE   r%   rB  rL  rM  s         r   rF  z&InterpreterBuilder.create_masked_store¤  s#   € Ü×!Ñ! $§)¡)¨U¯Z©Z¸¿¹ÓCÐCr   c                 óÚ  — |j                   j                  }|j                  }|t        j                  k(  r|t        j                  k(  s&|t        j                  k(  rY|t        j                  k(  rFt        |j                  ||d «      j                  t        |«      «      }t        ||j                  «      S t        |j                  j                  t        |«      «      |j                  «      S r   )r   rA   r|   r‚   r€   r´   r   Úviewr‹   r   r?   )r   ÚsrcÚdst_typeÚsrc_element_typeÚdst_element_typer   s         r   Ú	cast_implzInterpreterBuilder.cast_impl¨  s©   € ØŸ9™9×+Ñ+ÐØ#Ÿ?™?ÐØ¤§¡Ò+Ð0@ÄBÇJÁJÒ0NØ¤§
¡
Ò*Ð/?Ä2Ç;Á;Ò/NÜ! #§(¡(Ð,<Ð>NÐPTÓU×ZÑZÔ[hÐiqÓ[rÓsˆDÜ  h§o¡oÓ6Ð6ä §¡§¡´¸hÓ0GÓ HÈ(Ï/É/ÓZÐZr   c                 ó&   — | j                  ||«      S r   ©rX  ©r   rT  rU  s      r   rÓ   zInterpreterBuilder.<lambda>²  ó   € °$·.±.ÀÀhÓ2O€ r   c                 ó&   — | j                  ||«      S r   rZ  r[  s      r   rÓ   zInterpreterBuilder.<lambda>³  r\  r   c                 ó&   — | j                  ||«      S r   rZ  r[  s      r   rÓ   zInterpreterBuilder.<lambda>´  r\  r   c                 ó&   — | j                  ||«      S r   rZ  r[  s      r   rÓ   zInterpreterBuilder.<lambda>µ  r\  r   c                 ó&   — | j                  ||«      S r   rZ  r[  s      r   rÓ   zInterpreterBuilder.<lambda>¶  s   € °·±¸sÀHÓ0M€ r   c                 ó&   — | j                  ||«      S r   rZ  r[  s      r   rÓ   zInterpreterBuilder.<lambda>·  r\  r   c                 ó&   — | j                  ||«      S r   rZ  )r   rT  rU  Ú	is_signeds       r   rÓ   zInterpreterBuilder.<lambda>¸  s   € ¸T¿^¹^ÈCÐQYÓ=Z€ r   c                 óÒ   — |j                   j                  }|j                  }t        |j                  |||«      j	                  t        |«      «      }t        ||j                  «      S r   )r   rA   r´   r   rS  r‹   r   )r   rT  rU  r�   rV  rW  r   s          r   rÅ   z"InterpreterBuilder.create_fp_to_fpº  sU   € ØŸ9™9×+Ñ+ÐØ#Ÿ?™?ÐÜ˜cŸh™hÐ(8Ð:JÈMÓZ×_Ñ_Ô`mÐnvÓ`wÓxˆÜ˜D (§/¡/Ó2Ð2r   c                 ór   — t        |j                  j                  t        |«      «      |j                  «      S r   )r   r   rS  r‹   rA   r[  s      r   Úcreate_bitcastz!InterpreterBuilder.create_bitcastÀ  s%   € Ü˜CŸH™HŸM™M¬-¸Ó*AÓBÀHÇOÁOÓTÐTr   c                 óx   — t         ||j                  |j                  «      |j                  j                  «      S r   ©r   r   r   rA   )r   ÚlhsÚrhsÚops       r   Ú	binary_opzInterpreterBuilder.binary_opÄ  s(   € Ü™B˜sŸx™x¨¯©Ó2°C·I±I×4DÑ4DÓEÐEr   c                 óD   — | j                  ||t        j                  «      S r   ©rl  r8   Úadd©r   ri  rj  s      r   rÓ   zInterpreterBuilder.<lambda>Ç  s   € ¨¯©¸¸SÄ"Ç&Á&Ó)I€ r   c                 óD   — | j                  ||t        j                  «      S r   ©rl  r8   Úmultiplyrp  s      r   rÓ   zInterpreterBuilder.<lambda>È  ó   € ¨¯©¸¸SÄ"Ç+Á+Ó)N€ r   c                 óD   — | j                  ||t        j                  «      S r   ©rl  r8   Údividerp  s      r   rÓ   zInterpreterBuilder.<lambda>É  s   € ¨¯©¸¸SÄ"Ç)Á)Ó)L€ r   c                 óD   — | j                  ||t        j                  «      S r   ©rl  r8   Úfmodrp  s      r   rÓ   zInterpreterBuilder.<lambda>Ê  ó   € ¨¯©¸¸SÄ"Ç'Á'Ó)J€ r   c                 óD   — | j                  ||t        j                  «      S r   ©rl  r8   Úsubtractrp  s      r   rÓ   zInterpreterBuilder.<lambda>Ë  rt  r   c                 óD   — | j                  ||t        j                  «      S r   rr  rp  s      r   rÓ   zInterpreterBuilder.<lambda>Ì  ó   € ¨¯©°s¸CÄÇÁÓ(M€ r   c                 óD   — | j                  ||t        j                  «      S r   rv  rp  s      r   rÓ   zInterpreterBuilder.<lambda>Í  s   € °·±ÀÀSÌ"Ï)É)Ó1T€ r   c                 ó&   — | j                  ||«      S r   ©Úcreate_idivrp  s      r   rÓ   zInterpreterBuilder.<lambda>Î  ó   € ¨×)9Ñ)9¸#¸sÓ)C€ r   c                 ó&   — | j                  ||«      S r   rƒ  rp  s      r   rÓ   zInterpreterBuilder.<lambda>Ï  r…  r   c                 óD   — | j                  ||t        j                  «      S r   ry  rp  s      r   rÓ   zInterpreterBuilder.<lambda>Ñ  r{  r   c                 óD   — | j                  ||t        j                  «      S r   ry  rp  s      r   rÓ   zInterpreterBuilder.<lambda>Ò  r{  r   c                 óD   — | j                  ||t        j                  «      S r   rn  rp  s      r   rÓ   zInterpreterBuilder.<lambda>Ó  s   € ¨¯©°s¸CÄÇÁÓ(H€ r   c                 óD   — | j                  ||t        j                  «      S r   r}  rp  s      r   rÓ   zInterpreterBuilder.<lambda>Ô  r€  r   c                 óD   — | j                  ||t        j                  «      S r   )rl  r8   Ú
left_shiftrp  s      r   rÓ   zInterpreterBuilder.<lambda>Õ  s   € ¨¯©°s¸CÄÇÁÓ(O€ r   c                 óD   — | j                  ||t        j                  «      S r   )rl  r8   Úright_shiftrp  s      r   rÓ   zInterpreterBuilder.<lambda>Ö  s   € ¨¯©¸¸SÄ"Ç.Á.Ó)Q€ r   c                 óD   — | j                  ||t        j                  «      S r   ©rl  r8   r–   rp  s      r   rÓ   zInterpreterBuilder.<lambda>×  ó   € ¨$¯.©.¸¸cÄ2Ç:Á:Ó*N€ r   c                 óD   — | j                  ||t        j                  «      S r   r�  rp  s      r   rÓ   zInterpreterBuilder.<lambda>Ø  r‘  r   c                 óD   — | j                  ||t        j                  «      S r   r�  rp  s      r   rÓ   zInterpreterBuilder.<lambda>Ù  ó   € ¨T¯^©^¸CÀÄbÇjÁjÓ-Q€ r   c                 óD   — | j                  ||t        j                  «      S r   r�  rp  s      r   rÓ   zInterpreterBuilder.<lambda>Ú  ó   € ¨D¯N©N¸3ÀÄRÇZÁZÓ,P€ r   c                 óD   — | j                  ||t        j                  «      S r   ©rl  r8   r•   rp  s      r   rÓ   zInterpreterBuilder.<lambda>Û  r‘  r   c                 óD   — | j                  ||t        j                  «      S r   r˜  rp  s      r   rÓ   zInterpreterBuilder.<lambda>Ü  r‘  r   c                 óD   — | j                  ||t        j                  «      S r   r˜  rp  s      r   rÓ   zInterpreterBuilder.<lambda>Ý  r”  r   c                 óD   — | j                  ||t        j                  «      S r   r˜  rp  s      r   rÓ   zInterpreterBuilder.<lambda>Þ  r–  r   c                 óD   — | j                  ||t        j                  «      S r   ©rl  r8   Ú
less_equalrp  s      r   rÓ   zInterpreterBuilder.<lambda>ß  ó   € ¨D¯N©N¸3ÀÄRÇ]Á]Ó,S€ r   c                 óD   — | j                  ||t        j                  «      S r   ©rl  r8   Úlessrp  s      r   rÓ   zInterpreterBuilder.<lambda>à  ó   € ¨D¯N©N¸3ÀÄRÇWÁWÓ,M€ r   c                 óD   — | j                  ||t        j                  «      S r   ©rl  r8   Úgreater_equalrp  s      r   rÓ   zInterpreterBuilder.<lambda>á  ó   € ¨D¯N©N¸3ÀÄR×EUÑEUÓ,V€ r   c                 óD   — | j                  ||t        j                  «      S r   ©rl  r8   Úgreaterrp  s      r   rÓ   zInterpreterBuilder.<lambda>â  r–  r   c                 óD   — | j                  ||t        j                  «      S r   r�  rp  s      r   rÓ   zInterpreterBuilder.<lambda>ã  rŸ  r   c                 óD   — | j                  ||t        j                  «      S r   r¡  rp  s      r   rÓ   zInterpreterBuilder.<lambda>ä  r£  r   c                 óD   — | j                  ||t        j                  «      S r   r¥  rp  s      r   rÓ   zInterpreterBuilder.<lambda>å  r§  r   c                 óD   — | j                  ||t        j                  «      S r   r©  rp  s      r   rÓ   zInterpreterBuilder.<lambda>æ  r–  r   c                 óD   — | j                  ||t        j                  «      S r   ©rl  r8   Úequalrp  s      r   rÓ   zInterpreterBuilder.<lambda>ç  s   € ¨4¯>©>¸#¸sÄBÇHÁHÓ+M€ r   c                 óD   — | j                  ||t        j                  «      S r   ©rl  r8   Ú	not_equalrp  s      r   rÓ   zInterpreterBuilder.<lambda>è  s   € ¨4¯>©>¸#¸sÄBÇLÁLÓ+Q€ r   c                 óD   — | j                  ||t        j                  «      S r   r¡  rp  s      r   rÓ   zInterpreterBuilder.<lambda>é  r£  r   c                 óD   — | j                  ||t        j                  «      S r   r©  rp  s      r   rÓ   zInterpreterBuilder.<lambda>ê  r–  r   c                 óD   — | j                  ||t        j                  «      S r   r�  rp  s      r   rÓ   zInterpreterBuilder.<lambda>ë  rŸ  r   c                 óD   — | j                  ||t        j                  «      S r   r¥  rp  s      r   rÓ   zInterpreterBuilder.<lambda>ì  r§  r   c                 óD   — | j                  ||t        j                  «      S r   r°  rp  s      r   rÓ   zInterpreterBuilder.<lambda>í  ó   € ¨D¯N©N¸3ÀÄRÇXÁXÓ,N€ r   c                 óD   — | j                  ||t        j                  «      S r   r³  rp  s      r   rÓ   zInterpreterBuilder.<lambda>î  ó   € ¨D¯N©N¸3ÀÄRÇ\Á\Ó,R€ r   c                 óD   — | j                  ||t        j                  «      S r   r¡  rp  s      r   rÓ   zInterpreterBuilder.<lambda>ï  r£  r   c                 óD   — | j                  ||t        j                  «      S r   r©  rp  s      r   rÓ   zInterpreterBuilder.<lambda>ð  r–  r   c                 óD   — | j                  ||t        j                  «      S r   r�  rp  s      r   rÓ   zInterpreterBuilder.<lambda>ñ  rŸ  r   c                 óD   — | j                  ||t        j                  «      S r   r¥  rp  s      r   rÓ   zInterpreterBuilder.<lambda>ò  r§  r   c                 óD   — | j                  ||t        j                  «      S r   r°  rp  s      r   rÓ   zInterpreterBuilder.<lambda>ó  rº  r   c                 óD   — | j                  ||t        j                  «      S r   r³  rp  s      r   rÓ   zInterpreterBuilder.<lambda>ô  r¼  r   c                 óD   — | j                  ||t        j                  «      S r   )rl  r8   Úbitwise_andrp  s      r   rÓ   zInterpreterBuilder.<lambda>õ  ó   € ¨¯©°s¸CÄÇÁÓ(P€ r   c                 óD   — | j                  ||t        j                  «      S r   )rl  r8   Úbitwise_xorrp  s      r   rÓ   zInterpreterBuilder.<lambda>ö  rÅ  r   c                 óD   — | j                  ||t        j                  «      S r   )rl  r8   Ú
bitwise_orrp  s      r   rÓ   zInterpreterBuilder.<lambda>÷  s   €  t§~¡~°c¸3ÄÇÁÓ'N€ r   c                 óÆ   — t        |j                  t        j                  |j                  |j                  «      z
  |j                  z  |j                  j
                  «      S r   )r   r   r8   rz  r   rA   rp  s      r   r„  zInterpreterBuilder.create_idivû  sC   € ô ˜SŸX™X¬¯©°·±¸#¿(¹(Ó(CÑCÈÏÉÑPÐRU×R[ÑR[×RbÑRbÓcÐcr   c                 ó@  — t        |j                  j                  «      }t        |j                  j                  «      }|j                  j                  |«      |_        |j                  j                  |«      |_        | j	                  ||t
        j                  «      S r   )ry   r   r   r?   rl  r8   rŽ  )r   ri  rj  Ú	lhs_dtypeÚ	rhs_dtypes        r   Úcreate_ashrzInterpreterBuilder.create_ashr  se   € ä(¨¯©¯©Ó8ˆ	Ü(¨¯©¯©Ó8ˆ	Ø—8‘8—?‘? 9Ó-ˆŒØ—8‘8—?‘? 9Ó-ˆŒØ�~‰~˜c 3¬¯©Ó7Ð7r   c                 óR  — |j                   j                  }|t        j                  k(  s|t        j                  k(  r>t        t        |j                   |j                   «      |j                  j                  «      S t        t        d|j                  dz  dz  › �«      }|j                   j                  |«      }|j                   j                  |«      }t        j                  ||«      |j                  dz  z	  }t        |j                  |«      |j                  j                  «      S )Nr�   r5   r	   )r   r   r8   rx   r@   r   Únp_umulhi_u64rA   rŽ   rW   r?   rs  )r   ri  rj  r   Úcompute_dtypeÚlhs_dataÚrhs_dataÚret_datas           r   Úcreate_umulhiz InterpreterBuilder.create_umulhi	  sÐ   € Ø—‘—‘ˆØ”B—H‘HÒ ¬¯©Ò 2Ü¤¨c¯h©h¸¿¹Ó AÀ3Ç9Á9×CSÑCSÓTÐTä#¤B¨$¨u¯~©~ÀÑ/AÀAÑ/EÐ.FÐ(GÓHˆMØ—x‘x—‘ }Ó5ˆHØ—x‘x—‘ }Ó5ˆHÜ—{‘{ 8¨XÓ6¸5¿>¹>ÈAÑ;MÑNˆHÜ §¡°Ó 6¸¿	¹	×8HÑ8HÓIÐIr   c                 óŽ   — t         ||j                  |j                  |j                  «      |j                  j                  «      S r   rh  )r   ri  rj  rC  rk  s        r   Ú
ternary_opzInterpreterBuilder.ternary_op  s.   € Ü™B˜sŸx™x¨¯©°5·:±:Ó>ÀÇÁ×@RÑ@RÓSÐSr   c                 óF   — | j                  |||t        j                  «      S r   )r×  r8   Úclip)r   ÚargÚloÚhiÚpropagate_nanss        r   rÓ   zInterpreterBuilder.<lambda>  s   € ¸d¿o¹oÈcÐSUÐWYÔ[]×[bÑ[bÓ>c€ r   c                 óF   — | j                  |||t        j                  «      S r   )r×  r8   Úwhere)r   Úcondri  rj  s       r   rÓ   zInterpreterBuilder.<lambda>  s   € °·±ÀÀsÈCÔQS×QYÑQYÓ1Z€ r   c                 óŠ   — t        |j                  |j                  z  |j                  z   |j                  j                  «      S r   rh  rÛ   s       r   Ú
create_fmazInterpreterBuilder.create_fma  s,   € Ü˜AŸF™F Q§V¡V™O¨a¯f©fÑ4°a·g±g·n±nÓEÐEr   c                 ób   — t         ||j                  «      |j                  j                  «      S r   rh  )r   rÚ  rk  s      r   Úunary_opzInterpreterBuilder.unary_op  s!   € Ü™B˜sŸx™x›L¨#¯)©)×*:Ñ*:Ó;Ð;r   c                 ó2  — |j                   }|j                  dz
  }t        t        d|j                  › �«      }|j                  j                  |«      }d|z  dz
  }||z  j                  t        |«      «      }t        ||j                   j                  «      S )Nr   r�   )	r   r7   rŽ   r8   r   rS  r‹   r   rA   )r   rÚ  rC   Úmask_bitwidthÚnp_uint_dtyper   rB  rO  s           r   Úcreate_fabszInterpreterBuilder.create_fabs"  s†   € à—9‘9ˆØ ×3Ñ3°aÑ7ˆÜ¤ d¨8×+FÑ+FÐ*GÐ$HÓIˆØ�x‰x�}‰}˜]Ó+ˆØ�]Ñ" aÑ'ˆØ�d‰{× Ñ ¤¨xÓ!8Ó9ˆÜ˜C §¡×!1Ñ!1Ó2Ð2r   c                 óB   — | j                  |t        j                  «      S r   )rä  r8   Úcos©r   rÚ  s     r   rÓ   zInterpreterBuilder.<lambda>,  ó   €  4§=¡=°´b·f±fÓ#=€ r   c                 óB   — | j                  |t        j                  «      S r   )rä  r8   Úexprë  s     r   rÓ   zInterpreterBuilder.<lambda>-  rì  r   c                 óB   — | j                  |t        j                  «      S r   )rä  r8   Úexp2rë  s     r   rÓ   zInterpreterBuilder.<lambda>.  ó   €  D§M¡M°#´r·w±wÓ$?€ r   c                 óB   — | j                  |t        j                  «      S r   )rä  r8   Úabsrë  s     r   rÓ   zInterpreterBuilder.<lambda>/  s   €  D§M¡M°#´r·v±vÓ$>€ r   c                 óB   — | j                  |t        j                  «      S r   )rä  r8   Úfloorrë  s     r   rÓ   zInterpreterBuilder.<lambda>0  s   €  T§]¡]°3¼¿¹Ó%A€ r   c                 óB   — | j                  |t        j                  «      S r   )rä  r8   Úceilrë  s     r   rÓ   zInterpreterBuilder.<lambda>1  rñ  r   c                 óB   — | j                  |t        j                  «      S r   )rä  r8   Úlogrë  s     r   rÓ   zInterpreterBuilder.<lambda>2  rì  r   c                 óB   — | j                  |t        j                  «      S r   )rä  r8   Úlog2rë  s     r   rÓ   zInterpreterBuilder.<lambda>3  rñ  r   c                 óB   — | j                  |t        j                  «      S r   ©rä  r8   Úsqrtrë  s     r   rÓ   zInterpreterBuilder.<lambda>4  s   € ¨D¯M©M¸#¼r¿w¹wÓ,G€ r   c                 óB   — | j                  |t        j                  «      S r   rý  rë  s     r   rÓ   zInterpreterBuilder.<lambda>5  rñ  r   c                 óB   — | j                  |t        j                  «      S r   )rä  r8   Úsinrë  s     r   rÓ   zInterpreterBuilder.<lambda>6  rì  r   c                 óä   — |j                   j                  t        j                  k(  rt	        |j                   «      nt        |j                   «      }t        ||j                  j                  «      S r   )r   r   r8   r€   Únp_erf_fp32Únp_erf_fp64r   rA   )r   rÚ  rO  s      r   Ú
create_erfzInterpreterBuilder.create_erf8  sH   € Ø'*§x¡x§~¡~¼¿¹Ò'CŒk˜#Ÿ(™(Ô#ÌÐUX×U]ÑU]ÓI^ˆÜ˜C §¡×!1Ñ!1Ó2Ð2r   c                 ó‚   — t        dt        j                  |j                  «      z  |j                  j
                  «      S ©Nr   )r   r8   rþ  r   r   rA   rë  s     r   Úcreate_rsqrtzInterpreterBuilder.create_rsqrt<  s+   € Ü˜A¤§¡¨¯©Ó 1Ñ1°3·9±9×3CÑ3CÓDÐDr   c                 ót   — t        |j                  j                  |«      |j                  j                  «      S r   )r   r   r>   r   rA   )r   rÚ  r/   Úallow_reorders       r   rÓ   zInterpreterBuilder.<lambda>@  s*   € ¼\È#Ï(É(×JZÑJZÐ[`ÓJaÐcf×clÑcl×csÑcsÓ=t€ r   c                 ó~   — t        t        j                  |j                  |«      |j                  j
                  «      S r   )r   r8   Ú	transposer   r   rA   )r   rÚ  Úperms      r   Úcreate_transzInterpreterBuilder.create_transB  s(   € ÜœBŸL™L¨¯©°4Ó8¸#¿)¹)×:JÑ:JÓKÐKr   c                 ó¸  — |j                   }|j                   }|j                  j                  dk(  r|j                  j                  «       s3|j                  j                  dk(  r |j                  j                  «       r†t	        ||j                  t
        j                  d «      j                  t        j                  «      }t	        ||j                  t
        j                  d «      j                  t        j                  «      }t        t        j                  |||j                   j                  ¬«      |j                   z   |j                  j                  «      S )Nr5   r6   )r   r   r7   Úis_floatingr´   r|   r   rS  r8   r   ÚmatmulrA   )r   r½   r¾   ÚdÚinput_precisionÚmax_num_imprecise_accÚa_dataÚb_datas           r   Ú
create_dotzInterpreterBuilder.create_dotE  sÜ   € Ø—‘ˆØ—‘ˆØ�G‰G×&Ñ&¨!Ò+°·±×0CÑ0CÔ0EØ�G‰G×&Ñ&¨!Ò+°·±×0CÑ0CÔ0EÜ# F¨A¯G©G´R·Z±ZÀÓF×KÑKÌBÏJÉJÓWˆFÜ# F¨A¯G©G´R·Z±ZÀÓF×KÑKÌBÏJÉJÓWˆFÜœBŸI™I f¨f¸A¿F¹F¿L¹LÔIÈAÏFÉFÑRÐTU×T[ÑT[×TbÑTbÓcÐcr   c                 ó~   — t        t        j                  ||t        j                  ¬«      t        j                  «      S r  )r   r8   r=   rw   r|   )r   ÚstartÚstops      r   Úcreate_make_rangez$InterpreterBuilder.create_make_rangeN  s$   € ÜœBŸI™I e¨T¼¿¹ÔBÄBÇHÁHÓMÐMr   c                 ó€   — t        t        j                  |j                  |d|f¬«      d   t        j
                  «      S )Nr   )Úbinsr;   )r   r8   Ú	histogramr   r|   rw   )r   r   r  s      r   Úcreate_histogramz#InterpreterBuilder.create_histogramQ  s1   € ÜœBŸL™L¨¯©¸ÀaÈÀYÔOÐPQÑRÔTV×T\ÑT\Ó]Ð]r   c                 ó–   — t        t        j                  |j                  |j                  |¬«      |j                  j
                  «      S )N©r7  )r   r8   Útake_along_axisr   r   rA   )r   rT  Úindicesr7  s       r   Úcreate_gatherz InterpreterBuilder.create_gatherT  s3   € ÜœB×.Ñ.¨s¯x©x¸¿¹ÈDÔQÐSV×S\ÑS\×ScÑScÓdÐdr   c                 óò   — |j                  «       }|j                  }t        d|dz  «      }t        |j                  ||j                  j                  t        j                  «      z  z   |j                  «      S )Nr   r5   )	r   r7   Úmaxr   r   r?   r8   r@   r   )r   r>  ÚoffsetrC   Úelement_bitwidthÚelement_bytewidths         r   Úcreate_addptrz InterpreterBuilder.create_addptrY  se   € Ø×%Ñ%Ó'ˆØ#×6Ñ6Ðä Ð#3°qÑ#8Ó9ÐÜ˜CŸH™HÐ'8¸6¿;¹;×;MÑ;MÌbÏiÉiÓ;XÑ'XÑXÐZ]×ZcÑZcÓdÐdr   c                 óê  — |j                  |«      \  }}|j                  «       }	t        |	«      }
|€d }nª|t        j                  j
                  k(  r,t        t        j                  |j                  |
¬«      |	«      }na|t        j                  j                  k(  r6t        t        j                  |j                  t        d«      |
¬«      |	«      }nt        d|› �«      ‚| j                  ||||||«      S )Nr6   Únanzunsupported padding option )rJ   r   r‹   r—   ÚPADDING_OPTIONÚPAD_ZEROr   r8   r”   r   ÚPAD_NANÚ	full_likeÚfloatrÙ   r=  )r   r>  rB   Úpadding_optionrL  rM  rA  rE   rF   rC   rN  rC  s               r   Úcreate_tensor_pointer_loadz-InterpreterBuilder.create_tensor_pointer_load`  sÎ   € à×.Ñ.¨~Ó>‰ˆˆeØ×&Ñ&Ó(ˆÜ  Ó*ˆØÐ!Ø‰EØœs×1Ñ1×:Ñ:Ò:Ü ¤§¡¨t¯y©yÀÔ!IÈ8ÓT‰EØœs×1Ñ1×9Ñ9Ò9Ü ¤§¡¨d¯i©i¼¸u»ÈXÔ!VÐX`Óa‰EäÐ:¸>Ð:JÐKÓLÐLØ×&Ñ& t¨U°E¸>È?Ð\gÓhÐhr   c                 óT   — |j                  |«      \  }}| j                  |||||«      S r   ©rJ   rF  )r   r>  r%   rB   rL  rM  rE   rF   s           r   Úcreate_tensor_pointer_storez.InterpreterBuilder.create_tensor_pointer_storeo  s/   € Ø×.Ñ.¨~Ó>‰ˆˆeØ×'Ñ'¨¨e°U¸NÈOÓ\Ð\r   c                 ó~   — t        t        j                  |j                  |«      |j                  j
                  «      S r   )r   r8   Úexpand_dimsr   r   rA   )r   rÚ  r7  s      r   Úcreate_expand_dimsz%InterpreterBuilder.create_expand_dimss  s(   € ÜœBŸN™N¨3¯8©8°TÓ:¸C¿I¹I×<LÑ<LÓMÐMr   c                 ó~   — t        t        j                  |j                  |«      |j                  j
                  «      S r   )r   r8   r9   r   r   rA   ©r   rÚ  r/   s      r   Úcreate_broadcastz#InterpreterBuilder.create_broadcastv  s(   € ÜœBŸO™O¨C¯H©H°eÓ<¸c¿i¹i×>NÑ>NÓOÐOr   c                 ó”   — t        t        j                  |j                  |j                  g«      |j                  j
                  «      S r   )r   r8   Úconcatenater   r   rA   rp  s      r   Ú
create_catzInterpreterBuilder.create_caty  s/   € ÜœBŸN™N¨C¯H©H°c·h±hÐ+?Ó@À#Ç)Á)×BRÑBRÓSÐSr   c                 ó˜   — t        t        j                  |j                  |j                  gd¬«      |j                  j
                  «      S )NrQ   r!  )r   r8   Ústackr   r   rA   rp  s      r   Úcreate_joinzInterpreterBuilder.create_join|  s1   € äœBŸH™H c§h¡h°·±Ð%9ÀÔCÀSÇYÁY×EUÑEUÓVÐVr   c                 ó¶   — t        |j                  d   |j                  j                  «      t        |j                  d   |j                  j                  «      fS )N).r   ).r   rh  )r   rG  s     r   Úcreate_splitzInterpreterBuilder.create_split€  sE   € ä˜SŸX™X fÑ-¨s¯y©y×/?Ñ/?Ó@Ä,ÈsÏxÉxÐX^ÑO_Ðad×ajÑaj×aqÑaqÓBrÐsÐsr   c           	      óœ  — t        |j                  t        j                  «      rVt	        t        j                  ||j                  d   t        |j                  «      ¬«      |j                  j                  «      S t	        t        j                  ||j                  t        |j                  «      ¬«      |j                  j                  «      S r2  )
r{   r   r|   rˆ   r   r8   Úfullr   r‹   rA   r;  s      r   Úcreate_splatzInterpreterBuilder.create_splat„  sƒ   € Ü�c—i‘i¤§¡Ô/Ü¤§¡¨¨s¯x©x¸©{Ä-ÐPS×PYÑPYÓBZÔ [Ð]`×]fÑ]f×]mÑ]mÓnÐnä¤§¡¨¨s¯x©x¼}ÈSÏYÉYÓ?WÔ XÐZ]×ZcÑZc×ZjÑZjÓkÐkr   c                 ó   — || j                   vrt        d|› �«      ‚| j                   |   }t        t        j                  |j
                  |j
                  |j
                  |«      |j                  j                  «      S )Núunsupported semantic )Úir_sem_to_interpreter_semrÙ   r   rJ  Ú
atomic_casr   r   rA   )r   r>  ÚcmprG  ÚsemÚscopes         r   Úcreate_atomic_casz$InterpreterBuilder.create_atomic_casŠ  sk   € Ø�d×4Ñ4Ñ4ÜÐ4°S°EÐ:Ó;Ð;Ø×,Ñ,¨SÑ1ˆÜœL×3Ñ3°C·H±H¸c¿h¹hÈÏÉÐRUÓVÐX[×XaÑXa×XhÑXhÓiÐir   c           	      óX  — || j                   vrt        d|› �«      ‚|| j                  vrt        d|› �«      ‚| j                   |   }| j                  |   }t        t	        j
                  ||j                  |j                  |j                  |«      |j                  j                  «      S )Nzunsupported rmwOp rI  )	Úir_rmw_op_to_interpreter_rmw_oprÙ   rJ  r   rJ  Ú
atomic_rmwr   r   rA   )r   ÚrmwOpr>  rG  rB  rM  rN  s          r   Úcreate_atomic_rmwz$InterpreterBuilder.create_atomic_rmw�  sž   € Ø˜×<Ñ<Ñ<ÜÐ1°%°Ð9Ó:Ð:Ø�d×4Ñ4Ñ4ÜÐ4°S°EÐ:Ó;Ð;Ø×4Ñ4°UÑ;ˆØ×,Ñ,¨SÑ1ˆÜœL×3Ñ3°E¸3¿8¹8ÀSÇXÁXÈtÏyÉyÐZ]Ó^Ð`c×`iÑ`i×`pÑ`pÓqÐqr   c                 ó   — t        d«      ‚)Nz4extern_elementwise not supported in interpreter mode©ÚNotImplementedError)r   ÚlibNameÚlibPathÚsymbolÚargListÚretTypeÚisPures          r   Úcreate_extern_elementwisez,InterpreterBuilder.create_extern_elementwise™  s   € Ü!Ð"XÓYÐYr   c                 ó   — t        d«      ‚)Nz,inline_asm not supported in interpreter moderV  )r   Ú	inlineAsmÚconstraintsÚvaluesr3  r]  Úpacks          r   Úcreate_inline_asmz$InterpreterBuilder.create_inline_asmœ  s   € Ü!Ð"PÓQÐQr   c                 ó*  — d| j                   d   › d| j                   d   › d| j                   d   › d�}|r|d|› �z  }|rt        j                  dd	„ i¬
«       |D ]  }t        |d|j                  › �z   «       Œ |rt        j                  d ¬
«       y y )Nú(r   z, r   r	   ú)ú r   c                 ó   — d| d›�S )NÚ0xÚ02xr*   r¸   s    r   rÓ   z1InterpreterBuilder.create_print.<locals>.<lambda>¨  s   € ¸bÀÀ3À¸L€ r   )Ú	formatter)rÚ   r8   Úset_printoptionsÚprintr   )r   ÚprefixÚhexrb  ÚisSignedÚmsgr%   s          r   Úcreate_printzInterpreterBuilder.create_printŸ  s£   € ð
 �$—-‘- Ñ"Ð# 2 d§m¡m°AÑ&6Ð%7°r¸$¿-¹-ÈÑ:JÐ9KÈ1ÐMˆÙØ�Q�v�h�<ÑˆCÙÜ×Ñ¨5Ñ2HÐ*IÕJØò 	*ˆEÜ�#˜!˜EŸJ™J˜<Ð(Ñ(Õ)ð	*áÜ×Ñ¨$Ö/ð r   c                 ó   — |sJ |› «       ‚y r   r*   )r   Ú	conditionÚmessages      r   Úcreate_assertz InterpreterBuilder.create_assert®  s   € áÐ&˜W˜IÓ&‰yr   c                 ó   — |sJ d«       ‚y )NzAssume failedr*   )r   ru  s     r   Úcreate_assumez InterpreterBuilder.create_assume²  s   € ÙÐ)˜/Ó)‰yr   c                  ó   — y r   r*   r   s    r   Úcreate_barrierz!InterpreterBuilder.create_barrierµ  s   € àr   c                 óf   — |D �cg c]  }|j                  «       ‘Œ }}t        ||||||«      S c c}w r   )r"   r,   )	r   r.   r/   r0   r1   r2   r3   r'  Únew_offsetss	            r   Úcreate_make_block_ptrz(InterpreterBuilder.create_make_block_ptr¹  s6   € à4;Ö<¨&�v—|‘|•~Ð<ˆÐ<Ü! $¨¨w¸À[ÐRWÓXÐXùò =s   ….c                 óÊ  — t        |j                  «      t        |«      k7  rt        d«      ‚|j                  D �cg c]  }|j                  «       ‘Œ }}t	        |j
                  |j                  |j                  ||j                  |j                  «      }t        t        |«      «      D ]1  }|j                  |   xj                  ||   j                  z  c_        Œ3 |S c c}w )Nz len(ptr.offsets) != len(offsets))r<   r1   rÙ   r"   r,   r.   r/   r0   r2   r3   r;   r   )r   r>  r1   r'  r}  rO  rª   s          r   Úcreate_advancez!InterpreterBuilder.create_advance¾  s²   € Üˆs�{‰{Óœs 7›|Ò+ÜÐ?Ó@Ð@à47·K±KÖ@¨&�v—|‘|•~Ð@ˆÐ@Ü  §¡¨3¯9©9°c·k±kÀ;ÐPS×P_ÑP_Ðad×ajÑajÓkˆÜ”s˜7“|Ó$ò 	3ˆAØ�K‰K˜‰N×Ò 7¨1¡:§?¡?Ñ2Öð	3àˆ
ùò	 As   »C r.   r/   r0   Útensor_shapec                 óB   — t        ||||«      }|j                  «        |S r   )rL   rT   )r   r.   r/   r0   r�  Údescs         r   Úcreate_make_tensor_descriptorz0InterpreterBuilder.create_make_tensor_descriptorÈ  s"   € ô    e¨W°lÓCˆØ�‰ŒØˆr   rƒ  r#  c                 ó|   — t        |t        «      sJ ‚|j                  |«      \  }}| j                  ||d ||d¬«      S )NF)rC  rL  rM  rA  )r{   rL   rJ   r=  )r   rƒ  r#  rL  rM  rE   rB  s          r   Úcreate_descriptor_loadz)InterpreterBuilder.create_descriptor_loadÓ  sO   € ä˜$Ô 0Ô1Ð1Ð1Ø×.Ñ.¨wÓ7‰
ˆˆdØ×&Ñ& t¨T¸ÈnØ7FÐTYð 'ó [ð 	[r   r%   c                 óT   — |j                  |«      \  }}| j                  |||d d «      S r   r5  )r   rƒ  r%   r#  rE   rB  s         r   Úcreate_descriptor_storez*InterpreterBuilder.create_descriptor_storeÚ  s/   € Ø×.Ñ.¨wÓ7‰
ˆˆdØ×'Ñ'¨¨e°T¸4ÀÓFÐFr   Ú	x_offsetsÚy_offsetc                 ó®  — |j                   j                  j                  }t        |«      }t	        j
                  |j                  j                  d   |j                  d   g|¬«      }d }d }	t        |j                  «      D ]F  \  }
}t        |t        j                  «      |g}| j                  ||||	«      j                  ||
d d …f<   ŒH t        ||«      S )Nr   rQ   r6   )r.   r   r   r‹   r8   Úzerosr   r/   r2   Ú	enumerater   r|   rw   r†  )r   rƒ  r‰  rŠ  r3  r   Únp_dtypeÚresultrL  rM  rª   Úx_offsetr#  s                r   Úcreate_descriptor_gatherz+InterpreterBuilder.create_descriptor_gatherÞ  sÀ   € Ø—	‘	—‘×*Ñ*ˆÜ  Ó'ˆÜ—‘˜9Ÿ>™>×/Ñ/°Ñ2°D×4DÑ4DÀRÑ4HÐIÐQYÔZˆØˆØˆÜ$ Y§^¡^Ó4ò 	l‰KˆAˆxÜ# H¬b¯h©hÓ7¸ÐBˆGØ×6Ñ6°t¸WÀnÐVeÓf×kÑkˆF�1’a�4ŠLð	lô ˜F EÓ*Ð*r   c                 óâ   — t        |j                  «      D ]W  \  }}t        |j                  |   |j                  «      }t        |t        j
                  «      |g}| j                  |||«       ŒY y r   )r�  r   r   r   r|   rw   rˆ  )	r   rƒ  r%   r‰  rŠ  rª   r�  Úslicer#  s	            r   Úcreate_descriptor_scatterz,InterpreterBuilder.create_descriptor_scatteré  s]   € ä$ Y§^¡^Ó4ò 	?‰KˆAˆxÜ  §¡¨A¡°·±Ó<ˆEÜ# H¬b¯h©hÓ7¸ÐBˆGØ×(Ñ(¨¨u°gÕ>ñ	?r   c                 ó¨   — t        |«      }d|j                  v r,t        t        j                  dd|¬«      |j
                  «      S t        d|› �«      ‚)NrX   r   rQ   r6   zunsupported type )r‹   Únamer   r8   rF  rA   Ú	TypeError)r   r3  Únp_types      r   Úget_all_ones_valuez%InterpreterBuilder.get_all_ones_valueð  sI   € Ü Ó%ˆØ�G—L‘LÑ Ü¤§¡¨¨2°WÔ =¸t¿{¹{ÓKÐKäÐ/°¨vÐ6Ó7Ð7r   ©ÚreturnN)·r'   r(   r)   r—   ÚMEM_SEMANTICÚACQUIRErJ  ÚRELEASEÚRELAXEDÚACQUIRE_RELEASErJ  Ú	ATOMIC_OPÚADDÚRMW_OPÚFADDÚMINÚUMINÚMAXÚUMAXÚANDÚORÚXORÚXCHGrQ  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.  r0  r4  r8  r:  rD  rH  r=  rF  rX  Úcreate_si_to_fpÚcreate_ui_to_fpÚcreate_fp_to_siÚcreate_fp_to_uiÚcreate_fp_extÚcreate_fp_truncÚcreate_int_castrÅ   rf  rl  Úcreate_faddÚcreate_fmulÚcreate_fdivÚcreate_fremÚcreate_fsubÚ
create_mulÚcreate_precise_divfÚcreate_sdivÚcreate_udivÚcreate_sremÚcreate_uremÚ
create_addÚ
create_subÚ
create_shlÚcreate_lshrÚcreate_minsiÚcreate_minuiÚcreate_minimumfÚcreate_minnumfÚcreate_maxsiÚcreate_maxuiÚcreate_maximumfÚcreate_maxnumfÚcreate_icmpSLEÚcreate_icmpSLTÚcreate_icmpSGEÚcreate_icmpSGTÚcreate_icmpULEÚcreate_icmpULTÚcreate_icmpUGEÚcreate_icmpUGTÚcreate_icmpEQÚcreate_icmpNEÚcreate_fcmpOLTÚcreate_fcmpOGTÚcreate_fcmpOLEÚcreate_fcmpOGEÚcreate_fcmpOEQÚcreate_fcmpONEÚcreate_fcmpULTÚcreate_fcmpUGTÚcreate_fcmpULEÚcreate_fcmpUGEÚcreate_fcmpUEQÚcreate_fcmpUNEÚ
create_andÚ
create_xorÚ	create_orÚcreate_int_to_ptrÚcreate_ptr_to_intr„  rÎ  rÕ  r×  Úcreate_clampfÚcreate_selectrâ  rä  rè  Ú
create_cosÚ
create_expÚcreate_exp2Úcreate_iabsÚcreate_floorÚcreate_ceilÚ
create_logÚcreate_log2Úcreate_precise_sqrtÚcreate_sqrtÚ
create_sinr  r  Úcreate_reshaper  r  r  r  r$  r*  r3  r6  r9  r<  r?  rB  rD  rG  rO  rT  r^  rd  rs  rw  ry  r{  r~  r€  r   r   rX   r„  rL   r†  rˆ  r‘  r”  r™  r*   r   r   rÍ   rÍ   ÿ   s  „ à×Ñ× Ñ  ,×";Ñ";×"CÑ"CØ×Ñ× Ñ  ,×";Ñ";×"CÑ"CØ×Ñ× Ñ  ,×";Ñ";×"CÑ"CØ×Ñ×(Ñ(¨,×*CÑ*C×*SÑ*Sð	!Ðð 	�‰×Ñ˜<×.Ñ.×2Ñ2Ø�‰×Ñ˜L×/Ñ/×4Ñ4Ø�‰×Ñ˜<×.Ñ.×2Ñ2Ø�‰×Ñ˜L×/Ñ/×4Ñ4Ø�‰×Ñ˜<×.Ñ.×2Ñ2Ø�‰×Ñ˜L×/Ñ/×4Ñ4Ø�‰×Ñ˜<×.Ñ.×2Ñ2Ø�‰×Ñ˜,×-Ñ-×0Ñ0Ø�‰×Ñ˜<×.Ñ.×2Ñ2Ø�‰×Ñ˜L×/Ñ/×4Ñ4ð'Ð#óNò"ò%ò
òòòòòòòòòòòòòòòòò3ò+òHòIòGòKòIòKòIòKòIòMòMòMòLòWò
WòNò
Dò+òDò[ñ P€OÙO€OÙO€OÙO€OÙM€MÙO€OÙZ€Oò3òUòFñ J€KÙN€KÙL€KÙJ€KÙN€KÙM€JÙTÐÙC€KÙC€KáJ€KÙJ€KÙH€JÙM€JÙO€JÙQ€KÙN€LÙN€LÙQ€OÙP€NÙN€LÙN€LÙQ€OÙP€NÙS€NÙM€NÙV€NÙP€NÙS€NÙM€NÙV€NÙP€NÙM€MÙQ€MÙM€NÙP€NÙS€NÙV€NÙN€NÙR€NÙM€NÙP€NÙS€NÙV€NÙN€NÙR€NÙP€JÙP€JÙN€IØ&ÐØ&Ðòdò8ò	JòTñ d€MÙZ€MòFò<ò3ñ >€JÙ=€JÙ?€KÙ>€KÙA€LÙ?€KÙ=€JÙ?€KÙGÐÙ?€KÙ=€Jò3òEñ u€NòLòdòNò^òeò
eòiò]òNòPòTòWòtòlòjòròZòRò0ò'ò*òòYò
ð	àð	ð �LÑ!ð	ð �lÑ#ð		ð
 ˜3‘ió	ð[Ð+;ð [ÀdÈ<ÑFXó [ðGÐ,<ð GÀ\ð GÐ\`ÐamÑ\nó Gð	+Ð-=ð 	+È,ð 	+Ðbnó 	+ð?Ð.>ð ?À|ð ?Ð`lð ?Ø,8ó?ó8r   rÍ   c                 ó0   ‡— |dœˆfd„
}t        | ||«       y )N)Úmemberc           
      óv   •—  | |i |j                  «       D ��ci c]  \  }}|dk7  r||“Œ c}}¤d‰i¤ŽS c c}}w )NrÉ   ©Úitems)rõ  ÚargsÚkwargsÚkÚvÚbuilders        €r   rÓ   z_patch_attr.<locals>.<lambda>ù  sZ   ø€ ¹Àð :hàMSÏ\É\Ë^÷AUÙEIÀQÈØDEÈÂOð BCÀAÁó AUñ:hð `gò:h€ ùóAUs   ™5)Úsetattr)Úobjr–  rõ  rý  Ú
new_members      ` r   Ú_patch_attrr  ø  s   ø€ Ø&,õ i€Jô ˆC��zÕ"r   c                 óš   — t        j                  | «      D ]3  \  }}t        j                  j	                  |«      sŒ&t        | |||«       Œ5 y r   )ÚinspectÚ
getmembersr|   ÚcoreÚ
is_builtinr  )Úpkgrý  r–  rõ  s       r   Ú_patch_builtinr     s@   € Ü×*Ñ*¨3Ó/ò 4‰ˆˆfÜ�7‰7×Ñ˜fÕ%Ü˜˜T 6¨7Õ3ñ4r   c                 óv   ‡— d„ Šd„ }d„ | _         ˆfd„| _        d„ | _        d„ | _        t	        |«      | _        y )Nc                 óf   — | j                   j                  }|j                  dk(  rt        |«      S dS )Nr   T)rÆ   r   Úsizer   )r   r   s     r   Ú	_get_boolz%_patch_lang_tensor.<locals>._get_bool  s,   € Ø�{‰{×Ñˆð "ŸY™Y¨!š^Œt�D‹zÐ5°Ð5r   c                 ó¼  — t        t        j                  | j                  j                  «      | j                  j
                  «      }| j                  j                  «       sJ ‚t        | j                  j                  «      }|d   |d   c|d<   |d<   t        j                  j                  | j
                  |«      }t        j                  j                  ||«      S )NéþÿÿÿrQ   )r   r8   r  rÆ   r   r   r3  Úis_blockÚlistr/   r|   r  rˆ   rÄ   )r   rÆ   r2   Úres_tys       r   Ú_get_transposez*_patch_lang_tensor.<locals>._get_transpose  sž   € ÜœbŸl™l¨4¯;©;×+;Ñ+;Ó<¸d¿k¹k×>OÑ>OÓPˆØ�y‰y×!Ñ!Ô#Ð#Ð#Ü˜4Ÿ9™9Ÿ?™?Ó+ˆØ+6°r©?¸KÈ¹OÐ(ˆ�B‰˜ R™Ü—‘×#Ñ# D§J¡J°Ó<ˆÜ�w‰w�~‰~˜f fÓ-Ð-r   c                 ó@   — t        | j                  j                  «      S r   )rX   rÆ   r   r   s    r   rÓ   z$_patch_lang_tensor.<locals>.<lambda>  s   € ¤C¨¯©×(8Ñ(8Ó$9€ r   c                 ó   •—  ‰| «      S r   r*   )r   r  s    €r   rÓ   z$_patch_lang_tensor.<locals>.<lambda>  s   ø€ ¡9¨T£?€ r   c                 ó@   — t        | j                  j                  «      S r   )ÚreprrÆ   r   r   s    r   rÓ   z$_patch_lang_tensor.<locals>.<lambda>  s   € ¤4¨¯©×(8Ñ(8Ó#9€ r   c                 ó@   — t        | j                  j                  «      S r   )rp   rÆ   r   r   s    r   rÓ   z$_patch_lang_tensor.<locals>.<lambda>  s   € ¤# d§k¡k×&6Ñ&6Ó"7€ r   )Ú	__index__r   Ú__repr__Ú__str__ÚpropertyÚT)rÄ   r  r  s     @r   Ú_patch_lang_tensorr    s8   ø€ ò6ò.ñ :€FÔÛ2€F„OÙ9€F„OÙ7€F„NÜ˜Ó'€F…Hr   c                   ó0   — e Zd Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zy)ÚReduceScanOpInterfacec                 ó    — || _         || _        y r   )r7  Ú
combine_fn)r   r7  r!  s      r   r   zReduceScanOpInterface.__init__  s   € ØˆŒ	Ø$ˆ�r   c                 óH   — |� |t        |«      k\  rt        d|› d|› �«      ‚y y )Nzaxis z out of bounds for shape )r<   rÙ   )r   r/   r7  s      r   Ú
check_axisz ReduceScanOpInterface.check_axis#  s4   € ØÐ ¬¨E«
Ò 2Ü˜u T FÐ*CÀEÀ7ÐKÓLÐLð !3Ðr   c                 óÔ   — |D ]c  }t        |t        j                  j                  «      st	        dt        |«      › �«      ‚| j                  |j                  | j                  «       Œe y )Nzinput must be a tensor, got )	r{   r|   r  rÄ   rÙ   r3  r#  r/   r7  )r   rš   rÚ  s      r   Úcheck_tensorz"ReduceScanOpInterface.check_tensor'  sP   € Øò 	2ˆCÜ˜c¤2§7¡7§>¡>Ô2Ü Ð#?ÄÀSÃ	¸{Ð!KÓLÐLØ�O‰O˜CŸI™I t§y¡yÕ1ñ	2r   c                 óZ  — t        |«      }t        |d«      rG|j                  r;|j                  |«      }t	        j
                  |t        |j                  «      «      }nt        j                  |g|¬«      }|}t        j                  j                  t        ||j                  «      |«      S )Nr/   r6   )r‹   r   r/   r?   r|   rˆ   r  r8   r  r  rÄ   r   rA   )r   rO  r   rŽ  Úret_types        r   Ú	to_tensorzReduceScanOpInterface.to_tensor-  sy   € Ü  Ó'ˆÜ�3˜Ô  S§Y¢YØ—*‘*˜XÓ&ˆCÜ—}‘} U¬D°·±«OÓ<‰Hä—(‘(˜C˜5¨Ô1ˆCØˆHÜ�w‰w�~‰~œl¨3°·±Ó=¸xÓHÐHr   c                 ól   — t        |t        «      s|f}| j                  |«       | j                  |«      S r   )r{   Útupler%  Ú
apply_impl©r   rš   s     r   ÚapplyzReduceScanOpInterface.apply7  s0   € Ü˜%¤Ô'Ø�IˆEØ×Ñ˜%Ô Ø�‰˜uÓ%Ð%r   c                 ó   — t        d«      ‚)Nzapply_impl not implementedrV  r,  s     r   r+  z ReduceScanOpInterface.apply_impl=  s   € Ü!Ð">Ó?Ð?r   N)	r'   r(   r)   r   r#  r%  r(  r-  r+  r*   r   r   r  r    s#   „ ò%òMò2òIò&ó@r   r  c                   ó>   ‡ — e Zd Zˆ fd„Zd„ Zd„ Zdd„Zd„ Zd„ Zˆ xZ	S )Ú	ReduceOpsc                 ó4   •— t         ‰| �  ||«       || _        y r   )Úsuperr   Ú	keep_dims)r   r7  r!  r3  Ú	__class__s       €r   r   zReduceOps.__init__C  s   ø€ Ü‰Ñ˜˜zÔ*Ø"ˆ�r   c                 óô   — g }|D ]e  }|�|j                  |«       Œd}|j                  | j                  |j                  j                  j	                  «       |j
                  «      «       Œg t        |«      |fS )Nr   )Úappendr(  rÆ   r   Úflattenr   r*  )r   rš   r7  rO  r   s        r   ÚunravelzReduceOps.unravelG  sn   € ØˆØò 	SˆDØÐØ—
‘
˜4Õ à�Ø—
‘
˜4Ÿ>™>¨$¯+©+×*:Ñ*:×*BÑ*BÓ*DÀdÇjÁjÓQÕRð	Sô �S‹z˜4ÐÐr   c                 ó$  ‡ ‡‡‡— ‰ j                   }‰ j                  ‰‰ j                   «      \  Š}g }g }‰d   j                  j                  j                  }|d| ||dz   d  z   }‰D ]k  }|j                  |j                  j                  «       |j                  t        j                  ||j                  j                  j                  ¬«      «       Œm t        |d   j                  «      D �]Y  }	t        j                  |	|«      Š‰d| ‰|dz   d  z   Št        ˆˆˆ fd„t        |«      D «       «      }
‰|   dk(  rGt        t        |«      «      D ]/  }|
|   j                  j                  j                  «       ||   ‰<   Œ1 Œ–t        ˆˆˆ fd„t        |«      D «       «      } ‰ j                   j"                  g |¢|
¢­Ž }t%        |t        «      s|fn|}t        t        |«      «      D ][  }t%        ||   t&        j(                  j*                  «      r'||   j                  j                  j                  «       n||   ||   ‰<   Œ] �Œ\ g }t        |«      D ]š  \  }	}‰ j,                  rI|�t        j.                  ||«      }nBt        t        |«      «      D ]  }t        j.                  |d«      }Œ n|€|j                  «       }|j                  ‰ j1                  |‰|	   j                  «      «       Œœ t        |«      dk(  r|d   S t        |«      S )Nr   r   r6   c              3   óh   •K  — | ])  \  }}‰j                  |‰   ‰|   j                  «      –— Œ+ y ­wr   ©r(  r   )Ú.0Úiir  rš   Úinput_indexr   s      €€€r   ú	<genexpr>z+ReduceOps.generic_reduce.<locals>.<genexpr>`  s/   øè ø€ ÒsÑTYÐTVÐXY §¡¨q°©~¸uÀR¹y¿¹× OÑsùó   ƒ/2c              3   óh   •K  — | ])  \  }}‰j                  |‰   ‰|   j                  «      –— Œ+ y ­wr   r;  )r<  ÚoiÚorš   Úoutput_indexr   s      €€€r   r?  z+ReduceOps.generic_reduce.<locals>.<genexpr>f  s/   øè ø€ Ò!wÑW\ÐWYÐ[\ $§.¡.°°<±À%ÈÁ)Ç/Á/×"RÑ!wùr@  )r7  r8  rÆ   r   r/   r6  r8   rŒ  r   r;   r  Úunravel_indexr*  r�  r<   rR   r!  Úfnr{   r|   r  rÄ   r3  r8  r(  )r   rš   Úoriginal_axisr7  Ú
input_dataÚoutput_dataÚinput_shapeÚoutput_shaperÚ  rª   Úinput_tupleÚjÚ	acc_tupleÚcombine_fn_retrO  r   Ú_r>  rD  s   ``               @@r   Úgeneric_reducezReduceOps.generic_reduceQ  sÎ  û€ ØŸ	™	ˆØ—l‘l 5¨$¯)©)Ó4‰ˆˆtØˆ
ØˆØ˜A‘h—o‘o×*Ñ*×0Ñ0ˆØ" 1 TÐ*¨[¸À¹¸Ð-CÑCˆØò 	TˆCØ×Ñ˜cŸj™jŸo™oÔ.Ø×ÑœrŸx™x¨¸C¿J¹J¿O¹O×<QÑ<QÔRÕSð	Tô �z !‘}×)Ñ)Ó*ó 	HˆAä×*Ñ*¨1¨kÓ:ˆKØ& q¨Ð.°¸TÀA¹X¸YÐ1GÑGˆLÜÕsÔ]fÐgqÓ]rÔsÓsˆKØ˜4Ñ  AÒ%äœs ;Ó/Ó0ò U�AØ3>¸q±>×3HÑ3H×3MÑ3M×3RÑ3RÓ3T�K ‘N <Ò0ñUô "Õ!wÔ`iÐjuÓ`vÔ!wÓw�	Ø!3 §¡×!3Ñ!3Ð!M°YÐ!MÀÒ!M�Ü6@ÀÔQVÔ6W˜^Ñ.Ð]k�	Üœs ;Ó/Ó0ò H�AÜV`Ø! !™¤b§g¡g§n¡nôW6°9¸Q±<×3FÑ3F×3KÑ3K×3PÑ3PÔ3RØ;DÀQ¹<ð   ‘N <Ò0òHð	Hð" ˆÜ  Ó-ò 	=‰GˆAˆtØ�~Š~Ø Ð,ÜŸ>™>¨$°Ó5‘Dä"¤3 {Ó#3Ó4ò 7˜Ü!Ÿ~™~¨d°AÓ6™ñ7ð Ð&à—y‘y“{�Ø�J‰J�t—~‘~ d¨E°!©H¯N©NÓ;Õ<ð	=ô ˜S› Qšˆs�1‰vÐ6¬E°#«JÐ6r   c                 ó²  — t        |t        «      r|d   n|}d }d }|rM| j                   ||j                  j                  | j
                  | j                  ¬«      |j                  «      }|rQ| j                   ||j                  j                  | j
                  | j                  ¬«      t        j                  «      }|�|�||fS |�|S |�|S t        d«      ‚)Nr   ©r7  Úkeepdimsz-val_reduce_op and idx_reduce_op are both None)r{   r*  r(  rÆ   r   r7  r3  r   r|   rw   rÙ   )r   rš   Úval_reduce_opÚidx_reduce_oprG  Úidxs         r   Úmin_maxzReduceOps.min_max|  sÂ   € ä& u¬eÔ4��a’¸%ˆØˆØˆÙØ—.‘.¡¨u¯|©|×/@Ñ/@ÀtÇyÁyÐ[_×[iÑ[iÔ!jÐlq×lwÑlwÓxˆCÙØ—.‘.¡¨u¯|©|×/@Ñ/@ÀtÇyÁyÐ[_×[iÑ[iÔ!jÔln×ltÑltÓuˆCØˆ?˜s˜Ø˜�8ˆOØˆ_ØˆJØˆ_ØˆJäÐLÓMÐMr   c                 ó¶   — | j                  t        j                  |j                  j                  | j
                  | j                  ¬«      |j                  «      S )NrS  )r(  r8   ÚsumrÆ   r   r7  r3  r   r,  s     r   rZ  zReduceOps.sumŽ  s<   € Ø�~‰~œbŸf™f U§\¡\×%6Ñ%6¸T¿Y¹YÐQU×Q_ÑQ_Ô`Ðbg×bmÑbmÓnÐnr   c                 ó2  — | j                   t        j                  j                  k(  r3| j	                  |d   t
        j                  t
        j                  ¬«      S | j                   t        j                  j                  k(  r3| j	                  |d   t
        j                  t
        j                  ¬«      S | j                   t        j                  j                  k(  r%| j	                  |d   t
        j                  d ¬«      S | j                   t        j                  j                  k(  r%| j	                  |d   t
        j                  d ¬«      S | j                   t        j                  j                  k(  r| j                  |d   «      S | j                  |«      S )Nr   )rU  rV  )r!  r|   ÚstandardÚ_argmin_combine_tie_break_leftrX  r8   ÚminÚargminÚ_argmax_combine_tie_break_leftr&  ÚargmaxÚ_elementwise_maxÚ_elementwise_minÚ_sum_combinerZ  rQ  r,  s     r   r+  zReduceOps.apply_impl‘  s  € Ø�?‰?œbŸk™k×HÑHÒHØ—<‘<  a¡¼¿¹ÌbÏiÉi�<ÓXÐXØ�_‰_¤§¡× JÑ JÒJØ—<‘<  a¡¼¿¹ÌbÏiÉi�<ÓXÐXØ�_‰_¤§¡× <Ñ <Ò<Ø—<‘<  a¡¼¿¹Èd�<ÓSÐSØ�_‰_¤§¡× <Ñ <Ò<Ø—<‘<  a¡¼¿¹Èd�<ÓSÐSØ�_‰_¤§¡× 8Ñ 8Ò8Ø—8‘8˜E !™HÓ%Ð%ð ×&Ñ& uÓ-Ð-r   r   )
r'   r(   r)   r   r8  rQ  rX  rZ  r+  Ú__classcell__©r4  s   @r   r0  r0  A  s$   ø„ ô#ò ò)7óVNò$oö.r   r0  c                   ó6   ‡ — e Zd Zˆ fd„Zd„ Zd„ Zd„ Zd„ Zˆ xZS )ÚScanOpsc                 ó4   •— t         ‰| �  ||«       || _        y r   )r2  r   Úreverse)r   r7  r!  rj  r4  s       €r   r   zScanOps.__init__£  s   ø€ Ü‰Ñ˜˜zÔ*Øˆ�r   c                 ó¤   — | j                  t        j                  |j                  j                  | j
                  ¬«      |j                  ¬«      gS ©Nr!  r6   )r(  r8   ÚcumsumrÆ   r   r7  r   r,  s     r   rm  zScanOps.cumsum§  s8   € Ø—‘œrŸy™y¨¯©×):Ñ):ÀÇÁÔKÐSX×S^ÑS^�Ó_Ð`Ð`r   c                 ó¤   — | j                  t        j                  |j                  j                  | j
                  ¬«      |j                  ¬«      gS rl  )r(  r8   ÚcumprodrÆ   r   r7  r   r,  s     r   ro  zScanOps.cumprodª  s8   € Ø—‘œrŸz™z¨%¯,©,×*;Ñ*;À$Ç)Á)ÔLÐTY×T_ÑT_�Ó`ÐaÐar   c           	      óò  ‡ ‡‡‡— g }g }‰d   j                   j                  j                  }‰D ]k  }|j                  |j                   j                  «       |j                  t	        j
                  ||j                   j                  j                  ¬«      «       Œm t        |d   j                  «      D �]|  }t	        j                  ||«      Št        ˆˆˆ fd„t        |«      D «       «      }‰‰ j                     dk(  rGt        t        |«      «      D ]/  }||   j                   j                  j                  «       ||   ‰<   Œ1 Œ’t        ˆˆ fd„t        t        ‰«      «      D «       «      Št        ˆˆˆ fd„t        |«      D «       «      }	 ‰ j                  j                   g |	¢|¢­Ž }
t#        |
t        «      s|
fn|
}	t        t        |«      «      D ][  }t#        |	|   t$        j&                  j(                  «      r'|	|   j                   j                  j                  «       n|	|   ||   ‰<   Œ] �Œ g }t        |«      D ]3  \  }}|j                  ‰ j+                  |‰|   j                  «      «       Œ5 |S )Nr   r6   c              3   óh   •K  — | ])  \  }}‰j                  |‰   ‰|   j                  «      –— Œ+ y ­wr   r;  )r<  r=  r  Úindexrš   r   s      €€€r   r?  z'ScanOps.generic_scan.<locals>.<genexpr>¸  s,   øè ø€ ÒfÁuÀrÈ1˜Ÿ™¨¨%©°%¸±)·/±/×BÑfùr@  c              3   óV   •K  — | ]   }|‰j                   k(  r‰|   d z
  n‰|   –— Œ" y­w)r   Nr!  )r<  rª   rr  r   s     €€r   r?  z'ScanOps.generic_scan.<locals>.<genexpr>¾  s-   øè ø€ Ò"kÐTU°1¸¿	¹	²> 5¨¡8¨a¢<ÀuÈQÁxÓ#OÑ"kùs   ƒ&)c              3   óh   •K  — | ])  \  }}‰j                  |‰   ‰|   j                  «      –— Œ+ y ­wr   r;  )r<  rB  rC  rš   Ú
prev_indexr   s      €€€r   r?  z'ScanOps.generic_scan.<locals>.<genexpr>¿  s/   øè ø€ Ò!uÑUZÐUWÐYZ $§.¡.°°:±ÀÀbÁ	ÇÁ×"PÑ!uùr@  )rÆ   r   r/   r6  r8   rŒ  r   r;   r  rE  r*  r�  r7  r<   rR   r!  rF  r{   r|   r  rÄ   r(  )r   rš   rH  rI  r/   rÚ  rª   r   rM  rN  rO  rO  rr  ru  s   ``          @@r   Úgeneric_scanzScanOps.generic_scan­  s!  û€ Øˆ
ØˆØ�a‘—‘×$Ñ$×*Ñ*ˆØò 	MˆCØ×Ñ˜cŸj™jŸo™oÔ.Ø×ÑœrŸx™x¨°S·Z±Z·_±_×5JÑ5JÔKÕLð	Mô �z !‘}×)Ñ)Ó*ó 	HˆAä×$Ñ$ Q¨Ó.ˆEÜÕfÔPYÐZdÓPeÔfÓfˆDØ�T—Y‘YÑ 1Ò$äœs ;Ó/Ó0ò G�AØ,0°©G¯N©N×,?Ñ,?×,DÑ,DÓ,F�K ‘N 5Ò)ñGô #Ô"kÔY^Ô_bÐchÓ_iÓYjÔ"kÓk�
Ü!Õ!uÔ^gÐhsÓ^tÔ!uÓu�	Ø!3 §¡×!3Ñ!3Ð!F°YÐ!FÀÒ!F�Ü6@ÀÔQVÔ6W˜^Ñ.Ð]k�	Üœs ;Ó/Ó0ò H�AÜOYØ! !™¤b§g¡g§n¡nôP6¨I°a©L×,?Ñ,?×,DÑ,D×,IÑ,IÔ,KØ;DÀQ¹<ð   ‘N 5Ò)òHð	Hð" ˆÜ  Ó-ò 	=‰GˆAˆtØ�J‰J�t—~‘~ d¨E°!©H¯N©NÓ;Õ<ð	=àˆ
r   c           	      óò  — g }| j                   rf|D ]`  }|j                  | j                  t        j                  |j
                  j                  | j                  ¬«      |j                  «      «       Œb n|}| j                  t        j                  j                  k(  r| j                  |d   «      }nM| j                  t        j                  j                  k(  r| j                  |d   «      }n| j!                  |«      }| j                   rK|D ]F  }t        j                  |j
                  j                  | j                  ¬«      |j
                  _        ŒH t#        |«      dk(  xr |d   xs t%        |«      S )Nr!  r   r   )rj  r6  r(  r8   ÚfliprÆ   r   r7  r   r!  r|   r\  rd  rm  Ú_prod_combinero  rv  r<   r*  )r   rš   Ú	new_inputrÚ  rO  s        r   r+  zScanOps.apply_implË  s  € Øˆ	Ø�<Š<Øò f�Ø× Ñ  §¡´·±¸¿
¹
¿¹ÈdÏiÉiÔ0XÐZ]×ZcÑZcÓ!dÕeñfð ˆIØ�?‰?œbŸk™k×6Ñ6Ò6Ø—+‘+˜i¨™lÓ+‰CØ�_‰_¤§¡× 9Ñ 9Ò9Ø—,‘,˜y¨™|Ó,‰Cð ×#Ñ# IÓ.ˆCØ�<Š<Øò K�Ü"$§'¡'¨#¯*©*¯/©/ÀÇ	Á	Ô"J�—
‘
•ðKä�3‹x˜1‰}Ò'  Q¡Ò5¬5°«:Ð5r   )	r'   r(   r)   r   rm  ro  rv  r+  re  rf  s   @r   rh  rh  ¡  s   ø„ ôòaòbòö<6r   rh  c                  ó”   — dd„} dd„}| t         _        |t         _        | t         j                  _        |t         j                  _        y )Nc                 ó:   — t        |||«      j                  | «      S r   )r0  r-  )rš   r7  r!  r3  rú  s        r   Ú_new_reducez'_patch_reduce_scan.<locals>._new_reduceã  s   € Ü˜˜z¨9Ó5×;Ñ;¸EÓBÐBr   c                 ó:   — t        |||«      j                  | «      S r   )rh  r-  )rš   r7  r!  rj  rú  s        r   Ú	_new_scanz%_patch_reduce_scan.<locals>._new_scanæ  s   € Ü�t˜Z¨Ó1×7Ñ7¸Ó>Ð>r   )F)r|   ÚreduceÚassociative_scanr  )r}  r  s     r   Ú_patch_reduce_scanr‚  ß  s5   € óCó?ð „B„IØ#„BÔØ „B‡G�G„NØ(„B‡G�GÕr   c                 ó  — d„ }d	d„}d
d„}d„ }|| _         || _        || _        t        | _        || j
                  _        t        |d¬«      | _        t        |d¬«      | _	        t        |d¬«      | _
        t        «        y )Nc                 ó>  — | j                   dk(  r|j                  «       S | j                   dk(  r|j                  «       S | j                   dk(  r|j                  «       S | j                   dk(  r|j	                  «       S | j                   dk(  r|j                  «       S | j                   dk(  r|j                  «       S | j                   dk(  r|j                  «       S | j                   dk(  r|j                  «       S | j                   d	k(  r|j                  «       S | j                   d
k(  r|j                  «       S | j                   dk(  r|j                  «       S | j                   dk(  r|j                  «       S | j                   dk(  r|j                  «       S | j                   dk(  r|j                  «       S | j                   dk(  r|j                  «       S | j                   dk(  r|j!                  «       S | j                   dk(  r|j#                  «       S t%        d| › d�«      ‚)NÚvoidr~   rs   rr   ru   rt   rw   rv   rx   r@   r`   rb   rd   Úfp16Úbf16Úfp32Úfp64zfail to convert z to ir type)r–  Úget_void_tyÚget_int1_tyrï   rñ   rô   rö   rù   rû   rý   rÿ   r	  r  r  rå   rè   rë   rí   rÙ   )r   rý  s     r   Ú
_new_to_irz$_patch_lang_core.<locals>._new_to_irñ  sâ  € à�9‰9˜ÒØ×&Ñ&Ó(Ð(Ø�Y‰Y˜&Ò Ø×&Ñ&Ó(Ð(Ø�Y‰Y˜&Ò Ø×&Ñ&Ó(Ð(Ø�Y‰Y˜'Ò!Ø×'Ñ'Ó)Ð)Ø�Y‰Y˜'Ò!Ø×'Ñ'Ó)Ð)Ø�Y‰Y˜(Ò"Ø×(Ñ(Ó*Ð*Ø�Y‰Y˜'Ò!Ø×'Ñ'Ó)Ð)Ø�Y‰Y˜(Ò"Ø×(Ñ(Ó*Ð*Ø�Y‰Y˜'Ò!Ø×'Ñ'Ó)Ð)Ø�Y‰Y˜(Ò"Ø×(Ñ(Ó*Ð*Ø�Y‰Y˜'Ò!Ø×'Ñ'Ó)Ð)Ø�Y‰Y˜)Ò#Ø×)Ñ)Ó+Ð+Ø�Y‰Y˜*Ò$Ø×*Ñ*Ó,Ð,Ø�Y‰Y˜&Ò Ø×&Ñ&Ó(Ð(Ø�Y‰Y˜&Ò Ø×&Ñ&Ó(Ð(Ø�Y‰Y˜&Ò Ø×'Ñ'Ó)Ð)Ø�Y‰Y˜&Ò Ø×(Ñ(Ó*Ð*ÜÐ+¨D¨6°Ð=Ó>Ð>r   c                 ó:   — |€d}|€d| }}n| |}}t        |||«      S )Nr   r   )r;   )Úarg1Úarg2Ústeprú  r  Úends         r   Ú
_new_rangez$_patch_lang_core.<locals>._new_range  s2   € Øˆ<ØˆDØˆ<Ø˜D�3‰Eà˜t�3ˆEÜ�U˜C Ó&Ð&r   c                 ó   — | sJ |«       ‚y r   r*   )rà  rr  s     r   Ú_new_static_assertz,_patch_lang_core.<locals>._new_static_assert"  s   € ÙÐ�SÓ‰tr   c                 ó’  — t        | t        j                  «      s| S t        |t        t        f«      s|gn|}|D �cg c]*  }t        |t        j
                  «      r|j                  n|‘Œ, }}t        |«      t        dt        | j                  «      «      k7  rt        d|› �«      ‚| j                  j                  ||«       | S c c}w )Nr   z$len(values) != len(input.shape) for )r{   r|   rÄ   r  r*  Ú	constexprr%   r<   r&  r/   rÙ   rÆ   r&   )rš   rb  r–  rü  s       r   Ú	_set_attrz#_patch_lang_core.<locals>._set_attr%  s¢   € ä˜%¤§¡Ô+ØˆLä!+¨F´T¼5°MÔ!B�&‘ÈˆØIOÖPÀAœZ¨¬2¯<©<Ô8�!—'’'¸aÑ?ÐPˆÐPÜˆv‹;œ#˜a¤ U§[¡[Ó!1Ó2Ò2ÜÐCÀDÀ6ÐJÓKÐKØ�‰×Ñ˜d FÔ+Øˆùò	 Qs   ¼/Cztt.divisibility)r–  ztt.contiguityztt.constancy)NN)Ú )r;   Ústatic_rangeÚstatic_assertrn  Ústatic_printr   Úto_irr   Úmultiple_ofÚmax_contiguousÚmax_constancyr‚  )ÚlangrŒ  r’  r”  r—  s        r   Ú_patch_lang_corer¡  ï  sx   € ò$?óP'óò
ð €D„JØ"€DÔØ+€DÔÜ€DÔØ!€D‡J�JÔÜ˜yÐ/@ÔA€DÔÜ! )°/ÔB€DÔÜ  °Ô@€DÔäÕr   c                 ó(  — | j                   j                  «       D ��cg c]6  \  }}t        j                  |«      sŒ|t        t        j
                  fv sŒ5|‘Œ8 }}}t        |«      dk\  sJ d«       ‚|D ]o  }t        |t        «       t        |j                  t        «       |t        k(  rt        |j                  t        «       t        |j                  «       t        |«       Œq t        t        j
                  j                  t        «       y c c}}w )Nr   z:triton.language must be visible from within jit'd function)Ú__globals__rø  r  Úismoduler|   r  r<   r  Úinterpreter_builderrÄ   r¶   r  r¡  Ú$_experimental_tensor_descriptor_base)rF  rP  r%   Úlangsr   s        r   Ú_patch_langr¨  =  sÎ   € Ø#%§>¡>×#7Ñ#7Ó#9×p‘x�q˜%¼W×=MÑ=MÈeÕ=TÐY^ÔceÔgi×gnÑgnÐboÒYoŠUÐp€EÑpÜˆu‹:˜Š?ÐXÐXÓXˆ?Øò ˆÜ�tÔ0Ô1Ü�t—{‘{Ô$7Ô8Ø”2Š:Ü˜4Ÿ9™9Ô&9Ô:Ü˜4Ÿ;™;Ô'Ü˜Õðô ”2—7‘7×?Ñ?ÔATÕUùó qs   žD¼DÁDc                 óX   — t        | d«      r t        | «      |Ž S  t        | «      |«      S )NÚ_fields)r   r3  )rÚ  Úcontentss     r   Ú_tuple_creater¬  J  s.   € ô $+¨3°	Ô#:ˆ9Œ4�‹9�hÐÐSÀ	ÄÀSÃ	È(Ó@SÐSr   c                 óÐ  — t        | t        «      �rt        j                  t        j
                  j                  j                  | «      «      }t        j                  }d| cxk  rdk  rn nt        j                  }nkd| cxk  rdk  rn nt        j                  }nLd| cxk  rdk  rn nt        j                  }n-d| cxk  rdk  rn nt        j                  }nt        d| › �«      ‚t        t        j                  | g|¬«      |«      }t        j                   ||«      S t#        | d	«      r�t        j                  t        j
                  j                  j                  | «      «      }t        t        j                  | j%                  «       gt        j                  ¬«      |«      }t        j                   ||«      S t        | t&        «      rt)        | t+        t,        | «      «      S | S )
Ni   €l        l        lûÿÿÿ         l            l            zUnsupported integer value r6   Údata_ptr)r{   rX   r|   Ú	str_to_tyÚtritonÚruntimeÚjitÚmangle_typer8   rw   rv   rx   r@   rÙ   r   r  rÄ   r   r®  r*  r¬  ÚmapÚ_implicit_cvt)rÚ  Útyr   rÆ   s       r   rµ  rµ  T  sR  € Ü�#”sÕÜ�\‰\œ&Ÿ.™.×,Ñ,×8Ñ8¸Ó=Ó>ˆÜ—‘ˆØ�SÔ ˜5Õ Ü—H‘H‰EØ�cÔ!˜EÕ!Ü—I‘I‰EØ�sÔ"˜UÕ"Ü—H‘H‰EØ�cÔ!˜EÕ!Ü—I‘I‰EäÐ9¸#¸Ð?Ó@Ð@ÜœbŸh™h¨ u°EÔ:¸BÓ?ˆÜ�y‰y˜ Ó$Ð$Üˆs�JÔÜ�\‰\œ&Ÿ.™.×,Ñ,×8Ñ8¸Ó=Ó>ˆÜœbŸh™h¨¯©«Ð'7¼r¿y¹yÔIÈ2ÓNˆÜ�y‰y˜ Ó$Ð$Ü	�CœÔ	Ü˜S¤#¤m°SÓ"9Ó:Ð:Ø€Jr   c                 óz   — t        | t        j                  j                  j                  «      r| j
                  S | S r   )r{   r°  r±  r²  ÚTensorWrapperr.   )Úts    r   Ú_unwrap_tensorrº  p  s*   € Ü�!”V—^‘^×'Ñ'×5Ñ5Ô6Ø�v‰vˆØ€Hr   c                 óÊ   — t        |t        j                  j                  j                  «      r4t        j                  j                  j	                  | |j
                  «      S | S r   )r{   r°  r±  r²  r¸  r   )r¹  Úoriginal_tensors     r   Ú_rewrap_tensorr½  v  sD   € Ü�/¤6§>¡>×#5Ñ#5×#CÑ#CÔDÜ�~‰~×!Ñ!×/Ñ/°°?×3HÑ3HÓIÐIØ€Hr   c                   ó$   — e Zd Zd„ Zd„ Zd„ Zd„ Zy)ÚGridExecutorc                 ó
  — ddl m} || _        || _        || _        |j
                  j                  «       D ��ci c]  \  }}| ||«      “Œ }}}|D �cg c]  }|j                  |«      dk(  sŒ|‘Œ c}| _        y c c}}w c c}w )Nr   )Ú_normalize_tyr–  )	r²  rÁ  rF  Ú	arg_namesÚgridro   rø  ÚgetÚ
constexprs)r   rF  rÂ  rÃ  rÁ  r–  r¶  ro   s           r   r   zGridExecutor.__init__~  sx   € Ý&àˆŒØ"ˆŒØˆŒ	ØCE×CUÑCU×C[ÑC[ÓC]×^±x°t¸R˜4¡¨rÓ!2Ñ2Ð^ˆÑ^Ø,5Öb D¸×9LÑ9LÈTÓ9RÐVaÓ9aš4Òbˆ�ùó _ùÚbs   ¹A:ÁB Á-B c                 óœ   ‡‡	— i Š	ˆˆ	fd„Š|D �cg c]
  } ‰|«      ‘Œ }}i }|j                  «       D ]  \  }} ‰|«      ||<   Œ ||fS c c}w )Nc                 ó&  •— t        | t        «      rt        | t        ‰| «      «      S t	        | d«      s| S t        | «      }|j                  «       j                  «       ‰vr1|j                  «       }|j                  «       ‰|j                  «       <   ‰|j                  «       j                  «          }|j                  dd¬«      }|j                  ||j                  «       |j                  «       |j                  «       «       t        || ¬«      }|S )Nr®  r   Úcpu)Údevice)r¼  )r{   r*  r¬  r´  r   rº  Úuntyped_storager®  rÈ  Ú	new_emptyÚset_Ústorage_offsetr  rS   r½  )rÚ  Úunwrapped_argÚstorageÚcpu_argÚ_to_cpuÚstoragess       €€r   rÑ  z,GridExecutor._init_args_hst.<locals>._to_cpuŠ  së   ø€ Ü˜#œuÔ%Ü$ S¬#¨g°sÓ*;Ó<Ð<Ü˜S *Ô-Ø�
ä*¨3Ó/ˆMØ×,Ñ,Ó.×7Ñ7Ó9ÀÑIØ'×7Ñ7Ó9�Ø/6¯{©{«}�˜×)Ñ)Ó+Ñ,à˜}×<Ñ<Ó>×GÑGÓIÑJˆGØ#×-Ñ-¨a¸Ð-Ó>ˆGØ�L‰L˜ -×">Ñ">Ó"@À-×BTÑBTÓBVÐXe×XlÑXlÓXnÔoÜ$ W¸cÔBˆGØˆNr   r÷  )
r   Úargs_devrú  rÚ  Úargs_hstÚ
kwargs_hstr$   r%   rÑ  rÒ  s
           @@r   Ú_init_args_hstzGridExecutor._init_args_hst‡  sc   ù€ Øˆõ	ð" -5Ö5 S‘G˜C•LÐ5ˆÐ5ð ˆ
Ø Ÿ,™,›.ò 	-‰JˆC�Ù% e›nˆJ�sŠOð	-à˜Ð#Ð#ùò 6s   �A	c                 óð   ‡
‡— i Šˆ
ˆfd„Š
t        ||«      D ]  \  }} ‰
||«       Œ |j                  «       D ]  \  }}||   }	 ‰
||	«       Œ ‰j                  «       D ]  \  }}|j                  |«       Œ y )Nc                 ó&  •— t        | d«      rVt        | «      t        |«      }} | j                  «       |j                  «       f‰| j                  «       j                  «       <   y t	        | t
        «      rt        | |«      D ]  \  } } ‰| |«       Œ y y )Nr®  )r   rº  rÊ  r®  r{   r*  Úzip)Úarg_devÚarg_hstÚ	_from_cpurÒ  s     €€r   rÜ  z1GridExecutor._restore_args_dev.<locals>._from_cpu¦  sˆ   ø€ Ü�w 
Ô+ä#1°'Ó#:¼NÈ7Ó<S˜�ØBI×BYÑBYÓB[Ð]d×]tÑ]tÓ]vÐAw�˜×0Ñ0Ó2×;Ñ;Ó=Ò>Ü˜G¤UÔ+Ü*-¨g°wÓ*?ò 0Ñ&�W˜gÙ˜g wÕ/ñ0ð ,r   )rÙ  rø  rb  Úcopy_)r   rÓ  rÔ  rú  rÕ  rÚ  rÛ  r$   Ú	kwarg_devÚ	kwarg_hstrÜ  rÒ  s             @@r   Ú_restore_args_devzGridExecutor._restore_args_dev£  s‹   ù€ Øˆõ	0ô !$ H¨hÓ 7ò 	(ÑˆG�WÙ�g˜wÕ'ð	(ð %Ÿl™l›nò 	,‰NˆC�Ø" 3™ˆIÙ�i Õ+ð	,ð #+§/¡/Ó"3ò 	#ÑˆW�gØ�M‰M˜'Õ"ñ	#r   c                 ó  — |j                  dd«      ry t        j                  | j                  «      }|j	                  «       D ��ci c]  \  }}||j
                  v sŒ||“Œ }}}| j                  ||«      \  }}t        | j                  «       t        j                  | j                  g|¢­i |¤Ž}|j	                  «       D �	�
ci c]!  \  }	}
|	|	| j                  v r|
n
t        |
«      “Œ# }}	}
t        | j                  «      r| j                  |«      n| j                  }t        |«      dk  sJ d«       ‚|ddt        |«      z
  z  z   }t        j                  |Ž  	 t!        |d   «      D ]Q  }t!        |d   «      D ]>  }t!        |d   «      D ]+  }t        j#                  |||«        | j                  d	i |¤Ž Œ- Œ@ ŒS 	 | j+                  ||||«       y c c}}w c c}
}	w # t$        $ r}t'        t)        |«      «      |‚d }~ww xY w)
NÚwarmupFé   z#grid must have at most 3 dimensions)r   r   r   r	   r*   )Úpopr  ÚgetfullargspecrF  rø  rù  rÖ  r¨  ÚgetcallargsrÅ  rµ  ÚcallablerÃ  r<   r¥  rã   r;   rÞ   Ú	Exceptionr   r  rà  )r   rÓ  rú  Úargspecrû  rü  rÔ  rÕ  rù  r–  rÚ  rÃ  r¹   rÜ   rÝ   Úes                   r   Ú__call__zGridExecutor.__call__º  sÛ  € Ø�:‰:�h Ô&Øô ×(Ñ(¨¯©Ó1ˆØ#)§<¡<£>×G™4˜1˜a°Q¸'¿,¹,Ò5F�!�Q‘$ÐGˆÑGà#×2Ñ2°8¸VÓDÑˆ�*ä�D—G‘GÔô ×"Ñ" 4§7¡7ÐD¨XÒD¸ÑDˆØ^b×^hÑ^hÓ^j×kÑQZÐQUÐWZ�˜T T§_¡_Ñ4‘c¼-ÈÓ:LÑLÐkˆÑkä"*¨4¯9©9Ô"5ˆt�y‰y˜Œ¸4¿9¹9ˆÜ�4‹y˜AŠ~ÐDÐDÓDˆ~Ø�e˜q¤3 t£9™}Ñ-Ñ-ˆÜ×(Ñ(¨$Ñ/ð	3Ü˜4 ™7“^ò (�Ü˜t A™w›ò (�AÜ" 4¨¡7›^ò (˜Ü+×8Ñ8¸¸A¸qÔAØ˜Ÿ™™ $›ñ(ñ(ñ(ð 	×Ñ˜x¨°6¸:ÕFùó/ Hùó løô ò 	3Ü"¤4¨£7Ó+°Ð2ûð	3ús+   ÁGÁGÃ&GÅA"G  Ç 	HÇ)G>Ç>HN)r'   r(   r)   r   rÖ  rà  rë  r*   r   r   r¿  r¿  |  s   „ òcò$ò8#ó.Gr   r¿  c                   ó   — e Zd Zd„ Zy)ÚASTTransformerc                 ó�  — g }|j                   D ]  }|| j                  |«      gz  }Œ t        |«      dkD  rt        d«      ‚t	        j
                  t	        j                  t	        j                  t	        j                  t	        j                  dt	        j                  «       ¬«      dt	        j                  «       ¬«      dt	        j                  «       ¬«      dt	        j                  «       ¬«      |j                  t	        j                  d	t	        j                  «       ¬«      t	        j                  d
¬«      gg ¬«      |_	        |S )Nr   z&Multiple assignments are not supportedr°  )ÚidÚctxÚlanguage)r%   r   rð  Úsemanticr(  r¥  F)r%   )Úfuncrù  Úkeywords)ÚtargetsÚvisitr<   rÙ   ÚastÚCallÚ	AttributeÚNameÚLoadr%   ÚConstant)r   ÚnodeÚnamesÚtargets       r   Úvisit_AssignzASTTransformer.visit_AssignÝ  sì   € ØˆØ—l‘lò 	*ˆFØ�d—j‘j Ó(Ð)Ñ)‰Eð	*äˆu‹:˜Š>ÜÐEÓFÐFô —X‘XÜ—‘Ü—m‘mÜŸ-™-¬c¯h©h¸(ÌÏÉË
Ô.SÐZdÔjm×jrÑjrÓjtÔuØ#¬¯©«ô5à;FÌCÏHÉHËJôXð —*‘*œcŸh™hÐ*?ÄSÇXÁXÃZÔPÜ—,‘, UÔ+ð-Ø79ô;ˆŒ
ð ˆr   N)r'   r(   r)   r   r*   r   r   rí  rí  Û  s   „ ór   rí  c                   óD   — e Zd Z e«       Zd„ Zd„ Zd„ Zd„ Zd„ Z	d„ Z
d„ Zy)	ÚFunctionRewriterc                 ó<   — || _         || _        d| _        d| _        y )Nr˜  r   )rF  rú  ÚfilenameÚdef_file_lineno)r   rF  rú  s      r   r   zFunctionRewriter.__init__ò  s   € ØˆŒØˆŒØˆŒà$%ˆÕr   c                 óJ  — 	 t        j                  | j                  «      \  }}| j	                  «       \  | _        | _        | j                  |«      | _        | j                  |«      }| j                  |«      }| j                  |«      S # t        $ r | j                  cY S w xY wr   )r  ÚgetsourcelinesrF  rè  Ú_get_jit_fn_file_liner  r  Ú	_find_defÚ
def_linenoÚ_prepare_sourceÚ_transform_astÚ_compile_and_exec)r   ÚlinesrP  rT  Útransformed_asts        r   Úrewrite_astzFunctionRewriter.rewrite_astù  s”   € ð	Ü×-Ñ-¨d¯g©gÓ6‰HˆE�1ð /3×.HÑ.HÓ.JÑ+ˆŒ�tÔ+ØŸ.™.¨Ó/ˆŒØ×"Ñ" 5Ó)ˆØ×-Ñ-¨cÓ2ˆØ×%Ñ% oÓ6Ð6øô ò 	Ø—7‘7ŠNð	ús   ‚"B
 Â
B"Â!B"c                 óB   — ddl m}m}  | || j                  «      «      S )Nr   )Úget_jit_fn_file_lineÚJITFunction)r²  r  r  rF  )r   r  r  s      r   r  z&FunctionRewriter._get_jit_fn_file_line  s   € ß:Ù#¡K°·±Ó$8Ó9Ð9r   c                 óz   — d}t        |«      D ]*  \  }}|j                  «       j                  d«      sŒ&|dz   }Œ, |S )Nr   zdef r   )r�  ÚstripÚ
startswith)r   r  r
  rª   Úlines        r   r	  zFunctionRewriter._find_def  sD   € Øˆ
ä  Ó'ò 	#‰GˆAˆtØ�z‰z‹|×&Ñ& vÕ.Ø ™U‘
ð	#ð Ðr   c                 ór   — || j                   dz
  d  }dj                  |«      }t        j                  |«      S )Nr   r˜  )r
  ÚjoinÚtextwrapÚdedent)r   r  rT  s      r   r  z FunctionRewriter._prepare_source  s4   € Ø�d—o‘o¨Ñ)Ð*Ð+ˆØ�g‰g�e‹nˆÜ�‰˜sÓ#Ð#r   c                 óÚ   — t        j                  |«      }| j                  j                  |«      }t        j                  |«       | j
                  dz
  }t        j                  ||«       |S r  )r÷  ÚparseÚast_transformerrö  Úfix_missing_locationsr  Úincrement_lineno)r   rT  Ú
parsed_astr  Ú
inc_linenos        r   r  zFunctionRewriter._transform_ast  s[   € ô —Y‘Y˜s“^ˆ
Ø×.Ñ.×4Ñ4°ZÓ@ˆÜ×!Ñ! /Ô2Ø×)Ñ)¨AÑ-ˆ
Ü×Ñ˜_¨jÔ9ØÐr   c                 ó  — t        || j                  d¬«      }i | j                  ¥}| j                  j                  }t        «       j                  «       D ]  \  }}||vsŒ|||<   Œ t        |||«       || j                  j                     S )NÚexec)r  Úmode)	Úcompiler  rú  rF  r£  Úglobalsrø  r$  r'   )r   r  Úcompiled_codeÚlocal_namespaceÚ
fn_globalsr$   r%   s          r   r  z"FunctionRewriter._compile_and_exec)  s€   € Ü ¸$¿-¹-ÈfÔUˆØ)˜TŸ[™[˜/ˆØ—W‘W×(Ñ(ˆ
Ü!›)Ÿ/™/Ó+ò 	(‰JˆC�Ø˜*Ò$Ø"'�
˜3’ð	(ô 	ˆ]˜J¨Ô8Ø˜tŸw™w×/Ñ/Ñ0Ð0r   N)r'   r(   r)   rí  r  r   r  r  r	  r  r  r  r*   r   r   r  r  ï  s-   „ Ù$Ó&€Oò&ò7ò(:òò$ò
	ó1r   r  c                   ó:   — e Zd Zi Zdd„Zd„ Zed„ «       Z d„ Zd„ Zy)ÚInterpretedFunctionNc                 óð   ‡ — |‰ _         t        |fi |¤Ž‰ _        ˆ fd„}|‰ _        t	        j
                  |«      }|j                  j                  «       D �cg c]  }|j                  ‘Œ c}‰ _	        y c c}w )Nc                  óh   •— |d   }‰j                  «       } t        |‰j                  |«      | i |¤ŽS )NrÃ  ©Úrewriter¿  rÂ  )rù  rú  rÃ  rF  r   s       €r   Úrunz)InterpretedFunction.__init__.<locals>.run<  s6   ø€ Ø˜&‘>ˆDØ—‘“ˆBØ9”<  D§N¡N°DÓ9¸4ÐJÀ6ÑJÐJr   )
rF  r  Úrewriterr1  r  Ú	signatureÚ
parametersrb  r–  rÂ  )r   rF  rú  r1  r3  rü  s   `     r   r   zInterpretedFunction.__init__8  sb   ø€ ØˆŒÜ(¨Ñ6¨vÑ6ˆŒô	Kð
 ˆŒÜ×%Ñ% bÓ)ˆ	Ø*3×*>Ñ*>×*EÑ*EÓ*GÖH Q˜!Ÿ&›&ÒHˆ�ùÒHs   ÁA3c                 óÆ   — | j                   | j                  vr1| j                  j                  «       | j                  | j                   <   | j                  | j                      S r   )rF  Úrewritten_fnr2  r  r   s    r   r0  zInterpretedFunction.rewriteE  sJ   € Ø�7‰7˜$×+Ñ+Ñ+Ø)-¯©×)BÑ)BÓ)DˆD×Ñ˜dŸg™gÑ&Ø× Ñ  §¡Ñ)Ð)r   c                 ó.   — | j                   j                  S r   )rF  r'   r   s    r   r'   zInterpretedFunction.__name__J  s   € à�w‰w×ÑÐr   c                 óP   — | j                  «       }t        || j                  |«      S r   r/  )r   rÃ  rF  s      r   Ú__getitem__zInterpretedFunction.__getitem__N  s    € Ø�\‰\‹^ˆÜ˜B §¡°Ó5Ð5r   c                 óª   — t        | j                  «       | j                  «       }	  ||i |¤ŽS # t        $ r}t	        t        |«      «      |‚d }~ww xY wr   )r¨  rF  r0  rè  r   r  )r   rù  rú  rF  rê  s        r   rë  zInterpretedFunction.__call__R  sO   € ä�D—G‘GÔØ�\‰\‹^ˆð	3Ù�tÐ&˜vÑ&Ð&øÜò 	3Ü"¤4¨£7Ó+°Ð2ûð	3ús   §/ ¯	A¸AÁArš  )	r'   r(   r)   r6  r   r0  r  r9  rë  r*   r   r   r,  r,  4  s0   „ à€LóIò*ð
 ñ ó ð ò6ó3r   r,  )=r÷  r  r  Útypingr   r   r¶   Únumpyr8   r°  Útriton.languagerñ  r|   Údataclassesr   Úerrorsr   Ú	functoolsr   Ú_C.libtritonr
   rJ  r   r—   r   r,   rL   r[   ry   r‹   r´   rº   r¿   Ú	vectorizer€   r  r�   r  r@   rÐ  rÂ   rÍ   r  r  r  r  r0  rh  r‚  r¡  r¨  r¬  rµ  r¥  rº  r½  r¿  ÚNodeTransformerrí  r  r,  r*   r   r   ú<module>rD     s‹  ðÛ 
Û Û ß ã Û ã Ý Ý !Ý $Ý Ý 6Ý $÷ñ ÷:ñ ÷4"ñ "ñJ �$Ô÷
&ð 
&ó ð
&ò	òò@='ò@ò
#ð ˆb�l‰l˜4¨¯©¨Ô5€Øˆb�l‰l˜4¨¯©¨Ô5€Ø�—‘˜Z°·±°Ô<€÷gñ g÷v8ñ v8òr#ò4ò(÷.!@ñ !@ôH].Ð%ô ].ô@;6Ð#ô ;6ò|)ò Kò\
VòTòñ2 )Ó*Ð òò÷\Gñ \Gô~�S×(Ñ(ô ÷(B1ñ B1÷J%3ò %3r   