Ë
    ÏÍ:j‹í  ã                  ó.  — d dl mZ d dlZd dlZd dlZd dlmZmZmZ d dl	Z	d dl
Zd dlZd dlmZ d dlZd dlmZ d dlmZ d dlmZ ddlmZ d d	lmZ d
dlmZ d
dlmZ e 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jP                  e&ejR                  g¬«      Z* ejP                  e&ejV                  g¬«      Z, ejP                  e'ejZ                  g¬«      Z. G d„ d«      Z/ G d„ d «      Z0d!„ Z1d"„ Z2d#„ Z3 G d$„ d%«      Z4 G d&„ d'e4«      Z5 G d(„ d)e4«      Z6d*„ Z7d+„ Z8d,„ Z9d-„ Z:d.„ Z; e0«       Z< ee<«      Z=d/„ Z>d0„ Z? G d1„ d2«      Z@ G d3„ d4ej‚                  «      ZB G d5„ d6«      ZC G d7„ d8«      ZDy)9é    )ÚannotationsN)ÚTupleÚListÚDict)Ú	dataclass)ÚTritonSemantic)ÚTensorDescriptoré   )ÚInterpreterError)Úpartialé   )Úinterpreter)Úirc                  ón   — e Zd ZU dZded<   ded<    ej                  e¬«      Zded<   d	„ Z	d
„ Z
d„ Zd„ Zy)ÚTensorHandlezô
        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
    znp.arrayÚdataztl.dtypeÚdtype)Údefault_factoryr   Úattrc                óH   — t        | j                  j                  «       «      S ©N)Úboolr   Úall©Úselfs    úo/home/mcse/projects/srt_converter/srt-converter-venv/lib/python3.12/site-packages/triton/runtime/interpreter.pyÚ__bool__zTensorHandle.__bool__#   s   € Ü�D—I‘I—M‘M“OÓ$Ð$ó    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__Ú__doc__Ú__annotations__ÚdataclassesÚfieldÚdictr   r   r"   r%   r)   © r   r   r   r      s>   … ñð ƒNØƒOØ"�×"Ñ"°4Ô8€Dˆ$Ó8ò%òò:ór   r   c                  ó   — e Zd Zd„ Zd„ Zy)ÚBlockPointerHandlec                óX   — || _         || _        || _        || _        || _        || _        y r   )ÚbaseÚshapeÚstridesÚoffsetsÚblock_shapeÚorder)r   r6   r7   r8   r9   r:   r;   s          r   Ú__init__zBlockPointerHandle.__init__5   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   )r6   r"   Úprimitive_bitwidthÚnpÚbroadcast_tor   r:   Úonesr   ÚrangeÚlenr9   ÚarangeÚreshaper8   ÚastypeÚuint64r7   r   r   Úscalar)	r   Úboundary_checkÚdtype_ttÚn_bytesÚptrsÚmasksÚdimÚ
bcast_dimsÚoffs	            r   Úmaterialize_pointersz'BlockPointerHandle.materialize_pointers=   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<   rS   r2   r   r   r4   r4   3   s   „ òór   r4   c                  ó&   — e Zd Z	 	 dd„Zd„ Zdd„Zy)ÚTensorDescHandlec                ó\   — || _         t        |«      | _        || _        || _        || _        y r   )r6   rE   Úndimr7   r8   r:   )r   r6   r7   r8   r:   s        r   r<   zTensorDescHandle.__init__O   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)r6   r   ÚitemrE   r8   rW   r:   )r   Ústrides     r   ÚvalidatezTensorDescHandle.validateW   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   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  }ŒÔ |j                  t        j"                  k(  sJ ‚t'        || j                  j                  j(                  «      }||fS )Nr>   rZ   rY   r   z*block offset start must be 16-byte alignedr?   r
   )rE   rW   r6   r   r    r@   r   rA   rB   r:   rC   r   rD   rF   rG   r8   rH   rI   r7   r   rJ   )	r   r9   Ú	scalar_tyÚitemsizerN   rO   rP   rQ   rR   s	            r   rS   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ð �z‰zœRŸY™YÒ&Ð&Ð&Ü˜D $§)¡)§/¡/×"8Ñ"8Ó9ˆØ�Uˆ{Ðr   N)r6   r   r7   úList[TensorHandle]r8   ra   r:   ú	List[int])r9   ra   )r*   r+   r,   r<   r]   rS   r2   r   r   rU   rU   M   s   „ ð'Ø'ó'òPôr   rU   T)Úfrozenc                  óš   — e Zd ZU dZded<   dZded<   dZded<   dZd	ed
<   dZded<   dZ	ded<   dZ
d	ed<   dZded<   dZded<   dZd	ed<   y)ÚInterpreterOptionsNr1   Úextern_libsFr   ÚdebugTÚsanitize_overflowÚstrÚarch)Úfp8e5Úfp8e5b16Úfp8e4nvÚfp8e4b8Úfp8e4b15z
Tuple[str]Úsupported_fp8_dtypesr2   Ú!deprecated_fp8_dot_operand_dtypesÚtf32Údefault_dot_input_precision)rr   Útf32x3ÚieeeÚallowed_dot_input_precisionsr   ÚintÚmax_num_imprecise_acc_defaultr   Úbackend_name)r*   r+   r,   rf   r.   rg   rh   rj   rp   rq   rs   rv   rx   ry   r2   r   r   re   re   s   sl   … à€K�ÓØ€Eˆ4ÓØ"Ð�tÓ"Ø€Dˆ#ÓØ'^Ð˜*Ó^Ø46Ð% zÓ6Ø'-Ð Ó-Ø/IÐ  *ÓIØ)*Ð! 3Ó*Ø%€L�#Ô%r   re   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   )	rA   Úuint8Úint8Úuint16Úint16Úuint32Úint32rI   Úint64r?   s    r   Ú_get_signed_np_dtyper‚   �   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_typerA   r   rI   Úint1r   Úfloat16Úfloat32Úfloat64r|   r{   r~   r}   r€   r   r�   Ú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Úuintr?   r
   r   )ÚgetattrrA   r@   Ú
