Ë
    7^(h2W ã                  ó¨  — d dl mZ d dlZd dlmZmZmZmZmZ d dl	Z	ddl
mZ ddlmZ  ed«      Z G d	„ d
e«      Zdtd„Zdtd„Zdud„Z	 	 	 	 dvd„Zdwdxd„Zdyd„Z	 	 dz	 d{d„Zd|d„Z	 	 	 	 d}d„Z	 	 	 	 d}d„Z	 	 	 	 d}d„Zd~d„Zd~d„Z	 	 	 	 dd„Zd~d„Z d€d„Z!d€d„Z"d�d„Z#	 	 	 	 d‚d„Z$dƒd„Z%dƒd„Z&dƒd „Z'dƒd!„Z(dƒd"„Z)d„d#„Z*dƒd$„Z+dƒd%„Z,dƒd&„Z-d…d'„Z.d†d(„Z/d‡d)„Z0dˆd*„Z1dƒd+„Z2dƒd,„Z3dƒd-„Z4dƒd.„Z5dƒd/„Z6dƒd0„Z7d‰d1„Z8dŠd2„Z9d‹d3„Z:dŒd4„Z;d�d5„Z<dŽd6„Z=d�d7„Z>d�d8„Z?d‘d9„Z@d’d:„ZAd“d;„ZBd”d<„ZCd•d=„ZD	 d–	 	 	 d—d>„ZEd?„ ZFd@„ ZGdA„ ZHdB„ ZIdC„ ZJdD„ ZKdE„ ZLdF„ ZMdG„ ZN	 	 	 	 	 	 	 	 	 	 	 	 d˜dH„ZOd™dI„ZPdJ„ ZQ	 	 	 	 dšdK„ZR	 	 	 	 d›dL„ZS	 	 	 	 dœdM„ZTd�dN„ZU	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 dždO„ZVdŸdP„ZWdQ„ ZXdR„ ZY	 	 	 	 	 	 d dS„ZZd¡dT„Z[	 	 	 	 d¢dU„Z\d£dV„Z]d£dW„Z^d£dX„Z_d£dY„Z`d£dZ„Zad£d[„Zb	 	 	 	 d£d\„Zcd]„ Zd	 	 	 	 	 	 d¤d^„Zed¥d_„Zfd¦d`„Zg	 	 	 	 	 	 	 	 	 	 	 	 d§da„Zhd¨db„Zidc„ Zjd©dd„Zk	 	 	 	 dªde„Zld«df„Zmd¬dg„Znd­dh„Zod­di„Zpd­dj„Zqd®dk„Zrd¯dl„Zsd°dm„Ztd®dn„Zudo„ Zvdwdp„Zwd±dq„Zxd±dr„Zy	 	 	 	 	 	 	 	 	 	 	 	 d²ds„Zzy)³é    )ÚannotationsN)ÚListÚOptionalÚSequenceÚTupleÚTypeVaré   )Úiré   )ÚcoreÚTc                  ó   ‡ — e Zd Zˆ fd„Zˆ xZS )ÚIncompatibleTypeErrorImplc                óÜ   •— || _         || _        d| j                   j                  «       z   dz   | j                  j                  «       z   | _        t        t
        | �  | j                  «       y )Nzinvalid operands of type ú and )Útype_aÚtype_bÚ__repr__ÚmessageÚsuperr   Ú__init__)Úselfr   r   Ú	__class__s      €úV/var/www/skyplay_api_hub/venv/lib/python3.12/site-packages/triton/language/semantic.pyr   z"IncompatibleTypeErrorImpl.__init__   sX   ø€ ØˆŒØˆŒØ2°T·[±[×5IÑ5IÓ5KÑKÈgÑUÐX\×XcÑXc×XlÑXlÓXnÑnˆŒÜÔ'¨Ñ7¸¿¹ÕEó    )Ú__name__Ú
__module__Ú__qualname__r   Ú__classcell__)r   s   @r   r   r      s   ø„ ÷Fð Fr   r   c                óŒ   — | dvrt        d| › �«      ‚t        j                  |j                  | «      t        j                  «      S )N©r   r   r	   z+program_id axis must be 0, 1, or 2 but got )Ú
ValueErrorÚtlÚtensorÚcreate_get_program_idÚint32©ÚaxisÚbuilders     r   Ú
program_idr*      s=   € Ø�9ÑÜÐFÀtÀfÐMÓNÐNÜ�9‰9�W×2Ñ2°4Ó8¼"¿(¹(ÓCÐCr   c                óŒ   — | dvrt        d| › �«      ‚t        j                  |j                  | «      t        j                  «      S )Nr!   z-num_programs axis must be 0, 1, or 2 but got )r"   r#   r$   Úcreate_get_num_programsr&   r'   s     r   Únum_programsr-   !   s=   € Ø�9ÑÜÐHÈÈÐOÓPÐPÜ�9‰9�W×4Ñ4°TÓ:¼B¿H¹HÓEÐEr   c                ó`  — | j                   }|j                   }| j                  }|j                  }||k(  r	||kD  r| S |S |t        j                  j                  j
                  k(  r	||k\  r| S |S |t        j                  j                  j
                  k(  r	||k\  r|S | S t        d|› d|› �«      ‚)Nzunexpected signedness r   )Úint_bitwidthÚint_signednessr#   ÚdtypeÚ
SIGNEDNESSÚUNSIGNEDÚ	TypeError)Úa_tyÚb_tyÚa_rankÚb_rankÚa_snÚb_sns         r   Úinteger_promote_implr;   ,   sµ   € Ø×Ñ€FØ×Ñ€FØ×Ñ€DØ×Ñ€Dð ˆt‚|Ø ’ˆtÐ0¨DÐ0Ø	”—‘×$Ñ$×-Ñ-Ò	-Ø Ò'ˆtÐ1¨TÐ1Ø	”—‘×$Ñ$×-Ñ-Ò	-Ø Ò'ˆtÐ1¨TÐ1Ü
Ð,¨T¨F°%¸°vÐ>Ó
?Ð?r   c                ó¨  — ||k7  rx|r| |fn|| f\  }}|j                  «       j                  |j                  «       j                  k  r6|r2|t        j                  t        j                  fv rt        j
                  S |S | j                  «       s|j                  «       rt        j                  S | j                  «       s|j                  «       rt        j
                  S | j                  «       s|j                  «       r"|rt        j
                  S t        j                  S | j                  «       r2|j                  «       r"|rt        j
                  S t        j                  S | j                  «       s|j                  «       rt        j
                  S | j                  «       r'|j                  «       r| |k(  r| S t        j                  S | j                  «       r|j                  «       st        d| › d|› �«      ‚|rL| j                  |j                  k7  r3t        d| j                  «       z   dz   |j                  «       z   dz   «      ‚t!        | |«      S )Núunexpected type r   zCannot use /, #, or % with úx because they have different signedness;this is unlikely to result in a useful answer. Cast them to the same signedness.)ÚkindÚvaluer#   Úfloat16Úbfloat16Úfloat32Úis_fp64Úfloat64Úis_fp32Úis_fp16Úis_bf16Úis_fp8Úis_intr4   r0   r   r;   )r5   Úa_is_scalarr6   Úb_is_scalarÚ
div_or_modÚ	scalar_tyÚ	tensor_tys          r   Úcomputation_type_implrP   <   sÆ  € ð
 �kÒ!Ù/:  d™|ÀÀtÀÑˆ	�9Ø�>‰>Ó×!Ñ! Y§^¡^Ó%5×%;Ñ%;Ò;á˜y¬R¯Z©Z¼¿¹Ð,EÑEÜ—z‘zÐ!ØÐð ‡|�|„~˜Ÿ™œÜ�z‰zÐð ‡|�|„~˜Ÿ™œÜ�z‰zÐð ‡|�|„~˜Ÿ™œÙÜ—:‘:Ðä—:‘:Ðà‡|�|„~˜$Ÿ,™,œ.ÙÜ—:‘:Ðä—;‘;ÐØ‡|�|„~˜Ÿ™œÜ�z‰zÐà‡{�{„}˜Ÿ™œØ˜t’|ˆtÐ3¬¯©Ð3Ø�;‰;Œ= §¡¤ÜÐ*¨4¨&°°d°VÐ<Ó=Ð=ñ �d×)Ñ)¨T×-@Ñ-@Ò@ÜÐ5¸¿¹»ÑGÈ'ÑQÐTX×TaÑTaÓTcÑcðkñkó lð 	lô    dÓ+Ð+r   c                ó®  — t        | t        «      r3t        j                  |j	                  | «      t        j
                  «      S t        | t        «      ršd| cxk  rdk  rn nt        j                  }nld| cxk  rdk  rn nt        j                  }nMd| 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| › d�«      ‚t        d	| ||¬
«      S t        | t        «      rnd}dddz  z  }t        d   | «      }|t        d«      k(  s|dk(  s| | k7  s||cxk  r|k  rn nt        j                  }nt        j                   }t        d	| ||¬
«      S t        | t        j"                  «      rt%        | j&                  |«      S t        | t        j                  «      r| S |rt)        d| › dt+        | «      › d�«      ‚| S )Né   €ì        l        ìûÿÿÿ         ì            l            zNonrepresentable integer ú.© ©r1   r)   g      8g   àÿÿÿ?r	   é   ÚabsÚinfç        zcannot convert z	 of type z
 to tensor)Ú
isinstanceÚboolr#   r$   Úget_int1Úint1Úintr&   Úuint32Úint64Úuint64r"   ÚfullÚfloatÚ__builtins__rC   rE   Ú	constexprÚ	to_tensorr@   r4   Útype)Úxr)   Ú
check_typer1   Úmin_float32Úmax_float32Úabs_xs          r   ri   ri   o   s‚  € Ü�!”TÔÜ�y‰y˜×)Ñ)¨!Ó,¬b¯g©gÓ6Ð6ä	�A”sÔ	Ø�QÔ˜ÕÜ—H‘H‰EØ�aÔ˜%ÕÜ—I‘I‰EØ�qÔ ˜5Õ Ü—H‘H‰EØ�aÔ˜%ÕÜ—I‘I‰EäÐ8¸¸¸1Ð=Ó>Ð>Ü�B˜ °Ô8Ð8Ü	�A”uÔ	ØˆØ! Q¨¡VÑ+ˆÜ˜UÑ# AÓ&ˆØ”E˜%“LÒ Ø�CŠ<Ø�Š6Ø˜%Ô. ;Õ.Ü—J‘J‰Eä—J‘JˆEÜ�B˜ °Ô8Ð8ä	�A”r—|‘|Ô	$Ü˜Ÿ™ 'Ó*Ð*Ü	�A”r—y‘yÔ	!ØˆÙÜ˜/¨!¨¨I´d¸1³g°Y¸jÐIÓJÐJØ€Hr   c                ó¼   — | j                  «       rL|st        | |«      ‚|j                  «       r| |k7  rt        | |«      ‚|j                  «       rt        | |«      ‚y y ©N)Úis_ptrr   Úis_floating)r   r   Úallow_ptr_as      r   Úcheck_ptr_type_implru   š   s[   € Ø‡}�}„ÙÜ+¨F°FÓ;Ð;à�=‰=Œ? ¨&Ò 0Ü+¨F°FÓ;Ð;à×ÑÔÜ+¨F°FÓ;Ð;ð  ð r   c                óÀ  — t        | t        j                  «      }t        |t        j                  «      }|r| }	t        | |«      } |r|}
t        ||«      }| j                  j
                  }|j                  j
                  }t        |||«       t        |||«       |�r.|j                  «       �s|j                  «       �st        |||||«      }|r	dk  r|j                  «       s|r 
dk  r|j                  «       rt        d«      ‚|j                  «       rx|r:|j                  «       	cxk  r|j                  «       k  sn t        d|	› d|› �«      ‚|r:|j                  «       
cxk  r|j                  «       k  sn t        d|
› d|› �«      ‚|rt        d	||¬«      nt        | ||«      } |rt        d
||¬«      nt        |||«      }t!        | ||«      \  } }| |fS )Nr   z{Cannot perform a binary operation between an unsigned tensor and a negative scalar. Perform a explicit cast on one of them.zScalar z is out of range for type rW   rX   )r]   ÚnumbersÚNumberri   rj   Úscalarru   rr   rP   Úis_int_unsignedr"   rJ   Úget_int_min_valueÚget_int_max_valuere   ÚcastÚbroadcast_impl_value)ÚlhsÚrhsr)   Úallow_lhs_ptrÚallow_rhs_ptrÚarithmetic_checkrM   Úlhs_is_scalarÚrhs_is_scalarÚ
lhs_scalarÚ
rhs_scalarÚ
lhs_sca_tyÚ
rhs_sca_tyÚ
ret_sca_tys                 r   Úbinary_op_type_checking_implr‹   ¦   så  € ô ˜s¤G§N¡NÓ3€MÜ˜s¤G§N¡NÓ3€MÙØˆ
Ü˜˜WÓ%ˆÙØˆ
Ü˜˜WÓ%ˆð —‘—‘€JØ—‘—‘€JÜ˜
 J°Ô>Ü˜
 J°Ô>Ú 
× 1Ñ 1Õ 3¸J×<MÑ<MÕ<OÜ*¨:°}ÀjÐR_ÐakÓlˆ
Ù˜j¨1šn°×1KÑ1KÔ1MÙ  Z°!¢^¸
×8RÑ8RÔ8TÜð Gó Hð Hà×ÑÔÙ j×&BÑ&BÓ&DÈ
Ô&tÐV`×VrÑVrÓVtÔ&tÜ  7¨:¨,Ð6PÐQ[ÐP\Ð!]Ó^Ð^Ù j×&BÑ&BÓ&DÈ
Ô&tÐV`×VrÑVrÓVtÔ&tÜ  7¨:¨,Ð6PÐQ[ÐP\Ð!]Ó^Ð^áBOô Ø�
 *°gõ?ÜUYÐZ]Ð_iÐkrÓUsð 	ñ CPô Ø�
 *°gõ?ÜUYÐZ]Ð_iÐkrÓUsð 	ô $ C¨¨gÓ6�H€CˆØ�ˆ8€Or   c                ó(  — | j                   j                  j                  dk\  s|j                  j                  sy | j                   j                  }|j                   j                  }||k(  sJ ‚|j                  «       sJ ‚t        | t        j                  |«      } t        |t        j                  |«      } || |d|«      }|j                  «       }t        j                  |j                  |«      t        j                  «      }|j                  «       }t        j                  |j                  |«      t        j                  «      }t        t        |||«      t        |||«      |«      }	d|j                  › d|j                   › �}
