Ë
    ©š¾jh ã                  ó¬  — U 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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! ddlm"Z# ddl$m%Z%  e
d«      Z&e G d„ d«      «       Z' G d„ d«      Z( ed¬«       G d„ d«      «       Z)d„ Z*d„ Z+d„ Z,d„ Z-d„ Z.d„ Z/d„ Z0d„ Z1d „ Z2d!„ Z3d"„ Z4 ejj                  e.ejl                  g¬#«      Z7 ejj                  e.ejp                  g¬#«      Z9 ejj                  e/ejt                  g¬#«      Z; G d$„ d%«      Z< G d&„ d'«      Z= e>«       Z? e=«       Z@ ee@«      ZAd(eBd)<    G d*„ d+«      ZCdCd,„ZDdCd-„ZEdCd.„ZF G d/„ d0«      ZG G d1„ d2eG«      ZH G d3„ d4eG«      ZIdCd5„ZJdCd6„ZKd7„ ZLd8„ ZMd9„ ZNd:„ ZO G d;„ d<«      ZP G d=„ d>ej¢                  «      ZR G d?„ d@«      ZS G dA„ dBee&   «      ZTy)Dé    )ÚannotationsN)ÚTupleÚListÚDictÚCallableÚTypeVarÚOptional)Ú	dataclass)ÚTritonSemantic)ÚKernelInterface)ÚTensorDescriptoré   )ÚInterpreterError)Úpartialé   )Úinterpreter)Úir)Ú_tuple_createÚTc                  ót   — 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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.ndarrayÚdataztl.dtypeÚdtype)Údefault_factoryr   Úattrc                óÖ   — t        | j                  | j                  «      sIt        d| j                  j                  dz  › d| j                  j
                  › d| j                  › �«      ‚y )Nznumpy data itemsize (é   z) bits) exceeds dtype primitive_bitwidth (z bits) for triton type )Ú_validate_np_data_sizer   r   Ú
ValueErrorÚitemsizeÚprimitive_bitwidth©Úselfs    úd/home/spazio/kingdom_calls/work5/transenv/lib/python3.12/site-packages/triton/runtime/interpreter.pyÚ__post_init__zTensorHandle.__post_init__'   sj   € Ü% d§i¡i°·±Ô<ÜÐ4°T·Y±Y×5GÑ5GÈ!Ñ5KÐ4Lð M!Ø!%§¡×!>Ñ!>Ð ?Ð?VÐW[×WaÑWaÐVbðdó eð eð =ó    c                óH   — t        | j                  j                  «       «      S ©N)Úboolr   Úallr"   s    r$   Ú__bool__zTensorHandle.__bool__,   s   € Ü�D—I‘I—M‘M“OÓ$Ð$r&   c                óh   — | j                   }t        |d«      r|j                  }t        |d«      rŒ|S )NÚ
element_ty)r   Úhasattrr-   )r#   r   s     r$   Úget_element_tyzTensorHandle.get_element_ty/   s1   € Ø—
‘
ˆÜ�e˜\Ô*Ø×$Ñ$ˆEô �e˜\Õ*àˆr&   c                ó^   — t        | j                  j                  «       | j                  «      S r(   )r   r   Úcopyr   r"   s    r$   ÚclonezTensorHandle.clone5   s   € Ü˜DŸI™IŸN™NÓ,¨d¯j©jÓ9Ð9r&   c                ó"   — || j                   |<   y r(   )r   )r#   ÚkeyÚvalues      r$   Úset_attrzTensorHandle.set_attr8   s   € Øˆ�	‰	�#Šr&   N)Ú__name__Ú
__module__Ú__qualname__Ú__doc__Ú__annotations__ÚdataclassesÚfieldÚdictr   r%   r+   r/   r2   r6   © r&   r$   r   r      sE   … ñð ÓØƒOØ"�×"Ñ"°4Ô8€Dˆ$Ó8òeò
%òò:ór&   r   c                  ó&   — e Zd Z	 	 dd„Zd„ Zdd„Zy)ÚTensorDescHandlec                ój   — || _         t        |«      | _        || _        || _        || _        || _        y r(   )ÚbaseÚlenÚndimÚshapeÚstridesÚblock_shapeÚpadding)r#   rC   rF   rG   rH   rI   s         r$   Ú__init__zTensorDescHandle.__init__>   s1   € àˆŒ	Ü˜“JˆŒ	ØˆŒ
ØˆŒØ&ˆÔØˆ�r&   c                ób  — | j                   j                  j                  «       dz  dk(  sJ d«       ‚t        | j                  «      | j
                  k(  sJ ‚t        | j                  «      | j
                  k(  sJ ‚| j
                  dk\  sJ d«       ‚| j                   j                  j                  }|j                  dz  }| j                  d d D ].  }|j                  j                  «       |z  }|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 alignedr   z"descriptor cannot be 0 dimensionalr   éÿÿÿÿzstride must be 16-byte alignedzlast dim must be contiguous)
rC   r   ÚitemrD   rG   rE   rH   r   r-   r!   )r#   Ú	scalar_tyr    ÚstrideÚbyte_strides        r$   ÚvalidatezTensorDescHandle.validateG   s  € Ø�y‰y�~‰~×"Ñ"Ó$ rÑ)¨QÒ.ÐNÐ0NÓNÐ.Ü�4—<‘<Ó  D§I¡IÒ-Ð-Ð-Ü�4×#Ñ#Ó$¨¯	©	Ò1Ð1Ð1Ø�y‰y˜AŠ~ÐCÐCÓCˆ~à—I‘I—O‘O×.Ñ.ˆ	Ø×/Ñ/°1Ñ4ˆØ—l‘l 3 BÐ'ò 	KˆFØ Ÿ+™+×*Ñ*Ó,¨xÑ7ˆKØ Ñ# qÓ(ÐJÐ*JÓJÐ(ð	Kð �|‰|˜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   rM   rL   r   z*block offset start must be 16-byte aligned©r   r   )rD   rE   rC   r   r-   r!   r   ÚnpÚbroadcast_torH   Úonesr)   ÚrangeÚarangeÚreshaperG   ÚastypeÚuint64rF   r   Úscalar)
r#   ÚoffsetsrO   r    Ú	ptrs_dataÚmasksÚdimÚ
bcast_dimsÚoffÚptrs_handles
             r$   Úmaterialize_pointersz%TensorDescHandle.materialize_pointersT   s®  € Ü�7‹|˜tŸy™yÒ(Ð(Ð(Ø—I‘I—O‘O×.Ñ.ˆ	Ø×/Ñ/°1Ñ4ˆØ˜‘× Ñ  8Ñ+¨rÑ1°QÒ6ÐdÐ8dÓdÐ6ä—O‘O D§I¡I§N¡N°D×4DÑ4DÓEˆ	Ü—‘˜×(Ñ(´Ô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Ø! X°¡^°d·l±lÀ3Ñ6G×6LÑ6LÑ%L×$TÑ$TÔUW×U^ÑU^Ó$_Ñ_ˆIØ˜Q #™XÑ&¨#°·
±
¸3±×0DÑ0DÑ*DÑE‰Eð	Fð �‰¤"§)¡)Ò+Ð+Ð+Ü" 9¨d¯i©i¯o©o×.DÑ.DÓEˆØ˜EÐ!Ð!r&   N)rC   r   rF   úList[TensorHandle]rG   rf   rH   ú	List[int])r^   rf   )r7   r8   r9   rJ   rR   re   r?   r&   r$   rA   rA   <   s   „ ðØ'óòPô"r&   rA   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)ÚInterpreterOptionsNzOptional[dict]Úextern_libsFr)   ÚdebugTÚsanitize_overflowzOptional[str]Úarch)Úfp8e5Úfp8e5b16Úfp8e4nvÚfp8e4b8Úfp8e4b15zTuple[str, ...]Úsupported_fp8_dtypesr?   Ú!deprecated_fp8_dot_operand_dtypesÚtf32ÚstrÚdefault_dot_input_precision)rv   Útf32x3ÚieeeÚallowed_dot_input_precisionsr   ÚintÚmax_num_imprecise_acc_defaultr   Úbackend_name)r7   r8   r9   rk   r;   rl   rm   rn   rt   ru   rx   r{   r}   r~   r?   r&   r$   rj   rj   g   sl   … à"&€K�Ó&Ø€Eˆ4ÓØ"Ð�tÓ"Ø€Dˆ-ÓØ,cÐ˜/ÓcØ9;Ð% Ó;Ø'-Ð Ó-Ø4NÐ  /ÓNØ)*Ð! 3Ó*Ø%€L�#Ô%r&   rj   c                óŠ   — t        |t        j                  «      ry| j                  dz  }|j                  }|dk  rd}||kD  ryy)NTr   F)Ú
isinstanceÚtlÚpointer_typer    r!   )Únp_arrayÚtl_dtypeÚnp_dtype_bitwidthÚtl_dtype_bitwidths       r$   r   r   u   sO   € Ü�(œBŸO™OÔ,Øà ×)Ñ)¨AÑ-ÐØ ×3Ñ3Ðð ˜1ÒØÐàÐ,Ò,ØØr&   c                ó  — | t         j                  k(  rt         j                  S | t         j                  k(  rt         j                  S | t         j
                  k(  rt         j                  S | t         j                  k(  rt         j                  S | S r(   )	rU   Úuint8Úint8Úuint16Úint16Úuint32Úint32r\   Úint64rT   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(   )r€   r�   r‚   rU   r   r\   Ú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ÚuintrT   r   r   )ÚgetattrrU   r!   Ú
frombufferÚtobytesÚfp_mantissa_widthÚexponent_biasr[   r�   ÚanyÚ
zeros_likerX   ÚmaximumÚminimumÚ_irÚROUNDING_MODEÚRTNErZ   rF   )Ú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é@   )r|   )ÚaÚbs     r$   Ú
_umulhi_64rÒ   ö   s   € ô �‹F”S˜“V‰O Ñ"Ð"r&   c                ó&  — | j                   t        j                  t        j                  fv sJ ‚| j	                  t        j                  «      } | j	                  t        j
                  «      } | dz  } | j                  t        j                  «      } | S )Né   )r   rU   rˆ   r‰   r[   r�   Úviewr“   )Úscales    r$   Ú_e8m0_to_f32r×   ü   se   € Ø�;‰;œ2Ÿ8™8¤R§W¡WÐ-Ñ-Ð-Ð-Ø�L‰LœŸ™Ó"€EØ�L‰LœŸ™Ó"€EØ�R‰K€EØ�J‰J”r—z‘zÓ"€EØ€Lr&   c                óB  — | j                   t        j                  k(  sJ ‚| t        j                  d«      z  }| t        j                  d«      z	  }| j                  d d | j                  d   dz  fz   }t        j                  |t        j                  ¬«      }||ddd d…f<   ||ddd d…f<   t        j
                  g d	¢t        j                  ¬«      }||t        j                  d