frombufferÚtobytesÚfp_mantissa_widthÚexponent_biasrH   r€   ÚanyÚ
zeros_likerD   ÚmaximumÚminimumÚ_irÚROUNDING_MODEÚRTNErG   r7   )Ú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é@   )rw   )ÚaÚbs     r   Ú
_umulhi_64rÈ   ò   s   € ô �‹F”S˜“V‰O Ñ"Ð"r   )Úotypesc                  ó   — e Zd Zed„ «       Zy)ÚExtraFunctionsc                óx   — t        j                  |j                  j                  | j                  ||«      |«      S r   )r…   ÚtensorÚbuilderÚcreate_fp_to_fpÚhandle)r£   Údst_tyÚfp_downcast_roundingÚ	_semantics       r   Ú_convert_custom_typesz$ExtraFunctions._convert_custom_typesÿ   s.   € ä�y‰y˜×*Ñ*×:Ñ:¸5¿<¹<ÈÐQeÓfÐhnÓoÐor   N)r*   r+   r,   ÚstaticmethodrÔ   r2   r   r   rË   rË   ý   s   „ àñpó ñpr   rË   c                  ót  — 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.„ ZEd/„ ZFd0„ ZGd1„ ZHd2„ ZId3„ ZJd4„ ZKd5„ ZLd6„ ZMd7„ ZNd8„ ZOd9„ ZPd:„ ZQd;„ ZRd<„ ZSd=„ ZTd>„ ZUd?„ ZVd@„ ZWdA„ ZXdB„ ZYdC„ ZZdD„ Z[dE„ Z\dF„ Z]dG„ Z^dH„ Z_dI„ Z`dJ„ ZadK„ ZbdL„ ZcdM„ ZddN„ ZedO„ ZfdP„ ZgdQ„ ZhdR„ ZidS„ ZjdT„ ZkdU„ ZldV„ ZmdW„ ZndX„ ZodY„ ZpdZ„ Zqd[„ Zrd\„ Zsd]„ Ztd^„ Zud_„ Zvd`„ Zwda„ Zxdb„ Zydc„ Zzdd„ Z{de„ Z|eKZ}eKZ~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�d•„Z®džd–„Z¯dŸd—„Z°d d˜„Z±	 	 d¡d™„Z²dš„ Z³y›)¢ÚInterpreterBuilderc                ó˜   — d | _         t        «       | _        i | _        t        j
                  | j                  d<   d„ | j                  d<   y )NÚconvert_custom_typesc                 ó   — y)N)r
   r
   r
   r2   )ÚlhsTypeÚrhsTypes     r   ú<lambda>z-InterpreterBuilder.__init__.<locals>.<lambda>  s   � r   Úmin_dot_size)rj   re   Ú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_ty1  ó   € Ü�{‰{Ðr   c                ó"   — t         j                  S r   )r…   r‰   r   s    r   Úget_float_tyzInterpreterBuilder.get_float_ty4  rð   r   c                ó"   — t         j                  S r   )r…   rŠ   r   s    r   Úget_double_tyz InterpreterBuilder.get_double_ty7  rð   r   c                ó"   — t         j                  S r   )r…   r‡   r   s    r   Úget_int1_tyzInterpreterBuilder.get_int1_ty:  ó   € Ü�w‰wˆr   c                ó"   — t         j                  S r   )r…   r|   r   s    r   Úget_int8_tyzInterpreterBuilder.get_int8_ty=  rú   r   c                ó"   — t         j                  S r   )r…   r{   r   s    r   Úget_uint8_tyzInterpreterBuilder.get_uint8_ty@  ó   € Ü�x‰xˆr   c                ó"   — t         j                  S r   )r…   r~   r   s    r   Úget_int16_tyzInterpreterBuilder.get_int16_tyC  rÿ   r   c                ó"   — t         j                  S r   )r…   r}   r   s    r   Úget_uint16_tyz InterpreterBuilder.get_uint16_tyF  ó   € Ü�y‰yÐr   c                ó"   — t         j                  S r   )r…   r€   r   s    r   Úget_int32_tyzInterpreterBuilder.get_int32_tyI  rÿ   r   c                ó"   — t         j                  S r   )r…   r   r   s    r   Úget_uint32_tyz InterpreterBuilder.get_uint32_tyL  r  r   c                ó"   — t         j                  S r   )r…   r�   r   s    r   Úget_int64_tyzInterpreterBuilder.get_int64_tyO  rÿ   r   c                ó"   — t         j                  S r   )r…   rI   r   s    r   Úget_uint64_tyz InterpreterBuilder.get_uint64_tyR  r  r   c                ó"   — t         j                  S r   )r…   rŽ   r   s    r   Úget_fp8e4nv_tyz!InterpreterBuilder.get_fp8e4nv_tyU  ó   € Ü�}‰}Ðr   c                ó"   — t         j                  S r   )r…   r�   r   s    r   Úget_fp8e4b15_tyz"InterpreterBuilder.get_fp8e4b15_tyX  ó   € Ü�~‰~Ðr   c                ó"   — t         j                  S r   )r…   r�   r   s    r   Úget_fp8e4b8_tyz!InterpreterBuilder.get_fp8e4b8_ty[  r  r   c                ó"   — t         j                  S r   )r…   rŒ   r   s    r   Úget_fp8e5_tyzInterpreterBuilder.get_fp8e5_ty^  ró   r   c                ó"   — t         j                  S r   )r…   r�   r   s    r   Úget_fp8e5b16_tyz"InterpreterBuilder.get_fp8e5b16_tya  r  r   c                ó.   — t        j                  ||«      S r   )r…   r†   )r   Úelt_tyÚ