t#        |	|
|«       y )Né@   Fra   z! overflow detected for operation )rj   ry   r/   ÚoptionsÚsanitize_overflowrJ   r}   r#   rc   r|   r$   Ú	get_int64r{   Úand_Ú
less_equalÚgreater_equalr   Údevice_assert)r   r€   r)   Ú	binary_oprˆ   r‰   ÚretÚ	max_valueÚ	min_valueÚcondÚmsgs              r   Ú binary_op_sanitize_overflow_implr›   Ì   s>  € Ø
‡x�x‡�×#Ñ# rÒ)°·±×1RÒ1RØØ—‘—‘€JØ—‘—‘€JØ˜Ò#Ð#Ð#Ø×ÑÔÐÐÜ
ˆs”B—H‘H˜gÓ
&€CÜ
ˆs”B—H‘H˜gÓ
&€CÙ
�C˜˜e WÓ
-€CØ×,Ñ,Ó.€IÜ—	‘	˜'×+Ñ+¨IÓ6¼¿¹ÓA€IØ×,Ñ,Ó.€IÜ—	‘	˜'×+Ñ+¨IÓ6¼¿¹ÓA€IÜ”
˜3 	¨7Ó3´]À3È	ÐSZÓ5[Ð]dÓe€DØ�
×'Ñ'Ð(Ð(IÈ)×J\ÑJ\ÐI]Ð
^€CÜ�$˜˜WÕ%r   c                ó`  — t        | ||dd«      \  } }| j                  j                  }|j                  j                  }|j                  «       r|j                  «       rt	        d«      ‚|j                  «       r@|j                  «       s0|| }} | j                  j                  }|j                  j                  }|j                  «       �r|j
                  }|j                  j                  «       r»|j                  j                  dk  r¢|j                  j                  «       rLt        j                  t        j                  |j                  j                  «       «      j                  |«      }nt        j                  j                  |«      }|j                  |j
                  |d«      }t        j                   |j#                  | j
                  |«      | j                  «      S |j%                  «       rDt        j                   |j'                  | j
                  |j
                  «      | j                  «      S |j)                  «       rX|rt+        | ||t,        «       t        j                   |j/                  | j
                  |j
                  «      | j                  «      S t	        d|› �«      ‚)NTzcannot add pointers togetherr�   Fr=   )r‹   rj   ry   rr   r4   Úhandler1   rz   r/   Úis_blockr#   Ú
block_typerc   Úget_block_shapesÚto_irÚcreate_int_castr$   Úcreate_addptrrs   Úcreate_faddrJ   r›   ÚaddÚ
create_add)ÚinputÚotherr�   r)   Úinput_scalar_tyÚother_scalar_tyÚother_handleÚi64_tys           r   r¥   r¥   ß   sí  € ä/°°u¸gÀtÈTÓR�L€Eˆ5Ø—j‘j×'Ñ'€OØ—j‘j×'Ñ'€OØ×ÑÔ O×$:Ñ$:Ô$<ÜÐ6Ó7Ð7ð ×ÑÔ¨×(>Ñ(>Ô(@Ø˜eˆuˆØŸ*™*×+Ñ+ˆØŸ*™*×+Ñ+ˆØ×ÑÕØ—|‘|ˆØ�;‰;×&Ñ&Ô(¨U¯[©[×-EÑ-EÈÒ-Jà�z‰z×"Ñ"Ô$ÜŸ™¤r§x¡x°·±×1LÑ1LÓ1NÓO×UÑUÐV]Ó^‘äŸ™Ÿ™¨Ó0�Ø"×2Ñ2°5·<±<ÀÈÓOˆLÜ�y‰y˜×.Ñ.¨u¯|©|¸\ÓJÈEÏJÉJÓWÐWà	×	$Ñ	$Ô	&Ü�y‰y˜×,Ñ,¨U¯\©\¸5¿<¹<ÓHÈ%Ï*É*ÓUÐUà	×	Ñ	Ô	!ÙÜ,¨U°E¸7ÄCÔHÜ�y‰y˜×+Ñ+¨E¯L©L¸%¿,¹,ÓGÈÏÉÓTÐTÜ
Ð& Ð&7Ð8Ó
9Ð9r   c           	     ó¢  — t        | ||dd«      \  } }| j                  j                  }|j                  «       rNt	        j
                  |j                  | j                  t        ||«      j                  «      | j                  «      S |j                  «       rDt	        j
                  |j                  | j                  |j                  «      | j                  «      S |j                  «       rX|rt        | ||t        «       t	        j
                  |j                  | j                  |j                  «      | j                  «      S t        d|› �«      ‚)NTFr=   )r‹   rj   ry   rr   r#   r$   r£   r�   Úminusrs   Úcreate_fsubrJ   r›   ÚsubÚ
create_subr4   ©r§   r¨   r�   r)   rN   s        r   r°   r°     sù   € ä/°°u¸gÀtÈUÓS�L€Eˆ5Ø—
‘
×!Ñ!€Ià×ÑÔÜ�y‰y˜×.Ñ.¨u¯|©|¼UÀ5È'Ó=R×=YÑ=YÓZÐ\a×\fÑ\fÓgÐgà×ÑÔÜ�y‰y˜×,Ñ,¨U¯\©\¸5¿<¹<ÓHÈ%Ï*É*ÓUÐUà	×	Ñ	Ô	ÙÜ,¨U°E¸7ÄCÔHÜ�y‰y˜×+Ñ+¨E¯L©L¸%¿,¹,ÓGÈÏÉÓTÐTÜ
Ð& y kÐ2Ó
3Ð3r   c                óâ  — t        | ||«      \  } }| j                  j                  }|j                  «       rDt	        j
                  |j                  | j                  |j                  «      | j                  «      S |j                  «       rX|rt        | ||t        «       t	        j
                  |j                  | j                  |j                  «      | j                  «      S t        d|› �«      ‚©Nr=   )r‹   rj   ry   rs   r#   r$   Úcreate_fmulr�   rJ   r›   ÚmulÚ
create_mulr4   r²   s        r   r¶   r¶     s³   € ä/°°u¸gÓF�L€Eˆ5Ø—
‘
×!Ñ!€Ià×ÑÔÜ�y‰y˜×,Ñ,¨U¯\©\¸5¿<¹<ÓHÈ%Ï*É*ÓUÐUà	×	Ñ	Ô	ÙÜ,¨U°E¸7ÄCÔHÜ�y‰y˜×+Ñ+¨E¯L©L¸%¿,¹,ÓGÈÏÉÓTÐTÜ
Ð& y kÐ2Ó
3Ð3r   c           	     ó6  — t        | ||dddd«      \  } }| j                  j                  }|j                  j                  }|j                  «       r|j	                  «       rt        |||«      }nè|j	                  «       r|j                  «       rt        | ||«      } nº|j	                  «       rG|j	                  «       r7t        | t        j                  |«      } t        |t        j                  |«      }nc|j                  «       rE|j                  «       r5|j                  |j                  kD  rt        |||«      }nt        | ||«      } nt        d|› �«      ‚t        j                  |j                  | j                  |j                  «      | j                  «      S ©NFTr=   )r‹   rj   ry   rs   rJ   r}   r#   rC   Úfp_mantissa_widthr4   r$   Úcreate_fdivr�   )r§   r¨   r)   r©   rª   s        r   Útruedivr¼   #  sH  € Ü/°°u¸gÀuÈeÐUYÐ[_Ó`�L€Eˆ5Ø—j‘j×'Ñ'€OØ—j‘j×'Ñ'€Oà×"Ñ"Ô$¨×)?Ñ)?Ô)AÜ�U˜O¨WÓ5‰à	×	Ñ	Ô	! o×&AÑ&AÔ&CÜ�U˜O¨WÓ5‰à	×	Ñ	Ô	! o×&<Ñ&<Ô&>Ü�UœBŸJ™J¨Ó0ˆÜ�UœBŸJ™J¨Ó0‰à	×	$Ñ	$Ô	&¨?×+FÑ+FÔ+HØ×,Ñ,¨×/PÑ/PÒPÜ˜ °Ó9‰Eä˜ °Ó9‰Eô Ð*¨?Ð*;Ð<Ó=Ð=Ü�9‰9�W×(Ñ(¨¯©°u·|±|ÓDÀeÇjÁjÓQÐQr   c           	     óZ  — t        | ||dddd«      \  } }| j                  j                  }|j                  j                  }|j                  «       rÎ|j                  «       r¾t	        ||«      }t        | ||«      } t        |||«      }|j                  «       rDt        j                  |j                  | j                  |j                  «      | j                  «      S t        j                  |j                  | j                  |j                  «      | j                  «      S t        d|› �«      ‚r¹   )r‹   rj   ry   rJ   r;   r}   Úis_int_signedr#   r$   Úcreate_sdivr�   Úcreate_udivr4   )r§   r¨   r)   r©   rª   Úret_tys         r   ÚfloordivrÂ   =  sï   € Ü/°°u¸gÀuÈeÐUYÐ[_Ó`�L€Eˆ5Ø—j‘j×'Ñ'€OØ—j‘j×'Ñ'€OØ×ÑÔ O×$:Ñ$:Ô$<Ü% o°ÓGˆÜ�U˜F GÓ,ˆÜ�U˜F GÓ,ˆØ×ÑÔ!Ü—9‘9˜W×0Ñ0°·±¸u¿|¹|ÓLÈeÏjÉjÓYÐYä—9‘9˜W×0Ñ0°·±¸u¿|¹|ÓLÈeÏjÉjÓYÐYÜ
Ð& Ð&7Ð8Ó
9Ð9r   c           	     ód  — | j                   j                  }|j                   j                  }|j                  «       r|j                  «       st        d«      ‚t	        | ||dddd«      \  } }|j                  | j                  |j                  «      }t        j                  || j                   «      S )Nz4both operands of fdiv must have floating scalar typeFT)	rj   ry   rs   r4   r‹   r»   r�   r#   r$   )r§   r¨   Úieee_roundingr)   r©   rª   r–   s          r   ÚfdivrÅ   L  sŽ   € à—j‘j×'Ñ'€OØ—j‘j×'Ñ'€OØ×&Ñ&Ô(°×0KÑ0KÔ0MÜÐNÓOÐOÜ/°°u¸gÀuÈeÐUZÐ\`Óa�L€Eˆ5Ø
×
Ñ
˜eŸl™l¨E¯L©LÓ
9€CÜ�9‰9�S˜%Ÿ*™*Ó%Ð%r   c           	     ó.  — t        | ||dddd«      \  } }| j                  j                  }|j                  j                  }|j                  «       rDt	        j
                  |j                  | j                  |j                  «      | j                  «      S |j                  «       rä|j                  |j                  k7  r3t        d|j                  «       z   dz   |j                  «       z   dz   «      ‚|j                  «       rDt	        j
                  |j                  | j                  |j                  «      | j                  «      S t	        j
                  |j                  | j                  |j                  «      | j                  «      S t        d|› �«      ‚)NFTzCannot mod z by r>   r=   )r‹   rj   ry   rs   r#   r$   Úcreate_fremr�   rJ   r0   r4   r   r¾   Úcreate_sremÚcreate_urem)r§   r¨   r)   rN   rª   s        r   ÚmodrÊ   W  sF  € Ü/°°u¸gÀuÈeÐUYÐ[_Ó`�L€Eˆ5Ø—
‘
×!Ñ!€IØ—j‘j×'Ñ'€Oà×ÑÔÜ�y‰y˜×,Ñ,¨U¯\©\¸5¿<¹<ÓHÈ%Ï*É*ÓUÐUà	×	Ñ	Ô	Ø×#Ñ# ×'EÑ'EÒEÜ˜M¨I×,>Ñ,>Ó,@Ñ@À6ÑIÈO×LdÑLdÓLfÑfð joñ oó pð pð ×"Ñ"Ô$Ü—9‘9˜W×0Ñ0°·±¸u¿|¹|ÓLÈeÏjÉjÓYÐYä—9‘9˜W×0Ñ0°·±¸u¿|¹|ÓLÈeÏjÉjÓYÐYÜ
Ð& y kÐ2Ó
3Ð3r   c                óf  — t        | ||«      \  } }| j                  }|j                  «       rÐ|t        j                  j
                  k(  rDt        j                  |j                  | j                  |j                  «      | j                  «      S |t        j                  j                  k(  rDt        j                  |j                  | j                  |j                  «      | j                  «      S t        d|› �«      ‚|j                  «       rDt        j                  |j                  | j                  |j                  «      | j                  «      S |j                  «       rDt        j                  |j!                  | j                  |j                  «      | j                  «      S t#        d|› �«      ‚©NzUnexpected propagate_nan úUnexpected dtype )r‹   r1   rs   r#   ÚPropagateNanÚALLr$   Úcreate_minimumfr�   rj   ÚNONEÚcreate_minnumfr"   r¾   Úcreate_minsirz   Úcreate_minuir4   ©rk   ÚyÚpropagate_nanr)   r1   s        r   ÚminimumrØ   p  ó3  € Ü'¨¨1¨gÓ6�D€A€qØ�G‰G€EØ×ÑÔØœBŸO™O×/Ñ/Ò/Ü—9‘9˜W×4Ñ4°Q·X±X¸q¿x¹xÓHÈ!Ï&É&ÓQÐQØœbŸo™o×2Ñ2Ò2Ü—9‘9˜W×3Ñ3°A·H±H¸a¿h¹hÓGÈÏÉÓPÐPäÐ8¸¸ÐHÓIÐIØ	×	Ñ	Ô	Ü�y‰y˜×-Ñ-¨a¯h©h¸¿¹ÓAÀ1Ç6Á6ÓJÐJØ	×	Ñ	Ô	 Ü�y‰y˜×-Ñ-¨a¯h©h¸¿¹ÓAÀ1Ç6Á6ÓJÐJäÐ+¨E¨7Ð3Ó4Ð4r   c                óf  — t        | ||«      \  } }| j                  }|j                  «       rÐ|t        j                  j
                  k(  rDt        j                  |j                  | j                  |j                  «      | j                  «      S |t        j                  j                  k(  rDt        j                  |j                  | j                  |j                  «      | j                  «      S t        d|› �«      ‚|j                  «       rDt        j                  |j                  | j                  |j                  «      | j                  «      S |j                  «       rDt        j                  |j!                  | j                  |j                  «      | j                  «      S t#        d|› �«      ‚rÌ   )r‹   r1   rs   r#   rÎ   rÏ   r$   Úcreate_maximumfr�   rj   rÑ   Úcreate_maxnumfr"   r¾   Úcreate_maxsirz   Úcreate_maxuir4   rÕ   s        r   Úmaximumrß   ‚  rÙ   r   c                óX  — t        |||«      \  }}t        | ||«      \  } }t        | ||«      \  } }| j                  }|j                  «       rPt        j                  |j                  | j                  |j                  |j                  |«      | j                  «      S t        d|› d�«      ‚)NrÍ   z(. Only floating point clamp is supported)	r‹   r1   rs   r#   r$   Úcreate_clampfr�   rj   r4   )rk   ÚminÚmaxr×   r)   r1   s         r   Úclamprä   ”  s™   € Ü+¨C°°gÓ>�H€CˆÜ)¨!¨S°'Ó:�F€A€sÜ)¨!¨S°'Ó:�F€A€sà�G‰G€EØ×ÑÔÜ�y‰y˜×.Ñ.¨q¯x©x¸¿¹ÀSÇZÁZÐQ^Ó_Ðab×agÑagÓhÐhäÐ+¨E¨7Ð2ZÐ[Ó\Ð\r   c                ó:  — t        | ||«      \  } }| j                  j                  }|j                  j                  }|j                  «       r|j                  «       st	        ||«      ‚t        ||«      }||k7  rt        | ||«      } ||k7  rt        |||«      }| |fS rq   )r‹   rj   ry   rJ   r   r;   r}   )r§   r¨   r)   Úinput_sca_tyÚother_sca_tyrŠ   s         r   Úbitwise_op_type_checking_implrè   ¥  s˜   € ä/°°u¸gÓF�L€Eˆ5Ø—:‘:×$Ñ$€LØ—:‘:×$Ñ$€LØ×ÑÔ ¨×(;Ñ(;Ô(=Ü'¨°lÓCÐCÜ% l°LÓA€JØ�\Ò!Ü�U˜J¨Ó0ˆØ�\Ò!Ü�U˜J¨Ó0ˆØ�%ˆ<Ðr   c                óª   — t        | ||«      \  } }t        j                  |j                  | j                  |j                  «      | j
                  «      S rq   )rè   r#   r$   Ú
create_andr�   rj   ©r§   r¨   r)   s      r   r‘   r‘   ´  ó>   € Ü0°¸¸wÓG�L€Eˆ5Ü�9‰9�W×'Ñ'¨¯©°e·l±lÓCÀUÇZÁZÓPÐPr   c                óª   — t        | ||«      \  } }t        j                  |j                  | j                  |j                  «      | j
                  «      S rq   )rè   r#   r$   Ú	create_orr�   rj   rë   s      r   Úor_rï   ¹  s>   € Ü0°¸¸wÓG�L€Eˆ5Ü�9‰9�W×&Ñ& u§|¡|°U·\±\ÓBÀEÇJÁJÓOÐOr   c                óª   — t        | ||«      \  } }t        j                  |j                  | j                  |j                  «      | j
                  «      S rq   )rè   r#   r$   Ú
create_xorr�   rj   rë   s      r   Úxor_rò   ¾  rì   r   c                ó  — | j                   j                  «       s t        | t        j                  d«      |«      } |j                   j                  «       s t        |t        j                  d«      |«      }t        | ||«      S ©Nr`   )rj   Úis_int1Úbitcastr#   r1   r‘   rë   s      r   Úlogical_andr÷   Ã  s_   € Ø�:‰:×ÑÔÜ˜œrŸx™x¨Ó/°Ó9ˆØ�:‰:×ÑÔÜ˜œrŸx™x¨Ó/°Ó9ˆÜ��u˜gÓ&Ð&r   c                ó  — | j                   j                  «       s t        | t        j                  d«      |«      } |j                   j                  «       s t        |t        j                  d«      |«      }t        | ||«      S rô   )rj   rõ   rö   r#   r1   rï   rë   s      r   Ú
logical_orrù   Ë  s_   € Ø�:‰:×ÑÔÜ˜œrŸx™x¨Ó/°Ó9ˆØ�:‰:×ÑÔÜ˜œrŸx™x¨Ó/°Ó9ˆÜˆu�e˜WÓ%Ð%r   c                óŽ   — | j                   j                  «       s t        | t        j                  d«      |«      } t        | |«      S rô   )rj   rõ   rö   r#   r1   Úinvert)r§   r)   s     r   Únot_rü   Ó  s6   € Ø�:‰:×ÑÔÜ˜œrŸx™x¨Ó/°Ó9ˆÜ�%˜Ó!Ð!r   c                óª   — t        | ||«      \  } }t        j                  |j                  | j                  |j                  «      | j
                  «      S rq   )rè   r#   r$   Úcreate_lshrr�   rj   rë   s      r   Úlshrrÿ   Ù  ó>   € Ü0°¸¸wÓG�L€Eˆ5Ü�9‰9�W×(Ñ(¨¯©°u·|±|ÓDÀeÇjÁjÓQÐQr   c                óª   — t        | ||«      \  } }t        j                  |j                  | j                  |j                  «      | j
                  «      S rq   )rè   r#   r$   Úcreate_ashrr�   rj   rë   s      r   Úashrr  Þ  r   r   c                óª   — t        | ||«      \  } }t        j                  |j                  | j                  |j                  «      | j
                  «      S rq   )rè   r#   r$   Ú
create_shlr�   rj   rë   s      r   Úshlr  ã  rì   r   c                ó   — | S rq   rW   )r§   s    r   Úplusr  í  s   € Ø€Lr   c                ó  — | j                   j                  }|j                  «       rt        d|j	                  «       z   dz   «      ‚t        j                  |j                  |j                  |«      «      |«      }t        || d|«      S )Nz$wrong type argument to unary minus (ú)T)
rj   ry   rr   r"   r   r#   r$   Úget_null_valuer¡   r°   )r§   r)   ræ   Ú_0s       r   r®   r®   ñ  st   € Ø—:‘:×$Ñ$€LØ×ÑÔÜÐ?À,×BWÑBWÓBYÑYÐ\_Ñ_Ó`Ð`Ü	�‰�7×)Ñ)¨,×*<Ñ*<¸WÓ*EÓFÈÓ	U€BÜˆr�5˜$ Ó(Ð(r   c                ó.  — | j                   j                  }|j                  «       s|j                  «       rt	        d|j                  «       z   dz   «      ‚t        j                  |j                  |j                  |«      «      |«      }t        | ||«      S )Nz%wrong type argument to unary invert (r
  )rj   ry   rr   rs   r"   r   r#   r$   Úget_all_ones_valuer¡   rò   )r§   r)   ræ   Ú_1s       r   rû   rû   ù  s}   € Ø—:‘:×$Ñ$€LØ×ÑÔ × 8Ñ 8Ô :ÜÐ@À<×CXÑCXÓCZÑZÐ]`Ñ`ÓaÐaÜ	�‰�7×-Ñ-¨l×.@Ñ.@ÀÓ.IÓJÈLÓ	Y€BÜ��r˜7Ó#Ð#r   c                óÊ   — | j                   j                  «       st        j                  S | j                   j                  }t        j
                  t        j                  |«      S rq   )rj   rž   r#   r`   ÚshaperŸ   )Úvr  s     r   Ú
_bool_liker    s;   € Ø�6‰6�?‰?ÔÜ�w‰wˆØ�F‰F�L‰L€EÜ�=‰=œŸ™ %Ó(Ð(r   c                ó\  — t        | ||«      \  } }| j                  j                  }|j                  «       rCt	        j
                  |j                  | j                  |j                  «      t        | «      «      S |j                  «       r–|j                  «       rCt	        j
                  |j                  | j                  |j                  «      t        | «      «      S t	        j
                  |j                  | j                  |j                  «      t        | «      «      S t        d|› �«      ‚r´   )r‹   rj   ry   rs   r#   r$   Úcreate_fcmpOGTr�   r  rJ   r¾   Úcreate_icmpSGTÚcreate_icmpUGTr4   ©r§   r¨   r)   rN   s       r   Úgreater_thanr    óà   € Ü/°°u¸gÓF�L€Eˆ5Ø—