«      z     }|t        j                  d«      z  dk7  }t        j                  || |«      S )Né   é   rM   r   rT   .r   r   )g        g      à?g      ð?g      ø?g       @g      @g      @g      @é   r   )r   rU   rˆ   rF   ÚemptyÚarrayr“   Úwhere)r5   ÚlowÚhighÚunpacked_shapeÚunpacked_valÚpositive_e2m1_lutÚ
abs_valuesÚsignss           r$   Ú_e2m1_to_f32ræ     sû   € Ø�;‰;œ"Ÿ(™(Ò"Ð"Ð"à
”"—(‘(˜4“.Ñ
 €CØ”B—H‘H˜Q“KÑ€Dà—[‘[  "Ð%¨¯©°R©¸1Ñ)<Ð(?Ñ?€NÜ—8‘8˜N´"·(±(Ô;€LØ!€L��a�d˜�d�ÑØ"€L��a�d˜�d�Ñô Ÿ™Ò!IÔQS×Q[ÑQ[Ô\ÐØ" L´2·8±8¸D³>Ñ$AÑC€JØœBŸH™H T›NÑ*¨qÑ0€EÜ�8‰8�E˜J˜;¨
Ó3Ð3r&   c                ó8  — | j                   t        j                  k(  rt        | j                  «      }n‚t        | j                  | j                   t        j                  t        j                  j                  «      j                  t        j                  «      j                  t        j                  «      }| j                   t        j                  k(  �rG| j                  t        j                  d«      k(  }t        j                  |t        j                  d«      |«      }| j                  t        j                  d«      k(  }t        j                  |t        j                  d«       |«      }t        j                   | j                  t        j                  d«      z  t        j                  d«      k(  | j                  t        j                  d«      z  t        j                  d«      k7  «      }t        j                  |t        j                  d«      |«      }|S | j                   t        j"                  k(  rb| j                  t        j                  d«      z  t        j                  d«      k(  }t        j                  |t        j                  d«      |«      }|S )Né|   Úinféü   é   r   Únané   )r   r�   rˆ   ræ   r   rÇ   r’   rª   r«   r¬   rÕ   rU   r[   r“   r–   rÞ   Úlogical_andr˜   )Úvalue_handleÚvalue_floatÚpos_inf_maskÚneg_inf_maskÚnan_masks        r$   Ú_mxfp_value_handle_to_float32rô     sÈ  € Ø×ÑœRŸX™XÒ%Ü" <×#4Ñ#4Ó5‰ô
 %Ø×ÑØ×ÑÜ�J‰JÜ×Ñ×"Ñ"ó	
÷
 ‰$Œr�z‰zÓ
Ÿ6™6¤"§*¡*Ó-ð 	ð ×ÑœRŸ[™[Ó(Ø#×(Ñ(¬B¯H©H°T«NÑ:ˆÜ—h‘h˜|¬R¯Z©Z¸Ó->ÀÓLˆà#×(Ñ(¬B¯H©H°T«NÑ:ˆÜ—h‘h˜|¬b¯j©j¸Ó.?Ð-?ÀÓMˆä—>‘> <×#4Ñ#4´r·x±xÀ³~Ñ#EÌ"Ï(É(ÐSWË.Ñ"XØ#/×#4Ñ#4´r·x±xÀ³{Ñ#BÄrÇxÁxÐPQÃ{Ñ"RóTˆä—h‘h˜x¬¯©°EÓ):¸KÓHˆð
 Ðð	 
×	Ñ	œrŸ}™}Ò	,Ø×$Ñ$¤r§x¡x°£~Ñ5¼¿¹À$»ÑGˆÜ—h‘h˜x¬¯©°EÓ):¸KÓHˆàÐr&   c                ót   — t        j                  | |d«      } t        | «      }t        j                  |d|«      S )NrM   )rU   Úmoveaxisræ   )r   ÚaxisÚunpackeds      r$   Ú_unpack_e2m1rù   7  s3   € ô �;‰;�t˜T 2Ó&€DÜ˜DÓ!€HÜ�;‰;�x  TÓ*Ð*r&   c                óê  — |t         j                  j                  k(  r&|r|rdnd}n|rdnd}t        | j                  |«      }nt        | «      }|€|S t        |j                  «      }|rQt        j                  ||j                  d   |j                  d   z  d¬«      }t        j                  |dd«      }||z  S t        j                  ||j                  d   |j                  d   z  d¬«      }||z  S )NéþÿÿÿrM   ©r÷   )rª   ÚScaleDotElemTypeTYÚE2M1rù   r   rô   r×   rU   ÚrepeatrF   Úswapaxes)rï   Úscale_handleÚformat_enumÚk_packÚis_rhsÚunpack_axisr5   rÖ   s           r$   Ú_prepare_dot_scaled_operandr  ?  sÝ   € Ø”c×,Ñ,×1Ñ1Ò1ÙÙ &™"¨B‰Ká &™"¨BˆKÜ˜\×.Ñ.°Ó<‰ä-¨lÓ;ˆàÐØˆä˜×*Ñ*Ó+€Eáä—	‘	˜% §¡¨R¡°E·K±KÀ±OÑ!CÈ"ÔMˆÜ—‘˜E 2 rÓ*ˆð �5‰=Ðô —	‘	˜% §¡¨R¡°E·K±KÀ±OÑ!CÈ"ÔMˆà�5‰=Ð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)r7   r8   r9   Ústaticmethodr  r?   r&   r$   r	  r	  ]  s   „ àñpó ñpr&   r	  c                  óŽ  — e Zd Zej                  j
                  ej                  j
                  ej                  j                  ej                  j                  ej                  j                  ej                  j                  ej                  j                  ej                  j                  iZ
ej                  j                  ej                  j                  ej                  j                  ej                  j                  ej                  j                  ej                  j                  ej                  j                   ej                  j                   ej                  j"                  ej                  j"                  ej                  j$                  ej                  j$                  ej                  j&                  ej                  j&                  ej                  j(                  ej                  j(                  ej                  j*                  ej                  j*                  ej                  j,                  ej                  j,                  i
Zd›d„Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd	„ Z d
„ Z!d„ Z"d„ Z#d„ Z$d„ Z%d„ Z&d„ Z'd„ Z(d„ Z)d„ Z*d„ Z+d„ Z,d„ Z-d„ Z.d„ Z/d„ Z0d„ Z1d„ Z2d„ Z3d„ Z4d„ Z5d„ Z6d „ Z7d!„ Z8d"„ Z9d#„ Z:d$„ Z;d%„ Z<d&„ Z=d'„ Z>d(„ Z?d)„ Z@d*„ ZAd+„ ZBd,„ ZCd-„ ZDd.„ 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|df„ Z}eKZ~eKZ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œ	 	 	 	 	 d�d“„Z¬džd”„Z­dŸd•„Z®d d–„Z¯	 	 d¡d—„Z°d˜„ Z±	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 d¢d™„Z²yš)£ÚInterpreterBuilderc                ó˜   — d | _         t        «       | _        i | _        t        j
                  | j                  d<   d„ | j                  d<   y )NÚconvert_custom_typesc                 ó   — y)N)r   r   r   r?   )ÚlhsTypeÚrhsTypes     r$   ú<lambda>z-InterpreterBuilder.__init__.<locals>.<lambda>~  s   � r&   Úmin_dot_size)rn   rj   ÚoptionsÚcodegen_fnsr	  r  r"   s    r$   rJ   zInterpreterBuilder.__init__y  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_dimr   Úgrid_idx©r#   rÌ   ÚyÚzs       r$   Úset_grid_idxzInterpreterBuilder.set_grid_idx€  sf   € Ø�4—=‘= Ñ#Ò#ÜÐ/Ó0Ð0Ø�4—=‘= Ñ#Ò#ÜÐ/Ó0Ð0Ø�4—=‘= Ñ#Ò#ÜÐ/Ó0Ð0Ø˜A˜q˜	ˆ�r&   c                ó   — |||f| _         y r(   )r   )r#   ÚnxÚnyÚnzs       r$   Úset_grid_dimzInterpreterBuilder.set_grid_dim‰  s   € Ø˜R ˜ˆ�r&   c                ó"   — t         j                  S r(   )r�   r’   r"   s    r$   Úget_half_tyzInterpreterBuilder.get_half_tyŽ  ó   € Ü�z‰zÐr&   c                ó"   — t         j                  S r(   )r�   r•   r"   s    r$   Úget_bf16_tyzInterpreterBuilder.get_bf16_ty‘  ó   € Ü�{‰{Ðr&   c                ó"   — t         j                  S r(   )r�   r“   r"   s    r$   Úget_float_tyzInterpreterBuilder.get_float_ty”  r-  r&   c                ó"   — t         j                  S r(   )r�   r”   r"   s    r$   Úget_double_tyz InterpreterBuilder.get_double_ty—  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�  r7  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_ty£  r<  r&   c                ó"   — t         j                  S r(   )r�   rŠ   r"   s    r$   Úget_uint16_tyz InterpreterBuilder.get_uint16_ty¦  ó   € Ü�y‰yÐr&   c                ó"   — t         j                  S r(   )r�   r�   r"   s    r$   Úget_int32_tyzInterpreterBuilder.get_int32_ty©  r<  r&   c                ó"   — t         j                  S r(   )r�   rŒ   r"   s    r$   Úget_uint32_tyz InterpreterBuilder.get_uint32_ty¬  rA  r&   c                ó"   — t         j                  S r(   )r�   rŽ   r"   s    r$   Úget_int64_tyzInterpreterBuilder.get_int64_ty¯  r<  r&   c                ó"   — t         j                  S r(   )r�   r\   r"   s    r$   Úget_uint64_tyz InterpreterBuilder.get_uint64_ty²  rA  r&   c                ó"   — t         j                  S r(   )r�   r˜   r"   s    r$   Úget_fp8e4nv_tyz!InterpreterBuilder.get_fp8e4nv_tyµ  ó   € Ü�}‰}Ðr&   c                ó"   — t         j                  S r(   )r�   rš   r"   s    r$   Úget_fp8e4b15_tyz"InterpreterBuilder.get_fp8e4b15_ty¸  ó   € Ü�~‰~Ðr&   c                ó"   — t         j                  S r(   )r�   r™   r"   s    r$   Úget_fp8e4b8_tyz!InterpreterBuilder.get_fp8e4b8_ty»  rL  r&   c                ó"   — t         j                  S r(   )r�   r–   r"   s    r$   Úget_fp8e5_tyzInterpreterBuilder.get_fp8e5_ty¾  r0  r&   c                ó"   — t         j                  S r(   )r�   r—   r"   s    r$   Úget_fp8e5b16_tyz"InterpreterBuilder.get_fp8e5b16_tyÁ  rO  r&   c                ó.   — t        j                  ||«      S r(   )r�   r‚   )r#   Úelt_tyÚ
addr_spaces      r$   Ú
get_ptr_tyzInterpreterBuilder.get_ptr_tyÄ  s   € Ü�‰˜v zÓ2Ð2r&   c                ó.   — t        j                  ||«      S r(   )r�   r›   )r#   r   rF   s      r$   Úget_block_tyzInterpreterBuilder.get_block_tyÇ  s   € Ü�}‰}˜U EÓ*Ð*r&   c                ó~   — t        t        j                  |gt        j                  ¬«      t        j
                  «      S ©NrT   )r   rU   rÝ   Úbool_r�   r‘   ©r#   r5   s     r$   Úget_int1zInterpreterBuilder.get_int1Ê  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   rU   rÝ   rˆ   r�   r_  s     r$   Ú	get_uint8zInterpreterBuilder.get_uint8Í  ó$   € Üœ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   rU   rÝ   r‰   r�   r_  s     r$   Úget_int8zInterpreterBuilder.get_int8Ð  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   rU   rÝ   rŠ   r�   r_  s     r$   Ú