addr_spaces      r   Ú
get_ptr_tyzInterpreterBuilder.get_ptr_tyd  s   € Ü�‰˜v zÓ2Ð2r   c                ó.   — t        j                  ||«      S r   )r…   r‘   )r   r   r7   s      r   Úget_block_tyzInterpreterBuilder.get_block_tyg  s   € Ü�}‰}˜U EÓ*Ð*r   c                ó~   — t        t        j                  |gt        j                  ¬«      t        j
                  «      S ©Nr?   )r   rA   ÚarrayÚbool_r…   r‡   ©r   r(   s     r   Úget_int1zInterpreterBuilder.get_int1j  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   rA   r!  r{   r…   r#  s     r   Ú	get_uint8zInterpreterBuilder.get_uint8m  ó$   € Üœ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   rA   r!  r|   r…   r#  s     r   Úget_int8zInterpreterBuilder.get_int8p  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   rA   r!  r}   r…   r#  s     r   Ú
get_uint16zInterpreterBuilder.get_uint16s  ó$   € ÜœBŸH™H e W´B·I±IÔ>ÄÇ	Á	ÓJÐJr   c                ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r   )r   rA   r!  r~   r…   r#  s     r   Ú	get_int16zInterpreterBuilder.get_int16v  r'  r   c                ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r   )r   rA   r!  r   r…   r#  s     r   Ú
get_uint32zInterpreterBuilder.get_uint32y  r,  r   c                ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r   )r   rA   r!  r€   r…   r#  s     r   Ú	get_int32zInterpreterBuilder.get_int32|  r'  r   c                ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r   )r   rA   r!  rI   r…   r#  s     r   Ú
get_uint64zInterpreterBuilder.get_uint64  r,  r   c                ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r   )r   rA   r!  r�   r…   r#  s     r   Ú	get_int64zInterpreterBuilder.get_int64‚  r'  r   c                ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r   )r   rA   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   rA   r!  r‰   r…   r#  s     r   Úget_fp32zInterpreterBuilder.get_fp32ˆ  r9  r   c                ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r   )r   rA   r!  rŠ   r…   r#  s     r   Úget_fp64zInterpreterBuilder.get_fp64‹  r9  r   c                óX   — t        t        j                  dgt        |«      ¬«      |«      S ©Nr   r?   )r   rA   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 Noner?   )rä   rã   r   rA   r!  r€   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   rA   r!  râ   r€   r…   rC  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   rA   Ú	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   rA   rI  r   r   r…   r‡   Úcreate_masked_store)r   rK  ÚvalrL  rM  rO  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   rA   r�   r   Ú_interpreterÚload)
r   rN   rO  rP  Úcache_modifierÚeviction_policyrN  rL   Údtype_npÚrets
             r   rJ  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   )rW  Ústorer   )r   rN   r(   rO  rY  rZ  s         r   rS  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   rJ   r…   r‹   r‰   r½   r   Úviewr”   r   rH   )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   ©re  ©r   ra  rb  s      r   rÝ   zInterpreterBuilder.<lambda>º  ó   € °$·.±.ÀÀhÓ2O€ r   c                ó&   — | j                  ||«      S r   rg  rh  s      r   rÝ   zInterpreterBuilder.<lambda>»  ri  r   c                ó&   — | j                  ||«      S r   rg  rh  s      r   rÝ   zInterpreterBuilder.<lambda>¼  ri  r   c                ó&   — | j                  ||«      S r   rg  rh  s      r   rÝ   zInterpreterBuilder.<lambda>½  ri  r   c                ó&   — | j                  ||«      S r   rg  rh  s      r   rÝ   zInterpreterBuilder.<lambda>¾  s   € °·±¸sÀHÓ0M€ r   c                ó&   — | j                  ||«      S r   rg  rh  s      r   rÝ   zInterpreterBuilder.<lambda>¿  ri  r   c                ó&   — | j                  ||«      S r   rg  )r   ra  rb  Ú	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   rJ   r½   r   r`  r”   r   )r   ra  rb  r¦   rc  rd  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   r`  r”   rJ   rh  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   rJ   )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   ©ry  rA   Úadd©r   rv  rw  s      r   rÝ   zInterpreterBuilder.<lambda>Ï  s   € ¨¯©¸¸SÄ"Ç&Á&Ó)I€ r   c                óD   — | j                  ||t        j                  «      S r   ©ry  rA   Úmultiplyr}  s      r   rÝ   zInterpreterBuilder.<lambda>Ð  ó   € ¨¯©¸¸SÄ"Ç+Á+Ó)N€ r   c                óD   — | j                  ||t        j                  «      S r   ©ry  rA   Údivider}  s      r   rÝ   zInterpreterBuilder.<lambda>Ñ  s   € ¨¯©¸¸SÄ"Ç)Á)Ó)L€ r   c                óD   — | j                  ||t        j                  «      S r   ©ry  rA   Úfmodr}  s      r   rÝ   zInterpreterBuilder.<lambda>Ò  ó   € ¨¯©¸¸SÄ"Ç'Á'Ó)J€ r   c                óD   — | j                  ||t        j                  «      S r   ©ry  rA   Úsubtractr}  s      r   rÝ   zInterpreterBuilder.<lambda>Ó  r�  r   c                óD   — | j                  ||t        j                  «      S r   r  r}  s      r   rÝ   zInterpreterBuilder.<lambda>Ô  ó   € ¨¯©°s¸CÄÇÁÓ(M€ r   c                óD   — | j                  ||t        j                  «      S r   rƒ  r}  s      r   rÝ   zInterpreterBuilder.<lambda>Õ  s   € °·±ÀÀSÌ"Ï)É)Ó1T€ r   c                ó&   — | j                  ||«      S r   ©Úcreate_idivr}  s      r   rÝ   zInterpreterBuilder.<lambda>Ö  ó   € ¨×)9Ñ)9¸#¸sÓ)C€ r   c                ó&   — | j                  ||«      S r   r�  r}  s      r   rÝ   zInterpreterBuilder.<lambda>×  r’  r   c                óD   — | j                  ||t        j                  «      S r   r†  r}  s      r   rÝ   zInterpreterBuilder.<lambda>Ù  rˆ  r   c                óD   — | j                  ||t        j                  «      S r   r†  r}  s      r   rÝ   zInterpreterBuilder.<lambda>Ú  rˆ  r   c                óD   — | j                  ||t        j                  «      S r   r{  r}  s      r   rÝ   zInterpreterBuilder.<lambda>Û  s   € ¨¯©°s¸CÄÇÁÓ(H€ r   c                óD   — | j                  ||t        j                  «      S r   rŠ  r}  s      r   rÝ   zInterpreterBuilder.<lambda>Ü  r�  r   c                óD   — | j                  ||t        j                  «      S r   )ry  rA   Ú