‘
×!Ñ!€Ià×ÑÔÜ�y‰y˜×/Ñ/°·±¸e¿l¹lÓKÌZÐX]ÓM^Ó_Ð_à	×	Ñ	Ô	Ø×"Ñ"Ô$Ü—9‘9˜W×3Ñ3°E·L±LÀ%Ç,Á,ÓOÔQ[Ð\aÓQbÓcÐcä—9‘9˜W×3Ñ3°E·L±LÀ%Ç,Á,ÓOÔQ[Ð\aÓQbÓcÐcÜ
Ð& y kÐ2Ó
3Ð3r   c                ó\  — t        | ||«      \  } }| j                  j                  }|j                  «       rCt	        j
                  |j                  | j                  |j                  «      t        | «      «      S |j                  «       r–|j                  «       rCt	        j
                  |j                  | j                  |j                  «      t        | «      «      S t	        j
                  |j                  | j                  |j                  «      t        | «      «      S t        d|› �«      ‚r´   )r‹   rj   ry   rs   r#   r$   Úcreate_fcmpOGEr�   r  rJ   r¾   Úcreate_icmpSGEÚcreate_icmpUGEr4   r  s       r   r“   r“     r  r   c                ó\  — t        | ||«      \  } }| j                  j                  }|j                  «       rCt	        j
                  |j                  | j                  |j                  «      t        | «      «      S |j                  «       r–|j                  «       rCt	        j
                  |j                  | j                  |j                  «      t        | «      «      S t	        j
                  |j                  | j                  |j                  «      t        | «      «      S t        d|› �«      ‚r´   )r‹   rj   ry   rs   r#   r$   Úcreate_fcmpOLTr�   r  rJ   r¾   Úcreate_icmpSLTÚcreate_icmpULTr4   r  s       r   Ú	less_thanr#  )  r  r   c                ó\  — t        | ||«      \  } }| j                  j                  }|j                  «       rCt	        j
                  |j                  | j                  |j                  «      t        | «      «      S |j                  «       r–|j                  «       rCt	        j
                  |j                  | j                  |j                  «      t        | «      «      S t	        j
                  |j                  | j                  |j                  «      t        | «      «      S t        d|› �«      ‚r´   )r‹   rj   ry   rs   r#   r$   Úcreate_fcmpOLEr�   r  rJ   r¾   Úcreate_icmpSLEÚcreate_icmpULEr4   r  s       r   r’   r’   8  r  r   c                ó¶  — t        | ||«      \  } }| j                  j                  }|j                  «       rCt	        j
                  |j                  | j                  |j                  «      t        | «      «      S |j                  «       rCt	        j
                  |j                  | j                  |j                  «      t        | «      «      S t        d|› �«      ‚r´   )r‹   rj   ry   rs   r#   r$   Úcreate_fcmpOEQr�   r  rJ   Úcreate_icmpEQr4   r  s       r   Úequalr+  G  ó¦   € Ü/°°u¸gÓF�L€Eˆ5Ø—
‘
×!Ñ!€Ià×ÑÔÜ�y‰y˜×/Ñ/°·±¸e¿l¹lÓKÌZÐX]ÓM^Ó_Ð_à	×	Ñ	Ô	Ü�y‰y˜×.Ñ.¨u¯|©|¸U¿\¹\ÓJÌJÐW\ÓL]Ó^Ð^Ü
Ð& y kÐ2Ó
3Ð3r   c                ó¶  — t        | ||«      \  } }| j                  j                  }|j                  «       rCt	        j
                  |j                  | j                  |j                  «      t        | «      «      S |j                  «       rCt	        j
                  |j                  | j                  |j                  «      t        | «      «      S t        d|› �«      ‚r´   )r‹   rj   ry   rs   r#   r$   Úcreate_fcmpUNEr�   r  rJ   Úcreate_icmpNEr4   r  s       r   Ú	not_equalr0  S  r,  r   c                óž  — t        | t        «      rt        |t        «      st        d«      ‚t        | dz	  «      }t        |dz	  «      }|s|rt        d«      ‚|| k  rt        d«      ‚|| z
  }||dz
  z  dk7  rt        d«      ‚|g}t	        j
                  t        j                  |«      }t	        j                  |j                  | |«      |«      S )Nz/arange's arguments must be of type tl.constexpré    zarange must fit in int32z=arange's end argument must be greater than the start argumentr   r   z#arange's range must be a power of 2)	r]   ra   r"   r^   r#   rŸ   r&   r$   Úcreate_make_range)ÚstartÚendr)   Úis_start_int64Úis_end_int64Úranger  rÁ   s           r   Úaranger9  d  s¿   € Ü�eœSÔ!¬°C¼Ô)=ÜÐJÓKÐKÜ˜% 2™+Ó&€NÜ˜˜r™	“?€LÙ™ÜÐ3Ó4Ð4Ø
ˆe‚|ÜÐXÓYÐYØ�%‰K€EØ�˜‘Ñ Ò!ÜÐ>Ó?Ð?ØˆG€EÜ�]‰]œ2Ÿ8™8 UÓ+€FÜ�9‰9�W×.Ñ.¨u°cÓ:¸FÓCÐCr   c                ó€  — t        |t        j                  «      r.|j                  j                  dk(  sJ d«       ‚t        |||«      }nj|€t        d«      ‚|dk(  r!|j                  |j                  |«      «      }n!t        |d|j                  › �«      } ||«      }t        j                  ||«      }t        || |«      S )Nr   zonly accepts size-1 tensorz2dtype must be specified when value is not a tensorr   Úget_)r]   r#   r$   Únumelr@   r}   r"   r  r¡   ÚgetattrÚnameÚsplat)r  r@   r1   r)   Úget_value_fns        r   re   re   u  s¯   € Ü�%œŸ™Ô#Ø�{‰{× Ñ  AÒ%ÐCÐ'CÓCÐ%Ü�U˜E 7Ó+‰ð ˆ=ÜÐQÓRÐRØ�AŠ:Ø×*Ñ*¨5¯;©;°wÓ+?Ó@‰Eä" 7¨d°5·:±:°,Ð,?Ó@ˆLÙ  Ó'ˆEÜ—	‘	˜% Ó'ˆä�˜˜wÓ'Ð'r   c                ó  — | j                   j                  «       rJ d«       ‚t        |«      dk(  r| S t        j                  | j
                  |«      }t        j                  |j                  | j                  |«      |«      S )NzCannot splat a block tensorr   )	rj   rž   Úlenr#   rŸ   r1   r$   Úcreate_splatr�   )r@   r  r)   rÁ   s       r   r?  r?  Œ  sd   € Ø�z‰z×"Ñ"Ô$ÐCÐ&CÓCÐ$Ü
ˆ5ƒz�Q‚ØˆÜ�]‰]˜5Ÿ;™;¨Ó.€FÜ�9‰9�W×)Ñ)¨%¯,©,¸Ó>ÀÓGÐGr   c                ó  — d}|D ]  }||z  }Œ	 | j                   j                  |k7  rt        d«      ‚t        j                  | j                   j
                  |«      }t        j                  |j                  | j                  ||«      |«      S )Nr   z:reshape() cannot change total number of elements in tensor)	rj   r<  r"   r#   rŸ   ry   r$   Úcreate_reshaper�   )r§   Ú	dst_shapeÚcan_reorderr)   r<  ÚsrÁ   s          r   ÚreshaperI  ”  s|   € Ø€EØò ˆØ�‰
‰ðà‡z�z×Ñ˜5Ò ÜÐUÓVÐVÜ�]‰]˜5Ÿ:™:×,Ñ,¨iÓ8€FÜ�9‰9�W×+Ñ+¨E¯L©L¸)À[ÓQÐSYÓZÐZr   c                óŒ  — | j                   D �cg c]  }t        j                  |«      ‘Œ }}|j                  |d«       | j                  j                  «       st        | ||¬«      S t        j                  | j                  j                  |«      }t        j                  |j                  | j                  |«      |«      S c c}w )Nr   )r  r)   )r  r#   Ú_constexpr_to_valueÚinsertrj   rž   r?  rŸ   ry   r$   Úcreate_expand_dimsr�   )r§   r(   r)   rk   rF  rÁ   s         r   Úexpand_dimsrN  ž  s•   € Ø49·K±KÖ@¨q”×'Ñ'¨Õ*Ð@€IÐ@Ø×Ñ�T˜1Ôà�:‰:×ÑÔ Ü�U )°WÔ=Ð=ä�]‰]˜5Ÿ:™:×,Ñ,¨iÓ8€FÜ�9‰9�W×/Ñ/°·±¸dÓCÀVÓLÐLùò As   �Cc                óL  — |sJ d«       ‚t        | j                  «      dk(  sJ ‚t        j                  | j                  j
                  | j                  d   |j                  d   z   g«      }t        j                  |j                  | j                  |j                  «      |«      S )Nz;current implementation of `cat` always may reorder elementsr   r   )	rB  r  r#   rŸ   rj   ry   r$   Ú
create_catr�   )r   r€   rG  r)   Úret_types        r   ÚcatrR  ©  s|   € ÙÐUÐUÓUˆ;Üˆs�y‰y‹>˜QÒÐÐÜ�}‰}˜SŸX™XŸ_™_¨s¯y©y¸©|¸c¿i¹iÈ¹lÑ/JÐ.KÓL€HÜ�9‰9�W×'Ñ'¨¯
©
°C·J±JÓ?ÀÓJÐJr   c                ó  — t        | ||«      \  } }| j                  g k(  }|rt        | d|«      } t        |d|«      }t        | j                  d   t        j
                  «      rt	        j
                  d«      }nd}| j                  |gz   }t	        j                  | j                  j                  |«      }t	        j                  |j                  | j                  |j                  «      |«      }|rt        |dgd|¬«      }|S )Nr   éÿÿÿÿr	   F©rG  r)   )r~   r  rN  r]   r#   rh   rŸ   rj   ry   r$   Úcreate_joinr�   rI  )ÚaÚbr)   Ú
was_rank_1ÚtwoÚ	new_shaperQ  r–   s           r   Újoinr\  °  sÐ   € Ü  1 gÓ.�D€A€qð —‘˜B‘€JÙÜ˜˜1˜gÓ&ˆÜ˜˜1˜gÓ&ˆä�!—'‘'˜"‘+œrŸ|™|Ô,Ü�l‰l˜1‹o‰àˆØ—‘˜3˜%‘€Iä�}‰}˜QŸV™VŸ]™]¨IÓ6€HÜ
�)‰)�G×'Ñ'¨¯©°!·(±(Ó;¸XÓ
F€CáÜ�c˜A˜3¨E¸7ÔCˆà€Jr   c                óŠ  — t        | j                  «      dkD  sJ ‚t        j                  | j                  d   «      dk(  sJ ‚| j                  d d }t        j                  | j
                  j                  |«      }|j                  | j                  «      \  }}t        j                  ||«      t        j                  ||«      fS )Nr   rT  r	   )
rB  r  r#   rK  rŸ   rj   ry   Úcreate_splitr�   r$   )rW  r)   r[  rQ  ÚoutLHSÚoutRHSs         r   Úsplitra  É  sž   € Ü�—‘‹L˜1ÒÐÐÜ×"Ñ" 1§7¡7¨2¡;Ó/°1Ò4Ð5Ð4à—‘˜˜�€IÜ�}‰}˜QŸV™VŸ]™]¨IÓ6€HØ×)Ñ)¨!¯(©(Ó3�N€FˆFä
�	‰	�&˜(Ó#Ü
�	‰	�&˜(Ó#ðð r   c                óÊ  — t        | j                  «      t        |«      k7  rt        d«      ‚t        d„ |D «       «      t	        t        t        |«      «      «      k7  rt        d|› �«      ‚t        j                  | j                  j                  |D �cg c]  }| j                  |   ‘Œ c}«      }t        j                  |j                  | j                  |«      |«      S c c}w )Nz5permute dims must have the same length as input shapec              3  óF   K  — | ]  }t        j                  |«      –— Œ y ­wrq   )r#   rK  )Ú.0Úds     r   ú	<genexpr>zpermute.<locals>.<genexpr>Ù  s   è ø€ Ò6¨AŒb×$Ñ$ Q×'Ñ6ùs   ‚!z?permute dims must be a permutation of 0, 1, ..., n-1, but were )rB  r  r"   ÚsortedÚlistr8  r#   rŸ   rj   ry   r$   Úcreate_transr�   )r§   Údimsr)   re  rQ  s        r   Úpermuterk  Ö  s®   € Ü
ˆ5�;‰;Óœ3˜t›9Ò$ÜÐPÓQÐQÜÑ6°Ô6Ó6¼$¼uÄSÈÃYÓ?OÓ:PÒPÜÐZÐ[_ÐZ`ÐaÓbÐbä�}‰}˜UŸZ™Z×.Ñ.ÈÖ0NÀA°·±¸Q³Ò0NÓO€HÜ�9‰9�W×)Ñ)¨%¯,©,¸Ó=¸xÓHÐHùò 1Os   ÂC 
c                ó˜  — | j                   j                  «       sPt        j                  | j                   |«      }t        j                  |j                  | j                  |«      |«      S | j                   j                  «       }t        |«      t        |«      k7  rt        d|› d|› �«      ‚||k(  r| S t        |«      D ]0  \  }}||   |k7  sŒ|dk7  sŒt        d||   › d|› d|› d|› d|› �
«      ‚ t        j                  | j                   j                  |«      }t        j                  |j                  | j                  |«      |«      S )Nz!Cannot broadcast, rank mismatch: z, r   z3Cannot broadcast, the expanded size of the tensor (z ) must match the existing size (z) at non-singleton dimension ú: )rj   rž   r#   rŸ   r$   rC  r�   r    rB  r"   Ú	enumeratery   Úcreate_broadcast)r§   r  r)   rÁ   Ú	src_shapeÚiÚitems          r   Úbroadcast_impl_shapers  à  s;  € Ø�:‰:×ÑÔ Ü—‘˜uŸz™z¨5Ó1ˆÜ�y‰y˜×-Ñ-¨e¯l©l¸EÓBÀFÓKÐKØ—
‘
×+Ñ+Ó-€IÜ
ˆ9ƒ~œ˜U›Ò#ÜÐ<¸Y¸KÀrÈ%ÈÐQÓRÐRØ�	ÒØˆÜ˜YÓ'ò <‰ˆˆ4Ø�‰8�tÓ ¨£	ÜÐRÐSXÐYZÑS[ÐR\ð ]?Ø?C¸fð E!Ø!"  2 i [°°5°'ð;ó <ð <ð<ô
 �]‰]˜5Ÿ:™:×,Ñ,¨eÓ4€FÜ�9‰9�W×-Ñ-¨e¯l©l¸EÓBÀFÓKÐKr   c           	     ó@  — | j                   }|j                   }|j                  «       r||j                  «       slt        j                  |j                  |j
                  «      }t        j                  |j                  |j                  |j                  «       «      |«      }| |fS |j                  «       s||j                  «       rlt        j                  |j                  |j
                  «      }t        j                  |j                  | j                  |j                  «       «      |«      } | |fS |j                  «       �rÚ|j                  «       �rÉ|j                  «       }|j                  «       }t        |«      t        |«      k  rœt        t        |«      t        |«      «      D ]z  }t        j                  |j                  | j                  d«      t        j                  |j                  dg|j                  z   «      «      } | j                   }|j                  «       }Œ| n²t        |«      t        |«      k  r›t        t        |«      t        |«      «      D ]z  }t        j                  |j                  |j                  d«      t        j                  |j                  dg|j                  z   «      «      }|j                   }|j                  «       }Œ| t        |«      t        |«      k(  sJ ‚g }t        |«      D ]q  \  }	}
||	   }|
dk(  r|j                  |«       Œ"|dk(  s||
k(  r|j                  |
«       Œ>t!        dt#        |	«      z   dz   t#        |
«      z   dz   t#        |«      z   «      ‚ ||k7  rPt        j                  |j                  |«      }t        j                  |j%                  | j                  |«      |«      } ||k7  rPt        j                  |j                  |«      }t        j                  |j%                  |j                  |«      |«      }| |fS )Nr   r   z?Cannot make_shape_compatible: incompatible dimensions at index rm  r   )rj   rž   r#   rŸ   ry   r  r$   rC  r�   r    rB  r8  rM  Úvaluesrn  Úappendr"   Ústrro  )r   r€   r)   Úlhs_tyÚrhs_tyÚ	lhs_shapeÚ	rhs_shapeÚ_Ú	ret_shaperq  ÚleftÚrightrÁ   s                r   r~   r~   ò  s_  € Ø�X‰X€FØ�X‰X€Fð ‡�Ô §¡Ô!2Ü—‘˜vŸ}™}¨f¯l©lÓ;ˆÜ�i‰i˜×,Ñ,¨S¯Z©Z¸×9PÑ9PÓ9RÓSÐU[Ó\ˆðV �ˆ8€OðS �_‰_Ô 6§?¡?Ô#4Ü—‘˜vŸ}™}¨f¯l©lÓ;ˆÜ�i‰i˜×,Ñ,¨S¯Z©Z¸×9PÑ9PÓ9RÓSÐU[Ó\ˆðN �ˆ8€OðK 
�‰Õ	˜vŸ™Õ0Ø×+Ñ+Ó-ˆ	Ø×+Ñ+Ó-ˆ	äˆy‹>œC 	›NÒ*äœ3˜y›>¬3¨y«>Ó:ò 6�Ü—i‘i × :Ñ :¸3¿:¹:ÀqÓ IÜ "§¡¨f¯m©m¸a¸SÀ9×CSÑCSÑ=SÓ TóV�àŸ™�Ø"×3Ñ3Ó5‘	ñ	6ô
 �‹^œc )›nÒ,äœ3˜y›>¬3¨y«>Ó:ò 6�Ü—i‘i × :Ñ :¸3¿:¹:ÀqÓ IÜ "§¡¨f¯m©m¸a¸SÀ9×CSÑCSÑ=SÓ TóV�àŸ™�Ø"×3Ñ3Ó5‘	ð	6ô
 �9‹~¤ Y£Ò/Ð/Ð/àˆ	Ü  Ó+ò 	a‰GˆAˆtØ˜a‘LˆEØ�qŠyØ× Ñ  Õ'Ø˜1’* %¨4¢-Ø× Ñ  Õ&ä ð "-Ü/2°1«vñ"6Ø8<ñ"=Ü?BÀ4»yñ"IØKRñ"SÜUXÐY^ÓU_ñ"`ó að að	að ˜	Ò!Ü—]‘] 6§=¡=°)Ó<ˆFÜ—)‘)˜G×4Ñ4°S·Z±ZÀÓKÈVÓTˆCØ˜	Ò!Ü—]‘] 6§=¡=°)Ó<ˆFÜ—)‘)˜G×4Ñ4°S·Z±ZÀÓKÈVÓTˆCà�ˆ8€Or   c                ó¢   — | €y | dk(  rt         j                  j                  S | dk(  rt         j                  j                  S t	        d| › d�«      ‚)NÚrtneÚrtzzInvalid rounding mode: z0. Supported rounding modes are 'rtne' and 'rtz'.)r
   ÚROUNDING_MODEÚRTNEÚRTZr"   )Úrounding_modes    r   Ú_str_to_rounding_moder‡  ,  sU   € ØÐØØ˜ÒÜ×Ñ×$Ñ$Ð$Ø˜ÒÜ×Ñ×#Ñ#Ð#Ü