get_uint16zInterpreterBuilder.get_uint16Ó  ó$   € ÜœBŸH™H e W´B·I±IÔ>ÄÇ	Á	ÓJÐJr&   c                ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r]  )r   rU   rÝ   r‹   r�   r_  s     r$   Ú	get_int16zInterpreterBuilder.get_int16Ö  rc  r&   c                ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r]  )r   rU   rÝ   rŒ   r�   r_  s     r$   Ú
get_uint32zInterpreterBuilder.get_uint32Ù  rh  r&   c                ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r]  )r   rU   rÝ   r�   r�   r_  s     r$   Ú	get_int32zInterpreterBuilder.get_int32Ü  rc  r&   c                ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r]  )r   rU   rÝ   r\   r�   r_  s     r$   Ú
get_uint64zInterpreterBuilder.get_uint64ß  rh  r&   c                ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r]  )r   rU   rÝ   rŽ   r�   r_  s     r$   Ú	get_int64zInterpreterBuilder.get_int64â  rc  r&   c                ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r]  )r   rU   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   rU   rÝ   r“   r�   r_  s     r$   Úget_fp32zInterpreterBuilder.get_fp32è  ru  r&   c                ó~   — t        t        j                  |gt        j                  ¬«      t        j                  «      S r]  )r   rU   rÝ   r”   r�   r_  s     r$   Úget_fp64zInterpreterBuilder.get_fp64ë  ru  r&   c                óX   — t        t        j                  dgt        |«      ¬«      |«      S ©Nr   rT   )r   rU   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 NonerT   )r!  r   r   rU   rÝ   r�   r�   ©r#   r÷   s     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   rU   rÝ   r   r�   r�   r  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   rU   Ú	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   rU   r„  r   r)   r�   r‘   Úcreate_masked_store)r#   r†  Úvalr‡  rˆ  rŠ  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   rU   r§   r   Ú_interpreterÚload)
r#   ÚptrsrŠ  r‹  Úcache_modifierÚeviction_policyr‰  Údtype_ttÚdtype_npÚrets
             r$   r…  z%InterpreterBuilder.create_masked_load  si   € Ø×&Ñ&Ó(ˆÜ  Ó*ˆØˆ=Ü ¤§¡¨t¯y©yÀÔ!IÈ8ÓTˆEÜ×Ñ §	¡	¨4¯9©9°e·j±jÀ(ÓKˆÜ˜C Ó*Ð*r&   c                ól   — t        j                  |j                  |j                  |j                  «      S r(   )r’  Ústorer   )r#   r”  r5   rŠ  r•  r–  s         r$   rŽ  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   r]   r�   r•   r“   rÇ   r   rÕ   rž   r   r[   )r#   ÚsrcÚdst_typeÚsrc_element_typeÚdst_element_typer   s         r$   Ú	cast_implzInterpreterBuilder.cast_impl  s©   € ØŸ9™9×+Ñ+ÐØ#Ÿ?™?ÐØ¤§¡Ò+Ð0@ÄBÇJÁJÒ0NØ¤§
¡
Ò*Ð/?Ä2Ç;Á;Ò/NÜ! #§(¡(Ð,<Ð>NÐPTÓU×ZÑZÔ[hÐiqÓ[rÓsˆDÜ  h§o¡oÓ6Ð6ä §¡§¡´¸hÓ0GÓ HÈ(Ï/É/ÓZÐZr&   c                ó&   — | j                  ||«      S r(   ©r¡  ©r#   r�  rž  s      r$   r  zInterpreterBuilder.<lambda>  ó   € °$·.±.ÀÀhÓ2O€ r&   c                ó&   — | j                  ||«      S r(   r£  r¤  s      r$   r  zInterpreterBuilder.<lambda>  r¥  r&   c                ó&   — | j                  ||«      S r(   r£  r¤  s      r$   r  zInterpreterBuilder.<lambda>  r¥  r&   c                ó&   — | j                  ||«      S r(   r£  r¤  s      r$   r  zInterpreterBuilder.<lambda>  r¥  r&   c                ó&   — | j                  ||«      S r(   r£  r¤  s      r$   r  zInterpreterBuilder.<lambda>  s   € °·±¸sÀHÓ0M€ r&   c                ó&   — | j                  ||«      S r(   r£  r¤  s      r$   r  zInterpreterBuilder.<lambda>  r¥  r&   c                ó&   — | j                  ||«      S r(   r£  )r#   r�  rž  Ú	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   r]   rÇ   r   rÕ   rž   r   )r#   r�  rž  r°   rŸ  r   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ž   r]   r¤  s      r$   Úcreate_bitcastz!InterpreterBuilder.create_bitcast(  s%   € Ü˜CŸH™HŸM™M¬-¸Ó*AÓBÀHÇOÁOÓTÐTr&   c                óÌ   —  ||j                   |j                   «      }|j                  j                  }t        ||«      s|j	                  t        |«      «      }t        ||«      S r(   ©r   r   r]   r   r[   rž   r   )r#   ÚlhsÚrhsÚoprÆ   r„   s         r$   Ú	binary_opzInterpreterBuilder.binary_op,  sO   € Ù�C—H‘H˜cŸh™hÓ'ˆØ—9‘9×#Ñ#ˆä% f¨hÔ7Ø—]‘]¤=°Ó#:Ó;ˆFä˜F HÓ-Ð-r&   c                óD   — | j                  ||t        j                  «      S r(   ©rµ  rU   Úadd©r#   r²  r³  s      r$   r  zInterpreterBuilder.<lambda>5  s   € ¨¯©¸¸SÄ"Ç&Á&Ó)I€ r&   c                óD   — | j                  ||t        j                  «      S r(   ©rµ  rU   Úmultiplyr¹  s      r$   r  zInterpreterBuilder.<lambda>6  ó   € ¨¯©¸¸SÄ"Ç+Á+Ó)N€ r&   c                óD   — | j                  ||t        j                  «      S r(   ©rµ  rU   Údivider¹  s      r$   r  zInterpreterBuilder.<lambda>7  s   € ¨¯©¸¸SÄ"Ç)Á)Ó)L€ r&   c                óD   — | j                  ||t        j                  «      S r(   ©rµ  rU   Úfmodr¹  s      r$   r  zInterpreterBuilder.<lambda>8  ó   € ¨¯©¸¸SÄ"Ç'Á'Ó)J€ r&   c                óD   — | j                  ||t        j                  «      S r(   ©rµ  rU   Úsubtractr¹  s      r$   r  zInterpreterBuilder.<lambda>9  r½  r&   c                óB   — | j                  |t        j                  «      S r(   )Úunary_oprU   Únegative©r#   r­   s     r$   r  zInterpreterBuilder.<lambda>:  s   €  d§m¡m°E¼2¿;¹;Ó&G€ 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>A  rÄ  r&   c                óD   — | j                  ||t        j                  «      S r(   r·  r¹  s      r$   r  zInterpreterBuilder.<lambda>B  s   € ¨¯©°s¸CÄÇÁÓ(H€ r&   c                óD   — | j                  ||t        j                  «      S r(   rÆ  r¹  s      r$   r  zInterpreterBuilder.<lambda>C  rÍ  r&   c                óD   — | j                  ||t        j                  «      S r(   )rµ  rU   Ú
left_shiftr¹  s      r$   r  zInterpreterBuilder.<lambda>D  s   € ¨¯©°s¸CÄÇÁÓ(O€ r&   c                óD   — | j                  ||t        j                  «      S r(   )rµ  rU   Úright_shiftr¹  s      r$   r  zInterpreterBuilder.<lambda>E  s   € ¨¯©¸¸SÄ"Ç.Á.Ó)Q€ r&   c                óD   — | j                  ||t        j                  «      S r(   ©rµ  rU   r©   r¹  s      r$   r  zInterpreterBuilder.<lambda>F  ó   € ¨$¯.©.¸¸cÄ2Ç:Á:Ó*N€ r&   c                óD   — | j                  ||t        j                  «      S r(   rÝ  r¹  s      r$   r  zInterpreterBuilder.<lambda>G  rÞ  r&   c                óD   — | j                  ||t        j                  «      S r(   rÝ  r¹  s      r$   r  zInterpreterBuilder.<lambda>H  ó   € ¨T¯^©^¸CÀÄbÇjÁjÓ-Q€ r&   c                óD   — | j                  ||t        j                  «      S r(   )rµ  rU   Úfminr¹  s      r$   r  zInterpreterBuilder.<lambda>I  ó   € ¨D¯N©N¸3ÀÄRÇWÁWÓ,M€ r&   c                óD   — | j                  ||t        j                  «      S r(   ©rµ  rU   r¨   r¹  s      r$   r  zInterpreterBuilder.<lambda>J  rÞ  r&   c                óD   — | j                  ||t        j                  «      S r(   ræ  r¹  s      r$   r  zInterpreterBuilder.<lambda>K  rÞ  r&   c                óD   — | j                  ||t        j                  «      S r(   ræ  r¹  s      r$   r  zInterpreterBuilder.<lambda>L  rá  r&   c                óD   — | j                  ||t        j                  «      S r(   )rµ  rU   Úfmaxr¹  s      r$   r  zInterpreterBuilder.<lambda>M  rä  r&   c                óD   — | j                  ||t        j                  «      S r(   ©rµ  rU   Ú
less_equalr¹  s      r$   r  zInterpreterBuilder.<lambda>N  ó   € ¨D¯N©N¸3ÀÄRÇ]Á]Ó,S€ r&   c                óD   — | j                  ||t        j                  «      S r(   ©rµ  rU   Úlessr¹  s      r$   r  zInterpreterBuilder.<lambda>O  rä  r&   c                óD   — | j                  ||t        j                  «      S r(   ©rµ  rU   Úgreater_equalr¹  s      r$   r  zInterpreterBuilder.<lambda>P  ó   € ¨D¯N©N¸3ÀÄR×EUÑEUÓ,V€ r&   c                óD   — | j                  ||t        j                  «      S r(   ©rµ  rU   Úgreaterr¹  s      r$   r  zInterpreterBuilder.<lambda>Q  ó   € ¨D¯N©N¸3ÀÄRÇZÁZÓ,P€ r&   c                óD   — | j                  ||t        j                  «      S r(   rì  r¹  s      r$   r  zInterpreterBuilder.<lambda>R  rî  r&   c                óD   — | j                  ||t        j                  «      S r(   rð  r¹  s      r$   r  zInterpreterBuilder.<lambda>S  rä  r&   c                óD   — | j                  ||t        j                  «      S r(   ró  r¹  s      r$   r  zInterpreterBuilder.<lambda>T  rõ  r&   c                óD   — | j                  ||t        j                  «      S r(   r÷  r¹  s      r$   r  zInterpreterBuilder.<lambda>U  rù  r&   c                óD   — | j                  ||t        j                  «      S r(   ©rµ  rU   Úequalr¹  s      r$   r  zInterpreterBuilder.<lambda>V  s   € ¨4¯>©>¸#¸sÄBÇHÁHÓ+M€ r&   c                óD   — | j                  ||t        j                  «      S r(   ©rµ  rU   Ú	not_equalr¹  s      r$   r  zInterpreterBuilder.<lambda>W  s   € ¨4¯>©>¸#¸sÄBÇLÁLÓ+Q€ r&   c                óD   — | j                  ||t        j                  «      S r(   rð  r¹  s      r$   r  zInterpreterBuilder.<lambda>X  rä  r&   c                óD   — | j                  ||t        j                  «      S r(   r÷  r¹  s      r$   r  zInterpreterBuilder.<lambda>Y  rù  r&   c                óD   — | j                  ||t        j                  «      S r(   rì  r¹  s      r$   r  zInterpreterBuilder.<lambda>Z  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>a  rõ  r&   c                óD   — | j                  ||t        j                  «      S r(   rÿ  r¹  s      r$   r  zInterpreterBuilder.<lambda>b  r	  r&   c                óD   — | j                  ||t        j                  «      S r(   r  r¹  s      r$   r  zInterpreterBuilder.<lambda>c  r  r&   c                óD   — | j                  ||t        j                  «      S r(   )rµ  rU   Úbitwise_andr¹  s      r$   r  zInterpreterBuilder.<lambda>d  ó   € ¨¯©°s¸CÄÇÁÓ(P€ r&   c                óD   — | j                  ||t        j                  «      S r(   )rµ  rU   Úbitwise_xorr¹  s      r$   r  zInterpreterBuilder.<lambda>e  r  r&   c                óD   — | j                  ||t        j                  «      S r(   )rµ  rU   Ú
bitwise_orr¹  s      r$   r  zInterpreterBuilder.<lambda>f  s   €  t§~¡~°c¸3ÄÇÁÓ'N€ r&   c                óÆ   — t        |j                  t        j                  |j                  |j                  «      z
  |j                  z  |j                  j
                  «      S r(   )r   r   rU   rÃ  r   r]   r¹  s      r$   rÑ  zInterpreterBuilder.create_idivj  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                  |«      }t	        ||j                  «      }t	        ||j                  «      }| j                  ||t        j                  «      S r(   )r�   r   r   r[   r   rµ  rU   rÛ  )r#   r²  r³  Ú	lhs_dtypeÚ	rhs_dtypeÚlhs_dataÚrhs_datas          r$   Úcreate_ashrzInterpreterBuilder.create_ashrp  s�   € ä(¨¯©¯©Ó8ˆ	Ü(¨¯©¯©Ó8ˆ	Ø—8‘8—?‘? 9Ó-ˆØ—8‘8—?‘? 9Ó-ˆÜ˜8 S§Y¡YÓ/ˆÜ˜8 S§Y¡YÓ/ˆØ�~‰~˜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   rU   rŽ   r\   r   Únp_umulhi_u64r]   r¡   r    r[   r¼  )r#   r²  r³  r   Úcompute_dtyper  r  Úret_datas           r$   Úcreate_umulhiz InterpreterBuilder.create_umulhiz  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                óâ   —  ||j                   |j                   |j                   «      }|j                  j                  }t        ||«      s|j	                  t        |«      «      }t        ||«      S r(   r±  )r#   r²  r³  r‹  r´  rÆ   r„   s          r$   Ú
ternary_opzInterpreterBuilder.ternary_op†  sU   € Ù�C—H‘H˜cŸh™h¨¯
©
Ó3ˆØ—;‘;×%Ñ%ˆä% f¨hÔ7Ø—]‘]¤=°Ó#:Ó;ˆFä˜F HÓ-Ð-r&   c                ó   — |t         j                  j                  k(  r@| j                  | j                  ||t        j
                  «      |t        j                  «      S | j                  |||t        j                  «      S r(   )	rª   ÚPROPAGATE_NANÚNONErµ  rU   rê  rã  r&  Úclip)r#   ÚargÚloÚhiÚpropagate_nanss        r$   Úcreate_clampfz InterpreterBuilder.create_clampf�  sW   € ØœS×.Ñ.×3Ñ3Ò3Ø—>‘> $§.¡.°°b¼"¿'¹'Ó"BÀBÌÏÉÓPÐPà—?‘? 3¨¨B´·±Ó8Ð8r&   c                óF   — | j                  |||t        j                  «      S r(   )r&  rU   rÞ   )r#   Úcondr²  r³  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(   ©r   r   r   r]   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(   r3  )r#   r+  r´  s      r$   rÉ  z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¡   rU   r   rÕ   rž   r   r]   )r#   r+  r—  Úmask_bitwidthÚnp_uint_dtyper   rŠ  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É  rU   Úcos©r#   r+  s     r$   r  zInterpreterBuilder.<lambda>¨  ó   €  4§=¡=°´b·f±fÓ#=€ r&   c                óB   — | j                  |t        j                  «      S r(   )rÉ  rU   Úexpr<  s     r$   r  zInterpreterBuilder.<lambda>©  r=  r&   c                óB   — | j                  |t        j                  «      S r(   )rÉ  rU   Úexp2r<  s     r$   r  zInterpreterBuilder.<lambda>ª  ó   €  D§M¡M°#´r·w±wÓ$?€ r&   c                óB   — | j                  |t        j                  «      S r(   )rÉ  rU   Úabsr<  s     r$   r  zInterpreterBuilder.<lambda>«  s   €  D§M¡M°#´r·v±vÓ$>€ r&   c                óB   — | j                  |t        j                  «      S r(   )rÉ  rU   Úfloorr<  s     r$   r  zInterpreterBuilder.<lambda>¬  s   €  T§]¡]°3¼¿¹Ó%A€ r&   c                óB   — | j                  |t        j                  «      S r(   )rÉ  rU   Úceilr<  s     r$   r  zInterpreterBuilder.<lambda>­  rB  r&   c                óB   — | j                  |t        j                  «      S r(   )rÉ  rU   Úlogr<  s     r$   r  zInterpreterBuilder.<lambda>®  r=  r&   c                óB   — | j                  |t        j                  «      S r(   )rÉ  rU   Úlog2r<  s     r$   r  zInterpreterBuilder.<lambda>¯  rB  r&   c                óB   — | j                  |t        j                  «      S r(   ©rÉ  rU   Úsqrtr<  s     r$   r  zInterpreterBuilder.<lambda>°  s   € ¨D¯M©M¸#¼r¿w¹wÓ,G€ r&   c                óB   — | j                  |t        j                  «      S r(   rN  r<  s     r$   r  zInterpreterBuilder.<lambda>±  rB  r&   c                óB   — | j                  |t        j                  «      S r(   )rÉ  rU   Ú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   rU   r“   Únp_erf_fp32Únp_erf_fp64r   r]   )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   rU   rO  r   r   r]   r<  s     r$   Úcreate_rsqrtzInterpreterBuilder.create_rsqrt¸  s+   € Ü˜A¤§¡¨¯©Ó 1Ñ1°3·9±9×3CÑ3CÓDÐDr&   c                ót   — t        |j                  j                  |«      |j                  j                  «      S r(   )r   r   rZ   r   r]   )r#   r+  rF   Úallow_reorders       r$   r  zInterpreterBuilder.<lambda>¼  s*   € ¼\È#Ï(É(×JZÑJZÐ[`ÓJaÐcf×clÑcl×csÑcsÓ=t€ r&   c                ó~   — t        t        j                  |j                  |«      |j                  j
                  «      S r(   )r   rU   Ú	transposer   r   r]   )r#   r+  Úperms      r$   Úcreate_transzInterpreterBuilder.create_trans¾  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   rT   )r   r   r!   Úis_floatingrÇ   r�   r’   rÕ   rU   r   Úmatmulr]   )r#   rÐ   rÑ   ÚdÚinput_precisionÚmax_num_imprecise_accÚa_dataÚb_datas           r$   Ú