left_shiftr}  s      r   rÝ   zInterpreterBuilder.<lambda>Ý  s   € ¨¯©°s¸CÄÇÁÓ(O€ r   c                óD   — | j                  ||t        j                  «      S r   )ry  rA   Úright_shiftr}  s      r   rÝ   zInterpreterBuilder.<lambda>Þ  s   € ¨¯©¸¸SÄ"Ç.Á.Ó)Q€ r   c                óD   — | j                  ||t        j                  «      S r   ©ry  rA   rŸ   r}  s      r   rÝ   zInterpreterBuilder.<lambda>ß  ó   € ¨$¯.©.¸¸cÄ2Ç:Á:Ó*N€ r   c                óD   — | j                  ||t        j                  «      S r   r�  r}  s      r   rÝ   zInterpreterBuilder.<lambda>à  rž  r   c                óD   — | j                  ||t        j                  «      S r   r�  r}  s      r   rÝ   zInterpreterBuilder.<lambda>á  ó   € ¨T¯^©^¸CÀÄbÇjÁjÓ-Q€ r   c                óD   — | j                  ||t        j                  «      S r   r�  r}  s      r   rÝ   zInterpreterBuilder.<lambda>â  ó   € ¨D¯N©N¸3ÀÄRÇZÁZÓ,P€ r   c                óD   — | j                  ||t        j                  «      S r   ©ry  rA   rž   r}  s      r   rÝ   zInterpreterBuilder.<lambda>ã  rž  r   c                óD   — | j                  ||t        j                  «      S r   r¥  r}  s      r   rÝ   zInterpreterBuilder.<lambda>ä  rž  r   c                óD   — | j                  ||t        j                  «      S r   r¥  r}  s      r   rÝ   zInterpreterBuilder.<lambda>å  r¡  r   c                óD   — | j                  ||t        j                  «      S r   r¥  r}  s      r   rÝ   zInterpreterBuilder.<lambda>æ  r£  r   c                óD   — | j                  ||t        j                  «      S r   ©ry  rA   Ú
less_equalr}  s      r   rÝ   zInterpreterBuilder.<lambda>ç  ó   € ¨D¯N©N¸3ÀÄRÇ]Á]Ó,S€ r   c                óD   — | j                  ||t        j                  «      S r   ©ry  rA   Úlessr}  s      r   rÝ   zInterpreterBuilder.<lambda>è  ó   € ¨D¯N©N¸3ÀÄRÇWÁWÓ,M€ r   c                óD   — | j                  ||t        j                  «      S r   ©ry  rA   Úgreater_equalr}  s      r   rÝ   zInterpreterBuilder.<lambda>é  ó   € ¨D¯N©N¸3ÀÄR×EUÑEUÓ,V€ r   c                óD   — | j                  ||t        j                  «      S r   ©ry  rA   Úgreaterr}  s      r   rÝ   zInterpreterBuilder.<lambda>ê  r£  r   c                óD   — | j                  ||t        j                  «      S r   rª  r}  s      r   rÝ   zInterpreterBuilder.<lambda>ë  r¬  r   c                óD   — | j                  ||t        j                  «      S r   r®  r}  s      r   rÝ   zInterpreterBuilder.<lambda>ì  r°  r   c                óD   — | j                  ||t        j                  «      S r   r²  r}  s      r   rÝ   zInterpreterBuilder.<lambda>í  r´  r   c                óD   — | j                  ||t        j                  «      S r   r¶  r}  s      r   rÝ   zInterpreterBuilder.<lambda>î  r£  r   c                óD   — | j                  ||t        j                  «      S r   ©ry  rA   Úequalr}  s      r   rÝ   zInterpreterBuilder.<lambda>ï  s   € ¨4¯>©>¸#¸sÄBÇHÁHÓ+M€ r   c                óD   — | j                  ||t        j                  «      S r   ©ry  rA   Ú	not_equalr}  s      r   rÝ   zInterpreterBuilder.<lambda>ð  s   € ¨4¯>©>¸#¸sÄBÇLÁLÓ+Q€ r   c                óD   — | j                  ||t        j                  «      S r   r®  r}  s      r   rÝ   zInterpreterBuilder.<lambda>ñ  r°  r   c                óD   — | j                  ||t        j                  «      S r   r¶  r}  s      r   rÝ   zInterpreterBuilder.<lambda>ò  r£  r   c                óD   — | j                  ||t        j                  «      S r   rª  r}  s      r   rÝ   zInterpreterBuilder.<lambda>ó  r¬  r   c                óD   — | j                  ||t        j                  «      S r   r²  r}  s      r   rÝ   zInterpreterBuilder.<lambda>ô  r´  r   c                óD   — | j                  ||t        j                  «      S r   r½  r}  s      r   rÝ   zInterpreterBuilder.<lambda>õ  ó   € ¨D¯N©N¸3ÀÄRÇXÁXÓ,N€ r   c                óD   — | j                  ||t        j                  «      S r   rÀ  r}  s      r   rÝ   zInterpreterBuilder.<lambda>ö  ó   € ¨D¯N©N¸3ÀÄRÇ\Á\Ó,R€ r   c                óD   — | j                  ||t        j                  «      S r   r®  r}  s      r   rÝ   zInterpreterBuilder.<lambda>÷  r°  r   c                óD   — | j                  ||t        j                  «      S r   r¶  r}  s      r   rÝ   zInterpreterBuilder.<lambda>ø  r£  r   c                óD   — | j                  ||t        j                  «      S r   rª  r}  s      r   rÝ   zInterpreterBuilder.<lambda>ù  r¬  r   c                óD   — | j                  ||t        j                  «      S r   r²  r}  s      r   rÝ   zInterpreterBuilder.<lambda>ú  r´  r   c                óD   — | j                  ||t        j                  «      S r   r½  r}  s      r   rÝ   zInterpreterBuilder.<lambda>û  rÇ  r   c                óD   — | j                  ||t        j                  «      S r   rÀ  r}  s      r   rÝ   zInterpreterBuilder.<lambda>ü  rÉ  r   c                óD   — | j                  ||t        j                  «      S r   )ry  rA   Úbitwise_andr}  s      r   rÝ   zInterpreterBuilder.<lambda>ý  ó   € ¨¯©°s¸CÄÇÁÓ(P€ r   c                óD   — | j                  ||t        j                  «      S r   )ry  rA   Úbitwise_xorr}  s      r   rÝ   zInterpreterBuilder.<lambda>þ  rÒ  r   c                óD   — | j                  ||t        j                  «      S r   )ry  rA   Ú
bitwise_orr}  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   rA   r‡  r   rJ   r}  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   )r‚   r   r   rH   ry  rA   r›  )r   rv  rw  Ú	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–   r>   r   )r   r   rA   r�   rI   r   Únp_umulhi_u64rJ   r—   r`   rH   r€  )r   rv  rw  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   ru  )r   rv  rw  rP  rx  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ä  rA   Ú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ä  rA   Úwhere)r   Úcondrv  rw  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   ru  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   ru  )r   rç  rx  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   r@   r—   rA   r   r`  r”   r   rJ   )r   rç  rL   Úmask_bitwidthÚnp_uint_dtyper   rO  r\  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ñ  rA   Úcos©r   rç  s     r   rÝ   zInterpreterBuilder.<lambda>4  ó   €  4§=¡=°´b·f±fÓ#=€ r   c                óB   — | j                  |t        j                  «      S r   )rñ  rA   Úexprø  s     r   rÝ   zInterpreterBuilder.<lambda>5  rù  r   c                óB   — | j                  |t        j                  «      S r   )rñ  rA   Úexp2rø  s     r   rÝ   zInterpreterBuilder.<lambda>6  ó   €  D§M¡M°#´r·w±wÓ$?€ r   c                óB   — | j                  |t        j                  «      S r   )rñ  rA   Úabsrø  s     r   rÝ   zInterpreterBuilder.<lambda>7  s   €  D§M¡M°#´r·v±vÓ$>€ r   c                óB   — | j                  |t        j                  «      S r   )rñ  rA   Úfloorrø  s     r   rÝ   zInterpreterBuilder.<lambda>8  s   €  T§]¡]°3¼¿¹Ó%A€ r   c                óB   — | j                  |t        j                  «      S r   )rñ  rA   Úceilrø  s     r   rÝ   zInterpreterBuilder.<lambda>9  rþ  r   c                óB   — | j                  |t        j                  «      S r   )rñ  rA   Úlogrø  s     r   rÝ   zInterpreterBuilder.<lambda>:  rù  r   c                óB   — | j                  |t        j                  «      S r   )rñ  rA   Úlog2rø  s     r   rÝ   zInterpreterBuilder.<lambda>;  rþ  r   c                óB   — | j                  |t        j                  «      S r   ©rñ  rA   Úsqrtrø  s     r   rÝ   zInterpreterBuilder.<lambda><  s   € ¨D¯M©M¸#¼r¿w¹wÓ,G€ r   c                óB   — | j                  |t        j                  «      S r   r
  rø  s     r   rÝ   zInterpreterBuilder.<lambda>=  rþ  r   c                óB   — | j                  |t        j                  «      S r   )rñ  rA   Úsinrø  s     r   rÝ   zInterpreterBuilder.<lambda>>  rù  r   c                óä   — |j                   j                  t        j                  k(  rt	        |j                   «      nt        |j                   «      }t        ||j                  j                  «      S r   )r   r   rA   r‰   Únp_erf_fp32Únp_erf_fp64r   rJ   )r   rç  r\  s      r   Ú