Ð.¨}¨oÐ=mÐnÓ
oÐor   c                óF  — | j                   }|j                  «       r8t        j                  |j                  | j                   j                  «       «      }||k(  r| S |j                  }|j                  }|j                  «       s|j                  «       rt        | ||«      S |j                  }|j                  }||k7  r&t        dt        |«      z   dz   t        |«      z   «      ‚t        j                  |j                  | j                  |j                  |«      «      |«      S )Nz!Cannot bitcast data-type of size z to data-type of size )rj   rž   r#   rŸ   ry   r    rr   r}   Úprimitive_bitwidthr"   rw  r$   Úcreate_bitcastr�   r¡   )r§   Údst_tyr)   Úsrc_tyÚ
src_sca_tyÚ
dst_sca_tyÚsrc_bitsÚdst_bitss           r   rö   rö   6  sõ   € Ø�Z‰Z€FØ‡�ÔÜ—‘˜vŸ}™}¨e¯j©j×.IÑ.IÓ.KÓLˆØ�ÒØˆØ—‘€JØ—‘€JØ×ÑÔ˜j×/Ñ/Ô1Ü�E˜6 7Ó+Ð+à×,Ñ,€HØ×,Ñ,€HØ�8ÒÜÐ<¼sÀ8»}ÑLð P.ñ .Ü03°H³ñ>ó ?ð 	?ä�9‰9�W×+Ñ+¨E¯L©L¸&¿,¹,ÀwÓ:OÓPÐRXÓYÐYr   c                óÎ  — | j                   }|j                  «       r8t        j                  |j                  | j                   j                  «       «      }||k(  r| S |j                  }|j                  }t        |«      }d}|j                  «       rf|j                  «       rV|j                  |j                  k  r=|€t        j                  j                  }nH|t        j                  j                  k7  r+d}n(|�&t        dt        |«      z   dz   t        |«      z   «      ‚|j                  «       s|j                  «       r<|j                  j!                  d«      	 €J d«       ‚ |j                  d   | |||¬«      S |j#                  «       r|j                  «       s"|j                  «       r|j#                  «       s|r@t        j$                  |j'                  | j(                  |j+                  |«      |«      |«      S |j-                  «       r|j/                  «       r |j1                  «       r6|j/                  «       s&t3        t3        | t        j4                  |«      ||«      S |j                  «       xr+ |j                  «       xr |j                  |j                  kD  }|r?t        j$                  |j7                  | j(                  |j+                  |«      «      |«      S |j                  «       xr+ |j                  «       xr |j                  |j                  k  }	|	r?t        j$                  |j9                  | j(                  |j+                  |«      «      |«      S |j;                  «       �r|j;                  «       rü|j<                  |j<                  k7  s|j>                  |j>                  k7  rÊ|jA                  «       xr |jC                  «        }
|jC                  «       rW| jD                  j+                  |«      }t        j$                  |jG                  |«      | jD                  «      }tI        | ||«      S t        j$                  |jK                  | j(                  |j+                  |«      |
«      |«      S |jM                  «       �r|j;                  «       rõ|jC                  «       rW| jD                  j+                  |«      }t        j$                  |jG                  |«      | jD                  «      }tI        | ||«      S |jA                  «       r?t        j$                  |jO                  | j(                  |j+                  |«      «      |«      S t        j$                  |jQ                  | j(                  |j+                  |«      «      |«      S |j;                  «       r®|jM                  «       rž|jC                  «       s|jA                  «       s?t        j$                  |jS                  | j(                  |j+                  |«      «      |«      S t        j$                  |jU                  | j(                  |j+                  |«      «      |«      S |jW                  «       r¼|j;                  «       r¬|j<                  }|dk(  r?t        j$                  |jY                  | j(                  |j+                  |«      «      |«      S |d	k(  rWtI        t3        | t        jZ                  |«      t        j$                  |j]                  d
«      t        jZ                  «      |«      S |j;                  «       rO|jW                  «       r?t        j$                  |j_                  | j(                  |j+                  |«      «      |«      S |jW                  «       rO|jW                  «       r?t        j$                  |ja                  | j(                  |j+                  |«      «      |«      S J d| › d|› �«       ‚)NFTz]fp_downcast_rounding should be set only for truncating fp conversions. Source scalar type is z and destination type is Úconvert_custom_typesz0target doesn't provide conversion for this type.©Ú_builderr�   r   r   zcannot cast z to )1rj   rž   r#   rŸ   ry   r    r‡  rs   r‰  r
   rƒ  r„  r"   rw  Úis_fp8e4b15Úcodegen_fnsÚgetrI   r$   Úcreate_fp_to_fpr�   r¡   rG   rF   rH   r}   rC   Úcreate_fp_truncÚcreate_fp_extrJ   r/   r0   r¾   Úis_boolr1   r  r0  r¢   Úis_standard_floatingÚcreate_fp_to_siÚcreate_fp_to_uiÚcreate_ui_to_fpÚcreate_si_to_fprr   Úcreate_ptr_to_intrc   r�   Úcreate_int_to_ptrrŠ  )r§   r‹  r)   Úfp_downcast_roundingrŒ  r�  rŽ  Úuse_custom_roundingÚtruncate_fpÚext_fpÚsign_extendÚtyr  Úbitwidths                 r   r}   r}   I  sq  € à�Z‰Z€FØ‡�ÔÜ—‘˜vŸ}™}¨e¯j©j×.IÑ.IÓ.KÓLˆØ�ÒØˆà—‘€JØ—‘€Jô 1Ð1EÓFÐØÐØ×ÑÔ J×$:Ñ$:ô %Ø
×
'Ñ
'¨*×*GÑ*GÒ
GØÐ'Ä×@PÑ@P×@UÑ@UÑ)=Ø!¤R×%5Ñ%5×%:Ñ%:Ò:ÐRVÑ<OàÐ+Üð 6Ü8;¸J»ñHØJeñfÜhkÐlvÓhwñxó yð yð 	×ÑÔ  J×$:Ñ$:Ô$<Ø×"Ñ"×&Ñ&Ø"ó$Ø+/ð0ð 	dØ1có	dð 0à:ˆw×"Ñ"Ð#9Ñ:¸5À&ÐJ^ÐipÔqÐqð 	×ÑÔ 
× 6Ñ 6Ô 8Ø×ÑÔ  Z×%6Ñ%6Ô%8ÙÜ�y‰y˜×0Ñ0°·±¸v¿|¹|ÈGÓ?TÐVjÓkÐmsÓtÐtð 	×ÑÔ Z×%7Ñ%7Ô%9Ø×ÑÔ Z×%7Ñ%7Ô%9Ü”D˜¤§
¡
¨GÓ4°jÀ'ÓJÐJð
 ×(Ñ(Ó*ò FØ×ÑÓ òFà×%Ñ%¨
×(EÑ(EÑEð ñ Ü�y‰y˜×0Ñ0°·±¸v¿|¹|ÈGÓ?TÓUÐW]Ó^Ð^ð ×#Ñ#Ó%ò FØ×ÑÓ òFà×%Ñ%¨
×(EÑ(EÑEð ñ Ü�y‰y˜×.Ñ.¨u¯|©|¸V¿\¹\È'Ó=RÓSÐU[Ó\Ð\ð ×ÑÕ˜z×0Ñ0Ô2Ø×Ñ :×#:Ñ#:Ò:¸j×>WÑ>WÐ[e×[tÑ[tÒ>tØ ×.Ñ.Ó0ÒM¸×9KÑ9KÓ9MÐ5MˆØ×ÑÔØ—‘×"Ñ" 7Ó+ˆBÜ—‘˜7×1Ñ1°"Ó5°u·{±{ÓCˆBÜ˜U B¨Ó0Ð0ä—9‘9˜W×4Ñ4°U·\±\À6Ç<Á<ÐPWÓCXÐZeÓfÐhnÓoÐoð ×&Ñ&Õ(¨Z×->Ñ->Ô-@Ø×ÑÔØ—‘×"Ñ" 7Ó+ˆBÜ—‘˜7×1Ñ1°"Ó5°u·{±{ÓCˆBÜ˜U B¨Ó0Ð0Ø×%Ñ%Ô'Ü—9‘9˜W×4Ñ4°U·\±\À6Ç<Á<ÐPWÓCXÓYÐ[aÓbÐbä—9‘9˜W×4Ñ4°U·\±\À6Ç<Á<ÐPWÓCXÓYÐ[aÓbÐbð ×ÑÔ˜z×>Ñ>Ô@Ø×ÑÔ z×'?Ñ'?Ô'AÜ—9‘9˜W×4Ñ4°U·\±\À6Ç<Á<ÐPWÓCXÓYÐ[aÓbÐbä—9‘9˜W×4Ñ4°U·\±\À6Ç<Á<ÐPWÓCXÓYÐ[aÓbÐbð ×ÑÔ˜z×0Ñ0Ô2Ø×*Ñ*ˆØ�rŠ>Ü—9‘9˜W×6Ñ6°u·|±|ÀVÇ\Á\ÐRYÓEZÓ[Ð]cÓdÐdØ�qŠ=ÜœT %¬¯©°7Ó;¼R¿Y¹YÀw×GXÑGXÐYZÓG[Ô]_×]eÑ]eÓ=fÐhoÓpÐpð ×ÑÔ˜z×0Ñ0Ô2Ü�y‰y˜×2Ñ2°5·<±<ÀÇÁÈgÓAVÓWÐY_Ó`Ð`ð ×ÑÔ˜z×0Ñ0Ô2Ü�y‰y˜×/Ñ/°·±¸f¿l¹lÈ7Ó>SÓTÐV\Ó]Ð]à4�L   t¨F¨8Ð4Ó4ˆ5r   c                ó"  — t         j                  j                  }| rr| dk(  rt         j                  j                  }|S | dk(  rt         j                  j                  }|S | dk(  rt         j                  j
                  }|S t        d| › d�«      ‚|S )Nz.caú.cgz.cvúCache modifier ú not supported)r
   ÚCACHE_MODIFIERrÑ   ÚCAÚCGÚCVr"   ©Úcache_modifierÚcaches     r   Ú_str_to_load_cache_modifierrµ  ¹  s–   € Ü×Ñ×"Ñ"€EÙØ˜UÒ"Ü×%Ñ%×(Ñ(ˆEð €Lð ˜uÒ$Ü×%Ñ%×(Ñ(ˆEð
 €Lð	 ˜uÒ$Ü×%Ñ%×(Ñ(ˆEð €Lô ˜¨~Ð.>¸nÐMÓNÐNØ€Lr   c                ód  — t         j                  j                  }| r“| dk(  rt         j                  j                  }|S | dk(  rt         j                  j                  }|S | dk(  rt         j                  j
                  }|S | dk(  rt         j                  j                  }|S t        d| › d�«      ‚|S )Nz.wbr«  z.csz.wtr¬  r­  )r
   r®  rÑ   ÚWBr°  ÚCSÚWTr"   r²  s     r   Ú_str_to_store_cache_modifierrº  Ç  s¸   € Ü×Ñ×"Ñ"€EÙØ˜UÒ"Ü×%Ñ%×(Ñ(ˆEð €Lð ˜uÒ$Ü×%Ñ%×(Ñ(ˆEð €Lð ˜uÒ$Ü×%Ñ%×(Ñ(ˆEð
 €Lð	 ˜uÒ$Ü×%Ñ%×(Ñ(ˆEð €Lô ˜¨~Ð.>¸nÐMÓNÐNØ€Lr   c                óà   — t         j                  j                  }| rQ| dk(  rt         j                  j                  }|S | dk(  rt         j                  j                  }|S t        d| › d�«      ‚|S )NÚ
evict_lastÚevict_firstzEviction policy r­  )r
   ÚEVICTION_POLICYÚNORMALÚ
EVICT_LASTÚEVICT_FIRSTr"   )Úeviction_policyÚevictions     r   Ú_str_to_eviction_policyrÄ  ×  su   € Ü×!Ñ!×(Ñ(€HÙØ˜lÒ*Ü×)Ñ)×4Ñ4ˆHð
 €Oð	  Ò-Ü×)Ñ)×5Ñ5ˆHð €Oô Ð/°Ð/@ÀÐOÓPÐPØ€Or   c                ó°   — d }| rQ| dk(  rt         j                  j                  }|S | dk(  rt         j                  j                  }|S t	        d| › d�«      ‚|S )NÚzeroÚnanzPadding option r­  )r
   ÚPADDING_OPTIONÚPAD_ZEROÚPAD_NANr"   )Úpadding_optionÚpaddings     r   Ú_str_to_padding_optionrÍ  ã  sh   € Ø€GÙØ˜VÒ#Ü×'Ñ'×0Ñ0ˆGð
 €Nð	 ˜uÒ$Ü×'Ñ'×/Ñ/ˆGð €Nô ˜¨~Ð.>¸nÐMÓNÐNØ€Nr   c                ód  — t         j                  j                  }| r“| dk(  rt         j                  j                  }|S | dk(  rt         j                  j                  }|S | dk(  rt         j                  j                  }|S | dk(  rt         j                  j
                  }|S t        d| › d�«      ‚|S )NÚacquireÚreleaseÚacq_relÚrelaxedúMemory semantic r­  )r
   ÚMEM_SEMANTICÚACQUIRE_RELEASEÚACQUIREÚRELEASEÚRELAXEDr"   )Ú
sem_optionÚsems     r   Ú_str_to_semrÛ  ï  s®   € Ü
�/‰/×
)Ñ
)€CÙØ˜Ò"Ü—/‘/×)Ñ)ˆCð €Jð ˜9Ò$Ü—/‘/×)Ñ)ˆCð €Jð ˜9Ò$Ü—/‘/×1Ñ1ˆCð
 €Jð	 ˜9Ò$Ü—/‘/×)Ñ)ˆCð €Jô Ð/°
¨|¸>ÐJÓKÐKØ€Jr   c                ó"  — t         j                  j                  }| rr| dk(  rt         j                  j                  }|S | dk(  rt         j                  j                  }|S | dk(  rt         j                  j                  }|S t        d| › d�«      ‚|S )NÚgpuÚctaÚsysrÓ  r­  )r
   ÚMEM_SYNC_SCOPEÚGPUÚCTAÚSYSTEMr"   )Úscope_optionÚscopes     r   Ú_str_to_scoperæ  ÿ  s–   € Ü×Ñ×!Ñ!€EÙØ˜5Ò Ü×%Ñ%×)Ñ)ˆEð €Lð ˜UÒ"Ü×%Ñ%×)Ñ)ˆEð
 €Lð	 ˜UÒ"Ü×%Ñ%×,Ñ,ˆEð €Lô Ð/°¨~¸^ÐLÓMÐMØ€Lr   c                ó~  — | r¶t        | d«      s| g} | D �cg c]*  }t        |t        j                  «      r|j                  n|‘Œ, } }| D ]+  }t        |t
        «      rd|cxk  rt        |«      k  rŒ(J ‚ J ‚ t        | «      dkD  sJ ‚t        | «      t        t        | «      «      k(  sJ d«       ‚t        | «      S yc c}w )NÚ__iter__r   z'Duplicate dimension in `boundary_check`rW   )	Úhasattrr]   r#   rh   r@   ra   rB  Úsetrg  )Úboundary_checkÚblock_shapeÚelemÚdims       r   Ú_canonicalize_boundary_checkrï    sÆ   € ÙÜ�~ zÔ2Ø,Ð-ˆNØ]kÖlÐUY¬
°4¼¿¹Ô(F˜$Ÿ*š*ÈDÑPÐlˆÐlØ!ò 	HˆCÜ˜c¤3Ô'¨A°Ô,G´s¸;Ó7GÓ,GÐGÐGÐ,GÐGÐGð	Hä�>Ó" QÒ&Ð&Ð&Ü�>Ó"¤c¬#¨nÓ*=Ó&>Ò>ÐiÐ@iÓiÐ>Ü�nÓ%Ð%Øùò ms   –/B:c	           
     óÌ  — |€|�t        d«      ‚| j                  j                  j                  }	|	t        j                  k7  sJ d«       ‚|	j                  «       r(|t        j                  j                  k(  rt        d«      ‚| j                  j                  }
t        ||
j                  «       «      }t        j                  |j                  | j                  |||||«      |
«      S )NúK`mask` and `other` arguments cannot be specified for loading block pointersú4`tl.int1` should be rewritten in `tl.make_block_ptr`z@Padding option `nan` is not supported for integer block pointers)r"   rj   Ú
element_tyr#   r`   rJ   r
   rÈ  rÊ  rï  r    r$   Úcreate_tensor_pointer_loadr�   )ÚptrÚmaskr¨   rë  rÌ  r´  rÃ  Úis_volatiler)   Úelt_tyr‹  s              r   Ú_load_block_pointerrù    sÏ   € ð Ð˜5Ð,ÜÐfÓgÐgà�X‰X× Ñ ×+Ñ+€FØ”R—W‘WÒÐTÐTÓTÐØ‡}�}„˜7¤b×&7Ñ&7×&?Ñ&?Ò?ÜÐ[Ó\Ð\ð �X‰X× Ñ €Fô 2°.À&×BYÑBYÓB[Ó\€Nô �9‰9Ø×*Ñ*¨3¯:©:°~ÀwÐPUÐW_ÐalÓmÐouówð wr   c	           
     óT  — | j                   j                  j                  «       s't        d| j                   j	                  «       › d�«      ‚|€|�t        d«      ‚|s|rt        d«      ‚| j                   j                  «       sN|r%|j                   j                  «       rt        d«      ‚|r%|j                   j                  «       rt        d«      ‚| j                   j                  «       rN|�%t        || j                   j                  «       |«      }|�%t        || j                   j                  «       |«      }| j                   j                  }	|	j                  }
|
t        j                  k(  }|r=t        j                  }
t        j                  |
|	j                  «      }	t        | |	|«      } |�t        ||
|«      }| j                   j                  «       r1| j                   j                  «       }t        j                  |
|«      }n|
}|€3t        j                   |j#                  | j$                  |||«      |«      }nLt        j                   |j'                  | j$                  |j$                  |r|j$                  nd |||«      |«      }|rt        |t        j                  |«      }|S )NúUnsupported ptr type z in `tl.load`z)`other` cannot be provided without `mask`zä`padding_option` or `boundary_check` argument is not supported for loading a tensor ofpointers or loading a scalar. Because the compiler does not know the boundary; please use block pointers (defined by `make_block_ptr`) insteadúEMask argument cannot be block type if pointer argument is not a blockzFOther argument cannot be block type if pointer argument is not a block)rj   ry   rr   r"   r   rž   rs  r    ró  r#   r`   Úint8Úpointer_typeÚaddress_spacer}   rŸ   r$   Úcreate_loadr�   Úcreate_masked_load)rõ  rö  r¨   rë  rÌ  r´  rÃ  r÷  r)   Úptr_tyrø  r›  r  r‹  r–   s                  r   Ú_load_legacyr  0  sA  € à�8‰8�?‰?×!Ñ!Ô#ÜÐ0°·±×1BÑ1BÓ1DÐ0EÀ]ÐSÓTÐTð €|˜Ð)ÜÐDÓEÐEÙ‘.Üð Tó Uð 	Uð
 �8‰8×ÑÔÙ�D—I‘I×&Ñ&Ô(ÜÐdÓeÐeÙ�U—Z‘Z×(Ñ(Ô*ÜÐeÓfÐfð ‡x�x×ÑÔØÐÜ'¨¨c¯h©h×.GÑ.GÓ.IÈ7ÓSˆDØÐÜ(¨°·±×0IÑ0IÓ0KÈWÓUˆEð �X‰X�_‰_€FØ×Ñ€Fð œŸ™Ñ€GÙÜ—‘ˆÜ—‘ ¨×)=Ñ)=Ó>ˆÜ�3˜ Ó(ˆð ÐÜ�U˜F GÓ,ˆð ‡x�x×ÑÔØ—‘×)Ñ)Ó+ˆÜ—‘˜v uÓ-‰ð ˆð €|Ü�i‰i˜×+Ñ+¨C¯J©J¸¸xÈÓUÐW]Ó^‰ä�i‰iØ×&Ñ& s§z¡z°4·;±;ÑPUÀÇÂÐ[_ÐafÐhpØ'2ó4Ø5;ó=ˆñ Ü�3œŸ™ Ó)ˆØ€Jr   c	                ó  — t        |«      }	t        |«      }