create_dotzInterpreterBuilder.create_dotÁ  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   rU   rY   r�   r�   )r#   Úret_tyÚstartÚstops       r$   Úcreate_make_rangez$InterpreterBuilder.create_make_rangeÊ  s$   € ÜœBŸI™I e¨T¼¿¹ÔBÄBÇHÁHÓMÐMr&   c                óZ  — |€=t        t        j                  |j                  t        ¬«      t
        j                  «      }t        j                  |j                  |j                  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 )NrT   r   )ÚbinsrX   Úweights)r   rU   r„  r   r)   r�   r‘   r   rÞ   r§   Ú	histogramÚlogical_notÚsumr�   )r#   r   ro  rŠ  Údummy_weightsrq  s         r$   Úcreate_histogramz#InterpreterBuilder.create_histogramÍ  sÀ   € Øˆ<Ü¤§¡¨T¯Y©Y¼dÔ CÄRÇWÁWÓMˆDô Ÿ™ T§Y¡Y°d·i±i·o±oÔFˆô �x‰x˜Ÿ	™	 4§9¡9¬b¯m©m¸D¿I¹IÓ.FÓGˆÜ—L‘L ¨D¸¸D¸	È=ÔYÐZ[Ñ\ˆ	à�!‹œŸ™ t§y¡yÓ1×5Ñ5Ó7Ñ7‹Ü˜I¤r§x¡xÓ0Ð0r&   c                ó–   — t        t        j                  |j                  |j                  |¬«      |j                  j
                  «      S )Nrü   )r   rU   Útake_along_axisr   r   r]   )r#   r�  Úindicesr÷   s       r$   Úcreate_gatherz InterpreterBuilder.create_gatherß  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   r[   rU   r\   r   )r#   r†  Úoffsetr—  Úelement_bitwidthÚelement_bytewidths         r$   Úcreate_addptrz InterpreterBuilder.create_addpträ  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                ó~   — t        t        j                  |j                  |«      |j                  j
                  «      S r(   )r   rU   Úexpand_dimsr   r   r]   )r#   r+  r÷   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   rU   rV   r   r   r]   )r#   r+  rF   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   rU   Úconcatenater   r   r]   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 )NrM   rü   )r   rU   Ústackr   r   r]   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   r3  )r#   r�  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{  )rF   r€   r   r�   r›   r   rU   Úfullr   rž   r]   )r#   rj  r+  rF   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           	     ó®   — t        t        j                  d|j                  d   t	        |j
                  «      ¬«      |j
                  j                  «      S )N©r   r   rT   )r   rU   rŽ  r   rž   r   r]   r<  s     r$   Úcreate_unsplatz!InterpreterBuilder.create_unsplat  s:   € ÜœBŸG™G E¨3¯8©8°A©;¼mÈCÏIÉIÓ>VÔWÐY\×YbÑYb×YiÑYiÓjÐjr&   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   r’  Ú