create_erfzInterpreterBuilder.create_erf@  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   rA   r  r   r   rJ   rø  s     r   Úcreate_rsqrtzInterpreterBuilder.create_rsqrtD  s+   € Ü˜A¤§¡¨¯©Ó 1Ñ1°3·9±9×3CÑ3CÓDÐDr   c                ót   — t        |j                  j                  |«      |j                  j                  «      S r   )r   r   rG   r   rJ   )r   rç  r7   Úallow_reorders       r   rÝ   zInterpreterBuilder.<lambda>H  s*   € ¼\È#Ï(É(×JZÑJZÐ[`ÓJaÐcf×clÑcl×csÑcsÓ=t€ r   c                ó~   — t        t        j                  |j                  |«      |j                  j
                  «      S r   )r   rA   Ú	transposer   r   rJ   )r   rç  Úperms      r   Úcreate_transzInterpreterBuilder.create_transJ  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 )Nr>   r?   )r   r   r@   Úis_floatingr½   r…   rˆ   r`  rA   r   ÚmatmulrJ   )r   rÆ   rÇ   ÚdÚinput_precisionÚmax_num_imprecise_accÚa_dataÚb_datas           r   Ú
create_dotzInterpreterBuilder.create_dotM  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   rA   rF   r€   r…   )r   Úret_tyÚstartÚstops       r   Úcreate_make_rangez$InterpreterBuilder.create_make_rangeV  s$   € ÜœBŸI™I e¨T¼¿¹ÔBÄBÇHÁHÓMÐMr   c                óî  — |€=t        t        j                  |j                  t        ¬«      t
        j                  «      }t        j                  |j                  |j                  t        j                  |j                  «      «      }t        j                  ||d|f¬«      d   }|dxx   t        j                  |j                  «      j                  «       z  cc<   t        |t
        j                  «      S )Nr?   r   )ÚbinsrD   )r   rA   rI  r   r   r…   r‡   rì  r�   Ú	histogramÚlogical_notÚsumr€   )r   r   r+  rO  r,  s        r   Úcreate_histogramz#InterpreterBuilder.create_histogramY  s�   € Øˆ<Ü¤§¡¨T¯Y©Y¼dÔ CÄRÇWÁWÓMˆDä�x‰x˜Ÿ	™	 4§9¡9¬b¯m©m¸D¿I¹IÓ.FÓGˆÜ—L‘L ¨D¸¸D¸	ÔBÀ1ÑEˆ	à�!‹œŸ™ t§y¡yÓ1×5Ñ5Ó7Ñ7‹Ü˜I¤r§x¡xÓ0Ð0r   c                ó–   — t        t        j                  |j                  |j                  |¬«      |j                  j
                  «      S )N©rD  )r   rA   Útake_along_axisr   r   rJ   )r   ra  ÚindicesrD  s       r   Úcreate_gatherz InterpreterBuilder.create_gatherc  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
   r>   )	r"   r@   Úmaxr   r   rH   rA   rI   r   )r   rK  ÚoffsetrL   Úelement_bitwidthÚelement_bytewidths         r   Úcreate_addptrz InterpreterBuilder.create_addptrh  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 )Nr?   Únanzunsupported padding option )rS   r"   r”   r    ÚPADDING_OPTIONÚPAD_ZEROr   rA   r�   r   ÚPAD_NANÚ	full_likeÚfloatrã   rJ  )r   rK  rK   Úpadding_optionrY  rZ  rN  rN   rO   rL   r[  rP  s               r   Úcreate_tensor_pointer_loadz-InterpreterBuilder.create_tensor_pointer_loado  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   ©rS   rS  )r   rK  r(   rK   rY  rZ  rN   rO   s           r   Úcreate_tensor_pointer_storez.InterpreterBuilder.create_tensor_pointer_store~  s/   € Ø×.Ñ.¨~Ó>‰ˆˆeØ×'Ñ'¨¨e°U¸NÈOÓ\Ð\r   c                ó~   — t        t        j                  |j                  |«      |j                  j
                  «      S r   )r   rA   Úexpand_dimsr   r   rJ   )r   rç  rD  s      r   Úcreate_expand_dimsz%InterpreterBuilder.create_expand_dims‚  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   rA   rB   r   r   rJ   )r   rç  r7   s      r   Úcreate_broadcastz#InterpreterBuilder.create_broadcast…  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   rA   Úconcatenater   r   rJ   r}  s      r   Ú
create_catzInterpreterBuilder.create_catˆ  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 )NrZ   r1  )r   rA   Ústackr   r   rJ   r}  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
   ru  )r   rT  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           	     ó´  — |j                   }t        |j                  t        j                  «      rVt        t        j                  ||j                  d   t        |j                  «      ¬«      |j                  j                  «      S t        t        j                  ||j                  t        |j                  «      ¬«      |j                  j                  «      S r?  )r7   r„   r   r…   r‘   r   rA   Úfullr   r”   rJ   )r   r&  rç  r7   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   rW  Ú
atomic_casr   r   rJ   )r   rK  ÚcmprT  Ú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 rX  )	Úir_rmw_op_to_interpreter_rmw_oprã   rY  r   rW  Ú
atomic_rmwr   r   rJ   )r   ÚrmwOprK  rT  rO  r\  r]  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 modere  )r   Ú	inlineAsmÚconstraintsÚvaluesr@  rl  Ú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Ú02xr2   rÁ   s    r   rÝ   z1InterpreterBuilder.create_print.<locals>.<lambda>¸  s   € ¸bÀÀ3À¸L€ r   )Ú	formatter)rä   rA   Úset_printoptionsÚprintr   )r   ÚprefixÚhexrq  Ú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   r2   )r   Ú	conditionÚmessages      r   Úcreate_assertz InterpreterBuilder.create_assert¾  s   € áÐ&˜W˜IÓ&‰yr   c                ó   — |sJ d«       ‚y )NzAssume failedr2   )r   r„  s     r   Úcreate_assumez InterpreterBuilder.create_assumeÂ  s   € ÙÐ)˜/Ó)‰yr   c                 ó   — y r   r2   r   s    r   Úcreate_barrierz!InterpreterBuilder.create_barrierÅ  s   € àr   c                óf   — |D �cg c]  }|j                  «       ‘Œ }}t        ||||||«      S c c}w r   )r%   r4   )	r   r6   r7   r8   r9   r:   r;   r7  Ú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))rE   r9   rã   r%   r4   r6   r7   r8   r:   r;   rD   r   )r   rK  r9   r7  rŒ  r\  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 c                óB   — t        ||||«      }|j                  «        |S r   )rU   r]   )r   r6   r7   r8   Útensor_shaperp  Údescs          r   Úcreate_make_tensor_descriptorz0InterpreterBuilder.create_make_tensor_descriptorØ  s"   € ô    e¨W°lÓCˆØ�‰ŒØˆr   c                ó|   — t        |t        «      sJ ‚|j                  |«      \  }}| j                  ||d ||d¬«      S )NF)rP  rY  rZ  rN  )r„   rU   rS   rJ  )r   r’  r3  rY  rZ  rN   rO  s          r   Úcreate_descriptor_loadz)InterpreterBuilder.create_descriptor_loadä  sO   € ä˜$Ô 0Ô1Ð1Ð1Ø×.Ñ.¨wÓ7‰
ˆˆdØ×&Ñ& t¨T¸ÈnØ7FÐTYð 'ó [ð 	[r   c                óT   — |j                  |«      \  }}| j                  |||d d «      S r   rE  )r   r’  r(   r3  rN   rO  s         r   Úcreate_descriptor_storez*InterpreterBuilder.create_descriptor_storeë  s/   € Ø×.Ñ.¨wÓ7‰
ˆˆdØ×'Ñ'¨¨e°T¸4ÀÓFÐFr   c                ó®  — |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   rZ   r?   )r6   r   r    r”   rA   Úzerosr   r7   r:   Ú	enumerater   r…   r€   r•  )r   r’  Ú	x_offsetsÚy_offsetr@  r   Únp_dtypeÚresultrY  rZ  r³   Úx_offsetr3  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…   r€   r—  )	r   r’  r(   r›  rœ  r³   rŸ  Úslicer3  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        j                  k(  r,t        t        j                  dd|¬«      |j
                  «      S t        d|› �«      ‚)Nrw   r
   rZ   r?   Tzunsupported type )r”   Únamer   rA   rU  rJ   r"  Ú	TypeError)r   r@  Únp_types      r   Úget_all_ones_valuez%InterpreterBuilder.get_all_ones_value  ss   € Ü Ó%ˆØ�G—L‘LÑ Ü¤§¡¨¨2°WÔ =¸t¿{¹{ÓKÐKØœŸ™Ò Ü¤§¡¨¨4°wÔ ?ÀÇÁÓMÐMäÐ/°¨vÐ6Ó7Ð7r   N©ÚreturnÚNone)
r6   r   r7   ra   r8   ra   r‘  rb   rp  r   )r’  rU   r3  ra   )r’  rU   r(   r   r3  ra   )r’  rU   r›  r   rœ  r   )r’  rU   r(   r   r›  r   rœ  r   )´r*   r+   r,   r    ÚMEM_SEMANTICÚACQUIRErW  ÚRELEASEÚRELAXEDÚACQUIRE_RELEASErY  Ú	ATOMIC_OPÚADDÚRMW_OPÚFADDÚMINÚUMINÚMAXÚUMAXÚANDÚORÚXORÚXCHGr`  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  r2  r4  r6  r8  r;  r=  rA  rE  rG  rQ  rU  rJ  rS  re  Ú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Ï   rs  ry  Ú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/  r4  r:  rC  rF  rI  rK  rN  rQ  rS  rV  r^  rc  rm  rs  r‚  r†  rˆ  rŠ  r�  r�  r“  r•  r—  r   r£  r¨  r2   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ò1òeò
eòiò]òNòPòTòWòtòlòjòròZòRò0ò'ò*òòYò
ð
àð
ð "ð
ð $ð	
ð
  ð