t        |«      }| j                  j	                  «       r7| j                  j
                  j                  «       rt        | |||||	|
||«	      S t        | |||||	|
||«	      S rq   )	rµ  rÄ  rÍ  rj   rr   ró  rž   rù  r  )rõ  rö  r¨   rë  rË  r³  rÂ  r÷  r)   r´  rÃ  rÌ  s               r   Úloadr  n  s‚   € ô (¨Ó7€EÜ& Ó7€HÜ$ ^Ó4€Gà
‡x�x‡�Ô˜SŸX™X×0Ñ0×9Ñ9Ô;ä" 3¨¨e°^ÀWÈeÐU]Ð_jÐlsÓtÐtô ˜C  u¨n¸gÀuÈhÐXcÐelÓmÐmr   c                ó„   — |j                  | j                  |j                  |«      «      }t        j                  ||«      S rq   )Ú$create_reinterpret_tensor_descriptorr�   r¡   r#   Ú$_experimental_tensor_descriptor_base)Údesc_ptrÚblock_tyr)   r�   s       r   Úreinterpret_tensor_descriptorr  ~  s5   € Ø×9Ñ9¸(¿/¹/È8Ï>É>ÐZaÓKbÓc€FÜ×2Ñ2°6¸8ÓDÐDr   c                ó¤   — t        | «      dk7  ry | d   dk\  sJ d| d   › �«       ‚d|j                  z  dz  }| d   |k\  sJ |› d|› d| d   › �«       ‚y )	Nr	   r   é   zAtensor descriptor block shape must have at least 8 rows, but got r2  r   z2 tensor descriptor block shape must have at least ú columns, but got )rB  r‰  )r  r1   Úmin_colss      r   Úvalidate_descriptor_blockr  ƒ  s�   € Ü
ˆ5ƒz�Q‚Øð �‰8�qŠ=ÐhÐ]Ð^cÐdeÑ^fÐ]gÐhÓhˆ=Ø�U×-Ñ-Ñ-°Ñ1€HØØ	ñØòð {Ø ˜'Ð!SÐT\ÐS]Ð]oÐpuÐvwÑpxÐoyÐzó{ñ r   c                ó¨  — t        | t        j                  «      sJ ‚t        | j                  | j
                  «       t        | j                  «      }t        |«      |k(  sJ d|› dt        |«      › �«       ‚t        ||d¬«      }|j                  | j                  |t        |«      t        |«      «      }t        j                  || j                  «      S ©Nz	expected z offsets, but got F©Úrequire_i64)r]   r#   r  r  rì  r1   rB  Ú_convert_to_ir_valuesÚcreate_descriptor_loadr�   rµ  rÄ  r$   rŸ   )ÚdescÚoffsetsr³  rÂ  r)   Úndimrk   s          r   Údescriptor_loadr  Ž  s²   € ä�dœB×CÑCÔDÐDÐDÜ˜d×.Ñ.°·
±
Ô;Üˆt×ÑÓ €DÜˆw‹<˜4ÒÐS 9¨T¨FÐ2DÄSÈÃ\ÀNÐ!SÓSÐä# G¨WÀ%ÔH€GØ×&Ñ& t§{¡{°GÔ=XÐYgÓ=hÜ'>¸Ó'Oó	Q€Aä�9‰9�Q˜Ÿ™Ó(Ð(r   c                óÐ  — t        | t        j                  «      sJ ‚t        | j                  | j
                  «       t        | j                  «      }t        |«      |k(  sJ d|› dt        |«      › �«       ‚|j                  | j                  k(  sJ ‚t        ||d¬«      }t        j                  |j                  | j                  |j                  |«      t        j                  «      S r  )r]   r#   r  r  rì  r1   rB  r  r  r$   Úcreate_descriptor_storer�   Úvoid)r  r@   r  r)   r  s        r   Údescriptor_storer  ›  s»   € ä�dœB×CÑCÔDÐDÐDÜ˜d×.Ñ.°·
±
Ô;Üˆt×ÑÓ €DÜˆw‹<˜4ÒÐS 9¨T¨FÐ2DÄSÈÃ\ÀNÐ!SÓSÐØ�;‰;˜$×*Ñ*Ò*Ð*Ð*ä# G¨WÀ%ÔH€GÜ�9‰9�W×4Ñ4°T·[±[À%Ç,Á,ÐPWÓXÔZ\×ZaÑZaÓbÐbr   c                ó‚  — t        | t        j                  «      sJ ‚|dk(  sJ d«       ‚|dk(  sJ d«       ‚t        | j                  «      dk(  sJ d| j                  › �«       ‚| j                  d   dk(  sJ d| j                  › �«       ‚t        |j
                  «      dk(  sJ d	|j
                  › �«       ‚|j
                  d   d
k\  sJ d|j
                  › �«       ‚| j                  }d|j                  z  d
z  }| j                  d   |k\  sJ d|› d|› d| j                  d   › �«       ‚t        j                  | j                  |j
                  d   | j                  d   g«      }t        ||fd¬«      d   }|j                  | j                  |j                  ||j                  |«      «      }	t        j                  |	|«      S )NÚ z#cache modifier is not supported yetz$eviction policy is not supported yetr	   údescriptor must be 2D, but got r   r   ú*descriptor block must have 1 row, but got úx offsets must be 1D, but got r  z5descriptor gather must have at least 8 rows, but got r2  zdescriptor gather of ú must have at least r  Fr  )r]   r#   r  rB  rì  r  r1   r‰  rŸ   r  Úcreate_descriptor_gatherr�   r¡   r$   )
r  Ú	x_offsetsÚy_offsetr³  rÂ  r)   r1   r  rj   rk   s
             r   Údescriptor_gatherr(  §  sæ  € ä�dœB×CÑCÔDÐDÐDØ˜RÒÐFÐ!FÓFÐØ˜bÒ ÐHÐ"HÓHÐ ô ˆt×ÑÓ  AÒ%Ð[Ð)HÈ×IYÑIYÐHZÐ'[Ó[Ð%Ø×Ñ˜AÑ !Ò#ÐdÐ'QÐRV×RbÑRbÐQcÐ%dÓdÐ#ô ˆy�‰Ó 1Ò$ÐXÐ(FÀyÇÁÐFWÐ&XÓXÐ$ð �?‰?˜1Ñ Ò"ÐmÐ&[Ð\e×\kÑ\kÐ[lÐ$mÓmÐ"Ø�J‰J€EØ�U×-Ñ-Ñ-°Ñ1€HØ×ÑØ	ñØòð }Ø/°¨wÐ6JÈ8È*ÐTfÐgk×gwÑgwÐxyÑgzÐf{Ð|ó}ð ô �=‰=˜Ÿ™ i§o¡o°aÑ&8¸$×:JÑ:JÈ1Ñ:MÐ%NÓO€DÜ$ W¨x¨lÈÔNÈqÑQ€HØ×(Ñ(¨¯©°i×6FÑ6FÈÐRV×R\ÑR\Ð]dÓReÓf€AÜ�9‰9�Q˜ÓÐr   c                óì  — t        | t        j                  «      sJ ‚t        | j                  «      dk(  sJ d| j                  › �«       ‚| j                  d   dk(  sJ d| j                  › �«       ‚t        |j
                  «      dk(  sJ d|j                  › �«       ‚|j
                  d   dk\  sJ d|j
                  › �«       ‚| j                  }d	|j                  z  dz  }| j                  d   |k\  sJ d
|› d|› d| j                  d   › �«       ‚t        ||fd¬«      d   }|j                  | j                  |j                  |j                  |«       t        j                  d t        j                  «      S )Nr	   r!  r   r   r"  r#  r  z6descriptor scatter must have at least 8 rows, but got r2  zdescriptor scatter of r$  r  Fr  )r]   r#   r  rB  rì  r  Úshapaer1   r‰  r  Úcreate_descriptor_scatterr�   r$   r  )r  r@   r&  r'  r)   r1   r  s          r   Údescriptor_scatterr,  Á  s�  € Ü�dœB×CÑCÔDÐDÐDô ˆt×ÑÓ  AÒ%Ð[Ð)HÈ×IYÑIYÐHZÐ'[Ó[Ð%Ø×Ñ˜AÑ !Ò#ÐdÐ'QÐRV×RbÑRbÐQcÐ%dÓdÐ#ô ˆy�‰Ó 1Ò$ÐYÐ(FÀy×GWÑGWÐFXÐ&YÓYÐ$ð �?‰?˜1Ñ Ò"ÐnÐ&\Ð]f×]lÑ]lÐ\mÐ$nÓnÐ"Ø�J‰J€EØ�U×-Ñ-Ñ-°Ñ1€HØ×ÑØ	ñØòð ~Ø0°°Ð7KÈHÈ:ÐUgÐhl×hxÑhxÐyzÑh{Ðg|Ð}ó~ð ô % W¨x¨lÈÔNÈqÑQ€HØ×%Ñ% d§k¡k°5·<±<À×AQÑAQÐS[Ô\Ü�9‰9�Tœ2Ÿ7™7Ó#Ð#r   c                óÒ  — |r"|d   j                   t        j                  k(  sJ ‚t        j                  |
j	                  | j
                  |j
                  |D �cg c]  }|j
                  ‘Œ c}|D �cg c]  }|j
                  ‘Œ c}|D �cg c]  }|j
                  ‘Œ c}|D �cg c]  }|j
                  ‘Œ c}||||	«
      t        j                  «      S c c}w c c}w c c}w c c}w )Nr   )r1   r#   rc   r$   Úcreate_tensormap_creater�   r  )r	  Úglobal_addressÚbox_dimÚ
global_dimÚglobal_strideÚelement_strideÚ	elem_typeÚinterleave_layoutÚswizzle_modeÚ	fill_moder)   rk   s               r   Útensormap_creater8  ×  s½   € ñ  ¨aÑ 0× 6Ñ 6¼"¿(¹(Ò BÐBÐBÜ�9‰9Ø×'Ñ'Ø�O‰OØ×!Ñ!Ø&Ö'˜!ˆQ�X‹XÒ'Ø)Ö*˜!ˆQ�X‹XÒ*Ø,Ö-˜!ˆQ�X‹XÒ-Ø-Ö.˜!ˆQ�X‹XÒ.ØØØØó	
ô 	�‰óð ùò (ùÚ*ùÚ-ùÚ.s   ÁCÁ2CÂCÂ$C$c                ó|   — t        j                  |j                  | j                  «      t         j                  «      S rq   )r#   r$   Ú#create_tensormap_fenceproxy_acquirer�   r  )r	  r)   s     r   Útensormap_fenceproxy_acquirer;  ö  s)   € Ü�9‰9�W×@Ñ@ÀÇÁÓQÔSU×SZÑSZÓ[Ð[r   c           	     ó¤  — |�t        d«      ‚| j                  j                  j                  «       }|j                  j	                  «       st        |||«      }|j                  j	                  «       sJ d«       ‚||j                  j                  «       k(  s&J d|› d|j                  j                  «       › d�«       ‚| j                  j                  j                  |j                  j                  k(  s@J d| j                  j                  j                  › d|j                  j                  › d�«       ‚| j                  j                  j                  }|t        j                  k7  sJ d«       ‚t        ||«      }t        |||«      }t        j                  |j                  | j                  |j                  |||«      t        j                  «      S )	Nrñ  z-Value argument must be block type or a scalarzBlock shape(z) and value shape(z
) mismatchzBlock element type(z) and value element type(rò  )r"   rj   ró  r    rž   rs  r#   r`   rï  r}   r$   Úcreate_tensor_pointer_storer�   r  )	rõ  Úvalrö  rë  r´  rÃ  r)   rì  rø  s	            r   Ú_store_block_pointerr?  ú  sÆ  € ð ÐÜÐfÓgÐgð —(‘(×%Ñ%×6Ñ6Ó8€KØ�8‰8×ÑÔÜ" 3¨°WÓ=ˆØ�8‰8×ÑÔÐOÐ OÓOÐØ˜#Ÿ(™(×3Ñ3ó ò ð ]Ø	�k�]Ð"4°S·X±X×5NÑ5NÓ5PÐ4QÐQ[Ð\ó]ð à�8‰8×Ñ×)Ñ)¨S¯X©X×-@Ñ-@Ò@ð  qÐDWÐX[×X`ÑX`×XkÑXk×XvÑXvÐWwð  xQð  RU÷  RZñ  RZ÷  Reñ  Reð  Qfð  fpð  Cqó  qÐ@à�X‰X× Ñ ×+Ñ+€FØ”R—W‘WÒÐTÐTÓTÐô 2°.À+ÓN€Nô ˆs�F˜GÓ
$€Cô �9‰9�W×8Ñ8¸¿¹ÀSÇZÁZÐQ_ÐafÐhpÓqÜ—W‘Wóð r   c           	     óè  — | j                   j                  j                  «       s't        d| j                   j	                  «       › d�«      ‚|rt        d«      ‚| j                   j                  «       sL|j                   j                  «       rt        d«      ‚|r%|j                   j                  «       rt        d«      ‚| j                   j                  «       rLt        || j                   j                  «       |«      }|�%t        || j                   j                  «       |«      }| j                   j                  }|j                  }|t        j                  k(  r=t        j                  }t        j                  ||j                  «      }t        | ||«      } t        |||«      }|€Jt        j                  |j!                  | j"                  |j"                  ||«      t        j$                  «      S |j                   j                  j'                  «       st        d«      ‚t        j                  |j)                  | j"                  |j"                  |j"                  ||«      t        j$                  «      S )Nrû  z in `tl.store`zÑ`boundary_check` argument is not supported for storing a tensor of pointers or storing a scalar. Because the compiler does not know the boundary; please use block pointers (defined by `make_block_ptr`) insteadzFValue argument cannot be block type if pointer argument is not a blockrü  z"Mask must have boolean scalar type)rj   ry   rr   r"   r   rž   rs  r    ró  r#   r`   rý  rþ  rÿ  r}   r$   Úcreate_storer�   r  r›  Úcreate_masked_store)	rõ  r>  rö  rë  r´  rÃ  r)   r  rø  s	            r   Ú_store_legacyrC    sã  € à�8‰8�?‰?×!Ñ!Ô#ÜÐ0°·±×1BÑ1BÓ1DÐ0EÀ^ÐTÓUÐUñ Üð Aó Bð 	Bð
 �8‰8×ÑÔØ�8‰8×ÑÔÜÐeÓfÐfÙ�D—I‘I×&Ñ&Ô(ÜÐdÓeÐeð ‡x�x×ÑÔÜ" 3¨¯©×(AÑ(AÓ(CÀWÓMˆØÐÜ'¨¨c¯h©h×.GÑ.GÓ.IÈ7ÓSˆDà�X‰X�_‰_€FØ×Ñ€Fð ”—‘ÒÜ—‘ˆÜ—‘ ¨×)=Ñ)=Ó>ˆÜ�3˜ Ó(ˆô ˆs�F˜GÓ
$€Cð €|Ü�y‰y˜×-Ñ-¨c¯j©j¸#¿*¹*ÀeÈXÓVÔXZ×X_ÑX_Ó`Ð`Ø�9‰9×Ñ×#Ñ#Ô%ÜÐ=Ó>Ð>Ü�9‰9�W×0Ñ0°·±¸S¿Z¹ZÈÏÉÐV[Ð]eÓfÔhj×hoÑhoÓpÐpr   c           	     ó€  — t        |«      }t        |«      }| j                  j                  «       s$| j                  j                  j                  «       rt        d«      ‚| j                  j                  «       r5| j                  j                  j                  «       rt        | ||||||«      S t        | ||||||«      S )Nú"Cannot store to a constant pointer)rº  rÄ  rj   Úis_constry   r"   rr   ró  rž   r?  rC  )	rõ  r>  rö  rë  r³  rÂ  r)   r´  rÃ  s	            r   ÚstorerG  C  s›   € ô )¨Ó8€EÜ& Ó7€Hà
‡x�x×ÑÔ˜cŸh™hŸo™o×6Ñ6Ô8ÜÐ=Ó>Ð>à
‡x�x‡�Ô˜SŸX™X×0Ñ0×9Ñ9Ô;ä# C¨¨d°NÀEÈ8ÐU\Ó]Ð]ô ˜S # t¨^¸UÀHÈgÓVÐVr   c           	     óB  — t        |«      }t        |«      }| j                  j                  j                  }|j
                  dvrt        d«      ‚t        j                  |j                  | j                  |j                  |j                  ||«      |j                  «      S )N)é   r2  r�   z9atomic_cas only supports elements with width {16, 32, 64})rÛ  ræ  rj   ry   ró  r‰  r"   r#   r$   Úcreate_atomic_casr�   )rõ  Úcmpr>  rÚ  rå  r)   ró  s          r   Ú