atomic_casr   r   r]   )r#   r†  Úcmpr�  Ú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 r”  )	Úir_rmw_op_to_interpreter_rmw_opr   r•  r   r’  Ú
atomic_rmwr   r   r]   )r#   ÚrmwOpr†  r�  rŠ  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 moder¡  )r#   Ú	inlineAsmÚconstraintsÚvaluesr|  r¨  Úpacks          r$   Úcreate_inline_asmz$InterpreterBuilder.create_inline_asm  ó   € Ü!Ð"PÓQÐQr&   c                ó*  — d| j                   d   › d| j                   d   › d| j                   d   › d�}|r|d|› �z  }|rt        j                  dd	„ i¬
«       |D ]  }t        |d|j                  › �z   «       Œ |rt        j                  d ¬
«       y y )Nú(r   z, r   r   ú)ú r*   c                ó   — d| d›�S )NÚ0xÚ02xr?   rË   s    r$   r  z1InterpreterBuilder.create_print.<locals>.<lambda>$  s   € ¸bÀÀ3À¸L€ r&   )Ú	formatter)r!  rU   Úset_printoptionsÚprintr   )r#   ÚprefixÚhexr­  ÚisSignedÚmsgr5   s          r$   Úcreate_printzInterpreterBuilder.create_print  s£   € ð
 �$—-‘- Ñ"Ð# 2 d§m¡m°AÑ&6Ð%7°r¸$¿-¹-ÈÑ:JÐ9KÈ1ÐMˆÙØ�Q�v�h�<ÑˆCÙÜ×Ñ¨5Ñ2HÐ*IÕJØò 	*ˆEÜ�#˜!˜EŸJ™J˜<Ð(Ñ(Õ)ð	*áÜ×Ñ¨$Ö/ð r&   c                ó   — |sJ |› «       ‚y r(   r?   )r#   Ú	conditionÚmessages      r$   Úcreate_assertz InterpreterBuilder.create_assert*  s   € áÐ&˜W˜IÓ&‰yr&   c                ó   — |sJ d«       ‚y )NzAssume failedr?   )r#   rÁ  s     r$   Úcreate_assumez InterpreterBuilder.create_assume.  s   € ÙÐ)˜/Ó)‰yr&   c                 ó   — y r(   r?   r"   s    r$   Úcreate_barrierz!InterpreterBuilder.create_barrier1  s   € àr&   c                óD   — t        |||||«      }|j                  «        |S r(   )rA   rR   )r#   rC   rF   rG   Útensor_shaper¬  rI   Údescs           r$   Úcreate_make_tensor_descriptorz0InterpreterBuilder.create_make_tensor_descriptor5  s"   € ä  e¨W°lÀGÓLˆØ�‰ŒØˆr&   c                ó  — t        |t        «      sJ ‚|j                  |«      \  }}|j                  «       }t	        |«      }|j
                  }	|	t        j                  j                  k(  r,t        t        j                  |j                  |¬«      |«      }
na|	t        j                  j                  k(  r6t        t        j                  |j                  t        d«      |¬«      |«      }
nt!        d|	› �«      ‚| j#                  |||
||d¬«      S )NrT   rì   zunsupported padding F)r•  r–  r‰  )r€   rA   re   r/   rž   rI   rª   ÚPADDING_OPTIONÚPAD_ZEROr   rU   r§   r   ÚPAD_NANÚ	full_likeÚfloatr   r…  )r#   rÊ  rx  r•  r–  r”  rŠ  r—  r˜  rI   r‹  s              r$   Úcreate_descriptor_loadz)InterpreterBuilder.create_descriptor_load;  sè   € ä˜$Ô 0Ô1Ð1Ð1Ø×.Ñ.¨wÓ7‰
ˆˆdØ×&Ñ&Ó(ˆÜ  Ó*ˆØ—,‘,ˆØ”c×(Ñ(×1Ñ1Ò1Ü ¤§¡¨t¯y©yÀÔ!IÈ8ÓT‰EØœ×*Ñ*×2Ñ2Ò2Ü ¤§¡¨d¯i©i¼¸u»ÈXÔ!VÐX`Óa‰EäÐ3°G°9Ð=Ó>Ð>Ø×&Ñ& t¨T°5ÈØ7FÐTYð 'ó [ð 	[r&   c                óT   — |j                  |«      \  }}| j                  |||d d «      S r(   )re   rŽ  )r#   rÊ  r5   rx  r”  rŠ  s         r$   Úcreate_descriptor_storez*InterpreterBuilder.create_descriptor_storeK  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   rM   rT   )rC   r   r-   rž   rU   Úzerosr   rF   rH   Ú	enumerater   r�   r�   rÒ  )r#   rÊ  Ú	x_offsetsÚy_offsetr|  r   Únp_dtypeÚresultr•  r–  r½   Úx_offsetrx  s                r$   Úcreate_descriptor_gatherz+InterpreterBuilder.create_descriptor_gatherO  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Ê  r5   rØ  rÙ  r½   rÜ  Úslicerx  s	            r$   Úcreate_descriptor_scatterz,InterpreterBuilder.create_descriptor_scatterZ  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|› �«      ‚)Nr|   r   rM   rT   Tzunsupported type )rž   Únamer   rU   rŽ  r]   r^  Ú	TypeError)r#   r|  Únp_types      r$   Úget_all_ones_valuez%InterpreterBuilder.get_all_ones_valuea  ss   € Ü Ó%ˆØ�G—L‘LÑ Ü¤§¡¨¨2°WÔ =¸t¿{¹{ÓKÐKØœŸ™Ò Ü¤§¡¨¨4°wÔ ?ÀÇÁÓMÐMäÐ/°¨vÐ6Ó7Ð7r&   c                ó¼   — t        ||||d¬«      }t        ||||	d¬«      }t        j                  ||«      |
j                  z   }t	        |t
        j                  «      S )NF)r  T)r  rU   rb  r   r   r�   r“   )r#   r²  Úlhs_scale_handleÚlhs_format_enumr³  Úrhs_scale_handleÚrhs_format_enumÚ	fast_mathÚ
lhs_k_packÚ
rhs_k_packÚ
acc_handler  r  rÛ  s                 r$   Úcreate_dot_scaledz$InterpreterBuilder.create_dot_scaledj  sX   € ô
 /¨sÐ4DÀoÐWaÐjoÔpˆÜ.¨sÐ4DÀoÐWaÐjnÔoˆä—‘˜8 XÓ.°·±Ñ@ˆÜ˜F¤B§J¡JÓ/Ð/r&   N©ÚreturnÚNone)Úzero)rC   r   rF   rf   rG   rf   rÉ  rg   r¬  r)   rI   rw   )rÊ  rA   rx  rf   )rÊ  rA   r5   r   rx  rf   )rÊ  rA   rØ  r   rÙ  r   )rÊ  rA   r5   r   rØ  r   rÙ  r   )r²  r   rç  úOptional[TensorHandle]rè  ú_ir.ScaleDotElemTypeTYr³  r   ré  rô  rê  rõ  rë  r)   rì  r)   rí  r)   rî  r   rñ  r   )³r7   r8   r9   rª   ÚMEM_SEMANTICÚACQUIREr’  ÚRELEASEÚRELAXEDÚACQUIRE_RELEASEr•  Ú	ATOMIC_OPÚADDÚRMW_OPÚFADDÚMINÚUMINÚMAXÚUMAXÚANDÚORÚXORÚXCHGrœ  rJ   r%  r*  r,  r/  r2  r4  r6  r9  r;  r>  r@  rC  rE  rG  rI  rK  rN  rQ  rS  rU  rY  r[  r`  rb  re  rg  rj  rl  rn  rp  rr  rt  rw  ry  r}  r€  r‚  rŒ  r�  r…  rŽ  r¡  Ú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  r¯  rµ  Úcreate_faddÚcreate_fmulÚcreate_fdivÚcreate_fremÚcreate_fsubÚcreate_fnegÚ
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&  r/  Úcreate_selectr4  rÉ  r9  Ú
create_cosÚ
create_expÚcreate_exp2Úcreate_iabsÚcreate_floorÚcreate_ceilÚ
create_logÚcreate_log2Úcreate_precise_sqrtÚcreate_sqrtÚ
create_sinrV  rY  Úcreate_reshaper_  rh  rm  ru  ry  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  d  sV  „ à×Ñ× Ñ  ,×";Ñ";×"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ò.ñ J€KÙN€KÙL€KÙJ€KÙN€KÙG€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ÙM€NÙN€LÙN€LÙQ€OÙM€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ò.ò9ñ [€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òNòPòTòWòtòlòkòjòròZòRò0ò'ò*òð
 `fðØ4=ðØJNðØY\óó[ó Gó	+ð?Ø,8ó?ò8ð	0Ø+Að	0ØHTð	0à,Bð	0àUkð	0ð &*ð	0ð 8<ð	0ð JNð	0ð '3ð		0ð 8Dô		0r&   r  r   Úinterpreter_semanticc                  ó(   — e Zd ZdZdd„Zdd„Zdd„Zy)Ú_LangPatchScopez2Tracks patched attributes so they can be restored.c                ó   — g | _         y r(   )Ú_changesr"   s    r$   rJ   z_LangPatchScope.__init__~  s	   € Ø:<ˆ�r&   c                ó|   — t        ||t        «      }| j                  j                  |||f«       t	        |||«       y r(   )r¡   Ú_MISSINGrR  ÚappendÚsetattr)r#   Úobjrâ  r5   Úoriginals        r$   r6   z_LangPatchScope.set_attr�  s5   € Ü˜3 ¤hÓ/ˆØ�‰×Ñ˜c 4¨Ð2Ô3Ü��T˜5Õ!r&   c                ó¸   — | j                   rN| j                   j                  «       \  }}}|t        u rt        ||«       nt	        |||«       | j                   rŒMy y r(   )rR  ÚpoprT  ÚdelattrrV  )r#   rW  râ  rX  s       r$   Úrestorez_LangPatchScope.restore†  sI   € Ø�mŠmØ"&§-¡-×"3Ñ"3Ó"5ÑˆC��xØœ8Ñ#Ü˜˜TÕ"ä˜˜T 8Ô,ð �m�mr&   Nrð  )rW  Úobjectrâ  rw   r5   r]  rñ  rò  )r7   r8   r9   r:   rJ   r6   r\  r?   r&   r$   rP  rP  {  s   „ Ù<ó=ó"ô
-r&   rP  c                ó6   — |dœd„}|j                  | ||«       y )N)Úmemberc           
     ó|   —  | |i |j                  «       D ��ci c]  \  }}|dk7  r||“Œ c}}¤dt        i¤ŽS c c}}w )Nr  )ÚitemsrN  )r_  ÚargsÚkwargsÚkÚvs        r$   r  z_patch_attr.<locals>.<lambda>�  sZ   € ¹Àð :wàMSÏ\É\Ë^÷AVÙEIÀQÈØDEÈÒDTð BCÀAÁó AVñ:wô bvò:w€ ùóAVs   ˜8)r6   )rW  râ  r_  r  r™  Ú
new_members         r$   Ú_patch_attrrg  �  s   € Ø&,ó x€Jð 
‡N�N�3˜˜jÕ)r&   c                óœ   — t        j                  | «      D ]4  \  }}t        j                  j	                  |«      sŒ&t        | ||||«       Œ6 y r(   )ÚinspectÚ
getmembersr�   ÚcoreÚ
is_builtinrg  )Úpkgr  r™  râ  r_  s        r$   Ú_patch_builtinrn  —  sB   € Ü×*Ñ*¨3Ó/ò ;‰ˆˆfÜ�7‰7×Ñ˜fÕ%Ü˜˜T 6¨7°EÕ:ñ;r&   c                óî   ‡— d„ Šd„ }|j                  | dd„ «       |j                  | dˆfd„«       |j                  | dd„ «       |j                  | d	d
„ «       |j                  | 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 )Nrû   rM   )r   rU   r]  r  r   r   r|  Úis_blockÚlistrF   r�   rk  r›   r  )r#   r  rH   Ú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&   Ú	__index__c                ó\   — t        | j                  j                  j                  «       «      S r(   )r|   r  r   Úsqueezer"   s    r$   r  z$_patch_lang_tensor.<locals>.<lambda>­  s   € ´S¸¿¹×9IÑ9I×9QÑ9QÓ9SÓ5T€ r&   r+   c                ó   •—  ‰| «      S r(   r?   )r#   rr  s    €r$   r  z$_patch_lang_tensor.<locals>.<lambda>®  s   ø€ ±I¸d³O€ r&   Ú__repr__c                ó@   — t        | j                  j                  «      S r(   )Úreprr  r   r"   s    r$   r  z$_patch_lang_tensor.<locals>.<lambda>¯  s   € ´D¸¿¹×9IÑ9IÓ4J€ r&   Ú__str__c                ó@   — t        | j                  j                  «      S r(   )rw   r  r   r"   s    r$   r  z$_patch_lang_tensor.<locals>.<lambda>°  s   € ´3°t·{±{×7GÑ7GÓ3H€ r&   r   )r6   Úproperty)r  r™  rw  rr  s      @r$   Ú_patch_lang_tensorr‚  �  sh   ø€ ò6ò.ð 
‡N�N�6˜;Ñ(TÔUØ	‡N�N�6˜:Ó'CÔDØ	‡N�N�6˜:Ñ'JÔKØ	‡N�N�6˜9Ñ&HÔIØ	‡N�N�6˜3¤¨Ó 8Õ9r&   c                  ó0   — e Zd Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zy)ÚReduceScanOpInterfacec                ó    — || _         || _        y r(   )r÷   Ú
combine_fn)r#   r÷   r†  s      r$   rJ   zReduceScanOpInterface.__init__¶  s   € ØˆŒ	Ø$ˆ�r&   c                óH   — |� |t        |«      k\  rt        d|› d|› �«      ‚y y )Nzaxis z out of bounds for shape )rD   r   )r#   rF   r÷   s      r$   Ú
check_axisz ReduceScanOpInterface.check_axisº  s4   € ØÐ ¬¨E«
Ò 2Ü˜u T FÐ*CÀEÀ7ÐKÓLÐLð !3Ðr&   c                óÔ   — |D ]c  }t        |t        j                  j                  «      st	        dt        |«      › �«      ‚| j                  |j                  | j                  «       Œe y )Nzinput must be a tensor, got )	r€   r�   rk  r  r   r|  rˆ  rF   r÷   )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 )NrF   rT   )rž   r.   rF   r[   r�   r›   ru  rU   rÝ   rk  r  r   r]   )r#   r™  r   rÚ  Úret_types        r$   Ú	to_tensorzReduceScanOpInterface.to_tensorÄ  sy   € Ü  Ó'ˆÜ�3˜Ô  S§Y¢YØ—*‘*˜XÓ&ˆCÜ—}‘} U¬D°·±«OÓ<‰Hä—(‘(˜C˜5¨Ô1ˆCØˆHÜ�w‰w�~‰~œl¨3°·±Ó=¸xÓHÐHr&   c                ó   — t        d«      ‚)Nz,apply_impl must be implemented by subclassesr¡  rË  s     r$   Ú
apply_implz ReduceScanOpInterface.apply_implÎ  r°  r&   c                óØ   — t        |t        «      s| j                  |f«      d   S | j                  |«       | j	                  |«      }t        |t
        t        f«      rt        |«      S |fS ©Nr   )r€   ÚtupleÚapplyrŠ  r�  ru  )r#   r­   r™  s      r$   r“  zReduceScanOpInterface.applyÑ  s\   € Ü˜%¤Ô'Ø—:‘:˜u˜iÓ(¨Ñ+Ð+Ø×Ñ˜%Ô Ø�o‰o˜eÓ$ˆÜ'¨¬d´E¨]Ô;Œu�S‹zÐHÀ#ÀÐHr&   N)	r7   r8   r9   rJ   rˆ  rŠ  r�  r�  r“  r?   r&   r$   r„  r„  ´  s$   „ ò%òMò2òIòRó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(   )ÚsuperrJ   Ú	keep_dims)r#   r÷   r†  r˜  Ú	__class__s       €r$   rJ   zReduceOps.__init__Û  s   ø€ Ü‰Ñ˜˜zÔ*Ø"ˆ�r&   c                óô   — g }|D ]e  }|�|j                  |«       Œd}|j                  | j                  |j                  j                  j	                  «       |j
                  «      «       Œg t        |«      |fS r‘  )rU  r�  r  r   Úflattenr   r’  )r#   r­   r÷   r™  r   s        r$   ÚunravelzReduceOps.unravelß  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   rT   c              3  óh   •K  — | ])  \  }}‰j                  |‰   ‰|   j                  «      –— Œ+ y ­wr(   ©r�  r   )Ú.0Úiirc  r­   Úinput_indexr#   s      €€€r$   ú	<genexpr>z+ReduceOps.generic_reduce.<locals>.<genexpr>ø  s/   øè ø€ ÒsÑTYÐTVÐXY §¡¨q°©~¸uÀR¹y¿¹× OÑsùó   ƒ/2c              3  óh   •K  — | ])  \  }}‰j                  |‰   ‰|   j                  «      –— Œ+ y ­wr(   rŸ  )r   ÚoiÚor­   Úoutput_indexr#   s      €€€r$   r£  z+ReduceOps.generic_reduce.<locals>.<genexpr>þ  s/   øè ø€ Ò!wÑW\ÐWYÐ[\ $§.¡.°°<±À%ÈÁ)Ç/Á/×"RÑ!wùr¤  )r÷   rœ  r  r   rF   rU  rU   rÖ  r   rX   rq  Úunravel_indexr’  r×  rD   rN   r†  Úfnr€   r�   rk  r  r˜  r�  r�  )r#   r­   Úoriginal_axisr÷   Ú
input_dataÚoutput_dataÚinput_shapeÚoutput_shaper+  r½   Úinput_tupleÚjÚ	acc_tupleÚcombine_fn_retr™  r   Ú_r¢  r¨  s   ``               @@r$   Úgeneric_reducezReduceOps.generic_reduceé  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   ©r÷   Úkeepdimsz-val_reduce_op and idx_reduce_op are both None)r€   r’  r�  r  r   r÷   r˜  r   r�   r�   r   )r#   r­   Úval_reduce_opÚidx_reduce_opr�  Ú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 )Nr·  )r�  rU   rs  r  r   r÷   r˜  r   rË  s     r$   rs  zReduceOps.sum&  s<   € Ø�~‰~œbŸf™f U§\¡\×%6Ñ%6¸T¿Y¹YÐQU×Q_ÑQ_Ô`Ðbg×bmÑbmÓnÐnr&   c                óÀ  — | j                   t        j                  j                  u s&| j                   t        j                  j                  u r3| j                  |d   t        j                  t        j                  ¬«      S | j                   t        j                  j                  u s&| j                   t        j                  j                  u r3| j                  |d   t        j                  t        j                  ¬«      S | j                   t        j                  j                  u r%| j                  |d   t        j                  d ¬«      S | j                   t        j                  j                  u r%| j                  |d   t        j                  d ¬«      S | j                   t        j                  j                  u r| j!                  |d   «      S | j#                  |«      S )Nr   )r¹  rº  )r†  r�   ÚstandardÚ_argmin_combine_tie_break_leftÚ_argmin_combine_tie_break_fastr¼  rU   ÚnanminÚ	nanargminÚ_argmax_combine_tie_break_leftÚ_argmax_combine_tie_break_fastÚnanmaxÚ	nanargmaxÚ_elementwise_maxÚ_elementwise_minÚ_sum_combiners  rµ  rË  s     r$   r�  zReduceOps.apply_impl)  sH  € ð
 �O‰OœrŸ{™{×IÑIÑIØ—?‘?¤b§k¡k×&PÑ&PÑPØ—<‘<  a¡¼¿	¹	ÔQS×Q]ÑQ]�<Ó^Ð^Ø�o‰o¤§¡×!KÑ!KÑKØ—‘¤B§K¡K×$NÑ$NÑNØ—<‘<  a¡¼¿	¹	ÔQS×Q]ÑQ]�<Ó^Ð^Ø�_‰_¤§¡× <Ñ <Ñ<Ø—<‘<  a¡¼¿	¹	ÐQU�<ÓVÐVØ�_‰_¤§¡× <Ñ <Ñ<Ø—<‘<  a¡¼¿	¹	ÐQU�<ÓVÐVØ�_‰_¤§¡× 8Ñ 8Ñ8Ø—8‘8˜E !™HÓ%Ð%ð ×&Ñ& uÓ-Ð-r&   r(   )