ð ó
ó[óGó	+ð?Ø,8ó?ó8r   r×   c                óF   ‡— t        |«      Š|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Úsemantics        €r   rÝ   z_patch_attr.<locals>.<lambda>  s[   ø€ ¹Àð :kàMSÏ\É\Ë^÷AVÙEIÀQÈØDEÈÒDTð BCÀAÁó AVñ:kð bjò:k€ ùóAVs   ™5)r   Úsetattr)Úobjr¥  r  rÎ   Ú
new_memberr  s        @r   Ú_patch_attrr    s&   ø€ Ü˜gÓ&€HØ&,õ l€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éþÿÿÿrZ   )r   rA   r  rÐ   r   r   r@  Úis_blockÚlistr7   r…   r  r‘   rÍ   )r   rÐ   r:   Ú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   )rw   rÐ   r   r   s    r   rÝ   z$_patch_lang_tensor.<locals>.<lambda>*  s   € ¤C¨¯©×(8Ñ(8Ó$9€ r   c                ó   •—  ‰| «      S r   r2   )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   )ri   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                  ó*   — e Zd Zd„ Zd„ Zd„ Zd„ Zd„ Zy)ÚReduceScanOpInterfacec                ó    — || _         || _        y r   )rD  Ú
combine_fn)r   rD  r1  s      r   r<   zReduceScanOpInterface.__init__3  s   € ØˆŒ	Ø$ˆ�r   c                óH   — |� |t        |«      k\  rt        d|› d|› �«      ‚y y )Nzaxis z out of bounds for shape )rE   rã   )r   r7   rD  s      r   Ú
check_axisz ReduceScanOpInterface.check_axis7  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ã   r@  r3  r7   rD  )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 )Nr7   r?   )r”   r!   r7   rH   r…   r‘   r   rA   r!  r  rÍ   r   rJ   )r   r\  r   r�  Úret_types        r   Ú	to_tensorzReduceScanOpInterface.to_tensorA  sy   € Ü  Ó'ˆÜ�3˜Ô  S§Y¢YØ—*‘*˜XÓ&ˆCÜ—}‘} U¬D°·±«OÓ<‰Hä—(‘(˜C˜5¨Ô1ˆCØˆHÜ�w‰w�~‰~œl¨3°·±Ó=¸xÓHÐHr   c                óØ   — t        |t        «      s| j                  |f«      d   S | j                  |«       | j	                  |«      }t        |t
        t        f«      rt        |«      S |fS ©Nr   )r„   ÚtupleÚapplyr5  Ú
apply_implr   )r   r£   r\  s      r   r<  zReduceScanOpInterface.applyK  s\   € Ü˜%¤Ô'Ø—:‘:˜u˜iÓ(¨Ñ+Ð+Ø×Ñ˜%Ô Ø�o‰o˜eÓ$ˆÜ'¨¬d´E¨]Ô;Œu�S‹zÐHÀ#ÀÐHr   N)r*   r+   r,   r<   r3  r5  r8  r<  r2   r   r   r/  r/  1  s   „ ò%òMò2òIó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   rD  r1  rB  Ú	__class__s       €r   r<   zReduceOps.__init__U  s   ø€ Ü‰Ñ˜˜zÔ*Ø"ˆ�r   c                óô   — g }|D ]e  }|�|j                  |«       Œd}|j                  | j                  |j                  j                  j	                  «       |j
                  «      «       Œg t        |«      |fS r:  )Úappendr8  rÐ   r   Úflattenr   r;  )r   r£   rD  r\  r   s        r   ÚunravelzReduceOps.unravelY  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                  «      «       Œœ |S )Nr   r
   r?   c              3  óh   •K  — | ])  \  }}‰j                  |‰   ‰|   j                  «      –— Œ+ y ­wr   ©r8  r   )Ú.0Úiir  r£   Úinput_indexr   s      €€€r   ú	<genexpr>z+ReduceOps.generic_reduce.<locals>.<genexpr>r  s/   øè ø€ ÒsÑTYÐTVÐXY §¡¨q°©~¸uÀR¹y¿¹× OÑsùó   ƒ/2c              3  óh   •K  — | ])  \  }}‰j                  |‰   ‰|   j                  «      –— Œ+ y ­wr   rJ  )rK  ÚoiÚor£   Úoutput_indexr   s      €€€r   rN  z+ReduceOps.generic_reduce.<locals>.<genexpr>x  s/   øè ø€ Ò!wÑW\ÐWYÐ[\ $§.¡.°°<±À%ÈÁ)Ç/Á/×"RÑ!wùrO  )rD  rG  rÐ   r   r7   rE  rA   r™  r   rD   r  Úunravel_indexr;  rš  rE   r[   r1  Úfnr„   r…   r  rÍ   rB  rH  r8  )r   r£   Úoriginal_axisrD  Ú