atomic_casrL  Y  s|   € Ü
�cÓ
€CÜ˜%Ó €EØ—‘—‘×+Ñ+€JØ×$Ñ$¨LÑ8ÜÐTÓUÐUÜ�9‰9�W×.Ñ.¨s¯z©z¸3¿:¹:ÀsÇzÁzÐSVÐX]Ó^Ð`c×`hÑ`hÓiÐir   c                ó  — | j                   j                  j                  «       s&t        d| j                   j	                  «       z   «      ‚| j                   j                  «       s$| j                   j                  j                  «       rt        d«      ‚| j                   j                  j                  }|t        j                  u r|dk7  rt        d|z   dz   «      ‚|t        j                  t        j                  t        j                  t        j                  fv rt        d|z   dz   t        |«      z   «      ‚| j                   j                  «       rN|�%t        || j                   j!                  «       |«      }|�%t        || j                   j!                  «       |«      }t#        || j                   j                  j                  |«      }|€·|j%                  d«      }t        j                  }| j                   j                  «       rf|j'                  || j                   j!                  «       «      }t        j(                  t        j                  | j                   j!                  «       «      }t        j*                  ||«      }| ||fS )Nz)Pointer argument of store instruction is rE  r¥   Úatomic_z does not support fp16z does not support T)rj   ry   rr   r"   r   rF  ró  r#   rA   r`   rý  Úint16rB   rw  rž   rs  r    r}   r_   rC  rŸ   r$   )rõ  r>  rö  Úopr)   ró  Úmask_irÚmask_tys           r   Úatom_red_typechecking_implrS  b  sÕ  € à�8‰8�?‰?×!Ñ!Ô#ÜÐDÀsÇxÁx×GXÑGXÓGZÑZÓ[Ð[Ø
‡x�x×ÑÔ˜cŸh™h×1Ñ1×:Ñ:Ô<ÜÐ=Ó>Ð>Ø—‘—‘×+Ñ+€JØ”R—Z‘ZÑ B¨%¢KÜ˜ R™Ð*BÑBÓCÐCØ”b—g‘gœrŸw™w¬¯©´"·+±+Ð>Ñ>Ü˜ R™Ð*>Ñ>ÄÀZÃÑPÓQÐQØ
‡x�x×ÑÔØÐÜ'¨¨c¯h©h×.GÑ.GÓ.IÈ7ÓSˆDØˆ?Ü& s¨C¯H©H×,EÑ,EÓ,GÈÓQˆCÜ
ˆs�C—H‘H—O‘O×.Ñ.°Ó
8€CØ€|Ø×"Ñ" 4Ó(ˆÜ—'‘'ˆØ�8‰8×ÑÔØ×*Ñ*¨7°C·H±H×4MÑ4MÓ4OÓPˆGÜ—m‘m¤B§G¡G¨S¯X©X×-FÑ-FÓ-HÓIˆGÜ�y‰y˜ 'Ó*ˆØ��Tˆ>Ðr   c                ó¦  — t        | ||d|«      \  } }}t        |«      }t        |«      }|j                  j                  }|j                  «       rä|j                  «       rjt        j                  |j                  t        j                  j                  | j                  |j                  |j                  ||«      |j                  «      S t        j                  |j                  t        j                  j                  | j                  |j                  |j                  ||«      |j                  «      S |t        j                  t        j                   hvrt#        d|› �«      ‚t%        g d||«      }|t        j                  k(  rt        j&                  nt        j(                  }t+        |||«      }	t+        | t        j,                  |d«      |«      }
|t        j                  k(  rt        j.                  nt        j0                  }t+        |||«      }t+        | t        j,                  |d«      |«      }t3        |||«      }t5        |||«      }t        j                  |j                  t        j                  j                  |
j                  |	j                  t7        |||«      j                  ||«      |	j                  «      }t        j                  |j                  t        j                  j8                  |j                  |j                  t7        |||«      j                  ||«      |j                  «      }t;        ||||«      }t+        |||«      S )Nrã   z#atomic_max not supported for dtype r\   r   )rS  rÛ  ræ  rj   ry   rJ   r¾   r#   r$   Úcreate_atomic_rmwr
   Ú	ATOMIC_OPÚMAXr�   ÚUMAXrC   rE   r4   re   r&   rc   rö   rþ  rb   rd   r“   r#  r‘   ÚUMINÚwhere©rõ  r>  rö  rÚ  rå  r)   Úsca_tyrÆ  Úi_typeÚi_valÚi_ptrÚui_typeÚui_valÚui_ptrÚposÚnegÚpos_retÚneg_retr–   s                      r   Ú
atomic_maxrg  }  ó•  € Ü/°°S¸$ÀÀwÓO�N€CˆˆdÜ
�cÓ
€CÜ˜%Ó €EØ�X‰X�_‰_€Fà‡}�}„Ø×ÑÔ!Ü—9‘9Ø×)Ñ)¬"¯,©,×*:Ñ*:¸C¿J¹JÈÏ
É
ÐTX×T_ÑT_ÐadÐfkÓlÐnq×nvÑnvóxð xô —9‘9Ø×)Ñ)¬"¯,©,×*;Ñ*;¸S¿Z¹ZÈÏÉÐUY×U`ÑU`ÐbeÐglÓmÐor×owÑowóyð yð
 ”b—j‘j¤"§*¡*Ð-Ñ-ÜÐ=¸f¸XÐFÓGÐGä��C˜ Ó)€Dà¤2§:¡:Ò-ŒR�XŠX´2·8±8€FÜ�C˜ Ó)€EÜ�CœŸ™¨°Ó3°WÓ=€EØ!¤R§Z¡ZÒ/Œb�iŠi´R·Y±Y€GÜ�S˜' 7Ó+€FÜ�Sœ"Ÿ/™/¨'°1Ó5°wÓ?€FÜ
˜˜T 7Ó
+€CÜ
�C˜˜wÓ
'€CÜ�i‰iØ×!Ñ!¤"§,¡,×"2Ñ"2°E·L±LÀ%Ç,Á,Ü"& t¨S°'Ó":×"AÑ"AÀ3Èó	OØPU×PZÑPZó\€Gô �i‰iØ×!Ñ!¤"§,¡,×"3Ñ"3°V·]±]ÀFÇMÁMÜ"& t¨S°'Ó":×"AÑ"AÀ3Èó	OØPV×P[ÑP[ó]€Gô ��W˜g wÓ
/€CÜ�3˜ Ó(Ð(r   c                ó¦  — t        | ||d|«      \  } }}t        |«      }t        |«      }|j                  j                  }|j                  «       rä|j                  «       rjt        j                  |j                  t        j                  j                  | j                  |j                  |j                  ||«      |j                  «      S t        j                  |j                  t        j                  j                  | j                  |j                  |j                  ||«      |j                  «      S |t        j                  t        j                   hvrt#        d|› �«      ‚t%        g d||«      }|t        j                  k(  rt        j&                  nt        j(                  }t+        |||«      }	t+        | t        j,                  |d«      |«      }
|t        j                  k(  rt        j.                  nt        j0                  }t+        |||«      }t+        | t        j,                  |d«      |«      }t3        |||«      }t5        |||«      }t        j                  |j                  t        j                  j                  |
j                  |	j                  t7        |||«      j                  ||«      |	j                  «      }t        j                  |j                  t        j                  j8                  |j                  |j                  t7        |||«      j                  ||«      |j                  «      }t;        ||||«      }t+        |||«      S )Nrâ   z#atomic_min not supported for dtype r\   r   )rS  rÛ  ræ  rj   ry   rJ   r¾   r#   r$   rU  r
   rV  ÚMINr�   rY  rC   rE   r4   re   r&   rc   rö   rþ  rb   rd   r“   r#  r‘   rX  rZ  r[  s                      r   Ú
atomic_minrk  ¤  rh  r   c           
     ó¬  — t        | ||d|«      \  } }}t        |«      }t        |«      }|j                  j                  }|j                  «       rt        j                  j                  nt        j                  j                  }t        j                  |j                  || j                  |j                  |j                  ||«      |j                  «      S )Nr¥   )rS  rÛ  ræ  rj   ry   rs   r
   rV  ÚFADDÚADDr#   r$   rU  r�   )rõ  r>  rö  rÚ  rå  r)   r\  rP  s           r   Ú
atomic_addro  Ë  sœ   € Ü/°°S¸$ÀÀwÓO�N€CˆˆdÜ
�cÓ
€CÜ˜%Ó €EØ�X‰X�_‰_€FØ$×0Ñ0Ô2Œ�‰×	Ò	¼¿¹×8HÑ8H€BÜ�9‰9�W×.Ñ.¨r°3·:±:¸s¿z¹zÈ4Ï;É;ÐX[Ð]bÓcÐeh×emÑemÓnÐnr   c           
     ó(  — t        | ||d|«      \  } }}t        |«      }t        |«      }t        j                  |j                  t        j                  j                  | j                  |j                  |j                  ||«      |j                  «      S )NÚand)rS  rÛ  ræ  r#   r$   rU  r
   rV  ÚANDr�   rj   ©rõ  r>  rö  rÚ  rå  r)   s         r   Ú
atomic_andrt  Ô  óx   € Ü/°°S¸$ÀÀwÓO�N€CˆˆdÜ
�cÓ
€CÜ˜%Ó €EÜ�9‰9�W×.Ñ.¬r¯|©|×/?Ñ/?ÀÇÁÈSÏZÉZÐY]×YdÑYdÐfiÐkpÓqØ—X‘Xóð r   c           
     ó(  — t        | ||d|«      \  } }}t        |«      }t        |«      }t        j                  |j                  t        j                  j                  | j                  |j                  |j                  ||«      |j                  «      S )NÚor)rS  rÛ  ræ  r#   r$   rU  r
   rV  ÚORr�   rj   rs  s         r   Ú	atomic_orry  Ü  sv   € Ü/°°S¸$ÀÀgÓN�N€CˆˆdÜ
�cÓ
€CÜ˜%Ó €EÜ�9‰9�W×.Ñ.¬r¯|©|¯©ÀÇ
Á
ÈCÏJÉJÐX\×XcÑXcÐehÐjoÓpØ—X‘Xóð r   c           
     ó(  — t        | ||d|«      \  } }}t        |«      }t        |«      }t        j                  |j                  t        j                  j                  | j                  |j                  |j                  ||«      |j                  «      S )NÚxor)rS  rÛ  ræ  r#   r$   rU  r
   rV  ÚXORr�   rj   rs  s         r   Ú
atomic_xorr}  ä  ru  r   c           
     ó(  — t        | ||d|«      \  } }}t        |«      }t        |«      }t        j                  |j                  t        j                  j                  | j                  |j                  |j                  ||«      |j                  «      S )NÚxchg)rS  rÛ  ræ  r#   r$   rU  r
   rV  ÚXCHGr�   rj   rs  s         r   Úatomic_xchgr�  ì  sx   € ä/°°S¸$ÀÈÓP�N€CˆˆdÜ
�cÓ
€CÜ˜%Ó €EÜ�9‰9�W×.Ñ.¬r¯|©|×/@Ñ/@À#Ç*Á*ÈcÏjÉjÐZ^×ZeÑZeÐgjÐlqÓrØ—X‘Xóð r   c                óò   — | j                  «       |j                  j                  v s!J d|j                  j                  › d| › �«       ‚| j                  «       } | dk(  rd} t	        t
        j                  | «      S )Nzinput_precision must be one of ú. Got ÚTF32X3ÚTF32x3)ÚlowerrŽ   Úallowed_dot_input_precisionsÚupperr=  r
   ÚINPUT_PRECISION)Úinput_precisionr)   s     r   Ú_str_to_dot_input_precisionr‹  ú  sx   € Ø× Ñ Ó" g§o¡o×&RÑ&RÑRð pØ
)¨'¯/©/×*VÑ*VÐ)WÐW]Ð^mÐ]nÐoópÐRà%×+Ñ+Ó-€OØ˜(Ò"Ø"ˆÜ”2×%Ñ% Ó7Ð7r   c           
     ó  — | j                   j                  «       r|j                   j                  «       sJ ‚| j                  j                  «       r|j                  j                  «       r�n| j                  t        j
                  t        j                  t        j                  t        j                  t        j                  fv sJ d| j                  › �«       ‚|j                  t        j
                  t        j                  t        j                  t        j                  t        j                  fv sJ d|j                  › �«       ‚| j                  |j                  k(  s!J d| j                  › d|j                  › �«       ‚| j                  j                  «       s|j                  j                  «       r6t        | t        j                  |«      } t        |t        j                  |«      }|€|j                  j                  }t        ||«      }t        | j                   «      }t        |j                   «      }||cxk(  rdk(  s1n ||cxk(  rdk(  s$n J d| j                   › d|j                   › d	�«       ‚| j                   d
   j"                  |j                   d   j"                  k(  sVJ d| j                   › d|j                   › d| j                   d
   j"                  › d|j                   d   j"                  › d	�	«       ‚|j$                  j'                  d«      €J d«       ‚ |j$                  d   | j                   |j                   «      }	| j                   d   j"                  |	d   k\  r>| j                   d
   j"                  |	d   k\  r|j                   d
   j"                  |	d   k\  sJ d|	d   › d|	d   › d|	d   › �«       ‚| j                   j(                  j+                  «       rP| j                   j(                  t        j
                  k(  sJ d«       ‚|j-                  d«      }
t        j.                  }n¹|j1                  «       rt3        d«      ‚| j                   j(                  j5                  «       s$| j                   j(                  j1                  «       r"|j7                  d«      }
t        j                  }n4|j9                  «       r|j;                  d«      n|j7                  d«      }
|}| j                   j                   d   }|j                   j                   d
   }| j                   j                   d
   }|dk(  r| j                   j                   d   nd }t	        j<                  ||r|||gn||g«      }|€|j?                  |
|r|||gn||g«      }n|j@                  }|j                   |k(  sJ ‚|€N| j                  j                  «       r1|j                  j                  «       r|j                  jB                  }nNd}nK| j                  j                  «       r1|j                  j                  «       r||kD  rt3        d|› d|› d	�«      ‚t	        jD                  |jG                  | j@                  |j@                  |||«      |«      S )NzUnsupported lhs dtype zUnsupported rhs dtype z&Both operands must be same dtype. Got r   r	   é   ú+Both inputs must be either 2D or 3D; (lhs: ú	 vs rhs: r
  rT  éþÿÿÿzFirst input shape (z) and second input shape z= are not compatible for matmul (second index of first shape (z0) must be equal to first index of second shape (Úmin_dot_sizez2target doesn't provide lower shape bounds for dot.r   r   zInput shapes should have M >= z, N >= z
 and K >= zonly int8 supported!zhout_dtype=bfloat16 is unsupported. Please use out_dtype=float32/float16 and cast with `.to(tl.bfloat16)`zmax_num_imprecise_acc (z) must be <= K ()$rj   rž   r1   rI   r#   rý  Úuint8rA   rB   rC   r•  r}   rŽ   Údefault_dot_input_precisionr‹  rB  r  r@   r–  r—  ry   rJ   Ú	get_int32r&   rH   r"   rF   Úget_fp32rG   Úget_fp16rŸ   rC  r�   Úmax_num_imprecise_acc_defaultr$   Ú
create_dot)r   r€   ÚaccrŠ  Úmax_num_imprecise_accÚ	out_dtyper)   Úlhs_rankÚrhs_rankr‘  r  Úret_scalar_tyÚMÚNÚKÚBrÁ   Ú
acc_handles                     r   Údotr¤    sî  € à�8‰8×ÑÔ 3§8¡8×#4Ñ#4Ô#6Ð6Ð6à
‡y�y×ÑÔ˜cŸi™i×.Ñ.Ô0áà�y‰yœRŸW™W¤b§h¡h´·
±
¼B¿K¹KÜŸZ™Zð)ñ )ð 	OØ,BÀ3Ç9Á9À+Ð*Nó	Oð )à�y‰yœRŸW™W¤b§h¡h´·
±
¼B¿K¹KÜŸZ™Zð)ñ )ð 	OØ,BÀ3Ç9Á9À+Ð*Nó	Oð )à�y‰y˜CŸI™IÒ%ÐkÐ)OÐPS×PYÑPYÈ{ÐZ_Ð`c×`iÑ`iÐ_jÐ'kÓkÐ%à
‡y�y×ÑÔ #§)¡)×"7Ñ"7Ô"9ä�3œŸ
™
 GÓ,ˆÜ�3œŸ
™
 GÓ,ˆàÐØ!Ÿ/™/×EÑEˆä1°/À7ÓK€Oä�3—9‘9‹~€HÜ�3—9‘9‹~€HØ�xÔ$ 1Ô$¨°HÔ(AÀÔ(Að  RÐEpÐqt×qzÑqzÐp{ð  |Eð  FI÷  FOñ  FOð  EPð  PQð  DRó  RÐAØ�9‰9�R‰=×Ñ #§)¡)Ø
ñ#ß‰Eòð qØ(¨¯©¨Ð3LÈSÏYÉYÈKð  XUð  VY÷  V_ñ  V_ð  `bñ  Vc÷  Viñ  Við  Ujð  jZð  [^÷  [dñ  [dð  egñ  [h÷  [nñ  [nð  Zoð  opð  qóqð à×Ñ×"Ñ" >Ó2Ð>ÐtÐ@tÓtÐ>Ø6�7×&Ñ& ~Ñ6°s·x±xÀÇÁÓJ€LØ�9‰9�R‰=×Ñ ,¨q¡/Ò1°c·i±iÀ±m×6IÑ6IÈ\ÐZ[É_Ò6\Ø�I‰I�b‰M×Ñ <°¡?Ò2ðrà,¨\¸!©_Ð,=¸WÀ\ÐRSÁ_ÐDUÐU_Ð`lÐmnÑ`oÐ_pÐqórð 3ð ‡x�x‡�×ÑÔØ�x‰x�‰¤"§'¡'Ò)ÐAÐ+AÓAÐ)Ø×Ñ˜qÓ!ˆÜŸ™‰Ø	×	Ñ	Ô	ÜØvóxð 	xà	�‰�‰×	 Ñ	 Ô	" c§h¡h§o¡o×&=Ñ&=Ô&?Ø×Ñ˜aÓ ˆÜŸ
™
‰à$-×$5Ñ$5Ô$7ˆW×Ñ˜aÔ ¸W×=MÑ=MÈaÓ=PˆØ!ˆà�‰�‰�rÑ€AØ�‰�‰�rÑ€AØ�‰�‰�rÑ€AØ%¨š]ˆ�‰�‰�qÒ°€AÜ�]‰]˜=±q¨1¨a°©)¸qÀ!¸fÓE€FØ
€{Ø×)Ñ)¨"¹1¨q°!°Q©iÀ1ÀaÀ&ÓI‰
à—Z‘Zˆ
Ø�x‰x˜6Ò!Ð!Ð!ð Ð$Ø�9‰9×ÑÔ #§)¡)×"2Ñ"2Ô"4Ø$+§O¡O×$QÑ$QÑ!à$%Ñ!à�9‰9×ÑÔ #§)¡)×"2Ñ"2Ô"4Ð9NÐQRÒ9RÜÐ6Ð7LÐ6MÐM]Ð^_Ð]`Ð`aÐbÓcÐcä�9‰9�W×'Ñ'¨¯
©
°C·J±JÀ
ÈOÐ]rÓsØóð r   c                óz   — t        t        j                  | j                  «       d «      }|€t	        d| › d�«      ‚|S )NzInvalid float format: rV   )r=  r
   ÚScaleDotElemTypeTYrˆ  r"   )Úfloat_formatÚty_enums     r   Ú_str_to_fp_typer©  L  s>   € Ü”b×+Ñ+¨\×-?Ñ-?Ó-AÀ4ÓH€GØ€ÜÐ1°,°¸qÐAÓBÐBØ€Nr   c                ó2  — t         j                  t         j                  t         j                  t         j                  dœj                  |«      }|€B|dk(  s