r7   r8   r9   rJ   rœ  rµ  r¼  rs  r�  Ú__classcell__©r™  s   @r$   r•  r•  Ù  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(   )r—  rJ   Úreverse)r#   r÷   r†  rÐ  r™  s       €r$   rJ   zScanOps.__init__A  s   ø€ Ü‰Ñ˜˜zÔ*Øˆ�r&   c                ó¤   — | j                  t        j                  |j                  j                  | j
                  ¬«      |j                  ¬«      gS ©Nrü   rT   )r�  rU   Úcumsumr  r   r÷   r   rË  s     r$   rÓ  zScanOps.cumsumE  s8   € Ø—‘œrŸy™y¨¯©×):Ñ):ÀÇÁÔKÐSX×S^ÑS^�Ó_Ð`Ð`r&   c                ó¤   — | j                  t        j                  |j                  j                  | j
                  ¬«      |j                  ¬«      gS rÒ  )r�  rU   Úcumprodr  r   r÷   r   rË  s     r$   rÕ  zScanOps.cumprodH  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   rT   c              3  óh   •K  — | ])  \  }}‰j                  |‰   ‰|   j                  «      –— Œ+ y ­wr(   rŸ  )r   r¡  rc  Úindexr­   r#   s      €€€r$   r£  z'ScanOps.generic_scan.<locals>.<genexpr>V  s,   øè ø€ ÒfÁuÀrÈ1˜Ÿ™¨¨%©°%¸±)·/±/×BÑfùr¤  c              3  óV   •K  — | ]   }|‰j                   k(  r‰|   d z
  n‰|   –— Œ" y­w)r   Nrü   )r   r½   rØ  r#   s     €€r$   r£  z'ScanOps.generic_scan.<locals>.<genexpr>\  s-   øè ø€ Ò"kÐTU°1¸¿	¹	²> 5¨¡8¨a¢<ÀuÈQÁxÓ#OÑ"kùs   ƒ&)c              3  óh   •K  — | ])  \  }}‰j                  |‰   ‰|   j                  «      –— Œ+ y ­wr(   rŸ  )r   r¦  r§  r­   Ú