input_dataÚoutput_dataÚinput_shapeÚoutput_shaperç  r³   Úinput_tupleÚjÚ	acc_tupleÚcombine_fn_retr\  r   Ú_rM  rS  s   ``               @@r   Úgeneric_reducezReduceOps.generic_reducec  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Ó;Õ<ð	=ð ˆ
r   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   ©rD  Úkeepdimsz-val_reduce_op and idx_reduce_op are both None)r„   r;  r8  rÐ   r   rD  rB  r   r…   r€   rã   )r   r£   Úval_reduce_opÚidx_reduce_oprT  Ú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 )Nrb  )r8  rA   r.  rÐ   r   rD  rB  r   ©r   r£   s     r   r.  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   )rd  re  )r1  r…   ÚstandardÚ_argmin_combine_tie_break_leftrg  rA   ÚminÚargminÚ_argmax_combine_tie_break_leftr6  ÚargmaxÚ_elementwise_maxÚ_elementwise_minÚ_sum_combiner.  r`  ri  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<   rG  r`  rg  r.  r=  Ú__classcell__©rC  s   @r   r?  r?  S  s$   ø„ ô#ò ò)óVNò$oö.r   r?  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   )rA  r<   Úreverse)r   rD  r1  ry  rC  s       €r   r<   zScanOps.__init__µ  s   ø€ Ü‰Ñ˜˜zÔ*Øˆ�r   c                ó¤   — | j                  t        j                  |j                  j                  | j
                  ¬«      |j                  ¬«      gS ©Nr1  r?   )r8  rA   ÚcumsumrÐ   r   rD  r   ri  s     r   r|  zScanOps.cumsum¹  s8   € Ø—‘œrŸy™y¨¯©×):Ñ):ÀÇÁÔKÐSX×S^ÑS^�Ó_Ð`Ð`r   c                ó¤   — | j                  t        j                  |j                  j                  | j
                  ¬«      |j                  ¬«      gS r{  )r8  rA   ÚcumprodrÐ   r   rD  r   ri  s     r   r~  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   r?   c              3  óh   •K  — | ])  \  }}‰j                  |‰   ‰|   j                  «      –— Œ+ y ­wr   rJ  )rK  rL  r  Úindexr£   r   s      €€€r   rN  z'ScanOps.generic_scan.<locals>.<genexpr>Ê  s,   øè ø€ ÒfÁuÀrÈ1˜Ÿ™¨¨%©°%¸±)·/±/×BÑfùrO  c              3  óV   •K  — | ]   }|‰j                   k(  r‰|   d z
  n‰|   –— Œ" y­w)r
   Nr1  )rK  r³   r�  r   s     €€r   rN  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   rJ  )rK  rQ  rR  r£   Ú
prev_indexr   s      €€€r   rN  z'ScanOps.generic_scan.<locals>.<genexpr>Ñ  s/   øè ø€ Ò!uÑUZÐUWÐYZ $§.¡.°°:±ÀÀbÁ	ÇÁ×"PÑ!uùrO  )rÐ   r   r7   rE  rA   r™  r   rD   r  rT  r;  rš  rD  rE   r[   r1  rU  r„   r…   r  rÍ   r8  )r   r£   rW  rX  r7   rç  r³   r   r\  r]  r^  r\  r�  r„  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 |S )Nr1  r   )ry  rE  r8  rA   ÚfliprÐ   r   rD  r   r1  r…   rk  rs  r|  Ú_prod_combiner~  r…  )r   r£   Ú	new_inputrç  r\  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àˆ
r   )	r*   r+   r,   r<   r|  r~  r…  r=  rt  ru  s   @r   rw  rw  ³  s   ø„ ôòaòbòö<r   rw  c                 ó”   — dd„} dd„}| t         _        |t         _        | t         j                  _        |t         j                  _        y )Nc                ó:   — t        |||«      j                  | «      S r   )r?  r<  )r£   rD  r1  rB  r
  s        r   Ú_new_reducez'_patch_reduce_scan.<locals>._new_reduceõ  s   € Ü˜˜z¨9Ó5×;Ñ;¸EÓBÐBr   c                ó:   — t        |||«      j                  | «      S r   )rw  r<  )r£   rD  r1  ry  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‡   r|   r{   r~   r}   r€   r   r�   rI   rk   rm   ro   Úfp16Úbf16Úfp32Úfp64zfail to convert z to ir type)r¥  Úget_void_tyrù   rü   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   )rD   )Ú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   r2   )rí  r�  s     r   Ú_new_static_assertz,_patch_lang_core.<locals>._new_static_assert4  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(   rE   r6  r7   rã   rÐ   r)   )r£   rq  r¥  r  s       r   Ú	_set_attrz#_patch_lang_core.<locals>._set_attr7  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)Ú )rD   Ústatic_rangeÚstatic_assertr}  Ú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  rE   r  Úinterpreter_builderrÍ   r¿   r-  r¯  Útensor_descriptor_base)rU  r_  r(   Úlangsr®  s        r   Ú_patch_langr¶  O  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×1Ñ1Ô3FÕGùó qs   žD¼DÁDc                óX   — t        | d«      r t        | «      |Ž S  t        | «      |«      S )NÚ_fields)r!   r@  )rç  Úcontentss     r   Ú_tuple_createrº  \  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 t        | t.        «      rÕ| j0                  D �cg c]  }t-        |«      ‘Œ }}| j0                  d
   dk(  sJ ‚t        j2                  d«      |d
<   t5        t7        «       «      }|j9                  t-        | j:                  «      | j<                  D �cg c]  }t-        |«      ‘Œ c}|| j>                  D �cg c]  }t        j2                  |«      ‘Œ c}¬«      S | S c c}w c c}w c c}w )Ni   €l        l        lûÿÿÿ         l            l            zUnsupported integer value r?   Údata_ptrrZ   r
   )r6   r7   r8   r:   ) r„   rw   r…   Ú	str_to_tyÚtritonÚruntimeÚjitÚmangle_typerA   r€   r   r�   rI   rã   r   r!  rÍ   r!   r¼  r;  rº  ÚmapÚ_implicit_cvtr	   r8   r¤  r   r×   Úmake_tensor_descriptorr6   r7   r:   )rç  Útyr   rÐ   Úsr8   r  rÇ   s           r   rÃ  rÃ  f  s  € Ü�#”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Ó:Ð:Ü	�CÔ)Ô	*Ø-0¯[©[Ö9¨”= Õ#Ð9ˆÐ9Ø�{‰{˜2‰ !Ò#Ð#Ð#Ü—l‘l 1“oˆ�‰Ü!Ô"4Ó"6Ó7ˆØ×.Ñ.Ü˜sŸx™xÓ(Ø-0¯Y©YÖ7¨”= Õ#Ò7ØØ25·/±/ÖB¨QœŸ™ a�ÒBð	 /ó 
ð 	
ð €Jùò :ùò 8ùâBs   ÈKÊK
Ê'Kc                óz   — t        | t        j                  j                  j                  «      r| j
                  S | S r   )r„   r¾  r¿  rÀ  ÚTensorWrapperr6   )Úts    r   Ú_unwrap_tensorrÊ  Ž  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Í  ”  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Ñ  rU  Ú	arg_namesÚgridr.   r  ÚgetÚ