J d|› �«       ‚| j                  t         j                  k(  sJ d| j                  › �«       ‚| S | j                  |k(  r| S t         j                  t         j                  t         j                  t         j                  dœ|   }| j                  |k(  sJ d|› d| j                  › �«       ‚t        | ||«      S )z¦
    If float_format is subbyte, make sure it's packed as uint8 and return it.
    Otherwise, return a tensor (perhaps bitcasting) of the specified float format.
    )Úe5m2Úe4m3Úbf16Úfp16Úe2m1z)Internal Error: Unexpected float format: z)e2m1 format must be packed as uint8. Got zUnexpected dtype for rƒ  )
r#   Úfloat8e5Ú
float8e4nvrB   rA   r—  r1   r’  Úuint16rö   )r>  r§  r)   Ú	triton_tyÚunsigned_tys        r   Ú_bitcast_to_fp_typerµ  S  sø   € ô
 Ÿ™¬b¯m©mÄRÇ[Á[ÔZ\×ZdÑZdÑe×iÑiÐjvÓw€IØÐØ˜vÒ%ÐaÐ)RÐS_ÐR`Ð'aÓaÐ%Ø�y‰yœBŸH™HÒ$Ð]Ð(QÐRU×R[ÑR[ÐQ\Ð&]Ó]Ð$Øˆ
Ø
‡y�y�IÒØˆ
ä!Ÿx™x´·±Ä2Ç9Á9ÔVX×V_ÑV_Ñ`ÐamÑnˆØ�y‰y˜KÒ'Ð`Ð+@ÀÀÈfÐUX×U^ÑU^ÐT_Ð)`Ó`Ð'Ü�s˜I wÓ/Ð/r   c
                ó¶  — | j                   j                  «       r|j                   j                  «       sJ ‚t        | j                  «      }
t        |j                  «      }|
|cxk(  rdk(  s1n |
|cxk(  rdk(  s$n J d| j                  › d|j                  › d�«       ‚|j                  }|j                  }t        |«      }t        |«      }h d£}||v s
J d|› �«       ‚||v s
J d|› �«       ‚|d u xs* t        |t        j                  «      xr |j                  d u }|d u xs* t        |t        j                  «      xr |j                  d u }t        | ||	«      } t        |||	«      }| j                   j                  d	   }|j                   j                  d	d  \  }}|d
k(  rdnd}|d
k(  rdnd}||z  || j                   j                  d   z  k(  s"J d| j                  › d|j                  › d�«       ‚|
dk(  r| j                   j                  d   nd }t        j                  ||r|||gn||g«      }|	j                  d«      }|€|	j                  ||r|||gn||g«      }n|j                  }|j                   |k(  sJ ‚|rd n|j                  }|rd n|j                  }t        j                  |	j                  | j                  |||j                  ||||«      |«      S )Nr	   r�  rŽ  r�  r
  >   r­  r¯  r¬  r«  r®  zNYI: lhs_format zNYI: rhs_format r�  r¯  r   rT  zCReduction dimension should pack the same number of elements; (lhs: r   )rj   rž   rB  r  r@   r©  r]   r#   rh   rµ  rŸ   r•  rC  r�   r$   Úcreate_dot_scaled)r   Ú	lhs_scaleÚ
lhs_formatr€   Ú	rhs_scaleÚ
rhs_formatr™  Ú	fast_mathr›  r)   rœ  r�  Úlhs_format_enumÚrhs_format_enumÚallowed_formatsÚrhs_scale_is_noneÚlhs_scale_is_nonerŸ  r¡  r   ÚPACKED_AÚPACKED_Br¢  rÁ   r  r£  Úrhs_scale_handleÚlhs_scale_handles                               r   Ú
dot_scaledrÆ  e  s  € ð �8‰8×ÑÔ 3§8¡8×#4Ñ#4Ô#6Ð6Ð6ä�3—9‘9‹~€HÜ�3—9‘9‹~€HØ�xÔ$ 1Ô$¨°HÔ(AÀÔ(Að  RÐEpÐqt×qzÑqzÐp{ð  |Eð  FI÷  FOñ  FOð  EPð  PQð  DRó  RÐAØ ×&Ñ&€JØ ×&Ñ&€JÜ% jÓ1€OÜ% jÓ1€OÚ>€OØ˜Ñ(ÐIÐ,<¸Z¸LÐ*IÓIÐ(Ø˜Ñ(ÐIÐ,<¸Z¸LÐ*IÓIÐ(Ø! TÐ)Òn¬j¸ÄBÇLÁLÓ.QÒ.mÐV_×VeÑVeÐimÐVmÐØ! TÐ)Òn¬j¸ÄBÇLÁLÓ.QÒ.mÐV_×VeÑVeÐimÐVmÐÜ
˜c :¨wÓ
7€CÜ
˜c :¨wÓ
7€Cà�‰�‰�rÑ€AØ�8‰8�>‰>˜"˜#Ð�D€A€qØ &Ò(‰q¨a€HØ &Ò(‰q¨a€HØˆx‰<˜8 c§h¡h§n¡nØ
ñ'ñ ò ð tØRÐSV×S\ÑS\ÐR]Ð]fÐgj×gpÑgpÐfqÐqrÐsótð ð &¨š]ˆ�‰�‰�qÒ°€Aä�]‰]˜9±1 q¨!¨Q¡i¸1¸a¸&ÓA€FØ	×	Ñ	˜!Ó	€BØ
€{Ø×)Ñ)¨"¹1¨q°!°Q©iÀ1ÀaÀ&ÓI‰
à—Z‘Zˆ
Ø�x‰x˜6Ò!Ð!Ð!Ù0‘t°i×6FÑ6FÐÙ0‘t°i×6FÑ6FÐÜ�9‰9Ø×!Ñ! #§*¡*Ð.>ÀÐQT×Q[ÑQ[Ð]mØ"1°9¸jó	JØKQóSð Sr   c                ó  — | j                   t        j                  k7  r"t        j                  d| j                   › �«       t        | t        j                  |«      } t        |||dd«      \  }}| j                  j                  «       r!t        | ||«      \  } }t        |||«      \  }}nt        | ||«      \  } }|j                  }t        j                  |j                  | j                  |j                  |j                  «      |«      S )Nzgtl.where with a non-boolean condition is deprecated and will error out in a future triton release. Got T)r1   r#   r`   ÚwarningsÚwarnr}   r‹   rj   rž   r~   r$   Úcreate_selectr�   )Ú	conditionrk   rÖ   r)   r|  rÁ   s         r   rZ  rZ  •  sç   € Ø‡�œ"Ÿ'™'Ò!Ü�‰ØuÐv÷  wFñ  wFð  vGð  Hô	
ô �Y¤§¡¨Ó1€IÜ'¨¨1¨g°t¸TÓB�D€A€qà‡~�~×ÑÔ Ü+¨I°q¸'ÓB‰ˆ	�1Ü# A q¨'Ó2‰ˆ‰1ä+¨I°q¸'ÓB‰ˆ	�1Ø�V‰V€FÜ�9‰9�W×*Ñ*¨9×+;Ñ+;¸Q¿X¹XÀqÇxÁxÓPÐRXÓYÐYr   c                ód   — |rt        j                  ||«      }n|}t        j                  | |«      S rq   )r#   rŸ   r$   )rk   rN   r}  Úres_tys       r   Úwrap_tensorrÎ  «  s-   € ÙÜ—‘˜y¨)Ó4‰ð ˆÜ�9‰9�Q˜ÓÐr   c                ó  ‡ ‡‡‡	‡
— |€t        ˆfd„‰ D «       «      Š d}‰ d   j                  j                  Š
t        ‰
«      }||k  sJ d|› d�«       ‚t	        ‰
«      D ��cg c]  \  }}||k7  sŒ|‘Œ c}}Š	t        ˆ
fd„‰ D «       «      sJ d«       ‚‰j                  ‰ D �cg c]  }|j                  ‘Œ c}|«      Š |‰«       ‰j                  «        t        ˆ ˆˆ	fd„t        t        ‰ «      «      D «       «      S c c}}w c c}w )Nc              3  óf   •K  — | ](  }t        ||j                  j                  gd ‰¬«      –— Œ* y­w)TrU  N)rI  r<  r@   )rd  Útr)   s     €r   rf  zreduction.<locals>.<genexpr>¶  s*   øè ø€ ÒfÐZ[”w˜q 1§7¡7§=¡= /¸tÈW×UÐUÑfùs   ƒ.1r   z&reduction axis must be < inputs rank (r
  c              3  óP   •K  — | ]  }|j                   j                  ‰k(  –— Œ y ­wrq   )rj   r  )rd  rÑ  r  s     €r   rf  zreduction.<locals>.<genexpr>½  s   øè ø€ Ò5¨ˆq�v‰v�|‰|˜uÕ$Ñ5ùs   ƒ#&z-all reduction inputs must have the same shapec              3  ó„   •K  — | ]7  }t        ‰j                  |«      ‰|   j                  j                  ‰«      –— Œ9 y ­wrq   ©rÎ  Ú
get_resultrj   ry   )rd  rq  ÚinputsÚ	reduce_opr}  s     €€€r   rf  zreduction.<locals>.<genexpr>Ã  s4   øè ø€ ÒtÐ\]”˜Y×1Ñ1°!Ó4°f¸Q±i·n±n×6KÑ6KÈY×WÑtùó   ƒ=A )
Útuplerj   r  rB  rn  ÚallÚcreate_reducer�   Úverifyr8  )rÖ  r(   Úregion_builder_fnr)   Úrankrq  rH  rÑ  r×  r}  r  s   `  `    @@@r   Ú	reductionrß  ´  sì   ü€ Ø€|ÜÓfÐ_eÔfÓfˆØˆà�1‰I�N‰N× Ñ €EÜˆu‹:€DØ�$Š;ÐHÐ@ÀÀÀaÐHÓHˆ;Ü(¨Ó/×=‘t�q˜!°1¸³9’Ó=€IÜÓ5¨fÔ5Ô5ÐfÐ7fÓfÐ5à×%Ñ%¸Ö&@°A q§x£xÒ&@À$ÓG€IÙ�iÔ Ø×ÑÔäÕtÔafÔgjÐkqÓgrÓasÔtÓtÐtùó >ùò 'As   Á C;Á.C;Â Dc                óÊ  ‡ ‡‡— ‰ d   j                   j                  Št        ‰«      }| |cxk  r|k  sn J d|› d|› d�«       ‚|dk  r||z  }‰ D ]"  }|j                   j                  ‰k(  rŒJ d«       ‚ |j                  ‰ D �cg c]  }|j                  ‘Œ c}||«      Š |‰«       ‰j                  «        t        ˆ ˆˆfd„t        t        ‰ «      «      D «       «      S c c}w )Nr   z
scan axis z must be < inputs rank (r
  z(all scan inputs must have the same shapec              3  ó„   •K  — | ]7  }t        ‰j                  |«      ‰|   j                  j                  ‰«      –— Œ9 y ­wrq   rÔ  )rd  rq  rÖ  Úscan_opr  s     €€€r   rf  z#associative_scan.<locals>.<genexpr>Ü  s4   øè ø€ ÒnÐVW”˜W×/Ñ/°Ó2°F¸1±I·N±N×4IÑ4IÈ5×QÑnùrØ  )rj   r  rB  Úcreate_scanr�   rÜ  rÙ  r8  )	rÖ  r(   rÝ  Úreverser)   rÞ  rÑ  râ  r  s	   `      @@r   Úassociative_scanrå  Ë  sá   ú€ à�1‰I�N‰N× Ñ €EÜˆu‹:€Dàˆ5�DÔ˜4ÔÐS :¨d¨VÐ3KÈDÈ6ÐQRÐ!SÓSÐàˆa‚xØ�‰ˆàò QˆØ�v‰v�|‰|˜uÓ$ÐPÐ&PÓPÐ$ðQð ×!Ñ!°VÖ"<° 1§8£8Ò"<¸dÀGÓL€GÙ�gÔØ‡N�NÔäÕnÔ[`ÔadÐekÓalÓ[mÔnÓnÐnùò	 #=s   ÂC c                ó€  — |j                   j                  «       sJ d«       ‚t        | j                  j                  «      }t        |j                  j                  «      |k(  sJ d«       ‚| |cxk  r|k  sn J d|› d|› d�«       ‚|dk  r||z  }t        |«      D ]F  }||k(  rŒ	|j                  j                  |   | j                  j                  |   k(  rŒ=J d|› d�«       ‚ |j                  | j                  |j                  |«      }t        || j                  j                  |j                  j                  «      S )	Nzindex must be an integer tensorz0source and index tensors must have the same rankzgather axis z must be < source rank (r
  r   z
index dim z( must match the corresponding source dim)
r1   rJ   rB  rj   r  r8  Úcreate_gatherr�   rÎ  ry   )ÚsrcÚindexr(   r)   rÞ  re  Úgathers          r   rê  rê  ä  s%  € Ø�;‰;×ÑÔÐBÐ!BÓBÐäˆs�x‰x�~‰~Ó€DÜˆu�z‰z×ÑÓ  DÒ(Ð\Ð*\Ó\Ð(àˆ5�DÔ˜4ÔÐU <°¨vÐ5MÈdÈVÐSTÐ!UÓUÐØˆa‚xØ�‰ˆä�4‹[ò uˆØ�Š9ØØ�z‰z×Ñ Ñ" c§h¡h§n¡n°QÑ&7Ó7Ðt¸:ÀdÀVÐKsÐ9tÓtÐ7ðuð
 ×"Ñ" 3§:¡:¨u¯|©|¸TÓB€FÜ�v˜sŸx™xŸ™°·
±
×0@Ñ0@ÓAÐAr   c                ó(  — t        | j                  «      dk(  sJ d«       ‚| j                  j                  «       sJ d«       ‚t	        j
                  |j                  | j                  |«      t	        j                  t        j                  |g«      «      S )Nr   z histogram only supports 1D inputz%histogram only supports integer input)
rB  r  r1   rJ   r#   r$   Úcreate_histogramr�   rŸ   r&   )r§   Únum_binsr)   s      r   Ú	histogramrî  ü  ss   € Üˆu�{‰{Ó˜qÒ ÐDÐ"DÓDÐ Ø�;‰;×ÑÔÐHÐ!HÓHÐÜ�9‰9�W×-Ñ-¨e¯l©l¸HÓEÄrÇ}Á}ÔUW×U]ÑU]Ð`hÐ_iÓGjÓkÐkr   c                ó  — t        dt        | j                  «      «      t        |«      k7  rt        d«      ‚| j                  j                  dt        j                  || j                  j                  «       «      «       | S )Nr   zAShape of input to multiple_of does not match the length of valuesztt.divisibility)	rã   rB  r  r"   r�   Úset_attrr
   Ú	make_attrÚget_context©rk   ru  s     r   Úmultiple_ofrô    s[   € Ü
ˆ1Œc�!—'‘'‹lÓœs 6›{Ò*ÜÐ\Ó]Ð]Ø‡H�H×ÑÐ'¬¯©°f¸a¿h¹h×>RÑ>RÓ>TÓ)UÔVØ€Hr   c                óî   — t        | j                  «      t        |«      k7  rt        d«      ‚| j                  j	                  dt        j                  || j                  j                  «       «      «       | S )NzDShape of input to max_contiguous does not match the length of valuesztt.contiguity©rB  r  r"   r�   rð  r
   rñ  rò  ró  s     r   Úmax_contiguousr÷  	  sS   € Ü
ˆ1�7‰7ƒ|”s˜6“{Ò"ÜÐ_Ó`Ð`Ø‡H�H×Ñ�o¤r§|¡|°F¸A¿H¹H×<PÑ<PÓ<RÓ'SÔTØ€Hr   c                óî   — t        | j                  «      t        |«      k7  rt        d«      ‚| j                  j	                  dt        j                  || j                  j                  «       «      «       | S )NzCShape of input to max_constancy does not match the length of valuesztt.constancyrö  ró  s     r   Úmax_constancyrù    sS   € Ü
ˆ1�7‰7ƒ|”s˜6“{Ò"ÜÐ^Ó_Ð_Ø‡H�H×Ñ�n¤b§l¡l°6¸1¿8¹8×;OÑ;OÓ;QÓ&RÔSØ€Hr   c                óf   — t        j                  | j                  «       t         j                  «      S rq   )r#   r$   Úcreate_barrierr  )r)   s    r   Údebug_barrierrü    s    € Ü�9‰9�W×+Ñ+Ó-¬r¯w©wÓ7Ð7r   c           	     ó.  — | j                  d«      s|r| dz  } | j                  d«      s
|r| d d dz   } t        | «      dkD  r| j                  d«      sd| z   } |D �cg c]  }|j                  ‘Œ }}|D �cg c][  }|j                  t
        j                  t
        j                  t
        j                  t
        j                  t
        j                  fv ‘Œ] }}t        j                  |j                  | |||«      t
        j                  «      S c c}w c c}w )Nú rm  rT  r	   )ÚendswithrB  Ú