prev_indexr#   s      €€€r$   r£  z'ScanOps.generic_scan.<locals>.<genexpr>]  s/   øè ø€ Ò!uÑUZÐUWÐYZ $§.¡.°°:±ÀÀbÁ	ÇÁ×"PÑ!uùr¤  )r  r   rF   rU  rU   rÖ  r   rX   rq  r©  r’  r×  r÷   rD   rN   r†  rª  r€   r�   rk  r  r�  )r#   r­   r¬  r­  rF   r+  r½   r   r±  r²  r³  r™  rØ  rÛ  s   ``          @@r$   Úgeneric_scanzScanOps.generic_scanK  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 )Nrü   r   )rÐ  rU  r�  rU   Úflipr  r   r÷   r   r†  r�   r¿  rÊ  rÓ  Ú_prod_combinerÕ  rÜ  )r#   r­   Ú	new_inputr+  r™  s        r$   r�  zScanOps.apply_impli  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&   )	r7   r8   r9   rJ   rÓ  rÕ  rÜ  r�  rË  rÌ  s   @r$   rÎ  rÎ  ?  s   ø„ ôòaòbòö<r&   rÎ  c                óô   — dd„}dd„}| j                  t        d|«       | j                  t        d|«       | j                  t        j                  d|«       | j                  t        j                  d|«       y )Nc                ó:   — t        |||«      j                  | «      S r(   )r•  r“  )r­   r÷   r†  r˜  rc  s        r$   Ú_new_reducez'_patch_reduce_scan.<locals>._new_reduce�  s   € Ü˜˜z¨9Ó5×;Ñ;¸EÓBÐBr&   c                ó:   — t        |||«      j                  | «      S r(   )rÎ  r“  )r­   r÷   r†  rÐ  rc  s        r$   Ú	_new_scanz%_patch_reduce_scan.<locals>._new_scan„  s   € Ü�t˜Z¨Ó1×7Ñ7¸Ó>Ð>r&   ÚreduceÚassociative_scan)F)r6   r�   rk  )r™  rã  rå  s      r$   Ú_patch_reduce_scanrè  }  sY   € óCó?ð 
‡N�N”2�x Ô-Ø	‡N�N”2Ð)¨9Ô5Ø	‡N�N”2—7‘7˜H kÔ2Ø	‡N�N”2—7‘7Ð.°	Õ:r&   c                óÄ  — d„ }dd„}dd„}d„ }|j                  | d|«       |j                  | d|«       |j                  | d|«       |j                  | dt        «       |j                  | j                  d	|«       |j                  | d
t        |d¬«      «       |j                  | dt        |d¬«      «       |j                  | 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 | j                   dk(  r|j%                  «       S | j                   dk(  r|j'                  «       S t)        d| › d�«      ‚)NÚvoidr‘   r‰   rˆ   r‹   rŠ   r�   rŒ   rŽ   r\   ro   rq   rs   rp   rr   Úfp16Úbf16Úfp32Úfp64zfail to convert z to ir type)râ  Úget_void_tyr6  r9  r;  r>  r@  rC  rE  rG  rI  rS  rK  rN  rU  rQ  r,  r/  r2  r4  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˜&Ò Ø×&Ñ&Ó(Ð(Ø�Y‰Y˜&Ò Ø×'Ñ'Ó)Ð)Ø�Y‰Y˜&Ò Ø×(Ñ(Ó*Ð*ÜÐ+¨D¨6°Ð=Ó>Ð>r&   c                ó:   — |€d}|€d| }}n| |}}t        |||«      S )Nr   r   )rX   )Úarg1Úarg2Ústeprc  rk  Úends         r$   Ú
_new_rangez$_patch_lang_core.<locals>._new_range»  s2   € Øˆ<ØˆDØˆ<Ø˜D�3‰Eà˜t�3ˆEÜ�U˜C Ó&Ð&r&   c                ó   — | sJ |«       ‚y r(   r?   )r1  r¾  s     r$   Ú_new_static_assertz,_patch_lang_core.<locals>._new_static_assertÄ  s   € ÙÐ�SÓ‰tr&   c                ó’  — t        | t        j                  «      s| S t        |t        t        f«      s|gn|}|D �cg c]*  }t        |t        j
                  «      r|j                  n|‘Œ, }}t        |«      t        dt        | j                  «      «      k7  rt        d|› �«      ‚| j                  j                  ||«       | S c c}w )Nr   z$len(values) != len(input.shape) for )r€   r�   r  ru  r’  Ú	constexprr5   rD   r{  rF   r   r  r6   )r­   r­  râ  re  s       r$   Ú	_set_attrz#_patch_lang_core.<locals>._set_attrÇ  s¢   € ä˜%¤§¡Ô+ØˆLä!+¨F´T¼5°MÔ!B�&‘ÈˆØIOÖPÀAœZ¨¬2¯<©<Ô8�!—'’'¸aÑ?ÐPˆÐPÜˆv‹;œ#˜a¤ U§[¡[Ó!1Ó2Ò2ÜÐCÀDÀ6ÐJÓKÐKØ�‰×Ñ˜d FÔ+Øˆùò	 Qs   ¼/CrX   Ústatic_rangeÚstatic_assertÚstatic_printÚto_irÚmultiple_ofztt.divisibility)râ  Úmax_contiguousztt.contiguityÚmax_constancyztt.constancy)NN)Ú )r6   rº  r   r   rè  )Úlangr™  rñ  r÷  rù  rü  s         r$   Ú_patch_lang_corer  �  s½   € ò(?óX'óò
ð 
‡N�N�4˜ *Ô-Ø	‡N�N�4˜¨Ô4Ø	‡N�N�4˜Ð*<Ô=Ø	‡N�N�4˜¬Ô/Ø	‡N�N�4—:‘:˜w¨
Ô3Ø	‡N�N�4˜¬°	Ð@QÔ(RÔSØ	‡N�N�4Ð)¬7°9À?Ô+SÔTØ	‡N�N�4˜¬'°)À.Ô*QÔRä�uÕr&   c                óJ  — t        «       }| 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 ]t  }t        |t        |«       t        |j                  t        |«       |t
        k(  rt        |j                  t        |«       t        |j                  |«       t        ||«       Œv t        t
        j                  j                  t        |«       |S c c}}w )Nr   z:triton.language must be visible from within jit'd function)rP  Ú__globals__ra  ri  Úismoduler�   rk  rD   rn  Úinterpreter_builderr  rÉ   r‚  r  Útensor_descriptor_base)rª  r™  r´  r5   Úlangsr  s         r$   Ú_patch_langr  ß  sç   € ÜÓ€EØ#%§>¡>×#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°%Ô8Ü�t—{‘{Ô$7¸Ô?Ø”2Š:Ü˜4Ÿ9™9Ô&9¸5ÔAÜ˜4Ÿ;™;¨Ô.Ü˜˜uÕ%ð&ô ”2—7‘7×1Ñ1Ô3FÈÔNØ€Lùó qs   ¨DÁDÁDc                ó˜  — t        | t        «      rbt        t        j                  | gt        j
                  ¬«      t        j                  «      }t        j                  |t        j                  «      S t        | t        «      �rt        j                  t        j                  j                  j                  | «      d «      }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                  | «      d «      }t        t        j                  | j+                  «       gt        j$                  ¬«      |«      }t        j                  ||«      S t        | t,        «      rt/        | t1        t2        | «      «      S t        | t4        «      rÑ| j6                  D �cg c]  }t3        |«      ‘Œ }}| j6                  d
   dk(  sJ ‚t        j8                  d«      |d
<   t:        j=                  t3        | j>                  «      | j@                  D �cg c]  }t3        |«      ‘Œ c}|| jB                  D �cg c]  }t        j8                  |«      ‘Œ c}| jD                  ¬«      S | S c c}w c c}w c c}w )NrT   i   €l        l        lûÿÿÿ         l            l            zUnsupported integer value Údata_ptrrM   r   )rC   rF   rG   rH   Úpadding_option)#r€   r)   r   rU   rÝ   r^  r�   r‘   r  r|   Ú	str_to_tyÚtritonÚruntimeÚjitÚmangle_typer�   rŒ   rŽ   r\   r   r.   r  r’  r   ÚmapÚ_implicit_cvtr   rG   rû  rN  Úmake_tensor_descriptorrC   rF   rH   rI   )r+  r  Útyr   ÚsrG   rÑ   s          r$   r  r  ï  s]  € Ü�#”tÔÜœbŸh™h¨ u´B·H±HÔ=¼r¿w¹wÓGˆÜ�y‰y˜¤§¡Ó)Ð)Ü�#”sÕÜ�\‰\œ&Ÿ.™.×,Ñ,×8Ñ8¸Ó=¸tÓDˆÜ—‘ˆØ�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¸Ó=¸tÓDˆÜœ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ˆ�‰Ü#×:Ñ:ÄÈcÏhÉhÓ@WØ\_×\eÑ\eÖAfÐWXÄ-ÐPQÕBRÒAfÐpwØad×apÑapÖGqÐ\]ÌÏÉÐUVÍÒGqØJMÏ+É+ð ;ó Wð 	Wð €Jùò :ùò BgùÚGqs   É8L=Ë)M
ÌMc                óz   — t        | t        j                  j                  j                  «      r| j
                  S | S r(   )r€   r  r  r  ÚTensorWrapperrC   )Ú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g fd„Zd„ Zd„ Zd„ Zy)ÚGridExecutorc                ó*  — ddl m} || _        || _        || _        || _        t        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%  rª  Ú	arg_namesÚgridÚpre_run_hooksri  Úget_annotationsra  ÚgetÚ
constexprs)	r#   rª  r&  r'  r(  r%  râ  r  r;   s	            r$   rJ   zGridExecutor.__init__!  s†   € Ý&àˆŒØ"ˆŒØˆŒ	Ø*ˆÔÜCJ×CZÑCZÐ[]ÓC^×CdÑCdÓCf×g±x°t¸R˜4¡¨rÓ!2Ñ2ÐgˆÑgØ,5Öb D¸×9LÑ9LÈTÓ9RÐVaÓ9aš4Òbˆ�ùó hùÚbs   Á	B
Á#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        «      rRt	         ‰| j
                  «      | j                  | 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   rC   rF   rG   rH   rI   Úround_f32_to_tf32r.   r  Úuntyped_storager  r.  Ú	new_emptyÚset_Ústorage_offsetrq  rP   r!  )r+  Úunwrapped_argÚstorageÚcpu_argÚ_to_cpuÚstoragess       €€r$   r8  z,GridExecutor._init_args_hst.<locals>._to_cpu.  s6  ø€ Ü˜#œuÔ%Ü$ S¬#¨g°sÓ*;Ó<Ð<Ü˜CÔ!1Ô2Ü'Ù˜CŸH™HÓ%Ø—I‘IØ—K‘KØ—O‘OØ—K‘KØ×)Ñ)óð ô ˜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&   )ra  )