constexprs)r   rU  rÒ  rÓ  rÑ  r¥  rÅ  r.   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        | t        «      r<t	         ‰| j
                  «      | j                  | j                  | j                  «      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	   r6   r7   r8   r:   r!   rÊ  Úuntyped_storager¼  rØ  Ú	new_emptyÚset_Ústorage_offsetr  r\   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Ó*;Ó<Ð<Ü˜CÔ!1Ô2Ü'Ù˜CŸH™HÓ%Ø—I‘IØ—K‘KØ—O‘Oó	ð ô ˜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   ù€ Øˆõ	ð0 -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 t	        | t        «      r ‰| j                  |j                  «       y y )Nr¼  )	r!   rÊ  rÚ  r¼  r„   r;  Úzipr	   r6   )Ú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ä˜GÔ%5Ô6Ù˜'Ÿ,™,¨¯©Õ5ð 7r   )ré  r  rq  Úcopy_)r   rã  rä  r
  rå  rê  rë  r'   Ú	kwarg_devÚ	kwarg_hstrì  râ  s             @@r   Ú_restore_args_devzGridExecutor._restore_args_devÈ  s‹   ù€ Øˆõ		6ô !$ H¨hÓ 7ò 	(ÑˆG�WÙ�g˜wÕ'ð	(ð %Ÿl™l›nò 	,‰NˆC�Ø" 3™ˆIÙ�i Õ+ð	,ð #+§/¡/Ó"3ò 	#ÑˆW�gØ�M‰M˜'Õ"ñ	#r   c                óV  — |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 	 | j3                  ||||«       y c c}}w c c}
}	w # t$        $ r?}t&        j(                  j*                  j,                  r‚ t/        t1        |«      «      |‚d }~ww xY w)
NÚwarmupFé   z#grid must have at most 3 dimensions)r
   r   r
   r   r2   )Úpopr  ÚgetfullargspecrU  r  r	  ræ  r¶  ÚgetcallargsrÕ  rÃ  ÚcallablerÓ  rE   r³  rí   rD   rè   Ú	Exceptionr¾  ÚknobsÚcompilationÚfront_end_debuggingr   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ùó3 Hùó løô ò 	3Ü�|‰|×'Ñ'×;Ò;ØÜ"¤4¨£7Ó+°Ð2ûð	3ús+   ÁGÁGÃ&GÅA"G  Ç 	H(Ç):H#È#H(N)r*   r+   r,   r<   ræ  rð  rþ  r2   r   r   rÏ  rÏ  š  s   „ òcò!$òF#ó2 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                  dt	        j                  «       ¬«      dt	        j                  «       ¬«      |j                  t	        j                  d¬«      gg ¬	«      |_	        |S )
Nr
   z&Multiple assignments are not supportedÚinterpreter_semantic)ÚidÚctxr8  )r(   r   r  F)r(   )Úfuncr	  Úkeywords)ÚtargetsÚvisitrE   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Ü—‘¤S§X¡XÐ1GÌSÏXÉXËZÔ%XÐ_jÜ#&§8¡8£:ô/Ø6:·j±jÄ#Ç,Á,ÐUZÔB[Ð5\ÐgiôkˆŒ
ð ˆr   N)r*   r+   r,   r  r2   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   )rU  r
  ÚfilenameÚdef_file_lineno)r   rU  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  ÚgetsourcelinesrU  rø  Ú_get_jit_fn_file_liner  r  Ú	_find_defÚ
def_linenoÚ_prepare_sourceÚ_transform_astÚ_compile_and_exec)r   Úlinesr_  ra  Ú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%  rU  )r   r$  r%  s      r   r  z&FunctionRewriter._get_jit_fn_file_line2  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_def6  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   ra  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   ra  Ú
parsed_astr!  Ú
inc_linenos        r   r  zFunctionRewriter._transform_astC  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
  rU  r±  Úglobalsr  r6  r*   )r   r!  Úcompiled_codeÚlocal_namespaceÚ
fn_globalsr'   r(   s          r   r  z"FunctionRewriter._compile_and_execN  s€   € Ü ¸$¿-¹-ÈfÔUˆØ)˜TŸ[™[˜/ˆØ—W‘W×(Ñ(ˆ
Ü!›)Ÿ/™/Ó+ò 	(‰JˆC�Ø˜*Ò$Ø"'�
˜3’ð	(ô 	ˆ]˜J¨Ô8Ø˜tŸw™w×/Ñ/Ñ0Ð0r   N)r*   r+   r,   r   r0  r<   r"  r  r  r  r  r  r2   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)ÚInterpretedFunctionc                óð   ‡ — |‰ _         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Ó  rU  r   s       €r   Úrunz)InterpretedFunction.__init__.<locals>.runa  s6   ø€ Ø˜&‘>ˆDØ—‘“ˆBØ9”<  D§N¡N°DÓ9¸4ÐJÀ6ÑJÐJr   )
rU  r  ÚrewriterrC  r  Ú	signatureÚ
parametersrq  r¥  rÒ  )r   rU  r
  rC  rE  r  s   `     r   r<   zInterpretedFunction.__init__]  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   )rU  Úrewritten_fnrD  r"  r   s    r   rB  zInterpretedFunction.rewritej  sJ   € Ø�7‰7˜$×+Ñ+Ñ+Ø)-¯©×)BÑ)BÓ)DˆD×Ñ˜dŸg™gÑ&Ø× Ñ  §¡Ñ)Ð)r   c                ó.   — | j                   j                  S r   )rU  r*   r   s    r   r*   zInterpretedFunction.__name__o  s   € à�w‰w×ÑÐr   c                óP   — | j                  «       }t        || j                  |«      S r   rA  )r   rÓ  rU  s      r   Ú__getitem__zInterpretedFunction.__getitem__s  s    € Ø�\‰\‹^ˆÜ˜B §¡°Ó5Ð5r   c                óª   — t        | j                  «       | j                  «       }	  ||i |¤ŽS # t        $ r}t	        t        |«      «      |‚d }~ww xY wr   )r¶  rU  rB  rø  r   r&  )r   r	  r
  rU  rý  s        r   rþ  zInterpretedFunction.__call__w  sO   € ä�D—G‘GÔØ�\‰\‹^ˆð	3Ù�tÐ&˜vÑ&Ð&øÜò 	3Ü"¤4¨£7Ó+°Ð2ûð	3ús   §/ ¯	A¸AÁANr©  )	r*   r+   r,   rH  r<   rB  r+  rK  rþ  r2   r   r   r>  r>  Y  s0   „ à€LóIò*ð
 ñ ó ð ò6ó3r   r>  )EÚ
__future__r   r	  r,  r  Útypingr   r   r   r¿   ÚnumpyrA   r¾  Útriton.languageÚlanguager…   r/   r   Útriton.language.semanticr   Útriton.tools.tensor_descriptorr	   Úerrorsr   Ú	functoolsr   Ú_C.libtritonr   rW  r   r    r   r4   rU   re   r‚   r”   r½   rÃ   rÈ   Ú	vectorizer‰   r  rŠ   r  rI   rÝ  rË   r×   r  r  r-  r/  r?  rw  r‘  r¯  r¶  rº  rÃ  r³  r  rÊ  rÍ  rÏ  ÚNodeTransformerr   r  r>  r2   r   r   ú<module>rY     s¶  ðÝ "Û 
Û Û ß $Ñ $ã Û ã Ý Û Ý !å 3Ý ;Ý $Ý Ý 6Ý $ð ÷ð ó ð÷6ñ ÷4#ñ #ñL �$Ô÷
&ð 
&ó ð
&ò	òò@='ò@ò
#ð ˆb�l‰l˜4¨¯©¨Ô5€Øˆb�l‰l˜4¨¯©¨Ô5€Ø�—‘˜Z°·±°Ô<€÷pñ p÷D8ñ D8òN#ò4ò(÷.Iñ IôD].Ð%ô ].ô@;Ð#ô ;ò|)ò Kò\
HòTò!ñH )Ó*Ð Ù%Ð&9Ó:Ð òò÷gGñ gGôT�S×(Ñ(ô ÷ B1ñ B1÷J%3ò %3r   