startswithr�   r1   r#   r`   rý  rO  r&   rc   r$   Úcreate_printr  )ÚprefixÚargsÚhexr)   ÚargÚnew_argsÚ	is_signeds          r   Údevice_printr    sÛ   € ð �?‰?˜3Ô¡DØ�#‰ˆØ�?‰?˜4Ô ¡TØ˜˜�˜tÑ#ˆÜ
ˆ6ƒ{�Q‚˜v×0Ñ0°Ô5Ø�v‘ˆà&*Ö+˜s�—
“
Ð+€HÐ+ØZ^Ö_ÐSV�—‘œrŸw™w¬¯©´·±¼2¿8¹8ÄRÇXÁXÐNÒNÐ_€IÐ_Ü�9‰9�W×)Ñ)¨&°#°xÀÓKÌRÏWÉWÓUÐUùò ,ùÚ_s   ÁDÁ5A Dc                ó¬   — |j                   j                  sy t        j                  |j	                  | j
                  |«      t        j                  «      S rq   )rŽ   Údebugr#   r$   Úcreate_assertr�   r  )r™   rš   r)   s      r   r”   r”   *  s8   € Ø�?‰?× Ò ØÜ�9‰9�W×*Ñ*¨4¯;©;¸Ó<¼b¿g¹gÓFÐFr   c                ó|   — t        j                  |j                  | j                  «      t         j                  «      S rq   )r#   r$   Úcreate_assumer�   r  )r™   r)   s     r   Úassumer  0  s&   € Ü�9‰9�W×*Ñ*¨4¯;©;Ó7¼¿¹ÓAÐAr   c                ó¤  — t        |t        «      rt        j                  |«      }t        |t        j                  «      r�|rGd|j                  cxk  rdk  sn J d|j                  › d�«       ‚| j                  |j                  «      S d|j                  cxk  rdk  sn J d|j                  › d�«       ‚| j                  |j                  «      S t        |t        j                  «      rÕ|j                  j                  dk(  sJ d	«       ‚|j                  j                  «       sJ d
«       ‚|j                  t        j                  k7  rE|rC| j                  |j                  | j                  «       |j                  j                  «       «      S |j                  t        j                   k7  r	|sJ d«       ‚|j                  S J dt#        |«      › �«       ‚)NrT   rU   z@Block pointers only support 64 bit `shape/strides`, got a value z which is out of the rangerR   rS   zFBlock pointers only support 32 bit `offsets/block_shape`, got a value r   z*Expected a scalar in shape/strides/offsetsz8Expected an integer scalar type in shape/strides/offsetszzBlock pointers only support 32 bit `offsets/block_shape`, add a `.to(tl.int32)` or use regular indexing for 64 bit supportz3Unsupported element type in shape/strides/offsets: )r]   ra   r#   rh   r@   r�   r”  r$   r<  r1   rJ   rc   r¢   r�   Úget_int64_tyr¾   r&   rj   )r)   rí  r  s      r   Ú_convert_elem_to_ir_valuer  4  s•  € Ü�$œÔÜ�|‰|˜DÓ!ˆÜ�$œŸ™Ô%ÙØ˜TŸZ™ZÔ/¨%Ô/ð Fð 4Ø#Ÿz™z˜lÐ*Dð2Fó FÐ/à×$Ñ$ T§Z¡ZÓ0Ð0à˜TŸZ™ZÔ/¨%Ô/ð Fð 4Ø#Ÿz™z˜lÐ*Dð2Fó FÐ/à×$Ñ$ T§Z¡ZÓ0Ð0Ü	�Dœ"Ÿ)™)Ô	$Ø�z‰z×Ñ 1Ò$ÐRÐ&RÓRÐ$Ø�z‰z× Ñ Ô"Ð^Ð$^Ó^Ð"Ø�:‰:œŸ™Ò!¡kØ×*Ñ*¨4¯;©;¸×8LÑ8LÓ8NÐPT×PZÑPZ×PhÑPhÓPjÓkÐkØ�Z‰Zœ2Ÿ8™8Ò#©KðSð Só S�5à�{‰{ÐØTÐGÌÈTË
À|ÐTÓTˆ5r   c                óv   — t        |d«      r|D �cg c]  }t        | ||«      ‘Œ c}S t        | ||«      gS c c}w )Nrè  )ré  r  )r)   Ú	list_liker  rí  s       r   r  r  L  s?   € Üˆy˜*Ô%ØR[Ö\È$Ô)¨'°4¸ÕEÒ\Ð\Ü% g¨y¸+ÓFÐGÐGùò ]s   ‘6c           	     óŽ  ‡— t        ||«      }t        ||«      }t        ||d¬«      }| j                  j                  «       r$| j                  j                  j	                  «       rt        d«      ‚| j                  j                  t        j                  k(  rCt        | t        j                  t        j                  | j                  j                  «      |«      } t        ‰d«      s‰gŠ‰D �cg c]*  }t        |t        j                  «      r|j                  n|‘Œ, c}Št!        d„ ‰D «       «      sJ d«       ‚t        |d«      s|g}|D �cg c]*  }t        |t        j                  «      r|j                  n|‘Œ, }}t#        |«      t%        t'        t)        |«      «      «      k(  sJ d«       ‚t!        ˆfd„||||fD «       «      sJ d	«       ‚|j+                  | j,                  |||‰|«      }t        j.                  |t        j                  t        j0                  | j                  j                  ‰«      «      «      S c c}w c c}w )
NFr  zMExpected `base` to be a pointer type (but not a block pointer type or others)rè  c              3  ó`   K  — | ]&  }t        |t        «      xr d |cxk  xr dk  nc –— Œ( y­w)rR   rS   N)r]   ra   )rd  rí  s     r   rf  z!make_block_ptr.<locals>.<genexpr>e  s)   è ø€ ÒXÀDŒz˜$¤Ó$Ò?¨°4Ö)?¸%Ô)?Ó?ÑXùs   ‚,.zGExpected a list of constant integers (`int32_t` range) in `block_shape`z<Expected a permutation of (0, 1, ..., len(order)-1) in orderc              3  óL   •K  — | ]  }t        ‰«      t        |«      k(  –— Œ y ­wrq   )rB  )rd  r  rì  s     €r   rf  z!make_block_ptr.<locals>.<genexpr>o  s   øè ø€ Òd°iŒs�;Ó¤3 y£>Õ1Ñdùs   ƒ!$zBExpected shape/strides/offsets/block_shape to have the same length)r  rj   rr   ró  rž   r"   r#   r`   r}   rþ  rý  rÿ  ré  r]   rh   r@   rÚ  rg  rh  r8  rB  Úcreate_make_block_ptrr�   r$   rŸ   )	Úbaser  Ústridesr  rì  Úorderr)   rí  r�   s	       `    r   Úmake_block_ptrr  R  sã  ø€ ô " '¨5Ó1€EÜ# G¨WÓ5€GÜ# G¨WÀ%ÔH€Gð �9‰9×ÑÔ §¡×!5Ñ!5×!>Ñ!>Ô!@ÜÐhÓiÐið ‡y�y×ÑœrŸw™wÒ&Ü�Dœ"Ÿ/™/¬"¯'©'°4·9±9×3JÑ3JÓKÈWÓUˆô �; 
Ô+Ø"�mˆØVaÖbÈd¤¨D´"·,±,Ô!?�4—:’:ÀTÑIÒb€KÜÑXÈKÔXÔXð RØQóRÐXô �5˜*Ô%Ø�ˆØPUÖVÈœ: d¬B¯L©LÔ9ˆT�ZŠZ¸tÑCÐV€EÐVÜ�%‹=œD¤¤s¨5£zÓ!2Ó3Ò3ÐsÐ5sÓsÐ3ô ÓdÀEÈ7ÐT[Ð]bÐCcÔdÔdð MØLóMÐdð ×*Ñ*¨4¯;©;¸¸wÈÐQ\Ð^cÓd€FÜ�9‰9�VœRŸ_™_¬R¯]©]¸4¿9¹9×;OÑ;OÐQ\Ó-]Ó^Ó_Ð_ùò% cùò Ws   Ã./H=Å/Ic                ó’   — t        ||d¬«      }t        j                  |j                  | j                  |«      | j
                  «      S )NFr  )r  r#   r$   Úcreate_advancer�   rj   )r  r  r)   s      r   Úadvancer  y  s8   € ä# G¨WÀ%ÔH€Gô �9‰9�W×+Ñ+¨D¯K©K¸ÓAÀ4Ç9Á9ÓMÐMr   c                ó¸  — t        |«      }d|cxk  rdk  sn t        d|› d�«      ‚t        |«      |k7  rt        d|› dt        |«      › �«      ‚t        |«      |k7  rt        d|› dt        |«      › �«      ‚t        j                  |d	   «      |d	<   |d	   d
k7  rt        d|d	   › �«      ‚|D �cg c]  }t	        ||«      ‘Œ }}|D �cg c]-  }t	        ||«      j                  t        j                  |¬«      ‘Œ/ }}t        j                  |«      }t        | j                  t        j                  «      sJ ‚t        j                  | j                  j                  |«      }|j                  | j                  |D �cg c]  }|j                  ‘Œ c}|D �cg c]  }|j                  ‘Œ c}|«      }	t        j                  |	|||«      S c c}w c c}w c c}w c c}w )Nr	   é   z Expected 2 <= ndim <= 5 but got z dimensionsz	Expected z strides but got zExpected block_shape to have z dimensions but got rT  r   z-Tensor descriptor last dim must be 1 but got r“  )rB  r"   r#   rK  ri   Útorc   Ú_unwrap_shaper]   rj   rþ  rŸ   ró  Úcreate_make_tensor_descriptorr�   Ú_experimental_tensor_descriptor)
r  r  r  rì  r)   r  rk   rj   rH  r�   s
             r   Úmake_tensor_descriptorr%  �  s¸  € ô ˆu‹:€DØ�ŒN˜ŒNÜÐ;¸D¸6ÀÐMÓNÐNÜ
ˆ7ƒ|�tÒÜ˜9 T FÐ*;¼CÀ»L¸>ÐJÓKÐKÜ
ˆ;Ó˜4ÒÜÐ8¸¸Ð>RÔSVÐW^ÓS_ÐR`ÐaÓbÐbä×(Ñ(¨°©Ó5€GˆB�KØˆr�{�aÒÜÐHÈÐQSÉÈÐVÓWÐWà,1Ö2 qŒY�q˜'Õ"Ð2€EÐ2ØMTÖUÈŒy˜˜GÓ$×'Ñ'¬¯©¸7Ð'ÕCÐU€GÐUô ×"Ñ" ;Ó/€Kä�d—i‘i¤§¡Ô1Ð1Ð1Ü�=‰=˜Ÿ™×-Ñ-¨{Ó;€DØ×2Ñ2°4·;±;ÐSXÖ@YÈaÀÇÃÒ@YÐnuÖ[vÐijÐ\]×\dÓ\dÒ[vØ3>ó@€Fä×-Ñ-¨f°e¸WÀdÓKÐKùò 3ùÚUùò AZùÒ[vs   Â0GÃ	2GÅ<G
ÆG)r(   ra   r)   ú
ir.builderÚreturnú	tl.tensor)r5   útl.dtyper6   r)  r'  r)  )r5   r)  rK   r^   r6   r)  rL   r^   rM   r^   r'  r)  )T)rl   r^   )r   r)  r   r)  rt   r^   r'  ÚNone)FFTF)r   útl.tensor | numbers.Numberr€   r+  r)   r&  r'  úTuple[tl.tensor, tl.tensor])r   r(  r€   r(  r)   r&  r•   Úcallable)
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(  )rk   r(  rÖ   r(  r×   útl.PropagateNanr)   r&  )
rk   r(  râ   r(  rã   r(  r×   r.  r)   r&  )r§   r(  r¨   r(  r)   r&  r'  r,  )r§   r(  r¨   r(  r)   r&  r'  r(  )r§   r(  r)   r&  )r§   r(  r'  r(  )r§   r(  r)   r&  r'  r(  )r§   r(  r)   r(  r'  r(  )r  r(  r'  útl.block_type)r4  ra   r5  ra   r)   r&  r'  r(  )r  ú	List[int]r1   r)  r)   r&  r'  r(  )r@   r(  r  r0  r)   r&  r'  r(  )
r§   r(  rF  r0  rG  r^   r)   r&  r'  r(  )r§   r(  r(   ra   r)   r&  r'  r(  )
r   r(  r€   r(  rG  r^   r)   r&  r'  r(  )rW  r(  rX  r(  r)   r&  r'  r(  )rW  r(  r)   r&  r'  r,  )r§   r(  rj  z
Tuple[int]r)   r&  r'  r(  )r§   r(  r  r0  r)   r&  r'  r(  )r   r(  r€   r(  r)   r&  r'  r(  )r†  úOptional[str])r§   r(  r‹  r)  r)   r&  r'  r(  rq   )
r§   r(  r‹  r)  r)   r&  r£  r1  r'  r(  )rõ  r(  rö  úOptional[tl.tensor]r¨   r2  rë  r   rË  rw  r³  rw  rÂ  rw  r÷  r^   r)   r&  r'  r(  )r	  r(  r
  r/  r)   r&  )
r  z&tl._experimental_tensor_desciptor_baser³  rw  rÂ  rw  r)   r&  r'  r(  )r  z'tl._experimental_tensor_descriptor_baser@   r(  r)   r&  r'  r(  )r³  rw  rÂ  rw  r)   r&  r'  r(  )r@   r(  r)   r&  r'  r(  )r	  r(  r/  r(  r0  úList[tl.tensor]r1  r3  r2  r3  r3  r3  r4  ra   r5  ra   r6  ra   r7  ra   r)   r&  r'  r(  )r	  r(  r)   r&  r'  r(  )rõ  r(  r>  r(  rö  r2  r³  rw  rÂ  rw  r)   r&  r'  r(  )rõ  r(  rK  r(  r>  r(  rÚ  rw  rå  rw  r)   r&  r'  r(  )rõ  r(  r>  r(  rö  r(  rP  rw  r)   r&  r'  z&Tuple[tl.tensor, tl.tensor, tl.tensor])rõ  r(  r>  r(  rö  r(  rÚ  rw  rå  rw  r)   r&  r'  r(  )r   r(  r€   r(  r™  r(  rŠ  r1  rš  ra   r›  r)  r)   r&  r'  r(  )r§  rw  )r>  r(  r§  rw  r)   r&  )r   r(  r¸  r(  r¹  rw  r€   r(  rº  r2  r»  rw  r™  ztl.tensor | Noner¼  r^   r›  r)  r)   r&  r'  r(  )
rË  r(  rk   r(  rÖ   r(  r)   r&  r'  r(  )rÖ  úSequence[tl.tensor]r(   ra   r)   r&  r'  úTuple[tl.tensor, ...])
rÖ  r4  r(   ra   rä  r^   r)   r&  r'  r5  )
rè  r(  ré  r(  r(   ra   r)   r&  r'  r(  )r§   r(  rí  ra   r)   r&  r'  r(  )rk   r(  ru  r0  r'  r(  )r)   r&  r'  r(  )
r  rw  r  r3  r  r^   r)   r&  r'  r(  )r™   r(  rš   rw  r)   r&  r'  r(  )r  r(  r)   r&  r'  r(  )r  r(  r  r3  r  r3  rì  zList[tl.constexpr]r)   r&  r'  z"tl._experimental_tensor_descriptor){Ú
__future__r   rÈ  Útypingr   r   r   r   r   rw   Ú_C.libtritonr
   r   r   r#   r   Ú	Exceptionr   r*   r-   r;   rP   ri   ru   r‹   r›   r¥   r°   r¶   r¼   rÂ   rÅ   rÊ   rØ   rß   rä   rè   r‘   rï   rò   r÷   rù   rü   rÿ   r  r  r  r®   rû   r  r  r“   r#  r’   r+  r0  r9  re   r?  rI  rN  rR  r\  ra  rk  rs  r~   r‡  rö   r}   rµ  rº  rÄ  rÍ  rÛ  ræ  rï  rù  r  r  r  r  r  r  r(  r,  r8  r;  r?  rC  rG  rL  rS  rg  rk  ro  rt  ry  r}  r�  r‹  r¤  r©  rµ  rÆ  rZ  rÎ  rß  rå  rê  rî  rô  r÷  rù  rü  r  r”   r  r  r  r  r  r%  rW   r   r   ú<module>r:     sº  ðÝ "Û ç ;Õ ;Û å Ý áˆCƒL€ôF 	ô FóDóFó@ð 0,Ø&*ð0,Ø/7ó0,ôf#óV	<ð ]aØ,1ð#à6Qó#óL&ð& :Øð :Ø )ó :ðF4Øð4Ø )ó4ð$4Øð4Ø )ó4óRó4:ð&Øð&Ø!*ó&ó4ó25ó$5ó$	]ð"Ø+5ðØ:UóóQó
Pó
Qó
'ó&ó"óRó
Ró
Qóó)ó$ó)ó4ó4ó4ó4ó	4ó	4ó"Dó"(ó.Hó[óMóKóó2
óIóLó$2ótpóZð( 04ðh5Ø,ðh5Ø8Aóh5ò`òò 	ò	òò ò
òwò,;ð|nØðnØ.1ðnØDGðnØVZðnàðnà!*ónó Eò
{ð
)Ø'ð
)Ø,5ó
)ð	cØ(ð	cØ-6ó	cðØ)ðØ.7óó4$ð,Øðàðð ðð  ð	ð
 #ðð $ðð ðð ðð ðð ðð ðð óó>\òò:)qðXWØðWØ)3ðWØ8AóWó,jðØ(2ðØ7]óó6$)óN$)óNoóóóðØ#ðØ(1óò8ðFØðFØ&0ðFØ5>óFóRó0ð$(SØð(SØ%5ð(SØBFð(SØS[ð(Sà"ð(Sà'0ó(Só`Zò, óuð.oØ(ðoØ-Bóoó2Bó0lóóóó8óVóGóBòUó0Hó$`óNNðLØ
ðLàðLð ðLð $ð	Lð
 ðLð (ôLr   