r#   Úargs_devrc  r+  Úargs_hstÚ
kwargs_hstr4   r5   r8  r9  s
           @@r$   Ú_init_args_hstzGridExecutor._init_args_hst+  sc   ù€ Øˆõ	ð4 -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  r1  r  r€   r’  Úzipr   rC   )Úarg_devÚarg_hstÚ	_from_cpur9  s     €€r$   rC  z1GridExecutor._restore_args_dev.<locals>._from_cpuS  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@  ra  r­  Úcopy_)r#   r:  r;  rc  r<  rA  rB  r4   Ú	kwarg_devÚ	kwarg_hstrC  r9  s             @@r$   Ú_restore_args_devzGridExecutor._restore_args_devP  s‹   ù€ Øˆõ		6ô !$ H¨hÓ 7ò 	(ÑˆG�WÙ�g˜wÕ'ð	(ð %Ÿl™l›nò 	,‰NˆC�Ø" 3™ˆIÙ�i Õ+ð	,ð #+§/¡/Ó"3ò 	#ÑˆW�gØ�M‰M˜'Õ"ñ	#r&   c                ó®  — t        j                  | j                  «      }|j                  «       D ��ci c]  \  }}||j                  v sŒ||“Œ }}}| j                  ||«      \  }}| j                  D ]
  } ||i |¤Ž Œ 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                  «        | j5                  ||||«       y c c}}w c c}}w # t$        $ r?}t&        j(                  j*                  j,                  r‚ t/        t1        |«      «      |‚d }~ww xY w# |	j3                  «        w xY w)Nrë   z#grid must have at most 3 dimensionsr‘  r   r   r   r?   )ri  Úgetfullargspecrª  ra  rb  r=  r(  r  Úgetcallargsr+  r  Úcallabler'  rD   r
  r*  rX   r%  Ú	Exceptionr  ÚknobsÚcompilationÚfront_end_debuggingr   r~  r\  rG  )r#   r:  rc  Úargspecrd  re  r;  r<  ÚhookÚpatch_scoperb  râ  r+  r'  rÌ   r#  r$  Úes                     r$   Ú__call__zGridExecutor.__call__i  s&  € ô ×(Ñ(¨¯©Ó1ˆØ#)§<¡<£>×G™4˜1˜a°Q¸'¿,¹,Ò5F�!�Q‘$ÐGˆÑGà#×2Ñ2°8¸VÓDÑˆ�*à×&Ñ&ò 	*ˆDÙ�(Ð)˜jÓ)ð	*ô " $§'¡'Ó*ˆð	"ô ×&Ñ& t§w¡wÐH°ÒH¸ZÑHˆDØbf×blÑblÓbn×oÑU^ÐUYÐ[^�D ¨¯©Ñ!8™#¼mÈCÓ>PÑPÐoˆDÑoä&.¨t¯y©yÔ&9�4—9‘9˜T”?¸t¿y¹yˆDÜ�t“9 ’>ÐHÐ#HÓH�>Ø˜% 1¤s¨4£y¡=Ñ1Ñ1ˆDÜ×,Ñ,¨dÑ3ð	7Ü˜t A™w›ò ,�AÜ" 4¨¡7›^ò ,˜Ü!& t¨A¡w£ò ,˜AÜ/×<Ñ<¸QÀÀ1ÔEØ#˜DŸG™G™O d›Oñ,ñ,ñ,ð ×ÑÔ!à×Ñ˜x¨°6¸:ÕFùó? Hùó pøô ò 7Ü—<‘<×+Ñ+×?Ò?ØÜ&¤t¨A£wÓ/°QÐ6ûð7ûð
 ×ÑÕ!úsI   ³G+Á
G+Â6I Ã&G1Ã2A0I Å#A"G7 Ç1I Ç7	H?È :H:È:H?È?I ÉIN)r7   r8   r9   rJ   r=  rG  rT  r?   r&   r$   r#  r#    s   „ à:<ó cò#$òJ#ó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 supportedrN  )ÚidÚctxr�  )r5   r   rY  F)r5   )Úfuncrb  Úkeywords)ÚtargetsÚvisitrD   r   ÚastÚCallÚ	AttributeÚNameÚLoadr5   Ú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)r7   r8   r9   rg  r?   r&   r$   rV  rV  �  s   „ ór&   rV  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   )rª  rc  ÚfilenameÚdef_file_lineno)r#   rª  rc  s      r$   rJ   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(   )ri  Úgetsourcelinesrª  rL  Ú_get_jit_fn_file_linerk  rl  Ú	_find_defÚ
def_linenoÚ_prepare_sourceÚ_transform_astÚ_compile_and_exec)r#   Úlinesr´  r�  Ú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}  || j                  «      }|j                  |j                  fS )Nr   )ÚJITFunction)r  ry  rª  Ú	file_nameÚdef_file_line_number)r#   ry  Újit_funcs      r$   ro  z&FunctionRewriter._get_jit_fn_file_line¾  s*   € Ý$Ù˜tŸw™wÓ'ˆØ×!Ñ! 8×#@Ñ#@Ð@Ð@r&   c                óz   — d}t        |«      D ]*  \  }}|j                  «       j                  d«      sŒ&|dz   }Œ, |S )Nr   zdef r   )r×  ÚstripÚ
startswith)r#   ru  rq  r½   Úlines        r$   rp  zFunctionRewriter._find_defÃ  sD   € Øˆ
ä  Ó'ò 	#‰GˆAˆtØ�z‰z‹|×&Ñ& vÕ.Ø ™U‘
ð	#ð Ðr&   c                ór   — || j                   dz
  d  }dj                  |«      }t        j                  |«      S )Nr   r  )rq  ÚjoinÚtextwrapÚdedent)r#   ru  r�  s      r$   rr  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 rX  )r^  ÚparseÚast_transformerr]  Úfix_missing_locationsrl  Úincrement_lineno)r#   r�  Ú
parsed_astrv  Ú
inc_linenos        r$   rs  zFunctionRewriter._transform_astÐ  s[   € ô —Y‘Y˜s“^ˆ
Ø×.Ñ.×4Ñ4°ZÓ@ˆÜ×!Ñ! /Ô2Ø×)Ñ)¨AÑ-ˆ
Ü×Ñ˜_¨jÔ9ØÐr&   c                ó  — t        || j                  d¬«      }i | j                  ¥}| j                  j                  }t        «       j                  «       D ]  \  }}||vsŒ|||<   Œ t        |||«       || j                  j                     S )NÚexec)rk  Úmode)	Úcompilerk  rc  rª  r  Úglobalsra  r�  r7   )r#   rv  Úcompiled_codeÚlocal_namespaceÚ
fn_globalsr4   r5   s          r$   rt  z"FunctionRewriter._compile_and_execÛ  s€   € Ü ¸$¿-¹-ÈfÔUˆØ)˜TŸ[™[˜/ˆØ—W‘W×(Ñ(ˆ
Ü!›)Ÿ/™/Ó+ò 	(‰JˆC�Ø˜*Ò$Ø"'�
˜3’ð	(ô 	ˆ]˜J¨Ô8Ø˜tŸw™w×/Ñ/Ñ0Ð0r&   N)r7   r8   r9   rV  r‡  rJ   rw  ro  rp  rr  rs  rt  r?   r&   r$   ri  ri     s.   „ Ù$Ó&€Oò&ò7ò(Aò
ò$ò
	ó1r&   ri  c                  óL   — e Zd ZU i Zded<   d
d„Zd„ Zd„ Zd„ Ze	d„ «       Z d„ Z
y	)ÚInterpretedFunctionzDict[Callable, Callable]Úrewritten_fnc                óò   — || _         t        |fi |¤Ž| _        || _        g | _        t        j                  |«      }|j                  j                  «       D �cg c]  }|j                  ‘Œ c}| _
        y c c}w r(   )rª  ri  Úrewriterrc  r(  ri  Ú	signatureÚ
parametersr­  râ  r&  )r#   rª  rc  r™  re  s        r$   rJ   zInterpretedFunction.__init__ê  sa   € ØˆŒÜ(¨Ñ6¨vÑ6ˆŒØˆŒØˆÔä×%Ñ% bÓ)ˆ	Ø*3×*>Ñ*>×*EÑ*EÓ*GÖH Q˜!Ÿ&›&ÒHˆ�ùÒHs   ÁA4c               óx   — |ry | j                  «       } t        || j                  || j                  «      |i |¤ŽS r(   )Úrewriter#  r&  r(  )r#   r'  Úwarmuprb  rc  rª  s         r$   ÚrunzInterpretedFunction.runó  s;   € ÙØØ�\‰\‹^ˆØIŒ|˜B §¡°°d×6HÑ6HÓIÈ4ÐZÐSYÑZÐZr&   c                óT   — t        |«      sJ ‚| j                  j                  |«       y r(   )rK  r(  rU  )r#   rQ  s     r$   Úadd_pre_run_hookz$InterpretedFunction.add_pre_run_hookù  s"   € Ü˜Œ~Ðˆ~Ø×Ñ×!Ñ! $Õ'r&   c                óÆ   — | j                   | j                  vr1| j                  j                  «       | j                  | j                   <   | j                  | j                      S r(   )rª  r–  r˜  rw  r"   s    r$   rœ  zInterpretedFunction.rewriteý  sJ   € Ø�7‰7˜$×+Ñ+Ñ+Ø)-¯©×)BÑ)BÓ)DˆD×Ñ˜dŸg™gÑ&Ø× Ñ  §¡Ñ)Ð)r&   c                ó.   — | j                   j                  S r(   )rª  r7   r"   s    r$   r7   zInterpretedFunction.__name__  s   € à�w‰w×ÑÐr&   c                óª   — t        | j                  «       | j                  «       }	  ||i |¤ŽS # t        $ r}t	        t        |«      «      |‚d }~ww xY wr(   )r  rª  rœ  rL  r   r~  )r#   rb  rc  rª  rS  s        r$   rT  zInterpretedFunction.__call__  sO   € ä�D—G‘GÔØ�\‰\‹^ˆð	3Ù�tÐ&˜vÑ&Ð&øÜò 	3Ü"¤4¨£7Ó+°Ð2ûð	3ús   §/ ¯	A¸AÁANrð  )r7   r8   r9   r–  r;   rJ   rž  r   rœ  r�  rT  r?   r&   r$   r•  r•  æ  s<   … à-/€LÐ*Ó/óIò[ò(ò*ð
 ñ ó ð ó3r&   r•  )r™  rP  )UÚ
__future__r   r^  rƒ  ri  Útypingr   r   r   r   r   r	   rÉ   ÚnumpyrU   r  Útriton.languageÚlanguager�   r<   r
   Útriton.language.semanticr   Útriton.runtime.jitr   Útriton.tools.tensor_descriptorr   Úerrorsr   Ú	functoolsr   Ú_C.libtritonr   r’  r   rª   Ú_utilsr   r   r   rA   rj   r   r�   rž   rÇ   rÍ   rÒ   r×   ræ   rô   rù   r  Ú	vectorizer“   rT  r”   rU  r\   r!  r	  r  r]  rT  r
  rN  r;   rP  rg  rn  r‚  r„  r•  rÎ  rè  r  r  r  r  r!  r#  ÚNodeTransformerrV  ri  r•  r?   r&   r$   ú<module>r²     sñ  ðÞ "Û 
Û Û ß A× Aã Û ã Ý Û Ý !å 3Ý .Ý ;Ý $Ý Ý 6Ý $Ý "áˆCƒL€ð ÷ð ó ð÷@("ñ ("ñV �$Ô÷
&ð 
&ó ð
&òò 	òò@='ò@ò
#òò4ò$ò@+òð2 ˆb�l‰l˜4¨¯©¨Ô5€Øˆb�l‰l˜4¨¯©¨Ô5€Ø�—‘˜Z°·±°Ô<€÷pñ p÷O0ñ O0ñd ‹8€Ù(Ó*Ð Ù'5Ð6IÓ'JÐ �nÓ J÷-ñ -ó(*ó;ó:÷."Iñ "IôJc.Ð%ô c.ôL;Ð#ô ;ó|;ó Oòdò !òHò÷nGñ nGôb�S×(Ñ(ô ÷ C1ñ C1ôL'3˜/¨!Ñ,õ '3r&   