§
    ÕóÉjÚ| ã                  óâ   — d dl mZ d dlZd dlZd dlmZmZmZmZm	Z	m
Z
mZ d dlZd dlmZ ddlmZ ddlmZ  e	d	¦  «        Z e	d
¦  «        Z G d„ de¦  «        Z G d„ de
e         ¦  «        ZdS )é    )ÚannotationsN)ÚListÚOptionalÚSequenceÚTupleÚTypeVarÚGenericÚType)Údriveré   )Úiré   )ÚcoreÚTÚTensorTyc                  ó   ‡ — e Zd Zˆ fd„Zˆ xZS )ÚIncompatibleTypeErrorImplc                óú   •— || _         || _        d| j                              ¦   «         z   dz   | j                             ¦   «         z   | _        t	          t
          | ¦  «                             | j        ¦  «         d S )Nzinvalid operands of type ú and )Útype_aÚtype_bÚ__repr__ÚmessageÚsuperr   Ú__init__)Úselfr   r   Ú	__class__s      €ú}/home/lighthouse/badge-speaker-benchmark/.benchmark-venv-community/lib64/python3.11/site-packages/triton/language/semantic.pyr   z"IncompatibleTypeErrorImpl.__init__   sl   ø€ ØˆŒØˆŒØ2°T´[×5IÒ5IÑ5KÔ5KÑKÈgÑUÐX\ÔXc×XlÒXlÑXnÔXnÑnˆŒÝÕ'¨Ñ.Ô.×7Ò7¸¼ÑEÔEÐEÐEÐEó    )Ú__name__Ú
__module__Ú__qualname__r   Ú__classcell__)r   s   @r   r   r      sA   ø€ € € € € ðFð Fð Fð Fð Fð Fð Fð Fð Fr   r   c                  ó  — e Zd ZU ej        Zded<   eZded<   d„ Zdçd
„Zdçd„Z	dèd„Z
déd„Zdêd„Zd„ Zdëd„Z	 	 dì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ðd0„Zdòd5„Zdòd6„Zdód9„Zdôd:„Zdõd;„Zdõd<„Zdõd=„Zdõd>„Zdõd?„Z död@„Z!dõdA„Z"dõdB„Z#dõdC„Z$d÷dD„Z%d÷dE„Z&d÷dF„Z'dødI„Z(dõdJ„Z)dõdK„Z*dõdL„Z+dõdM„Z,dõdN„Z-dõdO„Z.dPdQœdùdU„Z/dúdW„Z0dúdX„Z1dûd[„Z2düd]„Z3dýd^„Z4dþda„Z5dÿdb„Z6�d dc„Z7�ddf„Z8�ddg„Z9�ddj„Z:�ddk„Z;�ddl„Z<�ddo„Z=�ddq„Z>�d�d	ds„Z?dt„ Z@du„ ZAdv„ ZBdw„ ZCdx„ ZDdy„ ZE�d
d„„ZF�dd‡„ZG�ddˆ„ZH�dd‰„ZI�ddŠ„ZJd‹„ ZKdŒ„ ZL�dd�„ZM�ddŽ„ZN�dd�„ZO�dd�„ZP�dd‘„ZQ�dd’„ZRdýd“„ZSd”„ ZT�dd–„ZU�ddš„ZV�dd�„ZW�ddž„ZX�ddŸ„ZY�dd „ZZ�dd¡„Z[�dd¢„Z\�dd£„Z]�dd¤„Z^�dd¥„Z_d¦„ Z`�dd¬„Za�dd®„Zb�dd¯„Zcd°„ Zdd±„ Ze�ddº„Zf�dd¼„Zgd½„ Zh�ddÁ„Zi�ddÃ„Zj�ddÆ„ZkdÇ„ Zl�ddÍ„Zm�ddÏ„Zn�ddÑ„Zo�ddÒ„Zp�ddÓ„Zq�ddÔ„Zr�d dÙ„Zs�d!dÜ„Zt�ddÝ„ZudÞ„ Zvdêdß„Zw	 �d"�d#dæ„ZxdPS ($  ÚTritonSemanticzType[TensorTy]Útensorz
ir.builderÚbuilderc                ó   — || _         d S ©N)r'   )r   r'   s     r   r   zTritonSemantic.__init__    s   € ØˆŒˆˆr   ÚaxisÚintÚreturnr   c                óž   — |dvrt          d|› �¦  «        ‚|                      | j                             |¦  «        t          j        ¦  «        S )N©r   r   r   z+program_id axis must be 0, 1, or 2 but got )Ú
ValueErrorr&   r'   Úcreate_get_program_idÚtlÚint32©r   r*   s     r   Ú
program_idzTritonSemantic.program_id'   sJ   € Ø�yÐ Ð ÝÐQÈ4ÐQÐQÑRÔRÐRØ�{Š{˜4œ<×=Ò=¸dÑCÔCÅRÄXÑNÔNÐNr   c                óž   — |dvrt          d|› �¦  «        ‚|                      | j                             |¦  «        t          j        ¦  «        S )Nr.   z-num_programs axis must be 0, 1, or 2 but got )r/   r&   r'   Úcreate_get_num_programsr1   r2   r3   s     r   Únum_programszTritonSemantic.num_programs,   sJ   € Ø�yÐ Ð ÝÐSÈTÐSÐSÑTÔTÐTØ�{Š{˜4œ<×?Ò?ÀÑEÔEÅrÄxÑPÔPÐPr   Úa_tyútl.dtypeÚb_tyc                ó  — |j         }|j         }|j        }|j        }||k    r
||k    r|n|S |t          j        j        j        k    r
||k    r|n|S |t          j        j        j        k    r
||k    r|n|S t          d|› d|› �¦  «        ‚)Nzunexpected signedness r   )Úint_bitwidthÚint_signednessr1   ÚdtypeÚ
SIGNEDNESSÚUNSIGNEDÚ	TypeError)r   r8   r:   Úa_rankÚb_rankÚa_snÚb_sns          r   Úinteger_promote_implz#TritonSemantic.integer_promote_impl5   s¬   € ØÔ"ˆØÔ"ˆØÔ"ˆØÔ"ˆð �4Š<ˆ<Ø! Fš?˜?�4�4°Ð4Ø•R”XÔ(Ô1Ò1Ð1Ø! VÒ+Ð+�4�4°Ð5Ø•R”XÔ(Ô1Ò1Ð1Ø! VÒ+Ð+�4�4°Ð5ÝÐB°ÐBÐB¸DÐBÐBÑCÔCÐCr   Úa_is_scalarÚboolÚb_is_scalarÚ
div_or_modc                óä  — ||k    rk|r||fn||f\  }}|                      ¦   «         j        |                      ¦   «         j        k    r*|r&|t          j        t          j        fv rt          j        S |S |                     ¦   «         s|                     ¦   «         rt          j        S |                     ¦   «         s|                     ¦   «         rt          j        S | 	                    ¦   «         s| 	                    ¦   «         r|rt          j        S t          j        S | 
                    ¦   «         r.| 
                    ¦   «         r|rt          j        S t          j        S | 
                    ¦   «         s| 
                    ¦   «         rt          j        S |                     ¦   «         r(|                     ¦   «         r||k    r|nt          j        S |                     ¦   «         r|                     ¦   «         st          d|› d|› �¦  «        ‚|rO|j        |j        k    r?t          d|                     ¦   «         z   dz   |                     ¦   «         z   dz   ¦  «        ‚|                      ||¦  «        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Úvaluer1   Úfloat16Úbfloat16Úfloat32Úis_fp64Úfloat64Úis_fp32Úis_fp16Úis_bf16Úis_fp8Úis_intrA   r=   r   rF   )r   r8   rG   r:   rI   rJ   Ú	scalar_tyÚ	tensor_tys           r   Úcomputation_type_implz$TritonSemantic.computation_type_implD   sO  € ð
 ˜+Ò%Ð%Ø3>Ð#P D¨$ < <ÀTÈ4ÀLÑ ˆI�yØ�~Š~ÑÔÔ%¨¯ªÑ)9Ô)9Ô)?Ò?Ð?àð & 9µ´½R¼[Ð0IÐ#IÐ#IÝœ:Ð%Ø Ð ð �<Š<‰>Œ>ð 	˜TŸ\š\™^œ^ð 	Ý”:Ðð �<Š<‰>Œ>ð 	˜TŸ\š\™^œ^ð 	Ý”:Ðð �<Š<‰>Œ>ð 	"˜TŸ\š\™^œ^ð 	"Øð "Ý”zÐ!å”zÐ!à�<Š<‰>Œ>ð 	#˜dŸlšl™nœnð 	#Øð #Ý”zÐ!å”{Ð"Ø�<Š<‰>Œ>ð 	˜TŸ\š\™^œ^ð 	Ý”:Ðà�;Š;‰=Œ=ð 	8˜TŸ[š[™]œ]ð 	8Ø 4š<˜<�4�4­R¬ZÐ7Ø�{Š{‰}Œ}ð 	B D§K¢K¡M¤Mð 	BÝÐ@¨tÐ@Ð@¸$Ð@Ð@ÑAÔAÐAð ð 	p˜$Ô-°Ô1DÒDÐDÝÐ9¸D¿MºM¹O¼OÑKÈgÑUÐX\×XeÒXeÑXgÔXgÑgðoñoñ pô pð pð ×(Ò(¨¨tÑ4Ô4Ð4r   Tc                ó`  — t          || j        ¦  «        r|S t          |t          j        ¦  «        r|j        n|}t          |t
          t          t          f¦  «        r,|                      |¦  «        }|  	                    ||¬¦  «        S |r#t          d|› dt          |¦  «        › d�¦  «        ‚|S )N©r>   úcannot convert ú	 of type ú
 to tensor)Ú
isinstancer&   r1   Ú	constexprrO   r+   ÚfloatrH   Úto_tensor_typeÚscalar_constantrA   Útype)r   ÚxÚ
check_typer>   s       r   Ú	to_tensorzTritonSemantic.to_tensorv   s³   € Ý�a˜œÑ%Ô%ð 	ØˆHÝ! !¥R¤\Ñ2Ô2Ð9ˆAŒGˆG¸ˆÝ�a�#�u¥dÐ+Ñ,Ô,ð 	OØ×'Ò'¨Ñ*Ô*ˆEØ×'Ò'¨°Ð'Ñ7Ô7Ð7Øð 	OÝÐM¨aÐMÐM½$¸q¹'¼'ÐMÐMÐMÑNÔNÐNØˆr   c                ó(  — t          |t          j        ¦  «        r|S t          |t          j        ¦  «        r|j        }t          |t
          ¦  «        rt          j        S t          |t          ¦  «        rƒd|cxk    rdk     rn nt          j        S d|cxk    rdk     rn nt          j	        S d|cxk    rdk     rn nt          j
        S d|cxk    rdk     rn nt          j        S t          d|› d�¦  «        ‚t          |t          ¦  «        red	}d
ddz  z  }t          j        |¦  «        }|t          d¦  «        k    s|dk    s||k    s||cxk    r|k    rn nt          j        S t          j        S t%          d|› dt'          |¦  «        › d�¦  «        ‚)Né   €ì        l        ìûÿÿÿ         ì            l            zNonrepresentable integer ú.g      8g   àÿÿÿ?r   é   Úinfg        r_   r`   ra   )rb   r1   r>   Úconstexpr_typerO   rH   Úint1r+   r2   Úuint32Úint64Úuint64r/   rd   ÚbuiltinsÚabsrR   rT   rA   rg   )r   rh   Úmin_float32Úmax_float32Úabs_xs        r   re   zTritonSemantic.to_tensor_type�   sÑ  € Ý�a�œÑ"Ô"ð 	ØˆHÝ˜�2Ô,Ñ-Ô-ð 	Ø”ˆAå�a�ÑÔð 	"Ý”7ˆNÝ˜�3ÑÔð 	"Ø˜Ð"Ð"Ò"Ð"˜UÒ"Ð"Ð"Ð"Ð"Ý”x�Ø˜!Ð#Ð#Ò#Ð#˜eÒ#Ð#Ð#Ð#Ð#Ý”yÐ Ø˜1Ð$Ð$Ò$Ð$˜uÒ$Ð$Ð$Ð$Ð$Ý”x�Ø˜!Ð#Ð#Ò#Ð#˜eÒ#Ð#Ð#Ð#Ð#Ý”yÐ ÝÐ=¸Ð=Ð=Ð=Ñ>Ô>Ð>Ý˜�5Ñ!Ô!ð 
	"Ø!ˆKØ%¨¨C©Ñ/ˆKÝ”L ‘O”OˆEØ�˜e™œÒ$Ð$Ø˜Š|ˆ|Ø�AŠvˆvØ˜eÐ2Ð2Ò2Ð2 {Ò2Ð2Ð2Ð2Ð2Ý”zÐ!å”zÐ!ÝÐI¨!ÐIÐIµd¸1±g´gÐIÐIÐIÑJÔJÐJr   r   r   Úallow_ptr_aÚNonec                óò   — |                      ¦   «         r`|st          ||¦  «        ‚|                      ¦   «         r||k    rt          ||¦  «        ‚|                     ¦   «         rt          ||¦  «        ‚d S d S r)   )Úis_ptrr   Úis_floating)r   r   r   r}   s       r   Úcheck_ptr_type_implz"TritonSemantic.check_ptr_type_impl¤   s”   € Ø�=Š=‰?Œ?ð 	@Øð @Ý/°¸Ñ?Ô?Ð?à�}Š}‰Œð @ F¨fÒ$4Ð$4Ý/°¸Ñ?Ô?Ð?à×!Ò!Ñ#Ô#ð @Ý/°¸Ñ?Ô?Ð?ð	@ð 	@ð@ð @r   FÚlhsúTensorTy | numbers.NumberÚrhsúTuple[TensorTy, TensorTy]c                óº  — t          |t          j        ¦  «        }t          |t          j        ¦  «        }|r|}	|                      |¦  «        }|r|}
|                      |¦  «        }|j        j        }|j        j        }|                      |||¦  «         |                      |||¦  «         |�r�|                     ¦   «         �s{|                     ¦   «         �sf|                      |||||¦  «        }|r|	dk     r| 	                    ¦   «         s|r)|
dk     r#| 	                    ¦   «         rt          d¦  «        ‚|                     ¦   «         r”|rH|                     ¦   «         |	cxk    r|                     ¦   «         k    sn t          d|	› d|› �¦  «        ‚|rH|                     ¦   «         |
cxk    r|                     ¦   «         k    sn t          d|
› d|› �¦  «        ‚|r|                      |	|¬¦  «        n|                      ||¦  «        }|r|                      |
|¬¦  «        n|                      ||¦  «        }|                      ||¦  «        \  }}||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 r^   )rb   ÚnumbersÚNumberrj   rg   Úscalarr‚   r€   r\   Úis_int_unsignedr/   rY   Úget_int_min_valueÚget_int_max_valuerf   ÚcastÚbroadcast_impl_value)r   rƒ   r…   Úallow_lhs_ptrÚallow_rhs_ptrÚarithmetic_checkrJ   Úlhs_is_scalarÚrhs_is_scalarÚ
lhs_scalarÚ
rhs_scalarÚ
lhs_sca_tyÚ
rhs_sca_tyÚ
ret_sca_tys                 r   Úbinary_op_type_checking_implz+TritonSemantic.binary_op_type_checking_impl¯   sä  € õ # 3­¬Ñ7Ô7ˆÝ" 3­¬Ñ7Ô7ˆØð 	&ØˆJØ—.’. Ñ%Ô%ˆCØð 	&ØˆJØ—.’. Ñ%Ô%ˆCð ”X”_ˆ
Ø”X”_ˆ
Ø× Ò  ¨Z¸ÑGÔGÐGØ× Ò  ¨Z¸ÑGÔGÐGØñ 	v J×$5Ò$5Ñ$7Ô$7ñ 	vÀ
×@QÒ@QÑ@SÔ@Sñ 	vØ×3Ò3°JÀÈzÐ[hÐjtÑuÔuˆJØð L *¨q¢. .°Z×5OÒ5OÑ5QÔ5Q .Ø$ð #1Ø)3°aª¨¸J×<VÒ<VÑ<XÔ<X¨Ý ð "Kñ Lô Lð Là× Ò Ñ"Ô"ð cØ ð c¨*×*FÒ*FÑ*HÔ*HÈJð +Ið +Iò +Ið +IØ*4×*FÒ*FÑ*HÔ*Hò+Ið +Ið +Ið +Iå$Ð%a¨zÐ%aÐ%aÐU_Ð%aÐ%aÑbÔbÐbØ ð c¨*×*FÒ*FÑ*HÔ*HÈJð +Ið +Iò +Ið +IØ*4×*FÒ*FÑ*HÔ*Hò+Ið +Ið +Ið +Iå$Ð%a¨zÐ%aÐ%aÐU_Ð%aÐ%aÑbÔbÐbØHUÐu�$×&Ò& z¸Ð&ÑDÔDÐDÐ[_×[dÒ[dÐehÐjtÑ[uÔ[uˆCØHUÐu�$×&Ò& z¸Ð&ÑDÔDÐDÐ[_×[dÒ[dÐehÐjtÑ[uÔ[uˆCð ×,Ò,¨S°#Ñ6Ô6‰ˆˆSØ�Cˆxˆr   Ú	binary_opÚcallablec                ó,  — |j         j        j        dk    s| j        j        j        sd S |j         j        }|j         j        }||k    sJ d|› d|› �¦   «         ‚|                     ¦   «         sJ d|› �¦   «         ‚|                      |t          j	        ¦  «        }|                      |t          j	        ¦  «        } |||d¦  «        }| 
                    ¦   «         }|                      |t          j	        ¦  «        }|                     ¦   «         }|                      |t          j	        ¦  «        }|                      |                      ||¦  «        |                      ||¦  «        ¦  «        }	d|j        › d|j        › �}
|                      |	|
d ¦  «         d S )Né@   z%expected matching operand types, got r   zexpected integer type, got Fr+   z! overflow detected for operation )rg   rŠ   r<   r'   ÚoptionsÚsanitize_overflowrY   rŽ   r1   rv   r�   rf   rŒ   Úand_Ú
less_equalÚgreater_equalr    Údevice_assert)r   rƒ   r…   r›   r—   r˜   ÚretÚ	max_valueÚ	min_valueÚcondÚmsgs              r   Ú binary_op_sanitize_overflow_implz/TritonSemantic.binary_op_sanitize_overflow_implÔ   s…  € ØŒ8Œ?Ô'¨2Ò-Ð-°T´\Ô5IÔ5[Ð-ØˆFØ”X”_ˆ
Ø”X”_ˆ
Ø˜ZÒ'Ð'Ð'Ð)nÐQ[Ð)nÐ)nÐblÐ)nÐ)nÑ'Ô'Ð'Ø× Ò Ñ"Ô"ÐNÐNÐ$NÀ*Ð$NÐ$NÑNÔNÐNØ�iŠi˜�RœXÑ&Ô&ˆØ�iŠi˜�RœXÑ&Ô&ˆØˆi˜˜S %Ñ(Ô(ˆØ×0Ò0Ñ2Ô2ˆ	Ø×(Ò(¨µB´HÑ=Ô=ˆ	Ø×0Ò0Ñ2Ô2ˆ	Ø×(Ò(¨µB´HÑ=Ô=ˆ	Ø�yŠy˜Ÿš¨¨iÑ8Ô8¸$×:LÒ:LÈSÐR[Ñ:\Ô:\Ñ]Ô]ˆØb�JÔ+ÐbÐbÈiÔN`ÐbÐbˆØ×Ò˜4  dÑ+Ô+Ð+Ð+Ð+r   ÚinputÚotherr    c                óº  — |                       ||dd¦  «        \  }}|j        j        }|j        j        }|                     ¦   «         r#|                     ¦   «         rt	          d¦  «        ‚|                     ¦   «         r0|                     ¦   «         s||}}|j        j        }|j        j        }|                     ¦   «         rÆ|j        }|j                             ¦   «         rm|j        j        dk     r]|j         	                    t          j        ¦  «                             | j        ¦  «        }| j                             |j        |d¦  «        }|                      | j                             |j        |¦  «        |j        ¦  «        S |                     ¦   «         r>|                      | j                             |j        |j        ¦  «        |j        ¦  «        S |                     ¦   «         r\|r|                      ||| j        ¦  «         |                      | j                             |j        |j        ¦  «        |j        ¦  «        S t	          d|› �¦  «        ‚)NTzcannot add pointers togetherrž   FrL   )rš   rg   rŠ   r€   rA   Úhandler>   r‹   r<   Úwith_element_tyr1   rv   Úto_irr'   Úcreate_int_castr&   Úcreate_addptrr�   Úcreate_faddrY   rª   ÚaddÚ
create_add)r   r«   r¬   r    Úinput_scalar_tyÚother_scalar_tyÚother_handleÚi64_tys           r   r´   zTritonSemantic.addæ   s   € à×8Ò8¸ÀÀtÈTÑRÔR‰ˆˆuØœ*Ô+ˆØœ*Ô+ˆØ×!Ò!Ñ#Ô#ð 	<¨×(>Ò(>Ñ(@Ô(@ð 	<ÝÐ:Ñ;Ô;Ð;ð ×!Ò!Ñ#Ô#ð 	0¨O×,BÒ,BÑ,DÔ,Dð 	0Ø  %�5ˆEØ#œjÔ/ˆOØ#œjÔ/ˆOØ×!Ò!Ñ#Ô#ð 	`Ø œ<ˆLØŒ{×*Ò*Ñ,Ô,ð Y°´Ô1IÈBÒ1NÐ1Nàœ×3Ò3µB´HÑ=Ô=×CÒCÀDÄLÑQÔQ�Ø#œ|×;Ò;¸E¼LÈ&ÐRWÑXÔX�Ø—;’;˜tœ|×9Ò9¸%¼,ÈÑUÔUÐW\ÔWaÑbÔbÐbà×(Ò(Ñ*Ô*ð 	`Ø—;’;˜tœ|×7Ò7¸¼ÀeÄlÑSÔSÐUZÔU_Ñ`Ô`Ð`à×#Ò#Ñ%Ô%ð 	`Ø ð NØ×5Ò5°e¸UÀDÄHÑMÔMÐMØ—;’;˜tœ|×6Ò6°u´|ÀUÄ\ÑRÔRÐTYÔT^Ñ_Ô_Ð_ÝÐ<¨?Ð<Ð<Ñ=Ô=Ð=r   c                óv  — |                       ||dd¦  «        \  }}|j        j        }|                     ¦   «         r+|                      ||                      |¦  «        d¬¦  «        S |                     ¦   «         r>|                      | j         	                    |j
        |j
        ¦  «        |j        ¦  «        S |                     ¦   «         r\|r|                      ||| j        ¦  «         |                      | j                             |j
        |j
        ¦  «        |j        ¦  «        S t          d|› �¦  «        ‚)NTF)r    rL   )rš   rg   rŠ   r€   r´   Úminusr�   r&   r'   Úcreate_fsubr®   rY   rª   ÚsubÚ
create_subrA   ©r   r«   r¬   r    rZ   s        r   r½   zTritonSemantic.sub  s"  € à×8Ò8¸ÀÀtÈUÑSÔS‰ˆˆuØ”JÔ%ˆ	à×ÒÑÔð 	OØ—8’8˜E 4§:¢:¨eÑ#4Ô#4È�8ÑNÔNÐNà× Ò Ñ"Ô"ð 	`Ø—;’;˜tœ|×7Ò7¸¼ÀeÄlÑSÔSÐUZÔU_Ñ`Ô`Ð`à×ÒÑÔð 	`Ø ð NØ×5Ò5°e¸UÀDÄHÑMÔMÐMØ—;’;˜tœ|×6Ò6°u´|ÀUÄ\ÑRÔRÐTYÔT^Ñ_Ô_Ð_ÝÐ6¨9Ð6Ð6Ñ7Ô7Ð7r   c                óô  — |                       ||¦  «        \  }}|j        j        }|                     ¦   «         r>|                      | j                             |j        |j        ¦  «        |j        ¦  «        S |                     ¦   «         r\|r|  	                    ||| j
        ¦  «         |                      | j                             |j        |j        ¦  «        |j        ¦  «        S t          d|› �¦  «        ‚©NrL   )rš   rg   rŠ   r�   r&   r'   Úcreate_fmulr®   rY   rª   ÚmulÚ
create_mulrA   r¿   s        r   rÃ   zTritonSemantic.mul  så   € à×8Ò8¸ÀÑFÔF‰ˆˆuØ”JÔ%ˆ	à× Ò Ñ"Ô"ð 	`Ø—;’;˜tœ|×7Ò7¸¼ÀeÄlÑSÔSÐUZÔU_Ñ`Ô`Ð`à×ÒÑÔð 	`Ø ð NØ×5Ò5°e¸UÀDÄHÑMÔMÐMØ—;’;˜tœ|×6Ò6°u´|ÀUÄ\ÑRÔRÐTYÔT^Ñ_Ô_Ð_ÝÐ6¨9Ð6Ð6Ñ7Ô7Ð7r   c                ó¨  — |                       ||dddd¦  «        \  }}|j        j        }|j        j        }|                     ¦   «         r,|                     ¦   «         r|                      ||¦  «        }�n |                     ¦   «         r+|                     ¦   «         r|                      ||¦  «        }ná|                     ¦   «         rU|                     ¦   «         rA|                      |t          j        ¦  «        }|                      |t          j        ¦  «        }nx|                     ¦   «         rR|                     ¦   «         r>|j        |j        k    r|                      ||¦  «        }n)|                      ||¦  «        }nt          d|› �¦  «        ‚|  
                    | j                             |j        |j        ¦  «        |j        ¦  «        S ©NFTrL   )rš   rg   rŠ   r�   rY   rŽ   r1   rR   Úfp_mantissa_widthrA   r&   r'   Úcreate_fdivr®   )r   r«   r¬   r¶   r·   s        r   ÚtruedivzTritonSemantic.truediv$  sµ  € Ø×8Ò8¸ÀÀuÈeÐUYÐ[_Ñ`Ô`‰ˆˆuØœ*Ô+ˆØœ*Ô+ˆà×&Ò&Ñ(Ô(ð 	B¨_×-CÒ-CÑ-EÔ-Eð 	BØ—I’I˜e _Ñ5Ô5ˆE‰Eà×#Ò#Ñ%Ô%ð 	B¨/×*EÒ*EÑ*GÔ*Gð 	BØ—I’I˜e _Ñ5Ô5ˆEˆEà×#Ò#Ñ%Ô%ð 	B¨/×*@Ò*@Ñ*BÔ*Bð 	BØ—I’I˜e¥R¤ZÑ0Ô0ˆEØ—I’I˜e¥R¤ZÑ0Ô0ˆEˆEà×(Ò(Ñ*Ô*ð 	B¨×/JÒ/JÑ/LÔ/Lð 	BØÔ0°?Ô3TÒTÐTØŸ	š	 %¨Ñ9Ô9��àŸ	š	 %¨Ñ9Ô9��õ Ð@¨Ð@Ð@ÑAÔAÐAØ�{Š{˜4œ<×3Ò3°E´LÀ%Ä,ÑOÔOÐQVÔQ[Ñ\Ô\Ð\r   c                ó„  — |                       ||dddd¦  «        \  }}|j        j        }|j        j        }|                     ¦   «         ræ|                     ¦   «         rÒ|                      ||¦  «        }|                      ||¦  «        }|                      ||¦  «        }|                     ¦   «         r>|                      | j         	                    |j
        |j
        ¦  «        |j        ¦  «        S |                      | j                             |j
        |j
        ¦  «        |j        ¦  «        S t          d|› �¦  «        ‚rÆ   )rš   rg   rŠ   rY   rF   rŽ   Úis_int_signedr&   r'   Úcreate_sdivr®   Úcreate_udivrA   )r   r«   r¬   r¶   r·   Úret_tys         r   ÚfloordivzTritonSemantic.floordiv=  s%  € Ø×8Ò8¸ÀÀuÈeÐUYÐ[_Ñ`Ô`‰ˆˆuØœ*Ô+ˆØœ*Ô+ˆØ×!Ò!Ñ#Ô#ð 	e¨×(>Ò(>Ñ(@Ô(@ð 	eØ×.Ò.¨ÀÑPÔPˆFØ—I’I˜e VÑ,Ô,ˆEØ—I’I˜e VÑ,Ô,ˆEØ×#Ò#Ñ%Ô%ð eØ—{’{ 4¤<×#;Ò#;¸E¼LÈ%Ì,Ñ#WÔ#WÐY^ÔYcÑdÔdÐdà—{’{ 4¤<×#;Ò#;¸E¼LÈ%Ì,Ñ#WÔ#WÐY^ÔYcÑdÔdÐdÝÐ<¨?Ð<Ð<Ñ=Ô=Ð=r   Úieee_roundingc                óª  — |j         j        }|j         j        }|                     ¦   «         r|                     ¦   «         st          d¦  «        ‚|                      ||dddd¦  «        \  }}|r&| j                             |j        |j        ¦  «        }n%| j                             |j        |j        ¦  «        }|  	                    ||j         ¦  «        S )Nz4both operands of fdiv must have floating scalar typeFT)
rg   rŠ   r�   rA   rš   r'   Úcreate_precise_divfr®   rÈ   r&   )r   r«   r¬   rÐ   r¶   r·   r¥   s          r   ÚfdivzTritonSemantic.fdivK  sÈ   € Øœ*Ô+ˆØœ*Ô+ˆØ×*Ò*Ñ,Ô,ð 	T°O×4OÒ4OÑ4QÔ4Qð 	TÝÐRÑSÔSÐSØ×8Ò8¸ÀÀuÈeÐUZÐ\`ÑaÔa‰ˆˆuØð 	GØ”,×2Ò2°5´<ÀÄÑNÔNˆCˆCà”,×*Ò*¨5¬<¸¼ÑFÔFˆCØ�{Š{˜3 ¤
Ñ+Ô+Ð+r   c                ó  — |                       ||dddd¦  «        \  }}|j        j        }|j        j        }|                     ¦   «         r>|                      | j                             |j        |j        ¦  «        |j        ¦  «        S |                     ¦   «         rß|j	        |j	        k    r?t          d|                     ¦   «         z   dz   |                     ¦   «         z   dz   ¦  «        ‚|                     ¦   «         r>|                      | j                             |j        |j        ¦  «        |j        ¦  «        S |                      | j                             |j        |j        ¦  «        |j        ¦  «        S t          d|› �¦  «        ‚)NFTzCannot mod z by rM   rL   )rš   rg   rŠ   r�   r&   r'   Úcreate_fremr®   rY   r=   rA   r   rË   Úcreate_sremÚcreate_urem)r   r«   r¬   rZ   r·   s        r   ÚmodzTritonSemantic.modW  sx  € Ø×8Ò8¸ÀÀuÈeÐUYÐ[_Ñ`Ô`‰ˆˆuØ”JÔ%ˆ	Øœ*Ô+ˆà× Ò Ñ"Ô"ð 	eØ—;’;˜tœ|×7Ò7¸¼ÀeÄlÑSÔSÐUZÔU_Ñ`Ô`Ð`à×ÒÑÔð 	eØÔ'¨?Ô+IÒIÐIÝ °	×0BÒ0BÑ0DÔ0DÑ DÀvÑ MÐP_×PhÒPhÑPjÔPjÑ jð nsñ !sñ tô tð tð ×&Ò&Ñ(Ô(ð eØ—{’{ 4¤<×#;Ò#;¸E¼LÈ%Ì,Ñ#WÔ#WÐY^ÔYcÑdÔdÐdà—{’{ 4¤<×#;Ò#;¸E¼LÈ%Ì,Ñ#WÔ#WÐY^ÔYcÑdÔdÐdÝÐ6¨9Ð6Ð6Ñ7Ô7Ð7r   rh   ÚyÚpropagate_nanútl.PropagateNanc                óF  — |                       ||¦  «        \  }}|j        }|                     ¦   «         r¸|t          j        j        k    r>|                      | j                             |j	        |j	        ¦  «        |j
        ¦  «        S |t          j        j        k    r>|                      | j                             |j	        |j	        ¦  «        |j
        ¦  «        S t          d|› �¦  «        ‚|                     ¦   «         r>|                      | j                             |j	        |j	        ¦  «        |j
        ¦  «        S |                     ¦   «         r>|                      | j                             |j	        |j	        ¦  «        |j
        ¦  «        S t%          d|› �¦  «        ‚©NzUnexpected propagate_nan úUnexpected dtype )rš   r>   r�   r1   ÚPropagateNanÚALLr&   r'   Úcreate_minimumfr®   rg   ÚNONEÚcreate_minnumfr/   rË   Úcreate_minsir‹   Úcreate_minuirA   ©r   rh   rÙ   rÚ   r>   s        r   ÚminimumzTritonSemantic.minimumn  ó`  € Ø×0Ò0°°AÑ6Ô6‰ˆˆ1Ø”ˆØ×ÒÑÔð 	9Ø¥¤Ô 3Ò3Ð3Ø—{’{ 4¤<×#?Ò#?ÀÄÈ!Ì(Ñ#SÔ#SÐUVÔU[Ñ\Ô\Ð\Ø¥"¤/Ô"6Ò6Ð6Ø—{’{ 4¤<×#>Ò#>¸q¼xÈÌÑ#RÔ#RÐTUÔTZÑ[Ô[Ð[å Ð!L¸]Ð!LÐ!LÑMÔMÐMØ× Ò Ñ"Ô"ð 	9Ø—;’;˜tœ|×8Ò8¸¼À1Ä8ÑLÔLÈaÌfÑUÔUÐUØ×"Ò"Ñ$Ô$ð 	9Ø—;’;˜tœ|×8Ò8¸¼À1Ä8ÑLÔLÈaÌfÑUÔUÐUåÐ7°Ð7Ð7Ñ8Ô8Ð8r   c                óF  — |                       ||¦  «        \  }}|j        }|                     ¦   «         r¸|t          j        j        k    r>|                      | j                             |j	        |j	        ¦  «        |j
        ¦  «        S |t          j        j        k    r>|                      | j                             |j	        |j	        ¦  «        |j
        ¦  «        S t          d|› �¦  «        ‚|                     ¦   «         r>|                      | j                             |j	        |j	        ¦  «        |j
        ¦  «        S |                     ¦   «         r>|                      | j                             |j	        |j	        ¦  «        |j
        ¦  «        S t%          d|› �¦  «        ‚rÝ   )rš   r>   r�   r1   rß   rà   r&   r'   Úcreate_maximumfr®   rg   râ   Úcreate_maxnumfr/   rË   Úcreate_maxsir‹   Úcreate_maxuirA   ræ   s        r   ÚmaximumzTritonSemantic.maximum  rè   r   ÚminÚmaxc                ó~  — |                       ||¦  «        \  }}|                       ||¦  «        \  }}|                       ||¦  «        \  }}|j        }|                     ¦   «         rE|                      | j                             |j        |j        |j        |¦  «        |j        ¦  «        S t          d|› d�¦  «        ‚)NrÞ   z(. Only floating point clamp is supported)	rš   r>   r�   r&   r'   Úcreate_clampfr®   rg   rA   )r   rh   rï   rð   rÚ   r>   s         r   ÚclampzTritonSemantic.clamp�  sº   € Ø×4Ò4°S¸#Ñ>Ô>‰ˆˆSØ×2Ò2°1°cÑ:Ô:‰ˆˆ3Ø×2Ò2°1°cÑ:Ô:‰ˆˆ3à”ˆØ×ÒÑÔð 	aØ—;’;˜tœ|×9Ò9¸!¼(ÀCÄJÐPSÔPZÐ\iÑjÔjÐlmÔlrÑsÔsÐsåÐ_°Ð_Ð_Ð_Ñ`Ô`Ð`r   c                óx  — |                       ||¦  «        \  }}|j        j        }|j        j        }|                     ¦   «         r|                     ¦   «         st	          ||¦  «        ‚|                      ||¦  «        }||k    r|                      ||¦  «        }||k    r|                      ||¦  «        }||fS r)   )rš   rg   rŠ   rY   r   rF   rŽ   )r   r«   r¬   Úinput_sca_tyÚother_sca_tyr™   s         r   Úbitwise_op_type_checking_implz,TritonSemantic.bitwise_op_type_checking_implŸ  sÃ   € Ø×8Ò8¸ÀÑFÔF‰ˆˆuØ”zÔ(ˆØ”zÔ(ˆØ×"Ò"Ñ$Ô$ð 	H¨L×,?Ò,?Ñ,AÔ,Að 	HÝ+¨L¸,ÑGÔGÐGØ×.Ò.¨|¸\ÑJÔJˆ
Ø˜Ò%Ð%Ø—I’I˜e ZÑ0Ô0ˆEØ˜Ò%Ð%Ø—I’I˜e ZÑ0Ô0ˆEØ�eˆ|Ðr   c                ó°   — |                       ||¦  «        \  }}|                      | j                             |j        |j        ¦  «        |j        ¦  «        S r)   )r÷   r&   r'   Ú
create_andr®   rg   ©r   r«   r¬   s      r   r¡   zTritonSemantic.and_¬  óI   € Ø×9Ò9¸%ÀÑGÔG‰ˆˆuØ�{Š{˜4œ<×2Ò2°5´<ÀÄÑNÔNÐPUÔPZÑ[Ô[Ð[r   c                ó°   — |                       ||¦  «        \  }}|                      | j                             |j        |j        ¦  «        |j        ¦  «        S r)   )r÷   r&   r'   Ú	create_orr®   rg   rú   s      r   Úor_zTritonSemantic.or_°  sG   € Ø×9Ò9¸%ÀÑGÔG‰ˆˆuØ�{Š{˜4œ<×1Ò1°%´,ÀÄÑMÔMÈuÌzÑZÔZÐZr   c                ó°   — |                       ||¦  «        \  }}|                      | j                             |j        |j        ¦  «        |j        ¦  «        S r)   )r÷   r&   r'   Ú
create_xorr®   rg   rú   s      r   Úxor_zTritonSemantic.xor_´  rû   r   c                ó  — |j                              ¦   «         s |                      |t          j        ¦  «        }|j                              ¦   «         s |                      |t          j        ¦  «        }|                      ||¦  «        S r)   )rg   Úis_int1Úbitcastr1   rt   r¡   rú   s      r   Úlogical_andzTritonSemantic.logical_and¸  sl   € ØŒz×!Ò!Ñ#Ô#ð 	1Ø—L’L ­¬Ñ0Ô0ˆEØŒz×!Ò!Ñ#Ô#ð 	1Ø—L’L ­¬Ñ0Ô0ˆEØ�yŠy˜ Ñ&Ô&Ð&r   c                ó  — |j                              ¦   «         s |                      |t          j        ¦  «        }|j                              ¦   «         s |                      |t          j        ¦  «        }|                      ||¦  «        S r)   )rg   r  r  r1   rt   rþ   rú   s      r   Ú
logical_orzTritonSemantic.logical_or¿  sl   € ØŒz×!Ò!Ñ#Ô#ð 	1Ø—L’L ­¬Ñ0Ô0ˆEØŒz×!Ò!Ñ#Ô#ð 	1Ø—L’L ­¬Ñ0Ô0ˆEØ�xŠx˜˜uÑ%Ô%Ð%r   c                óž   — |j                              ¦   «         s |                      |t          j        ¦  «        }|                      |¦  «        S r)   )rg   r  r  r1   rt   Úinvert©r   r«   s     r   Únot_zTritonSemantic.not_Æ  s?   € ØŒz×!Ò!Ñ#Ô#ð 	1Ø—L’L ­¬Ñ0Ô0ˆEØ�{Š{˜5Ñ!Ô!Ð!r   c                ó°   — |                       ||¦  «        \  }}|                      | j                             |j        |j        ¦  «        |j        ¦  «        S r)   )r÷   r&   r'   Úcreate_lshrr®   rg   rú   s      r   ÚlshrzTritonSemantic.lshrË  óI   € Ø×9Ò9¸%ÀÑGÔG‰ˆˆuØ�{Š{˜4œ<×3Ò3°E´LÀ%Ä,ÑOÔOÐQVÔQ[Ñ\Ô\Ð\r   c                ó°   — |                       ||¦  «        \  }}|                      | j                             |j        |j        ¦  «        |j        ¦  «        S r)   )r÷   r&   r'   Úcreate_ashrr®   rg   rú   s      r   ÚashrzTritonSemantic.ashrÏ  r  r   c                ó°   — |                       ||¦  «        \  }}|                      | j                             |j        |j        ¦  «        |j        ¦  «        S r)   )r÷   r&   r'   Ú
create_shlr®   rg   rú   s      r   ÚshlzTritonSemantic.shlÓ  rû   r   c                ó   — |S r)   © r
  s     r   ÚpluszTritonSemantic.plusÛ  s   € Øˆr   c                óâ  — |j         j        }|                     ¦   «         r't          d|                     ¦   «         z   dz   ¦  «        ‚|                     ¦   «         r8|                      | j                             |j	        ¦  «        |j         ¦  «        S |                      | j         
                    |                     | j        ¦  «        ¦  «        |¦  «        }|                      ||d¦  «        S )Nz$wrong type argument to unary minus (ú)T)rg   rŠ   r€   r/   r   r�   r&   r'   Úcreate_fnegr®   Úget_null_valuer°   r½   )r   r«   rõ   Ú_0s       r   r»   zTritonSemantic.minusÞ  sÎ   € Ø”zÔ(ˆØ×ÒÑ Ô ð 	eÝÐCÀl×F[ÒF[ÑF]ÔF]Ñ]Ð`cÑcÑdÔdÐdØ×#Ò#Ñ%Ô%ð 	SØ—;’;˜tœ|×7Ò7¸¼ÑEÔEÀuÄzÑRÔRÐRØ�[Š[˜œ×4Ò4°\×5GÒ5GÈÌÑ5UÔ5UÑVÔVÐXdÑeÔeˆØ�xŠx˜˜E 4Ñ(Ô(Ð(r   c                óp  — |j         j        }|                     ¦   «         s|                     ¦   «         r't	          d|                     ¦   «         z   dz   ¦  «        ‚|                      | j                             | 	                    | j        ¦  «        ¦  «        |¦  «        }|  
                    ||¦  «        S )Nz%wrong type argument to unary invert (r  )rg   rŠ   r€   r�   r/   r   r&   r'   Úget_all_ones_valuer°   r  )r   r«   rõ   Ú_1s       r   r	  zTritonSemantic.invertç  s£   € Ø”zÔ(ˆØ×ÒÑ Ô ð 	f L×$<Ò$<Ñ$>Ô$>ð 	fÝÐDÀ|×G\ÒG\ÑG^ÔG^Ñ^ÐadÑdÑeÔeÐeØ�[Š[˜œ×8Ò8¸×9KÒ9KÈDÌLÑ9YÔ9YÑZÔZÐ\hÑiÔiˆØ�yŠy˜ Ñ#Ô#Ð#r   Úvútl.block_typec                óJ   — |j                              t          j        ¦  «        S r)   )rg   r¯   r1   rt   )r   r!  s     r   Ú
_bool_likezTritonSemantic._bool_likeò  s   € ØŒv×%Ò%¥b¤gÑ.Ô.Ð.r   c                ó°  — |                       ||¦  «        \  }}|j        j        }|                     ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S | 	                    ¦   «         r¬| 
                    ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S |                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S t          d|› �¦  «        ‚rÁ   )rš   rg   rŠ   r�   r&   r'   Úcreate_fcmpOGTr®   r$  rY   rË   Úcreate_icmpSGTÚcreate_icmpUGTrA   ©r   r«   r¬   rZ   s       r   Úgreater_thanzTritonSemantic.greater_thanõ  ó)  € Ø×8Ò8¸ÀÑFÔF‰ˆˆuØ”JÔ%ˆ	à× Ò Ñ"Ô"ð 	tØ—;’;˜tœ|×:Ò:¸5¼<ÈÌÑVÔVÐX\×XgÒXgÐhmÑXnÔXnÑoÔoÐoà×ÒÑÔð 	tØ×&Ò&Ñ(Ô(ð tØ—{’{ 4¤<×#>Ò#>¸u¼|ÈUÌ\Ñ#ZÔ#ZÐ\`×\kÒ\kÐlqÑ\rÔ\rÑsÔsÐsà—{’{ 4¤<×#>Ò#>¸u¼|ÈUÌ\Ñ#ZÔ#ZÐ\`×\kÒ\kÐlqÑ\rÔ\rÑsÔsÐsÝÐ6¨9Ð6Ð6Ñ7Ô7Ð7r   c                ó°  — |                       ||¦  «        \  }}|j        j        }|                     ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S | 	                    ¦   «         r¬| 
                    ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S |                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S t          d|› �¦  «        ‚rÁ   )rš   rg   rŠ   r�   r&   r'   Úcreate_fcmpOGEr®   r$  rY   rË   Úcreate_icmpSGEÚcreate_icmpUGErA   r)  s       r   r£   zTritonSemantic.greater_equal  r+  r   c                ó°  — |                       ||¦  «        \  }}|j        j        }|                     ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S | 	                    ¦   «         r¬| 
                    ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S |                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S t          d|› �¦  «        ‚rÁ   )rš   rg   rŠ   r�   r&   r'   Úcreate_fcmpOLTr®   r$  rY   rË   Úcreate_icmpSLTÚcreate_icmpULTrA   r)  s       r   Ú	less_thanzTritonSemantic.less_than  r+  r   c                ó°  — |                       ||¦  «        \  }}|j        j        }|                     ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S | 	                    ¦   «         r¬| 
                    ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S |                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S t          d|› �¦  «        ‚rÁ   )rš   rg   rŠ   r�   r&   r'   Úcreate_fcmpOLEr®   r$  rY   rË   Úcreate_icmpSLEÚcreate_icmpULErA   r)  s       r   r¢   zTritonSemantic.less_equal  r+  r   c                óð  — |                       ||¦  «        \  }}|j        j        }|                     ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S | 	                    ¦   «         rL|                      | j         
                    |j        |j        ¦  «        |                      |¦  «        ¦  «        S t          d|› �¦  «        ‚rÁ   )rš   rg   rŠ   r�   r&   r'   Úcreate_fcmpOEQr®   r$  rY   Úcreate_icmpEQrA   r)  s       r   ÚequalzTritonSemantic.equal-  óÚ   € Ø×8Ò8¸ÀÑFÔF‰ˆˆuØ”JÔ%ˆ	à× Ò Ñ"Ô"ð 	oØ—;’;˜tœ|×:Ò:¸5¼<ÈÌÑVÔVÐX\×XgÒXgÐhmÑXnÔXnÑoÔoÐoà×ÒÑÔð 	oØ—;’;˜tœ|×9Ò9¸%¼,ÈÌÑUÔUÐW[×WfÒWfÐglÑWmÔWmÑnÔnÐnÝÐ6¨9Ð6Ð6Ñ7Ô7Ð7r   c                óð  — |                       ||¦  «        \  }}|j        j        }|                     ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S | 	                    ¦   «         rL|                      | j         
                    |j        |j        ¦  «        |                      |¦  «        ¦  «        S t          d|› �¦  «        ‚rÁ   )rš   rg   rŠ   r�   r&   r'   Úcreate_fcmpUNEr®   r$  rY   Úcreate_icmpNErA   r)  s       r   Ú	not_equalzTritonSemantic.not_equal8  r=  r   N)rÎ   ÚstartÚendrÎ   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k    rt          d¦  «        ‚|g}|€t	          j        t          j        |¦  «        }|                     | 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)rb   r+   r/   rH   r1   Ú
block_typer2   r°   r'   r&   Úcreate_make_range)	r   rB  rC  rÎ   Úis_start_int64Úis_end_int64ÚrangeÚshapeÚ	ret_ty_irs	            r   ÚarangezTritonSemantic.arangeG  s  € Ý˜%¥Ñ%Ô%ð 	P­Z¸½SÑ-AÔ-Að 	PÝÐNÑOÔOÐOÝ˜e r™kÑ*Ô*ˆÝ˜C 2™I‘”ˆØð 	9˜\ð 	9ÝÐ7Ñ8Ô8Ð8Ø�%Š<ˆ<ÝÐ\Ñ]Ô]Ð]Ø�e‘ˆØ�U˜Q‘YÑ AÒ%Ð%ÝÐBÑCÔCÐCØ�ˆØˆ>Ý”]¥2¤8¨UÑ3Ô3ˆFØ—L’L ¤Ñ.Ô.ˆ	Ø�{Š{˜4œ<×9Ò9¸)ÀUÈCÑPÔPÐRXÑYÔYÐYr   r>   c                óÖ  — |€t          d¦  «        ‚|dk    r3| j                             |                     | j        ¦  «        ¦  «        }nŠ|                     ¦   «         rN| j                             |¦  «        }| j                             ||                     | j        ¦  «        ¦  «        }n(t          | j        d|j        › �¦  «        } ||¦  «        }|  	                    ||¦  «        S )Nz2dtype must be specified when value is not a tensorr   Úget_)
r/   r'   r  r°   rX   Úget_fp32Úcreate_fp_truncÚgetattrÚnamer&   )r   rO   r>   Úget_value_fns       r   rf   zTritonSemantic.scalar_constantY  sÏ   € àˆ=ÝÐQÑRÔRÐRØ�AŠ:ˆ:Ø”L×/Ò/°·²¸D¼LÑ0IÔ0IÑJÔJˆEˆEØ�\Š\‰^Œ^ð 	(Ø”L×)Ò)¨%Ñ0Ô0ˆEØ”L×0Ò0°¸¿ºÀDÄLÑ8QÔ8QÑRÔRˆEˆEå" 4¤<Ð1D¸¼
Ð1DÐ1DÑEÔEˆLØ �L Ñ'Ô'ˆEØ�{Š{˜5 %Ñ(Ô(Ð(r   c                óÂ   — t          |t          j        ¦  «        r0|j        j        dk    s
J d¦   «         ‚|                      ||¦  «        S |                      ||¦  «        S )Nr   zonly accepts size-1 tensor)rb   r1   r&   ÚnumelrO   rŽ   rf   )r   rO   r>   s      r   Úmake_scalarzTritonSemantic.make_scalarg  s_   € Ý�e�RœYÑ'Ô'ð 	+Ø”;Ô$¨Ò)Ð)Ð)Ð+GÑ)Ô)Ð)Ø—9’9˜U EÑ*Ô*Ð*à×#Ò# E¨5Ñ1Ô1Ð1r   rK  ú	List[int]c                óV   — |                       |                      ||¦  «        |¦  «        S r)   )ÚsplatrW  )r   rK  rO   r>   s       r   ÚfullzTritonSemantic.fulln  s&   € Ø�zŠz˜$×*Ò*¨5°%Ñ8Ô8¸%Ñ@Ô@Ð@r   rO   c                ó>  — |j                              ¦   «         r
J d¦   «         ‚t          |¦  «        dk    r|S t          j        |j        |¦  «        }|                      | j                             | 	                    | j        ¦  «        |j
        ¦  «        |¦  «        S )NzCannot splat a block tensorr   )rg   Úis_blockÚlenr1   rF  r>   r&   r'   Úcreate_splatr°   r®   )r   rO   rK  rÎ   s       r   rZ  zTritonSemantic.splatu  s…   € Ø”:×&Ò&Ñ(Ô(ÐGÐGÐ*GÑGÔGÐGÝˆu‰:Œ:˜Š?ˆ?ØˆLÝ”˜uœ{¨EÑ2Ô2ˆØ�{Š{˜4œ<×4Ò4°V·\²\À$Ä,Ñ5OÔ5OÐQVÔQ]Ñ^Ô^Ð`fÑgÔgÐgr   c                ór   — |                       | j                             |j        ¦  «        |j        ¦  «        S r)   )r&   r'   Úcreate_unsplatr®   r>   )r   rO   s     r   ÚunsplatzTritonSemantic.unsplat|  s*   € Ø�{Š{˜4œ<×6Ò6°u´|ÑDÔDÀeÄkÑRÔRÐRr   Ú	dst_shapeÚcan_reorderc                ó   — d}|D ]}||z  }Œ|j         j        |k    rt          d¦  «        ‚t          j        |j         j        |¦  «        }|                      | j                             |j	        ||¦  «        |¦  «        S )Nr   z:reshape() cannot change total number of elements in tensor)
rg   rV  r/   r1   rF  rŠ   r&   r'   Úcreate_reshaper®   )r   r«   rc  rd  rV  ÚsrÎ   s          r   ÚreshapezTritonSemantic.reshape  s‚   € ØˆØð 	ð 	ˆAØ�Q‰JˆEˆEØŒ:Ô˜uÒ$Ð$ÝÐYÑZÔZÐZÝ”˜uœzÔ0°)Ñ<Ô<ˆØ�{Š{˜4œ<×6Ò6°u´|ÀYÐP[Ñ\Ô\Ð^dÑeÔeÐer   c                óV  — d„ |j         D ¦   «         }|                     |d¦  «         |j                             ¦   «         s|                      ||¬¦  «        S t          j        |j        j        |¦  «        }|                      | j	         
                    |j        |¦  «        |¦  «        S )Nc                ó6   — g | ]}t          j        |¦  «        ‘ŒS r  ©r1   Ú_unwrap_if_constexpr)Ú.0rh   s     r   ú
<listcomp>z.TritonSemantic.expand_dims.<locals>.<listcomp>‰  s#   € ÐEÐEÐE°A•RÔ,¨QÑ/Ô/ÐEÐEÐEr   r   ©rK  )rK  Úinsertrg   r]  rZ  r1   rF  rŠ   r&   r'   Úcreate_expand_dimsr®   )r   r«   r*   rc  rÎ   s        r   Úexpand_dimszTritonSemantic.expand_dimsˆ  s˜   € ØEÐE¸¼ÐEÑEÔEˆ	Ø×Ò˜˜qÑ!Ô!Ð!àŒz×"Ò"Ñ$Ô$ð 	6Ø—:’:˜e¨9�:Ñ5Ô5Ð5å”˜uœzÔ0°)Ñ<Ô<ˆØ�{Š{˜4œ<×:Ò:¸5¼<ÈÑNÔNÐPVÑWÔWÐWr   c                ón  — |s
J d¦   «         ‚t          |j        ¦  «        dk    s J dt          |j        ¦  «        › d�¦   «         ‚t          j        |j        j        |j        d         |j        d         z   g¦  «        }|                      | j                             |j	        |j	        ¦  «        |¦  «        S )Nz;current implementation of `cat` always may reorder elementsr   zexpected 1D input for cat, got ÚDr   )
r^  rK  r1   rF  rg   rŠ   r&   r'   Ú
create_catr®   )r   rƒ   r…   rd  Úret_types        r   ÚcatzTritonSemantic.cat’  sŸ   € ØÐYÐYÐYÑYÔYÐYÝ�3”9‰~Œ~ Ò"Ð"Ð"Ð$WÅcÈ#Ì)ÁnÄnÐ$WÐ$WÐ$WÑ"Ô"Ð"Ý”= ¤¤°3´9¸Q´<À#Ä)ÈAÄ,Ñ3NÐ2OÑPÔPˆØ�{Š{˜4œ<×2Ò2°3´:¸s¼zÑJÔJÈHÑUÔUÐUr   ÚaÚbc                ó  — |                       ||¦  «        \  }}|j        g k    }|r,|                      |d¦  «        }|                      |d¦  «        }t          |j        d         t          j        ¦  «        rt	          j        d¦  «        }nd}|j        |gz   }t	          j        |j        j        |¦  «        }|  	                    | j
                             |j        |j        ¦  «        |¦  «        }|r|                      |dgd¬¦  «        }|S )Nr   éÿÿÿÿr   F©rd  )r�   rK  rr  rb   r1   rc   rF  rg   rŠ   r&   r'   Úcreate_joinr®   rh  )r   rx  ry  Ú
was_rank_1ÚtwoÚ	new_shaperv  r¥   s           r   ÚjoinzTritonSemantic.join˜  sù   € Ø×(Ò(¨¨AÑ.Ô.‰ˆˆ1ð ”W ’]ˆ
Øð 	'Ø× Ò   AÑ&Ô&ˆAØ× Ò   AÑ&Ô&ˆAå�a”g˜b”k¥2¤<Ñ0Ô0ð 	Ý”,˜q‘/”/ˆCˆCàˆCØ”G˜s˜e‘Oˆ	å”= ¤¤°	Ñ:Ô:ˆØ�kŠk˜$œ,×2Ò2°1´8¸Q¼XÑFÔFÈÑQÔQˆàð 	<Ø—,’,˜s Q C°U�,Ñ;Ô;ˆCàˆ
r   c                óØ  — t          |j        ¦  «        dk    s
J d¦   «         ‚t          j        |j        d         ¦  «        dk    s*J dt          j        |j        d         ¦  «        › �¦   «         ‚|j        d d…         }t          j        |j        j        |¦  «        }| j                             |j	        ¦  «        \  }}|  
                    ||¦  «        |  
                    ||¦  «        fS )Nr   z"split requires a non-scalar tensorr{  r   z/expected last dimension to be 2 for split, got )r^  rK  r1   rl  rF  rg   rŠ   r'   Úcreate_splitr®   r&   )r   rx  r€  rv  ÚoutLHSÚoutRHSs         r   ÚsplitzTritonSemantic.split°  sÛ   € Ý�A”G‘”˜qÒ Ð Ð Ð#GÑ Ô Ð ÝÔ'¨¬°¬Ñ4Ô4¸Ò9Ð9Ð9Ød½bÔ>UÐVWÔV]Ð^`ÔVaÑ>bÔ>bÐdÐdñ :Ô9Ð9ð ”G˜C˜R˜C”Lˆ	Ý”= ¤¤°	Ñ:Ô:ˆØœ×2Ò2°1´8Ñ<Ô<‰ˆ�à�KŠK˜ Ñ)Ô)Ø�KŠK˜ Ñ)Ô)ð
ð 	
r   Údimsú
Tuple[int]c                ó   ‡— t          ‰j        ¦  «        t          |¦  «        k    r4t          dt          ‰j        ¦  «        › dt          |¦  «        › �¦  «        ‚t          d„ |D ¦   «         ¦  «        t	          t          t          |¦  «        ¦  «        ¦  «        k    rt          d|› �¦  «        ‚t          j        ‰j        j	        ˆfd„|D ¦   «         ¦  «        }|  
                    | j                             ‰j        |¦  «        |¦  «        S )Nz;permute dims must have the same length as input shape, got r   c              3  ó>   K  — | ]}t          j        |¦  «        V — Œd S r)   rk  )rm  Úds     r   ú	<genexpr>z)TritonSemantic.permute.<locals>.<genexpr>Á  s-   è è € Ð;Ð;°•"Ô)¨!Ñ,Ô,Ð;Ð;Ð;Ð;Ð;Ð;r   z?permute dims must be a permutation of 0, 1, ..., n-1, but were c                ó*   •— g | ]}‰j         |         ‘ŒS r  ro  )rm  r‹  r«   s     €r   rn  z*TritonSemantic.permute.<locals>.<listcomp>Ä  s   ø€ Ð4RÐ4RÐ4RÈ°U´[À´^Ð4RÐ4RÐ4Rr   )r^  rK  r/   ÚsortedÚlistrJ  r1   rF  rg   rŠ   r&   r'   Úcreate_transr®   )r   r«   r‡  rv  s    `  r   ÚpermutezTritonSemantic.permute½  sÿ   ø€ ÝˆuŒ{ÑÔ�s 4™yœyÒ(Ð(ÝØpÍcÐRWÔR]ÑN^ÔN^ÐpÐpÕehÐimÑenÔenÐpÐpñrô rð råÐ;Ð;°dÐ;Ñ;Ô;Ñ;Ô;½tÅEÍ#ÈdÉ)Ì)ÑDTÔDTÑ?UÔ?UÒUÐUÝÐeÐ_cÐeÐeÑfÔfÐfå”= ¤Ô!2Ð4RÐ4RÐ4RÐ4RÈTÐ4RÑ4RÔ4RÑSÔSˆØ�{Š{˜4œ<×4Ò4°U´\À4ÑHÔHÈ(ÑSÔSÐSr   c                óH  — |j                              ¦   «         s|                      ||¦  «        S |j                              ¦   «         }t	          |¦  «        t	          |¦  «        k    rt          d|› d|› �¦  «        ‚||k    r|S t          |¦  «        D ];\  }}||         |k    r*|dk    r$t          d||         › d|› d|› d|› d|› �
¦  «        ‚Œ<t          j        |j         j	        |¦  «        }|  
                    | j                             |j        |¦  «        |¦  «        S )Nz!Cannot broadcast, rank mismatch: ú, r   z3Cannot broadcast, the expanded size of the tensor (z ) must match the existing size (z) at non-singleton dimension ú: )rg   r]  rZ  Úget_block_shapesr^  r/   Ú	enumerater1   rF  rŠ   r&   r'   Úcreate_broadcastr®   )r   r«   rK  Ú	src_shapeÚiÚitemrÎ   s          r   Úbroadcast_impl_shapez#TritonSemantic.broadcast_impl_shapeÇ  s[  € ØŒz×"Ò"Ñ$Ô$ð 	,Ø—:’:˜e UÑ+Ô+Ð+Ø”J×/Ò/Ñ1Ô1ˆ	Ýˆy‰>Œ>�S ™ZœZÒ'Ð'ÝÐUÀÐUÐUÈeÐUÐUÑVÔVÐVØ�IÒÐØˆLÝ  Ñ+Ô+ð 	@ð 	@‰GˆAˆtØ�QŒx˜4ÒÐ D¨A¢I IÝ ð "?ÐW\Ð]^ÔW_ð "?ð "?ØCGð"?ð "?à%&ð"?ð "?à*3ð"?ð "?à7<ð"?ð "?ñ @ô @ð @øõ ”˜uœzÔ0°%Ñ8Ô8ˆØ�{Š{˜4œ<×8Ò8¸¼ÀuÑMÔMÈvÑVÔVÐVr   c           	     ó&	  — |j         }|j         }|                     ¦   «         r||                     ¦   «         sh|                     |j        ¦  «        }|                      | j                             |                     | j        ¦  «        |j        ¦  «        |¦  «        }�nð|                     ¦   «         s||                     ¦   «         rh|                     |j        ¦  «        }|                      | j                             |                     | j        ¦  «        |j        ¦  «        |¦  «        }�n`|                     ¦   «         �rK|                     ¦   «         �r6| 	                    ¦   «         }| 	                    ¦   «         }t          |¦  «        t          |¦  «        k     ržt          t          |¦  «        t          |¦  «        ¦  «        D ]r}|                      | j                             |j        d¦  «        t          j        |j        dg|j        z   ¦  «        ¦  «        }|j         }| 	                    ¦   «         }Œsn½t          |¦  «        t          |¦  «        k     r�t          t          |¦  «        t          |¦  «        ¦  «        D ]r}|                      | j                             |j        d¦  «        t          j        |j        dg|j        z   ¦  «        ¦  «        }|j         }| 	                    ¦   «         }Œst          |¦  «        t          |¦  «        k    s*J dt          |¦  «        › dt          |¦  «        › �¦   «         ‚g }t!          |¦  «        D ]�\  }	}
||	         }|
dk    r|                     |¦  «         Œ)|dk    s||
k    r|                     |
¦  «         ŒKt%          dt'          |	¦  «        z   dz   t'          |
¦  «        z   dz   t'          |¦  «        z   ¦  «        ‚||k    rNt          j        |j        |¦  «        }|                      | j                             |j        |¦  «        |¦  «        }||k    rNt          j        |j        |¦  «        }|                      | j                             |j        |¦  «        |¦  «        }||fS )Nr   r   z2expected tensors of equal rank for broadcast, got r   z?Cannot make_shape_compatible: incompatible dimensions at index r”  )rg   r]  r¯   rŠ   r&   r'   r_  r°   r®   r•  r^  rJ  rq  r1   rF  Úvaluesr–  Úappendr/   Ústrr—  )r   rƒ   r…   Úlhs_tyÚrhs_tyÚ	lhs_shapeÚ	rhs_shapeÚ_Ú	ret_shaper™  ÚleftÚrightrÎ   s                r   r�   z#TritonSemantic.broadcast_impl_value×  s/  € Ø”ˆØ”ˆð �?Š?ÑÔð ,	` V§_¢_Ñ%6Ô%6ð ,	`Ø×+Ò+¨F¬MÑ:Ô:ˆFØ—+’+˜dœl×7Ò7¸¿ºÀTÄ\Ñ8RÔ8RÐTWÔT^Ñ_Ô_ÐagÑhÔhˆC‰Cà—’Ñ"Ô"ð (	` v§¢Ñ'8Ô'8ð (	`Ø×+Ò+¨F¬MÑ:Ô:ˆFØ—+’+˜dœl×7Ò7¸¿ºÀTÄ\Ñ8RÔ8RÐTWÔT^Ñ_Ô_ÐagÑhÔhˆC‰Cà�_Š_ÑÔñ $	` 6§?¢?Ñ#4Ô#4ñ $	`Ø×/Ò/Ñ1Ô1ˆIØ×/Ò/Ñ1Ô1ˆIå�9‰~Œ~¥ I¡¤Ò.Ð.å�s 9™~œ~­s°9©~¬~Ñ>Ô>ð :ð :�AØŸ+š+ d¤l×&EÒ&EÀcÄjÐRSÑ&TÔ&TÝ&(¤m°F´MÀAÀ3ÈÔIYÑCYÑ&ZÔ&Zñ\ô \�Cà œX�FØ &× 7Ò 7Ñ 9Ô 9�I�Ið	:õ
 �Y‘”¥# i¡.¤.Ò0Ð0å�s 9™~œ~­s°9©~¬~Ñ>Ô>ð :ð :�AØŸ+š+ d¤l×&EÒ&EÀcÄjÐRSÑ&TÔ&TÝ&(¤m°F´MÀAÀ3ÈÔIYÑCYÑ&ZÔ&Zñ\ô \�Cà œX�FØ &× 7Ò 7Ñ 9Ô 9�I�IÝ�y‘>”>¥S¨¡^¤^Ò3Ð3Ð3ØjÅSÈÁ^Ä^ÐjÐjÕZ]Ð^gÑZhÔZhÐjÐjñ 4Ô3Ð3ð ˆIÝ$ YÑ/Ô/ð eð e‘��4Ø! !œ�Ø˜1’9�9Ø×$Ò$ UÑ+Ô+Ð+Ð+Ø˜q’j�j e¨t¢m mØ×$Ò$ TÑ*Ô*Ð*Ð*å$ð &1Ý36°q±6´6ñ&:Ø<@ñ&AÝCFÀtÁ9Ä9ñ&MØOVñ&WÝY\Ð]bÑYcÔYcñ&dñ eô eð eà˜IÒ%Ð%Ýœ v¤}°iÑ@Ô@�Ø—k’k $¤,×"?Ò"?ÀÄ
ÈIÑ"VÔ"VÐX^Ñ_Ô_�Ø˜IÒ%Ð%Ýœ v¤}°iÑ@Ô@�Ø—k’k $¤,×"?Ò"?ÀÄ
ÈIÑ"VÔ"VÐX^Ñ_Ô_�à�Cˆxˆr   Úrounding_modeúOptional[str]c                óŒ   — |€d S |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/   )r   r¨  s     r   Ú_str_to_rounding_modez$TritonSemantic._str_to_rounding_mode  sU   € ØÐ Ø�4Ø˜FÒ"Ð"ÝÔ#Ô(Ð(Ø˜EÒ!Ð!ÝÔ#Ô'Ð'ÝÐr°=ÐrÐrÐrÑsÔsÐsr   Údst_tyc                ó8  — |j         }|                     ¦   «         r|                     |j        ¦  «        }||k    r|S |j        }|j        }|                     ¦   «         s|                     ¦   «         r|                      ||¦  «        S |j        }|j        }||k    r2t          dt          |¦  «        z   dz   t          |¦  «        z   ¦  «        ‚|  	                    | j
                             |j        |                     | j
        ¦  «        ¦  «        |¦  «        S )Nz!Cannot bitcast data-type of size z to data-type of size )rg   r]  r¯   rŠ   r€   rŽ   Úprimitive_bitwidthr/   rŸ  r&   r'   Úcreate_bitcastr®   r°   )r   r«   r±  Úsrc_tyÚ
src_sca_tyÚ
dst_sca_tyÚsrc_bitsÚdst_bitss           r   r  zTritonSemantic.bitcast  s  € Ø”ˆØ�?Š?ÑÔð 	;Ø×+Ò+¨F¬MÑ:Ô:ˆFØ�VÒÐØˆLØ”]ˆ
Ø”]ˆ
Ø×ÒÑÔð 	, *×"3Ò"3Ñ"5Ô"5ð 	,Ø—9’9˜U FÑ+Ô+Ð+àÔ0ˆØÔ0ˆØ�xÒÐÝÐ@Å3ÀxÁ=Ä=ÑPð T2ñ 2Ý47¸±M´MñBñ Cô Cð Cà�{Š{˜4œ<×6Ò6°u´|ÀVÇ\Â\ÐRVÔR^ÑE_ÔE_Ñ`Ô`ÐbhÑiÔiÐir   Úfp_downcast_roundingc                óê  — |j         }|j        }|j        }||k    r|S |                     ¦   «         r|                     |¦  «        }|                      |¦  «        }d}|                     ¦   «         rP|                     ¦   «         r<|j        |j        k     r,|€t          j        j	        }nL|t          j        j	        k    rd}n4|�2t          dt          |¦  «        z   dz   t          |¦  «        z   ¦  «        ‚|                     ¦   «         s|                     ¦   «         rI| j        j                             d¦  «        	 €
J d¦   «         ‚ | j        j        d         |||| ¬¦  «        S |                     ¦   «         r|                     ¦   «         s*|                     ¦   «         r|                     ¦   «         s|rM|                      | j                             |j        |                     | j        ¦  «        |¦  «        |¦  «        S |                     ¦   «         r|                     ¦   «         r(|                     ¦   «         rH|                     ¦   «         s4|                      |                      |t2          j        ¦  «        |¦  «        S |                     ¦   «         o#|                     ¦   «         o|j        |j        k    }|rL|                      | j                             |j        |                     | j        ¦  «        ¦  «        |¦  «        S |                     ¦   «         o#|                     ¦   «         o|j        |j        k     }	|	rL|                      | j                             |j        |                     | j        ¦  «        ¦  «        |¦  «        S |                     ¦   «         �r'|                     ¦   «         �r|j        |j        k    s|j        |j        k    rò|                      ¦   «         o| !                    ¦   «          }
| !                    ¦   «         rh|j"                             | j        ¦  «        }|                      | j         #                    |¦  «        |j"        ¦  «        }|  $                    ||¦  «        S |                      | j         %                    |j        |                     | j        ¦  «        |
¦  «        |¦  «        S | &                    ¦   «         �r=|                     ¦   «         �r(| !                    ¦   «         rh|j"                             | j        ¦  «        }|                      | j         #                    |¦  «        |j"        ¦  «        }|  $                    ||¦  «        S |                      ¦   «         rL|                      | j         '                    |j        |                     | j        ¦  «        ¦  «        |¦  «        S |                      | j         (                    |j        |                     | j        ¦  «        ¦  «        |¦  «        S |                     ¦   «         rÔ| &                    ¦   «         rÀ| !                    ¦   «         s|                      ¦   «         sL|                      | j         )                    |j        |                     | j        ¦  «        ¦  «        |¦  «        S |                      | j         *                    |j        |                     | j        ¦  «        ¦  «        |¦  «        S | +                    ¦   «         rÝ|                     ¦   «         rÉ|j        }|dk    rL|                      | j         ,                    |j        |                     | j        ¦  «        ¦  «        |¦  «        S |d	k    rj|  $                    |                      |t2          j-        ¦  «        |                      | j         .                    d
¦  «        t2          j-        ¦  «        ¦  «        S |                     ¦   «         r`| +                    ¦   «         rL|                      | j         /                    |j        |                     | j        ¦  «        ¦  «        |¦  «        S | +                    ¦   «         r`| +                    ¦   «         rL|                      | j         0                    |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.)Ú	_semanticrž   r   r   zcannot cast z to )1rg   rŠ   r]  r¯   r°  r�   r³  r   r­  r®  r/   rŸ  Úis_fp8e4b15r'   Úcodegen_fnsÚgetrX   r&   Úcreate_fp_to_fpr®   r°   rV   rU   rW   rŽ   r1   rR   rQ  Úcreate_fp_extrY   r<   r=   rË   Úis_boolr>   r  rA  r±   Úis_standard_floatingÚcreate_fp_to_siÚcreate_fp_to_uiÚcreate_ui_to_fpÚcreate_si_to_fpr€   Úcreate_ptr_to_intrv   Ú	get_int64Úcreate_int_to_ptrr´  )r   r«   r±  rº  rµ  r¶  r·  Úuse_custom_roundingÚtruncate_fpÚext_fpÚsign_extendÚtyr  Úbitwidths                 r   rŽ   zTritonSemantic.cast+  sd  € Ø”ˆØ”]ˆ
Ø”]ˆ
Ø˜Ò#Ð#ØˆLØ�?Š?ÑÔð 	8Ø×+Ò+¨JÑ7Ô7ˆFð  $×9Ò9Ð:NÑOÔOÐØ#ÐØ×!Ò!Ñ#Ô#ð 	2¨
×(>Ò(>ñ )
ô )
ð 	2ØÔ+¨jÔ.KÒKÐKØ#Ð+ÅBÔDTÔDYÐ-AÐ-AØ%­Ô)9Ô)>Ò>Ð>ÐVZÐ@Søà#Ð/Ý ð ":Ý<?À
¹O¼Oñ"LØNiñ"jå!$ Z¡¤ñ"1ñ 2ô 2ð 2ð ×"Ò"Ñ$Ô$ð 	y¨
×(>Ò(>Ñ(@Ô(@ð 	yØ”<Ô+×/Ò/Ø&ñ(ô (Ø/3ð4ð 4Ø5gñ4ô 4ð 4àC�4”<Ô+Ð,BÔCÀEÈ6ÐSgÐswÐxÑxÔxÐxð ×ÒÑÔð 	v J×$:Ò$:Ñ$<Ô$<ð 	vØ×"Ò"Ñ$Ô$ð	vØ)3×):Ò):Ñ)<Ô)<ð	vàð	vð —;’;Ø”×,Ò,¨U¬\¸6¿<º<ÈÌÑ;UÔ;UÐWkÑlÔlÐntñvô vð vð ×ÒÑ Ô ð 	G¨×);Ò);Ñ)=Ô)=ð 	GØ×ÒÑ Ô ð	GØ)3×);Ò);Ñ)=Ô)=ð	Gà—9’9˜TŸYšY u­b¬jÑ9Ô9¸:ÑFÔFÐFð
 !×,Ò,Ñ.Ô.ð JØ×"Ò"Ñ$Ô$ðJàÔ)¨JÔ,IÒIð 	ð ð 	oØ—;’;˜tœ|×;Ò;¸E¼LÈ&Ï,Ê,ÐW[ÔWcÑJdÔJdÑeÔeÐgmÑnÔnÐnð ×'Ò'Ñ)Ô)ð JØ×"Ò"Ñ$Ô$ðJàÔ)¨JÔ,IÒIð 	ð ð 	mØ—;’;˜tœ|×9Ò9¸%¼,ÈÏÊÐUYÔUaÑHbÔHbÑcÔcÐekÑlÔlÐlð ×ÒÑÔñ 		+ :×#4Ò#4Ñ#6Ô#6ñ 		+ØÔ# zÔ'>Ò>Ð>À*ÔB[Ð_iÔ_xÒBxÐBxØ$×2Ò2Ñ4Ô4ÐQ¸Z×=OÒ=OÑ=QÔ=QÐ9QˆKØ×!Ò!Ñ#Ô#ð +Ø”[×&Ò& t¤|Ñ4Ô4�Ø—[’[ ¤×!<Ò!<¸RÑ!@Ô!@À%Ä+ÑNÔN�Ø—~’~ e¨RÑ0Ô0Ð0à—{’{ 4¤<×#?Ò#?ÀÄÈfÏlÊlÐ[_Ô[gÑNhÔNhÐjuÑ#vÔ#vØ#)ñ+ô +ð +ð ×*Ò*Ñ,Ô,ñ 	s°×1BÒ1BÑ1DÔ1Dñ 	sØ×!Ò!Ñ#Ô#ð sØ”[×&Ò& t¤|Ñ4Ô4�Ø—[’[ ¤×!<Ò!<¸RÑ!@Ô!@À%Ä+ÑNÔN�Ø—~’~ e¨RÑ0Ô0Ð0Ø×)Ò)Ñ+Ô+ð sØ—{’{ 4¤<×#?Ò#?ÀÄÈfÏlÊlÐ[_Ô[gÑNhÔNhÑ#iÔ#iÐkqÑrÔrÐrà—{’{ 4¤<×#?Ò#?ÀÄÈfÏlÊlÐ[_Ô[gÑNhÔNhÑ#iÔ#iÐkqÑrÔrÐrð ×ÒÑÔð 	s :×#BÒ#BÑ#DÔ#Dð 	sØ×!Ò!Ñ#Ô#ð s¨:×+CÒ+CÑ+EÔ+Eð sØ—{’{ 4¤<×#?Ò#?ÀÄÈfÏlÊlÐ[_Ô[gÑNhÔNhÑ#iÔ#iÐkqÑrÔrÐrà—{’{ 4¤<×#?Ò#?ÀÄÈfÏlÊlÐ[_Ô[gÑNhÔNhÑ#iÔ#iÐkqÑrÔrÐrð ×ÒÑÔð 	t :×#4Ò#4Ñ#6Ô#6ð 	tØ!Ô.ˆHØ˜2Š~ˆ~Ø—{’{ 4¤<×#AÒ#AÀ%Ä,ÐPV×P\ÒP\Ð]aÔ]iÑPjÔPjÑ#kÔ#kÐmsÑtÔtÐtØ˜1Š}ˆ}Ø—~’~ d§i¢i°µr´xÑ&@Ô&@À$Ç+Â+ÈdÌl×NdÒNdÐefÑNgÔNgÕikÔiqÑBrÔBrÑsÔsÐsð ×ÒÑÔð 	q :×#4Ò#4Ñ#6Ô#6ð 	qØ—;’;˜tœ|×=Ò=¸e¼lÈFÏLÊLÐY]ÔYeÑLfÔLfÑgÔgÐioÑpÔpÐpð ×ÒÑÔð 	n :×#4Ò#4Ñ#6Ô#6ð 	nØ—;’;˜tœ|×:Ò:¸5¼<ÈÏÊÐVZÔVbÑIcÔIcÑdÔdÐflÑmÔmÐmà8Ð8 UÐ8Ð8°Ð8Ð8Ñ8Ô8Ð8r   c                ó  — t           j        j        }|rs|dk    rt           j        j        }n[|dk    rt           j        j        }nC|dk    rt           j        j        }n+|dk    rt           j        j        }nt          d|› d�¦  «        ‚|S )Nz.caú.cgú.csz.cvúCache modifier ú not supported)r   ÚCACHE_MODIFIERrâ   ÚCAÚCGÚCSÚCVr/   ©r   Úcache_modifierÚcaches      r   Ú_str_to_load_cache_modifierz*TritonSemantic._str_to_load_cache_modifierš  ó•   € ÝÔ!Ô&ˆØð 
	SØ Ò&Ð&ÝÔ)Ô,��Ø 5Ò(Ð(ÝÔ)Ô,��Ø 5Ò(Ð(ÝÔ)Ô,��Ø 5Ò(Ð(ÝÔ)Ô,��å Ð!Q°>Ð!QÐ!QÐ!QÑRÔRÐRØˆr   c                ó  — t           j        j        }|rs|dk    rt           j        j        }n[|dk    rt           j        j        }nC|dk    rt           j        j        }n+|dk    rt           j        j        }nt          d|› d�¦  «        ‚|S )Nz.wbrÓ  rÔ  z.wtrÕ  rÖ  )r   r×  râ   ÚWBrÙ  rÚ  ÚWTr/   rÜ  s      r   Ú_str_to_store_cache_modifierz+TritonSemantic._str_to_store_cache_modifier©  rà  r   c                ó²   — t           j        j        }|rC|dk    rt           j        j        }n+|dk    rt           j        j        }nt          d|› d�¦  «        ‚|S )NÚ
evict_lastÚevict_firstzEviction policy rÖ  )r   ÚEVICTION_POLICYÚNORMALÚ
EVICT_LASTÚEVICT_FIRSTr/   )r   Úeviction_policyÚevictions      r   Ú_str_to_eviction_policyz&TritonSemantic._str_to_eviction_policy¸  se   € ÝÔ%Ô,ˆØð 	UØ ,Ò.Ð.ÝÔ-Ô8��Ø  MÒ1Ð1ÝÔ-Ô9��å Ð!S°OÐ!SÐ!SÐ!SÑTÔTÐTØˆr   c                ó”   — d }|rC|dk    rt           j        j        }n+|dk    rt           j        j        }nt	          d|› d�¦  «        ‚|S )NÚzeroÚnanzPadding option rÖ  )r   ÚPADDING_OPTIONÚPAD_ZEROÚPAD_NANr/   )r   Úpadding_optionÚpaddings      r   Ú_str_to_padding_optionz%TritonSemantic._str_to_padding_optionÃ  s_   € ØˆØð 	SØ Ò'Ð'ÝÔ+Ô4��Ø 5Ò(Ð(ÝÔ+Ô3��å Ð!Q°>Ð!QÐ!QÐ!QÑRÔRÐRØˆr   c                ó  — t           j        j        }|rs|dk    rt           j        j        }n[|dk    rt           j        j        }nC|dk    rt           j        j        }n+|dk    rt           j        j        }nt          d|› d�¦  «        ‚|S )NÚacquireÚreleaseÚacq_relÚrelaxedúMemory semantic rÖ  )r   ÚMEM_SEMANTICÚACQUIRE_RELEASEÚACQUIREÚRELEASEÚRELAXEDr/   )r   Ú
sem_optionÚsems      r   Ú_str_to_semzTritonSemantic._str_to_semÎ  s�   € ÝŒoÔ-ˆØð 
	PØ˜YÒ&Ð&Ý”oÔ-��Ø˜yÒ(Ð(Ý”oÔ-��Ø˜yÒ(Ð(Ý”oÔ5��Ø˜yÒ(Ð(Ý”oÔ-��å Ð!N°JÐ!NÐ!NÐ!NÑOÔOÐOØˆ
r   c                óâ   — t           j        j        }|r[|dk    rt           j        j        }nC|dk    rt           j        j        }n+|dk    rt           j        j        }nt          d|› d�¦  «        ‚|S )NÚgpuÚctaÚsysrý  rÖ  )r   ÚMEM_SYNC_SCOPEÚGPUÚCTAÚSYSTEMr/   )r   Úscope_optionÚscopes      r   Ú_str_to_scopezTritonSemantic._str_to_scopeÝ  s}   € ÝÔ!Ô%ˆØð 	RØ˜uÒ$Ð$ÝÔ)Ô-��Ø Ò&Ð&ÝÔ)Ô-��Ø Ò&Ð&ÝÔ)Ô0��å Ð!P°LÐ!PÐ!PÐ!PÑQÔQÐQØˆr   ÚptrÚmaskúOptional[TensorTy]Úboundary_checkr   rõ  rŸ  rÝ  rì  Úis_volatilec	           
     óx  — |                       |¦  «        }	|                      |¦  «        }
|                      |¦  «        }|j        j                             ¦   «         s*t          d|j                             ¦   «         › d�¦  «        ‚|€|�t          d¦  «        ‚|s|rt          d¦  «        ‚|j                             ¦   «         sT|r(|j                             ¦   «         rt          d¦  «        ‚|r(|j                             ¦   «         rt          d¦  «        ‚|j                             ¦   «         r6|�|  	                    ||¦  «        \  }}|�|  	                    ||¦  «        \  }}|j        j        }|j
        }|t          j        k    }|r<t          j        }t          j        ||j        ¦  «        }|                      ||¦  «        }|�|                      ||¦  «        }|j                             ¦   «         r|j                             |¦  «        }n|}|€7|                      | j                             |j        |	|
|¦  «        |¦  «        }nF|                      | j                             |j        |j        |r|j        nd |	|
|¦  «        |¦  «        }|r |                      |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)rß  rî  r÷  rg   rŠ   r€   r/   r   r]  r�   Ú
element_tyr1   rt   Úint8Úpointer_typeÚaddress_spacerŽ   r¯   r&   r'   Úcreate_loadr®   Úcreate_masked_load)r   r  r  r¬   r  rõ  rÝ  rì  r  rÞ  rí  rö  Úptr_tyÚelt_tyrÃ  r±  r¥   s                    r   ÚloadzTritonSemantic.loadê  sß  € à×0Ò0°Ñ@Ô@ˆØ×/Ò/°Ñ@Ô@ˆØ×-Ò-¨nÑ=Ô=ˆØŒxŒ×%Ò%Ñ'Ô'ð 	YÝÐW°S´X×5FÒ5FÑ5HÔ5HÐWÐWÐWÑXÔXÐXð ˆ<˜EÐ-ÝÐHÑIÔIÐIØð 	Y�nð 	YÝð Xñ Yô Yð Yð
 Œx× Ò Ñ"Ô"ð 	kØð j˜œ	×*Ò*Ñ,Ô,ð jÝ Ð!hÑiÔiÐiØð k˜œ×,Ò,Ñ.Ô.ð kÝ Ð!iÑjÔjÐjð Œ8×ÒÑÔð 	CØÐØ ×5Ò5°c¸4Ñ@Ô@‘	��TØÐ Ø!×6Ò6°s¸EÑBÔB‘
��Uð ””ˆØÔ"ˆð �BœGÒ#ˆØð 	)Ý”WˆFÝ”_ V¨VÔ-AÑBÔBˆFØ—)’)˜C Ñ(Ô(ˆCð ÐØ—I’I˜e VÑ,Ô,ˆEð Œ8×ÒÑÔð 	Ø”X×-Ò-¨fÑ5Ô5ˆFˆFð ˆFð ˆ<Ø—+’+˜dœl×6Ò6°s´zÀ5È(ÐT_Ñ`Ô`ÐbhÑiÔiˆCˆCà—+’+Ø”×/Ò/°´
¸D¼KÐY^ÐIhÈÌÈÐdhÐjoØ08¸+ñGô GØHNñPô PˆCð ð 	*Ø—)’)˜C¥¤Ñ)Ô)ˆCØˆ
r   Údescútl.tensor_descriptor_basec                óð  — t          |t          j        ¦  «        sJ dt          |¦  «        j        › �¦   «         ‚t          |j        ¦  «        }t          |¦  «        |k    sJ d|› dt          |¦  «        › �¦   «         ‚|                      |d¬¦  «        }| j         	                    |j
        ||                      |¦  «        |                      |¦  «        ¦  «        }|                      ||j        ¦  «        S )Nú"expected a tensor descriptor, got ú	expected ú offsets, but got F©Úrequire_i64)rb   r1   Útensor_descriptor_baserg   r    r^  Úblock_shapeÚ_convert_to_ir_valuesr'   Úcreate_descriptor_loadr®   rß  rî  r&   rF  )r   r"  ÚoffsetsrÝ  rì  Úndimrh   s          r   Údescriptor_loadzTritonSemantic.descriptor_load)  sù   € å˜$¥Ô 9Ñ:Ô:ð 	Gð 	GØFµ°d±´Ô1DÐFÐFñ	Gô 	Gð 	Gå�4Ô#Ñ$Ô$ˆÝ�7‰|Œ|˜tÒ#Ð#Ð#Ð%W°Ð%WÐ%WÍÈWÉÌÐ%WÐ%WÑ#Ô#Ð#à×,Ò,¨WÀ%Ð,ÑHÔHˆØŒL×/Ò/°´¸WÀd×FfÒFfÐguÑFvÔFvØ04×0LÒ0LÈ_Ñ0]Ô0]ñ_ô _ˆà�{Š{˜1˜dœoÑ.Ô.Ð.r   c                óT  — t          |t          j        ¦  «        sJ dt          |¦  «        j        › �¦   «         ‚t          |j        ¦  «        }t          |¦  «        |k    sJ d|› dt          |¦  «        › �¦   «         ‚|j        |j        k    sJ d|j        › d|j        › �¦   «         ‚d S )Nr%  r&  r'  zexpected value shape ú, got )rb   r1   r*  rg   r    r^  r+  rK  )r   r"  rO   r.  r/  s        r   Úvalidate_store_likez"TritonSemantic.validate_store_like5  sÆ   € Ý˜$¥Ô 9Ñ:Ô:ð 	Gð 	GØFµ°d±´Ô1DÐFÐFñ	Gô 	Gð 	Gå�4Ô#Ñ$Ô$ˆÝ�7‰|Œ|˜tÒ#Ð#Ð#Ð%W°Ð%WÐ%WÍÈWÉÌÐ%WÐ%WÑ#Ô#Ð#ØŒ{˜dÔ.Ò.Ð.Ð.ØI DÔ$4ÐIÐI¸E¼KÐIÐIñ /Ô.Ð.Ð.Ð.r   c                ó  — |                       |||¦  «         |                      ||j        ¦  «        }|                      |d¬¦  «        }|                      | j                             |j        |j        |¦  «        t          j	        ¦  «        S ©NFr(  )
r3  rŽ   r>   r,  r&   r'   Úcreate_descriptor_storer®   r1   Úvoid)r   r"  rO   r.  s       r   Údescriptor_storezTritonSemantic.descriptor_store=  sx   € Ø× Ò   u¨gÑ6Ô6Ð6à—	’	˜% ¤Ñ,Ô,ˆØ×,Ò,¨WÀ%Ð,ÑHÔHˆØ�{Š{˜4œ<×?Ò?ÀÄÈUÌ\Ð[bÑcÔcÕegÔelÑmÔmÐmr   c                ó´  — |                       |||¦  «         |j        t          j        t          j        t          j        t          j        t          j        t          j        hv s
J d¦   «         ‚|  	                    |d¬¦  «        }t          j        j        }|                      | j                             ||j        |j        |¦  «        t          j        ¦  «        S ©NúUnsupported dtypeFr(  )r3  r>   r1   ru   r2   rw   rR   rP   rQ   r,  r   ÚDESCRIPTOR_REDUCE_KINDÚADDr&   r'   Úcreate_descriptor_reducer®   r7  ©r   r"  rO   r.  rN   s        r   Údescriptor_atomic_addz$TritonSemantic.descriptor_atomic_addD  s©   € Ø× Ò   u¨gÑ6Ô6Ð6ØŒz�bœi­¬µ2´9½b¼jÍ"Ì*ÕVXÔVaÐbÐbÐbÐbÐdwÑbÔbÐbØ×,Ò,¨WÀ%Ð,ÑHÔHˆÝÔ(Ô,ˆØ�{Š{˜4œ<×@Ò@ÀÀtÄ{ÐTYÔT`ÐbiÑjÔjÕlnÔlsÑtÔtÐtr   c                ój   — t           j                             ¦   «         }|j        dk    o
|j        dk    S )NÚcudaéZ   )r   ÚactiveÚget_current_targetÚbackendÚarch)r   Útargets     r   Ú_has_native_tmazTritonSemantic._has_native_tmaK  s.   € Ý”×1Ò1Ñ3Ô3ˆØ” &Ò(Ð>¨V¬[¸BÒ->Ð?r   c                ó  — |t           j        t           j        t           j        t           j        t           j        t           j        hv s
J d¦   «         ‚|t           j        t           j        hv r|                      ¦   «         sJ d¦   «         ‚d S d S )Nr;  z-16-bit float types require native tma support)r1   ru   r2   rw   rv   rP   rQ   rI  )r   r>   s     r   Ú$_descriptor_atomic_min_max_supportedz3TritonSemantic._descriptor_atomic_min_max_supportedO  sx   € Ø�œ¥B¤H­b¬i½¼Å2Ä:ÍrÌ{Ð[Ð[Ð[Ð[Ð]pÑ[Ô[Ð[Ø•R”Z¥¤Ð-Ð-Ð-Ø×'Ò'Ñ)Ô)ÐZÐZÐ+ZÑZÔZÐZð .Ð-ØZÐZr   c                ó>  — |                       |||¦  «         |                      |j        ¦  «         |                      |d¬¦  «        }t          j        j        }|                      | j         	                    ||j
        |j
        |¦  «        t          j        ¦  «        S r5  )r3  rK  r>   r,  r   r<  ÚMINr&   r'   r>  r®   r1   r7  r?  s        r   Údescriptor_atomic_minz$TritonSemantic.descriptor_atomic_minT  óˆ   € Ø× Ò   u¨gÑ6Ô6Ð6Ø×1Ò1°$´*Ñ=Ô=Ð=Ø×,Ò,¨WÀ%Ð,ÑHÔHˆÝÔ(Ô,ˆØ�{Š{˜4œ<×@Ò@ÀÀtÄ{ÐTYÔT`ÐbiÑjÔjÕlnÔlsÑtÔtÐtr   c                ó>  — |                       |||¦  «         |                      |j        ¦  «         |                      |d¬¦  «        }t          j        j        }|                      | j         	                    ||j
        |j
        |¦  «        t          j        ¦  «        S r5  )r3  rK  r>   r,  r   r<  ÚMAXr&   r'   r>  r®   r1   r7  r?  s        r   Údescriptor_atomic_maxz$TritonSemantic.descriptor_atomic_max[  rO  r   c                óˆ  — |                       |||¦  «         |j        t          j        t          j        t          j        t          j        hv s
J d¦   «         ‚|                      |d¬¦  «        }t          j	        j
        }|                      | j                             ||j        |j        |¦  «        t          j        ¦  «        S r:  )r3  r>   r1   ru   r2   rw   rv   r,  r   r<  ÚANDr&   r'   r>  r®   r7  r?  s        r   Údescriptor_atomic_andz$TritonSemantic.descriptor_atomic_andb  óŸ   € Ø× Ò   u¨gÑ6Ô6Ð6ØŒz�bœi­¬µ2´9½b¼hÐGÐGÐGÐGÐI\ÑGÔGÐGØ×,Ò,¨WÀ%Ð,ÑHÔHˆÝÔ(Ô,ˆØ�{Š{˜4œ<×@Ò@ÀÀtÄ{ÐTYÔT`ÐbiÑjÔjÕlnÔlsÑtÔtÐtr   c                óˆ  — |                       |||¦  «         |j        t          j        t          j        t          j        t          j        hv s
J d¦   «         ‚|                      |d¬¦  «        }t          j	        j
        }|                      | j                             ||j        |j        |¦  «        t          j        ¦  «        S r:  )r3  r>   r1   ru   r2   rw   rv   r,  r   r<  ÚORr&   r'   r>  r®   r7  r?  s        r   Údescriptor_atomic_orz#TritonSemantic.descriptor_atomic_ori  sŸ   € Ø× Ò   u¨gÑ6Ô6Ð6ØŒz�bœi­¬µ2´9½b¼hÐGÐGÐGÐGÐI\ÑGÔGÐGØ×,Ò,¨WÀ%Ð,ÑHÔHˆÝÔ(Ô+ˆØ�{Š{˜4œ<×@Ò@ÀÀtÄ{ÐTYÔT`ÐbiÑjÔjÕlnÔlsÑtÔtÐtr   c                óˆ  — |                       |||¦  «         |j        t          j        t          j        t          j        t          j        hv s
J d¦   «         ‚|                      |d¬¦  «        }t          j	        j
        }|                      | j                             ||j        |j        |¦  «        t          j        ¦  «        S r:  )r3  r>   r1   ru   r2   rw   rv   r,  r   r<  ÚXORr&   r'   r>  r®   r7  r?  s        r   Údescriptor_atomic_xorz$TritonSemantic.descriptor_atomic_xorp  rV  r   c                ó
  — t          |t          j        ¦  «        sJ d|j        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        t          j	        t          j
        hv 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¦  «        }|                      |fd¬¦  «        d         }| j                             |j        |j        ||                     | j        ¦  «        ¦  «        }	|                      |	|¦  «        S )Nr%  Ú 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 ú2x offsets must have dtype int16 or int32, but got é   z5descriptor gather must have at least 8 rows, but got rE  zdescriptor gather of ú must have at least ú columns, but got Fr(  )rb   r1   r*  r   r    r^  r+  rK  r>   Úint16r2   r³  rF  r,  r'   Úcreate_descriptor_gatherr®   r°   r&   )
r   r"  Ú	x_offsetsÚy_offsetrÝ  rì  r>   Úmin_colsrg   rh   s
             r   Údescriptor_gatherz TritonSemantic.descriptor_gatherw  s…  € Ý˜$¥Ô 9Ñ:Ô:ð 	Kð 	KØJ°´Ô1HÐJÐJñ	Kô 	Kð 	Kà Ò#Ð#Ð#Ð%JÑ#Ô#Ð#Ø "Ò$Ð$Ð$Ð&LÑ$Ô$Ð$õ �4Ô#Ñ$Ô$¨Ò)Ð)Ð)Ð+_ÈTÔM]Ð+_Ð+_Ñ)Ô)Ð)ØÔ Ô" aÒ'Ð'Ð'Ð)hÐVZÔVfÐ)hÐ)hÑ'Ô'Ð'õ �9”?Ñ#Ô# qÒ(Ð(Ð(Ð*\È9Ì?Ð*\Ð*\Ñ(Ô(Ð(ØŒ¥2¤8Ý#%¤8ð#-ð -ð -ð -Ø.tÐclÔcrÐ.tÐ.tñ-ô -ð -ð Œ˜qÔ! QÒ&Ð&Ð&Ð(qÐ`iÔ`oÐ(qÐ(qÑ&Ô&Ð&Ø”
ˆØ˜Ô1Ñ1°AÑ5ˆØÔØôØòð ð ð A°Eð  Að  AÈxð  Að  AÐkoÔk{Ð|}Ôk~ð  Að  Añô ð õ Œ}˜TœZ¨)¬/¸!Ô*<¸dÔ>NÈqÔ>QÐ)RÑSÔSˆØ×-Ò-¨x¨lÈÐ-ÑNÔNÈqÔQˆØŒL×1Ò1°$´+¸yÔ?OÐQYÐ[_×[eÒ[eÐfjÔfrÑ[sÔ[sÑtÔtˆØ�{Š{˜1˜dÑ#Ô#Ð#r   c                ód  — t          |t          j        ¦  «        sJ dt          |¦  «        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        t          j	        t          j
        hv 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         › �¦   «         ‚|                      |fd¬¦  «        d         }| j                             |j        |j        |j        |¦  «         |                      d t          j        ¦  «        S )Nr%  r   r_  r   r   r`  ra  rb  rc  z6descriptor scatter must have at least 8 rows, but got rE  zdescriptor scatter of rd  re  Fr(  )rb   r1   r*  rg   r    r^  r+  rK  r>   rf  r2   r³  r,  r'   Úcreate_descriptor_scatterr®   r&   r7  )r   r"  rO   rh  ri  r>   rj  s          r   Údescriptor_scatterz!TritonSemantic.descriptor_scatter’  s!  € Ý˜$¥Ô 9Ñ:Ô:ð 	Gð 	GØFµ°d±´Ô1DÐFÐFñ	Gô 	Gð 	Gõ �4Ô#Ñ$Ô$¨Ò)Ð)Ð)Ð+_ÈTÔM]Ð+_Ð+_Ñ)Ô)Ð)ØÔ Ô" aÒ'Ð'Ð'Ð)hÐVZÔVfÐ)hÐ)hÑ'Ô'Ð'õ �9”?Ñ#Ô# qÒ(Ð(Ð(Ð*\È9Ì?Ð*\Ð*\Ñ(Ô(Ð(ØŒ¥2¤8Ý#%¤8ð#-ð -ð -ð -Ø.tÐclÔcrÐ.tÐ.tñ-ô -ð -ð Œ˜qÔ! QÒ&Ð&Ð&Ð(rÐajÔapÐ(rÐ(rÑ&Ô&Ð&Ø”
ˆØ˜Ô1Ñ1°AÑ5ˆØÔØôØòð ð ð B°Uð  Bð  BÐPXð  Bð  BÐlpÔl|Ð}~Ôlð  Bð  Bñô ð ð ×-Ò-¨x¨lÈÐ-ÑNÔNÈqÔQˆØŒ×.Ò.¨t¬{¸E¼LÈ)ÔJZÐ\dÑeÔeÐeØ�{Š{˜4¥¤Ñ)Ô)Ð)r   c                óü   — |j         }|€|                      ||¦  «        \  }}n0|                      |¦  «        }|                      |||¦  «        \  }}}||j         k    rt          d|j         › d|› �¦  «        ‚|||fS )Nz(Expected pointer argument to have shape z	 but got )rK  Úbroadcast_tensorsrj   r/   )r   r  Úvalr  Ú	ptr_shapes        r   Ú_broadcast_ptr_val_maskz&TritonSemantic._broadcast_ptr_val_maskª  s’   € Ø”Iˆ	Øˆ<Ø×-Ò-¨c°3Ñ7Ô7‰HˆC��à—>’> $Ñ'Ô'ˆDØ!×3Ò3°C¸¸dÑCÔC‰NˆC��dØ˜œ	Ò!Ð!ÝÐgÈÌ	ÐgÐgÐ\eÐgÐgÑhÔhÐhØ�C˜ˆ~Ðr   rq  c           	     ó:  — |                       |¦  «        }|                      |¦  «        }|j                             ¦   «         s|j        j                             ¦   «         rt          d¦  «        ‚|j        j                             ¦   «         s*t          d|j                             ¦   «         › d�¦  «        ‚|rt          d¦  «        ‚|j                             ¦   «         sR|j                             ¦   «         rt          d¦  «        ‚|r(|j                             ¦   «         rt          d¦  «        ‚|j                             ¦   «         r|  	                    |||¦  «        \  }}}|j        j        }	|	j
        }
|
t          j        k    r<t          j        }
t          j        |
|	j        ¦  «        }	|                      ||	¦  «        }|                      ||
¦  «        }|€E|                      | j                             |j        |j        ||¦  «        t          j        ¦  «        S |j        j                             ¦   «         st          d¦  «        ‚|                      | j                             |j        |j        |j        ||¦  «        t          j        ¦  «        S )Nú"Cannot store to a constant pointerr  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  ú"Mask must have boolean scalar type)rä  rî  rg   Úis_constrŠ   r/   r€   r   r]  rs  r  r1   rt   r  r  r  rŽ   r&   r'   Úcreate_storer®   r7  rÃ  Úcreate_masked_store)r   r  rq  r  r  rÝ  rì  rÞ  rí  r  r   s              r   ÚstorezTritonSemantic.storeµ  sx  € à×1Ò1°.ÑAÔAˆØ×/Ò/°Ñ@Ô@ˆØŒ8×ÒÑÔð 	C #¤(¤/×":Ò":Ñ"<Ô"<ð 	CÝÐAÑBÔBÐBð ŒxŒ×%Ò%Ñ'Ô'ð 	ZÝÐX°S´X×5FÒ5FÑ5HÔ5HÐXÐXÐXÑYÔYÐYð ð 	FÝð Eñ Fô Fð Fð
 Œx× Ò Ñ"Ô"ð 	jØŒx× Ò Ñ"Ô"ð kÝ Ð!iÑjÔjÐjØð j˜œ	×*Ò*Ñ,Ô,ð jÝ Ð!hÑiÔiÐið Œ8×ÒÑÔð 	JØ!×9Ò9¸#¸sÀDÑIÔI‰NˆC��dà””ˆØÔ"ˆð •R”WÒÐÝ”WˆFÝ”_ V¨VÔ-AÑBÔBˆFØ—)’)˜C Ñ(Ô(ˆCð �iŠi˜˜VÑ$Ô$ˆð ˆ<Ø—;’;˜tœ|×8Ò8¸¼ÀSÄZÐQVÐX`ÑaÔaÕceÔcjÑkÔkÐkØŒyÔ×'Ò'Ñ)Ô)ð 	CÝÐAÑBÔBÐBØ�{Š{˜4œ<×;Ò;¸C¼JÈÌ
ÐTXÔT_ÐafÐhpÑqÔqÝœ7ñ$ô $ð 	$r   Úcmpr  r  c           	     ó4  — |                       |¦  «        }|                      |¦  «        }|j        j        j        }|j        dvrt          d¦  «        ‚|                      | j         	                    |j
        |j
        |j
        ||¦  «        |j        ¦  «        S )N)é   rE  rž   z9atomic_cas only supports elements with width {16, 32, 64})r  r  rg   rŠ   r  r³  r/   r&   r'   Úcreate_atomic_casr®   )r   r  r{  rq  r  r  r  s          r   Ú
atomic_caszTritonSemantic.atomic_casé  s‹   € Ø×Ò˜sÑ#Ô#ˆØ×"Ò" 5Ñ)Ô)ˆØ”X”_Ô/ˆ
ØÔ(°Ð<Ð<ÝÐXÑYÔYÐYØ�{Š{˜4œ<×9Ò9¸#¼*ÀcÄjÐRUÔR\Ð^aÐchÑiÔiÐknÔksÑtÔtÐtr   Úopú#Tuple[TensorTy, TensorTy, TensorTy]c                óê  — |j         j                             ¦   «         s)t          d|j                              ¦   «         z   ¦  «        ‚|j                              ¦   «         s|j         j                             ¦   «         rt          d¦  «        ‚|j         j        j        }|t          j        u r|dk    rt          d|z   dz   ¦  «        ‚|t          j	        u r|dk    rt          d|z   dz   ¦  «        ‚|t          j
        t          j        fv s|j        dk     r%t          d|z   dz   t          |¦  «        z   ¦  «        ‚|j                              ¦   «         r|                      |||¦  «        \  }}}|                      ||j         j        j        ¦  «        }|€­| j                             d	¦  «        }t          j        }|j                              ¦   «         rW|j                              t          j        ¦  «        }| j                             |                     | j        ¦  «        |¦  «        }|                      ||¦  «        }n-|j         j                             ¦   «         st          d
¦  «        ‚|||fS )Nz)Pointer argument of store instruction is ru  r´   Úatomic_z does not support fp16z does not support bf16r}  z does not support Trv  )rg   rŠ   r€   r/   r   rw  r  r1   rP   rQ   rf  Úuint16r³  rŸ  r]  rs  rŽ   r'   Úget_int1rt   r¯   r_  r°   r&   rÃ  )r   r  rq  r  r€  r  Úmask_irÚmask_tys           r   Úatom_red_typechecking_implz)TritonSemantic.atom_red_typechecking_implñ  s4  € àŒxŒ×%Ò%Ñ'Ô'ð 	`ÝÐHÈ3Ì8×K\ÒK\ÑK^ÔK^Ñ^Ñ_Ô_Ð_ØŒ8×ÒÑÔð 	C #¤(Ô"5×">Ò">Ñ"@Ô"@ð 	CÝÐAÑBÔBÐBØ”X”_Ô/ˆ
Ø�œÐ#Ð#¨¨eª¨Ý˜Y¨™^Ð.FÑFÑGÔGÐGØ�œÐ$Ð$¨¨uª¨Ý˜Y¨™^Ð.FÑFÑGÔGÐGØ�"œ(¥B¤IÐ.Ð.Ð.°*Ô2OÐRTÒ2TÐ2TÝ˜Y¨™^Ð.BÑBÅSÈÁ_Ä_ÑTÑUÔUÐUØŒ8×ÒÑÔð 	JØ!×9Ò9¸#¸sÀDÑIÔI‰NˆC��dØ�iŠi˜˜SœXœ_Ô7Ñ8Ô8ˆØˆ<Ø”l×+Ò+¨DÑ1Ô1ˆGÝ”gˆGØŒx× Ò Ñ"Ô"ð ZØœ(×2Ò2µ2´7Ñ;Ô;�Øœ,×3Ò3°G·M²MÀ$Ä,Ñ4OÔ4OÐQXÑYÔY�Ø—;’;˜w¨Ñ0Ô0ˆDˆDØ”Ô!×)Ò)Ñ+Ô+ð 	CÝÐAÑBÔBÐBØ�C˜ˆ~Ðr   c                óä   — |j         j        }t          j        |d¬¦  «        }|                      ||¦  «        }|                      ||dz
  ¦  «        }|                      |t          j        ¦  «        S )NF)rÑ  Úsignedr   )r>   r³  r1   Úget_int_dtyper  r  rŽ   rt   )r   rh   rÑ  ÚidtypeÚixÚsignbits         r   Ú_signbitzTritonSemantic._signbit  sa   € Ø”7Ô-ˆÝÔ!¨8¸EÐBÑBÔBˆØ�\Š\˜!˜VÑ$Ô$ˆØ—)’)˜B ¨1¡Ñ-Ô-ˆØ�yŠy˜¥"¤'Ñ*Ô*Ð*r   c                ó˜  — |                       |||d¦  «        \  }}}|                      |¦  «        }|                      |¦  «        }|j        j        }|                     ¦   «         rÀ|                     ¦   «         rV|                      | j         	                    t          j        j        |j        |j        |j        ||¦  «        |j        ¦  «        S |                      | j         	                    t          j        j        |j        |j        |j        ||¦  «        |j        ¦  «        S |t          j        t          j        hvrt%          d|› �¦  «        ‚|t          j        k    rt          j        nt          j        }|                      ||¦  «        }|                      |t          j        |d¦  «        ¦  «        }	|t          j        k    rt          j        nt          j        }
|                      ||
¦  «        }|                      |t          j        |
d¦  «        ¦  «        }|                      |¦  «        }|                      |¦  «        }|                      | j         	                    t          j        j        |	j        |j        |                      ||¦  «        j        ||¦  «        |j        ¦  «        }|                      | j         	                    t          j        j        |j        |j        |                      ||¦  «        j        ||¦  «        |j        ¦  «        }|                      |||¦  «        }|                      ||¦  «        S )Nrð   z#atomic_max not supported for dtype r   )rˆ  r  r  rg   rŠ   rY   rË   r&   r'   Úcreate_atomic_rmwr   Ú	ATOMIC_OPrQ  r®   ÚUMAXr1   rR   rT   rA   r2   rv   r  r  ru   rw   r�  r  r¡   ÚUMINÚwhere©r   r  rq  r  r  r  Úsca_tyÚi_typeÚi_valÚi_ptrÚui_typeÚui_valÚui_ptrÚnegÚposÚpos_retÚneg_retr¥   s                     r   Ú
atomic_maxzTritonSemantic.atomic_max  óÊ  € Ø×8Ò8¸¸cÀ4ÈÑOÔO‰ˆˆS�$Ø×Ò˜sÑ#Ô#ˆØ×"Ò" 5Ñ)Ô)ˆØ””ˆà�=Š=‰?Œ?ð 	Ø×#Ò#Ñ%Ô%ð Ø—{’{Ø”L×2Ò2µ2´<Ô3CÀSÄZÐQTÔQ[Ð]aÔ]hÐjmÐotÑuÔuØ”Hñô ð ð —{’{Ø”L×2Ò2µ2´<Ô3DÀcÄjÐRUÔR\Ð^bÔ^iÐknÐpuÑvÔvØ”Hñô ð ð �"œ*¥b¤jÐ1Ð1Ð1ÝÐJÀ&ÐJÐJÑKÔKÐKà#¥r¤zÒ1Ð1•”�µr´xˆØ—’˜S &Ñ)Ô)ˆØ—’˜S¥"¤/°&¸!Ñ"<Ô"<Ñ=Ô=ˆØ%­¬Ò3Ð3•"”)�)½¼ˆØ—’˜c 7Ñ+Ô+ˆØ—’˜c¥2¤?°7¸AÑ#>Ô#>Ñ?Ô?ˆØ�mŠm˜CÑ Ô ˆØ�iŠi˜‰nŒnˆØ—+’+ØŒL×*Ò*­2¬<Ô+;¸U¼\È5Ì<Ø+/¯9ª9°T¸3Ñ+?Ô+?Ô+FÈÈUñTô TØUZÔU_ñaô aˆð —+’+ØŒL×*Ò*­2¬<Ô+<¸f¼mÈVÌ]Ø+/¯9ª9°T¸3Ñ+?Ô+?Ô+FÈÈUñTô TØU[ÔU`ñbô bˆð �jŠj˜˜g wÑ/Ô/ˆØ�|Š|˜C Ñ(Ô(Ð(r   c                ó˜  — |                       |||d¦  «        \  }}}|                      |¦  «        }|                      |¦  «        }|j        j        }|                     ¦   «         rÀ|                     ¦   «         rV|                      | j         	                    t          j        j        |j        |j        |j        ||¦  «        |j        ¦  «        S |                      | j         	                    t          j        j        |j        |j        |j        ||¦  «        |j        ¦  «        S |t          j        t          j        hvrt%          d|› �¦  «        ‚|t          j        k    rt          j        nt          j        }|                      ||¦  «        }|                      |t          j        |d¦  «        ¦  «        }	|t          j        k    rt          j        nt          j        }
|                      ||
¦  «        }|                      |t          j        |
d¦  «        ¦  «        }|                      |¦  «        }|                      |¦  «        }|                      | j         	                    t          j        j        |	j        |j        |                      ||¦  «        j        ||¦  «        |j        ¦  «        }|                      | j         	                    t          j        j        |j        |j        |                      ||¦  «        j        ||¦  «        |j        ¦  «        }|                      |||¦  «        }|                      ||¦  «        S )Nrï   z#atomic_min not supported for dtype r   )rˆ  r  r  rg   rŠ   rY   rË   r&   r'   r‘  r   r’  rM  r®   r”  r1   rR   rT   rA   r2   rv   r  r  ru   rw   r�  r  r¡   r“  r•  r–  s                     r   Ú
atomic_minzTritonSemantic.atomic_min9  r£  r   c           
     ó   — |                       |||d¦  «        \  }}}|                      |¦  «        }|                      |¦  «        }|j        j        }|                     ¦   «         rt          j        j        nt          j        j	        }|  
                    | j                             ||j        |j        |j        ||¦  «        |j        ¦  «        S )Nr´   )rˆ  r  r  rg   rŠ   r�   r   r’  ÚFADDr=  r&   r'   r‘  r®   )r   r  rq  r  r  r  r—  r€  s           r   Ú
atomic_addzTritonSemantic.atomic_add_  s¶   € Ø×8Ò8¸¸cÀ4ÈÑOÔO‰ˆˆS�$Ø×Ò˜sÑ#Ô#ˆØ×"Ò" 5Ñ)Ô)ˆØ””ˆØ"(×"4Ò"4Ñ"6Ô"6ÐL�RŒ\ÔÐ½B¼LÔ<LˆØ�{Š{˜4œ<×9Ò9¸"¸c¼jÈ#Ì*ÐVZÔVaÐcfÐhmÑnÔnØœ8ñ%ô %ð 	%r   c           
     ó:  — |                       |||d¦  «        \  }}}|                      |¦  «        }|                      |¦  «        }|                      | j                             t          j        j        |j	        |j	        |j	        ||¦  «        |j
        ¦  «        S )NÚand)rˆ  r  r  r&   r'   r‘  r   r’  rT  r®   rg   ©r   r  rq  r  r  r  s         r   Ú
atomic_andzTritonSemantic.atomic_andh  ó’   € Ø×8Ò8¸¸cÀ4ÈÑOÔO‰ˆˆS�$Ø×Ò˜sÑ#Ô#ˆØ×"Ò" 5Ñ)Ô)ˆØ�{Š{ØŒL×*Ò*­2¬<Ô+;¸S¼ZÈÌÐUYÔU`ÐbeÐglÑmÔmÐorÔowñyô yð 	yr   c           
     ó:  — |                       |||d¦  «        \  }}}|                      |¦  «        }|                      |¦  «        }|                      | j                             t          j        j        |j	        |j	        |j	        ||¦  «        |j
        ¦  «        S )NÚor)rˆ  r  r  r&   r'   r‘  r   r’  rX  r®   rg   r«  s         r   Ú	atomic_orzTritonSemantic.atomic_oro  s‘   € Ø×8Ò8¸¸cÀ4ÈÑNÔN‰ˆˆS�$Ø×Ò˜sÑ#Ô#ˆØ×"Ò" 5Ñ)Ô)ˆØ�{Š{ØŒL×*Ò*­2¬<¬?¸C¼JÈÌ
ÐTXÔT_ÐadÐfkÑlÔlÐnqÔnvñxô xð 	xr   c           
     ó:  — |                       |||d¦  «        \  }}}|                      |¦  «        }|                      |¦  «        }|                      | j                             t          j        j        |j	        |j	        |j	        ||¦  «        |j
        ¦  «        S )NÚxor)rˆ  r  r  r&   r'   r‘  r   r’  r[  r®   rg   r«  s         r   Ú
atomic_xorzTritonSemantic.atomic_xorv  r­  r   c           
     ó:  — |                       |||d¦  «        \  }}}|                      |¦  «        }|                      |¦  «        }|                      | j                             t          j        j        |j	        |j	        |j	        ||¦  «        |j
        ¦  «        S )NÚxchg)rˆ  r  r  r&   r'   r‘  r   r’  ÚXCHGr®   rg   r«  s         r   Úatomic_xchgzTritonSemantic.atomic_xchg}  sŽ   € Ø×8Ò8¸¸cÀ4ÈÑPÔP‰ˆˆS�$Ø×Ò˜sÑ#Ô#ˆØ×"Ò" 5Ñ)Ô)ˆØ�{Š{ØŒL×*Ò*­2¬<Ô+<¸c¼jÈ#Ì*ÐVZÔVaÐcfÐhmÑnÔnØŒHñô ð 	r   c                ó  — |                      ¦   «         | j        j        j        v sJ d| j        j        j        › d|› �¦   «         ‚|                     ¦   «         }|dk    rd}|dk    rd}|dk    rd}t          t          j        |¦  «        S )	Nzinput_precision must be one of ú. Got ÚTF32X3ÚTF32x3ÚBF16X3ÚBF16x3ÚBF16X6ÚBF16x6)Úlowerr'   rŸ   Úallowed_dot_input_precisionsÚupperrR  r   ÚINPUT_PRECISION)r   Úinput_precisions     r   Ú_str_to_dot_input_precisionz*TritonSemantic._str_to_dot_input_precision‰  s    € Ø×$Ò$Ñ&Ô&¨$¬,Ô*>Ô*[Ð[Ð[Ð[Øx¨d¬lÔ.BÔ._ÐxÐxÐgvÐxÐxñ \Ô[Ð[à)×/Ò/Ñ1Ô1ˆØ˜hÒ&Ð&Ø&ˆOØ˜hÒ&Ð&Ø&ˆOØ˜hÒ&Ð&Ø&ˆOÝ•rÔ)¨?Ñ;Ô;Ð;r   ÚaccrÄ  Úmax_num_imprecise_accÚ	out_dtypeútl.dtype | Nonec           
     ó4  — |j                              ¦   «         r|j                              ¦   «         s
J d¦   «         ‚|j                             ¦   «         r|j                             ¦   «         rnä|j        t          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	        t          j
        fv sJ d|j        › �¦   «         ‚|j        |j        k    sJ d|j        › d|j        › �¦   «         ‚|j                             ¦   «         s|j                             ¦   «         rgd| j        j        j        v rt          j        d¦  «         |                      |t          j        ¦  «        }|                      |t          j        ¦  «        }|j                             ¦   «         p|j                             ¦   «         }|j                             ¦   «         p|j                             ¦   «         }|s|r‡|rdnd	}	|	| j        j        j        v rn| j        j        j        }
t          j        |	› d
|
› d|
› d�¦  «         |                      |t          j        ¦  «        }|                      |t          j        ¦  «        }|€| j        j        j        }|€|€t          j	        n|j         j        }|                      |¦  «        }t1          |j        ¦  «        }t1          |j        ¦  «        }||cxk    rdk    s,n ||cxk    rdk    sn J d|j        › d|j        › d�¦   «         ‚|j        d         j        |j        d         j        k    sAJ d|j        › d|j        › d|j        d         j        › d|j        d         j        › d�	¦   «         ‚| j        j                             d¦  «        	 €
J d¦   «         ‚ | j        j        d         |j         |j         ¦  «        }|j        d         j        |d         k    r8|j        d         j        |d         k    r|j        d         j        |d         k    s%J d|d         › d|d         › d|d         › �¦   «         ‚|j         j                             ¦   «         rL|j         j        t          j        k    s
J d¦   «         ‚| j                             d¦  «        }t          j         }�n| !                    ¦   «         rtE          d ¦  «        ‚|j         j         #                    ¦   «         s|j         j         !                    ¦   «         r'| j         $                    d¦  «        }t          j	        }n�|j         j         %                    ¦   «         r'| j         &                    d¦  «        }t          j
        }nJ| '                    ¦   «         r| j         (                    d¦  «        n| j         $                    d¦  «        }|}|j         j        d         }|j         j        d         }|j         j        d         }|dk    r|j         j        d         nd }t	          j)        ||r|||gn||g¦  «        }|€4| j         *                    | +                    | j        ¦  «        |¦  «        }ns|j,        }|j         j        |j        k    sJ d!|j        › d"|j         j        › �¦   «         ‚|j         j        |k    s(J d#|› d"|j         j        › d$|j         j        › d%�¦   «         ‚|€G|j                             ¦   «         r+|j                             ¦   «         r| j        j        j-        }nQd}nN|j                             ¦   «         r5|j                             ¦   «         r||k    rtE          d&|› d'|› d�¦  «        ‚|  .                    | j         /                    |j,        |j,        |||¦  «        |¦  «        S )(Nz0dot operands must be block tensors (not scalars)zUnsupported lhs dtype zUnsupported rhs dtype z&Both operands must be same dtype. Got r   Úfp8e4b15z–the use of fp8e4b15 is deprecated on Hopper and later architectures and can cause significant slow down. It will be removed in a future triton releaseÚfp8e4b8Úfp8e5b16z- is AMD gfx942 specific and not supported on z^ so it's upcasted to fp16 and can cause significant slow down. Please use OCP fp8 variants on z for performancer   é   ú+Both inputs must be either 2D or 3D; (lhs: ú	 vs rhs: r  r{  éþÿÿÿ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)`úexpected accumulator shape r2  úexpected accumulator dtype ú; pass out_dtype=ú to use this accumulator dtypezmax_num_imprecise_acc (z) must be <= K ()0rg   r]  r>   rX   r1   r  Úuint8rP   rQ   rR   rT   r¾  r'   rŸ   Ú!deprecated_fp8_dot_operand_dtypesÚwarningsÚwarnrŽ   Ú
is_fp8e4b8Úis_fp8e5b16rG  Údefault_dot_input_precisionr  rÅ  r^  rK  rO   r¿  rÀ  rŠ   rY   Ú	get_int32r2   rW   r/   rU   rP  rS   Úget_fp64rV   Úget_fp16rF  r_  r°   r®   Úmax_num_imprecise_acc_defaultr&   Ú
create_dot)r   rƒ   r…   rÆ  rÄ  rÇ  rÈ  Úuses_fp8e4b8Úuses_fp8e5b16Ú	type_namerG  Úlhs_rankÚrhs_rankrÒ  r  Úret_scalar_tyÚMÚNÚKÚBrÎ   Ú
acc_handles                         r   ÚdotzTritonSemantic.dot•  sœ  € àŒx× Ò Ñ"Ô"Ðn s¤x×'8Ò'8Ñ':Ô':ÐnÐnÐ<nÑnÔnÐnàŒ9×ÒÑÔð 	p #¤)×"2Ò"2Ñ"4Ô"4ð 	pàà”9¥¤­"¬(µB´JÅÄÍRÌZÝ!#¤ð!-ð -ð -ð -Ø.RÀsÄyÐ.RÐ.Rñ-ô -ð -à”9¥¤­"¬(µB´JÅÄÍRÌZÝ!#¤ð!-ð -ð -ð -Ø.RÀsÄyÐ.RÐ.Rñ-ô -ð -à”9 ¤	Ò)Ð)Ð)Ð+oÐTWÔT]Ð+oÐ+oÐdgÔdmÐ+oÐ+oÑ)Ô)Ð)àŒ9× Ò Ñ"Ô"ð 	- c¤i×&;Ò&;Ñ&=Ô&=ð 	-Ø˜Tœ\Ô1ÔSÐSÐSÝ”ð mñô ð ð —)’)˜C¥¤Ñ,Ô,ˆCØ—)’)˜C¥¤Ñ,Ô,ˆCà”y×+Ò+Ñ-Ô-ÐG°´×1EÒ1EÑ1GÔ1GˆØœ	×-Ò-Ñ/Ô/ÐJ°3´9×3HÒ3HÑ3JÔ3JˆØð 	1˜=ð 	1Ø%1ÐA˜	˜	°zˆIØ˜DœLÔ0ÔRÐRÐRØ”|Ô+Ô0�Ý”Ø ð Mð MÈtð Mð MØ6:ðMð Mð MñNô Nð Nð —i’i ¥R¤ZÑ0Ô0�Ø—i’i ¥R¤ZÑ0Ô0�àÐ"Ø"œlÔ2ÔNˆOàÐØ&) k�œ
˜
°s´xÔ7JˆIà×:Ò:¸?ÑKÔKˆå�s”y‘>”>ˆÝ�s”y‘>”>ˆØ˜8Ð(Ð(Ò(Ð( qÒ(Ð(Ð(Ð(¨H¸Ð,EÐ,EÒ,EÐ,EÀAÒ,EÐ,EÐ,EÐ,EÐ,Eð  HVÐuxÔu~ð  HVð  HVð  JMô  JSð  HVð  HVð  HVÑ,EÔ,EÐ,EØŒy˜Œ}Ô" c¤iØô'Üòð ð ð u¨S¬Yð  uð  uÐQTÔQZð  uð  uð  Z]ô  Zcð  dfô  Zgô  Zmð  uð  uð  _bô  _hð  ikô  _lô  _rð  uð  uð  uñô ð àŒ|Ô'×+Ò+Øñô Ø#'ð(ð (Ø)]ñ(ô (ð (à?�t”|Ô/°Ô?ÀÄÈ#Ì(ÑSÔSˆØŒy˜Œ}Ô" l°1¤oÒ5Ð5¸#¼)ÀB¼-Ô:MÐQ]Ð^_ÔQ`Ò:`Ð:`Ø”	˜"”Ô# |°A¤Ò6Ð6Ð6Øu°¸a´ÐuÐuÈÐVWÌÐuÐuÐdpÐqrÔdsÐuÐuñ 7Ô6Ð6àŒ8Œ?×!Ò!Ñ#Ô#ð 	&Ø”8”?¥b¤gÒ-Ð-Ð-Ð/EÑ-Ô-Ð-Ø”×'Ò'¨Ñ*Ô*ˆBÝœHˆM‰MØ×ÒÑ Ô ð 	&ÝØzñô ð ð ŒXŒ_×$Ò$Ñ&Ô&ð 	&¨#¬(¬/×*AÒ*AÑ*CÔ*Cð 	&Ø”×&Ò& qÑ)Ô)ˆBÝœJˆMˆMØŒXŒ_×$Ò$Ñ&Ô&ð 	&Ø”×&Ò& qÑ)Ô)ˆBÝœJˆMˆMà-6×->Ò->Ñ-@Ô-@Ð^�”×&Ò& qÑ)Ô)Ð)ÀdÄl×F[ÒF[Ð\]ÑF^ÔF^ˆBØ%ˆMàŒHŒN˜2ÔˆØŒHŒN˜2ÔˆØŒHŒN˜2ÔˆØ!)¨Q¢ ˆCŒHŒN˜1ÔÐ°DˆÝ”˜}¸1Ð.H¨q°!°Q¨i¨iÀ1ÀaÀ&ÑIÔIˆØˆ;Øœ×2Ò2°6·<²<ÀÄÑ3MÔ3MÈrÑRÔRˆJˆJàœˆJØ”8”> V¤\Ò1Ð1Ð1ØR¨f¬lÐRÐRÀ#Ä(Ä.ÐRÐRñ 2Ô1Ð1à”8Ô&¨)Ò3Ð3Ð3ðV¨ið Vð V¸s¼xÔ?Rð Vð VØ"%¤(Ô"5ðVð Vð Vñ 4Ô3Ð3ð
 !Ð(ØŒy×ÒÑ!Ô!ð * c¤i×&6Ò&6Ñ&8Ô&8ð *Ø(,¬Ô(<Ô(ZÐ%Ð%à()Ð%Ð%àŒy×ÒÑ!Ô!ð h c¤i×&6Ò&6Ñ&8Ô&8ð hÐ=RÐUVÒ=VÐ=VÝ Ð!fÐ;PÐ!fÐ!fÐbcÐ!fÐ!fÐ!fÑgÔgÐgà�{Š{ØŒL×#Ò# C¤J°´
¸JÈÐYnÑoÔoÐqwñyô yð 	yr   Úfloat_formatc                óŠ   — t          t          j        |                     ¦   «         d ¦  «        }|€t	          d|› d�¦  «        ‚|S )NzInvalid float format: rp   )rR  r   ÚScaleDotElemTypeTYrÂ  r/   )r   rï  Úty_enums      r   Ú_str_to_fp_typezTritonSemantic._str_to_fp_typeù  sF   € Ý�"Ô/°×1CÒ1CÑ1EÔ1EÀtÑLÔLˆØˆ?ÝÐE°lÐEÐEÐEÑFÔFÐFØˆr   c                óò  — t           j        t           j        t           j        t           j        dœ                     |¦  «        }|€<|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        › �¦   «         ‚|  	                    ||¦  «        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Úfp16NÚe2m1z)Internal Error: Unexpected float format: z)e2m1 format must be packed as uint8. Got zUnexpected dtype for r¹  )
r1   Úfloat8e5Ú
float8e4nvrQ   rP   rÀ  r>   r×  r„  r  )r   rq  rï  Ú	triton_tyÚunsigned_tys        r   Ú_bitcast_to_fp_typez"TritonSemantic._bitcast_to_fp_typeÿ  s  € õ
  œ[µ"´-ÍÌÝ”Zð!ð !ß!$¢ \Ñ!2Ô!2ð 	àÐØ 6Ò)Ð)Ð)Ð+eÐWcÐ+eÐ+eÑ)Ô)Ð)Ø”9¥¤Ò(Ð(Ð(Ð*aÐVYÔV_Ð*aÐ*aÑ(Ô(Ð(ØˆJØŒ9˜	Ò!Ð!ØˆJå#%¤8µR´XÅrÄyÕZ\ÔZcÐdÐdÐeqÔrˆKØ”9 Ò+Ð+Ð+Ð-dÀ\Ð-dÐ-dÐY\ÔYbÐ-dÐ-dÑ+Ô+Ð+Ø—<’<  YÑ/Ô/Ð/r   c	                óÖ   — d„ }	t          j        |j        j         |	|¦  «        |                      |¦  «        ||j        j         |	|¦  «        |                      |¦  «        |¦  «        S )Nc                ór   — t          | t          j        ¦  «        r| j        j        dk    r| j        j        S d S )Nr   )rb   r1   r&   rV  rO   rg   rK  )Úscales    r   Ú_to_scale_handlez<TritonSemantic.deduce_scale_factor.<locals>._to_scale_handle  s5   € Ý˜%¥¤Ñ+Ô+ð (°´Ô0AÀQÒ0FÐ0FØ”zÔ'Ð'Ø�4r   )r   Údeduce_scale_factorrg   rK  ró  )
r   rƒ   Ú	lhs_scaleÚ
lhs_formatÚ
lhs_k_packr…   Ú	rhs_scaleÚ
rhs_formatÚ
rhs_k_packr  s
             r   r  z"TritonSemantic.deduce_scale_factor  s|   € ð	ð 	ð 	õ
 Ô% c¤h¤nÐ6FÐ6FÀyÑ6QÔ6QÐSW×SgÒSgÐhrÑSsÔSsØ&0°#´(´.ÐBRÐBRÐS\ÑB]ÔB]Ø&*×&:Ò&:¸:Ñ&FÔ&FÈ
ñTô Tð 	Tr   c                óæ   — |�5|j         j        }|dd …         |||z  gk    sJ d|› d||z  › d|› �¦   «         ‚|�5|j         j        }|dd …         |||z  gk    sJ d|› d||z  › d|› �¦   «         ‚d S d S )NrÑ  z*lhs_scale must be a tensor of shape [..., r“  z]. Got z*rhs_scale must be a tensor of shape [..., ©rg   rK  )	r   ré  rê  rë  r  r  Úscale_factorÚlhs_scale_shapeÚrhs_scale_shapes	            r   Úverify_scaled_shapez"TritonSemantic.verify_scaled_shape  sê   € ØÐ Ø'œnÔ2ˆOØ" 2 3 3Ô'Ø�1˜Ñ$ð,ò ð ð àl¸AÐlÐlÀÀlÑARÐlÐlÐ[jÐlÐlñô ð ð Ð Ø'œnÔ2ˆOØ" 2 3 3Ô'Ø�1˜Ñ$ð,ò ð ð àl¸AÐlÐlÀÀlÑARÐlÐlÐ[jÐlÐlñô ð ð !Ð ðð r   r  r  r  r  úTensorTy | NoneÚ	fast_mathr  r	  c                óN  — |j                              ¦   «         r|j                              ¦   «         s
J d¦   «         ‚t          |j        ¦  «        }t          |j        ¦  «        }||cxk    rdk    s,n ||cxk    rdk    sn J d|j        › d|j        › d�¦   «         ‚|                      |¦  «        }|                      |¦  «        }h d£}||v sJ d|› �¦   «         ‚||v sJ d	|› �¦   «         ‚|d u p"t          |t          j        ¦  «        o|j        d u }|d u p"t          |t          j        ¦  «        o|j        d u }|  	                    ||¦  «        }|  	                    ||¦  «        }|	s|d
k    s
J d¦   «         ‚|
s|d
k    s
J d¦   «         ‚|j         j        dd …         \  }}|j         j        dd …         \  }}|d
k    rdnd}|d
k    rdnd}|	r||z  n|}|
r||z  n|}||k    sJ d|j        › d|j        › d�¦   «         ‚|dk    r|j         j        d         nd }|}|	s||z  }n||z  }|
s||z  }t          j
        ||r|||gn||g¦  «        }| j                             d¦  «        }|€4| j                             |                     | j        ¦  «        |¦  «        }ns|j        }|j         j        |j        k    sJ d|j        › d|j         j        › �¦   «         ‚|j         j        |k    s(J d|› d|j         j        › d|j         j        › d�¦   «         ‚|rd n|j        } |rd n|j        }!|                      ||||	||||
¦  «        }"|                      ||||rd n||rd n||"¦  «         |                      | j                             |j        |!||j        | |||	|
|¦
  «
        |¦  «        S )Nz7dot_scaled operands must be block tensors (not scalars)r   rÎ  rÏ  rÐ  r  >   r÷  rù  rö  rõ  rø  zNYI: lhs_format zNYI: rhs_format rù  zBonly mxfp4 inputs can be packed along a dimension different than KrÑ  r   zCReduction dimension should pack the same number of elements; (lhs: r   rÓ  r2  rÔ  rÕ  rÖ  )rg   r]  r^  rK  ró  rb   r1   rc   rO   rþ  rF  r'   rP  r_  r°   r®   r  r  r  r&   Úcreate_dot_scaled)#r   rƒ   r  r  r…   r  r  rÆ  r  r  r	  rÈ  ræ  rç  Úlhs_format_enumÚrhs_format_enumÚallowed_formatsÚrhs_scale_is_noneÚlhs_scale_is_noneré  ÚK_LHSÚK_RHSrê  ÚPACKED_AÚPACKED_BÚPACKED_A_DIMÚPACKED_B_DIMrì  rë  rÎ   r  rí  Úrhs_scale_handleÚlhs_scale_handler  s#                                      r   Ú
dot_scaledzTritonSemantic.dot_scaled(  s  € ð Œx× Ò Ñ"Ô"Ðu s¤x×'8Ò'8Ñ':Ô':ÐuÐuÐ<uÑuÔuÐuå�s”y‘>”>ˆÝ�s”y‘>”>ˆØ˜8Ð(Ð(Ò(Ð( qÒ(Ð(Ð(Ð(¨H¸Ð,EÐ,EÒ,EÐ,EÀAÒ,EÐ,EÐ,EÐ,EÐ,Eð  HVÐuxÔu~ð  HVð  HVð  JMô  JSð  HVð  HVð  HVÑ,EÔ,EÐ,EØ×.Ò.¨zÑ:Ô:ˆØ×.Ò.¨zÑ:Ô:ˆØBÐBÐBˆØ˜_Ð,Ð,Ð,Ð.MÀÐ.MÐ.MÑ,Ô,Ð,Ø˜_Ð,Ð,Ð,Ð.MÀÐ.MÐ.MÑ,Ô,Ð,Ø%¨Ð-Ðrµ*¸YÍÌÑ2UÔ2UÐ2qÐZcÔZiÐmqÐZqÐØ%¨Ð-Ðrµ*¸YÍÌÑ2UÔ2UÐ2qÐZcÔZiÐmqÐZqÐØ×&Ò& s¨JÑ7Ô7ˆØ×&Ò& s¨JÑ7Ô7ˆàÐw˜Z¨6Ò1Ð1Ð1Ð3wÑ1Ô1Ð1ØÐw˜Z¨6Ò1Ð1Ð1Ð3wÑ1Ô1Ð1Ø”8”> " # #Ô&‰ˆˆ5Ø”8”> " # #Ô&‰ˆˆqØ" fÒ,Ð,�1�1°!ˆØ" fÒ,Ð,�1�1°!ˆØ+5Ð@�x %Ñ'Ð'¸5ˆØ+5Ð@�x %Ñ'Ð'¸5ˆØ˜|Ò+Ð+Ð+ð  .TÐsvÔs|ð  .Tð  .Tð  HKô  HQð  .Tð  .Tð  .TÑ+Ô+Ð+à!)¨Q¢ ˆCŒHŒN˜1ÔÐ°DˆØˆØð 	Ø�H‘ˆAˆAà�H‘ˆAØð 	Ø�H‘ˆAÝ”˜y°qÐ*D¨1¨a°¨)¨)¸qÀ!¸fÑEÔEˆØŒ\×"Ò" 1Ñ%Ô%ˆØˆ;Øœ×2Ò2°6·<²<ÀÄÑ3MÔ3MÈrÑRÔRˆJˆJàœˆJØ”8”> V¤\Ò1Ð1Ð1ØR¨f¬lÐRÐRÀ#Ä(Ä.ÐRÐRñ 2Ô1Ð1à”8Ô&¨)Ò3Ð3Ð3ðV¨ið Vð V¸s¼xÔ?Rð Vð VØ"%¤(Ô"5ðVð Vð Vñ 4Ô3Ð3ð $5ÐJ˜4˜4¸)Ô:JÐØ#4ÐJ˜4˜4¸)Ô:JÐà×/Ò/°°YÀ
ÈJÐX[Ð]fÐhrØ0:ñ<ô <ˆà× Ò   A qÐ2CÐ*R¨$¨$ÈØ):Ð!I  À	È<ñ	Yô 	Yð 	Yà�{Š{ØŒL×*Ò*¨3¬:Ð7GÈÐZ]ÔZdÐfvØ+:¸IÀzÐS]Ð_iñkô kØlrñtô tð 	tr   Ú	conditionc                ó0  — |j         t          j        k    rt          j        d|j         › �¦  «         |                      |t          j        ¦  «        }|                      ||dd¦  «        \  }}|j                             ¦   «         r3|  	                    ||¦  «        \  }}|  	                    ||¦  «        \  }}n|  	                    ||¦  «        \  }}|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)r>   r1   rt   rÙ  rÚ  rŽ   rš   rg   r]  r�   r&   r'   Úcreate_selectr®   )r   r"  rh   rÙ   r¤  rÎ   s         r   r•  zTritonSemantic.wheref  s  € ØŒ?�bœgÒ%Ð%ÝŒMð Lð  {Dô  {Jð  Lð  Lñô ð ð —I’I˜i­¬Ñ1Ô1ˆ	Ø×0Ò0°°A°t¸TÑBÔB‰ˆˆ1àŒ>×"Ò"Ñ$Ô$ð 	CØ×4Ò4°YÀÑBÔB‰LˆI�qØ×,Ò,¨Q°Ñ2Ô2‰DˆAˆqˆqà×4Ò4°YÀÑBÔB‰LˆI�qØ”ˆØ�{Š{˜4œ<×5Ò5°iÔ6FÈÌÐRSÔRZÑ[Ô[Ð]cÑdÔdÐdr   c                ób   — |rt          j        ||¦  «        }n|}|                      ||¦  «        S r)   )r1   rF  r&   )r   rh   rZ   r¥  Úres_tys        r   Úwrap_tensorzTritonSemantic.wrap_tensorz  s8   € Øð 	Ý”] 9¨iÑ8Ô8ˆFˆFð ˆFØ�{Š{˜1˜fÑ%Ô%Ð%r   ÚinputsúSequence[TensorTy]úTuple[TensorTy, ...]c                ó2  ‡ ‡‡‡‡‡— ‰€t          ˆ fd„‰D ¦   «         ¦  «        ŠdŠ‰d         j        j        Št          ‰¦  «        }‰|k     sJ d|› d�¦   «         ‚ˆfd„t	          ‰¦  «        D ¦   «         Št          ˆfd„‰D ¦   «         ¦  «        s
J d¦   «         ‚‰ j                             d„ ‰D ¦   «         ‰¦  «        Š |‰¦  «         ‰                     ¦   «         sJ ‚t          ˆˆˆˆ fd	„t          t          ‰¦  «        ¦  «        D ¦   «         ¦  «        S )
Nc              3  ó^   •K  — | ]'}‰                      ||j        j        gd ¬¦  «        V — Œ(dS )Tr|  N)rh  rV  rO   )rm  Útr   s     €r   rŒ  z+TritonSemantic.reduction.<locals>.<genexpr>„  s<   øè è € Ð^Ð^ÐRS˜4Ÿ<š<¨¨A¬G¬M¨?È˜<ÑMÔMÐ^Ð^Ð^Ð^Ð^Ð^r   r   z&reduction axis must be < inputs rank (r  c                ó&   •— g | ]\  }}|‰k    ¯|‘ŒS r  r  )rm  r™  rg  r*   s      €r   rn  z,TritonSemantic.reduction.<locals>.<listcomp>Š  s"   ø€ ÐAÐAÐA™4˜1˜a°q¸D²y°y�Q°y°y°yr   c              3  ó8   •K  — | ]}|j         j        ‰k    V — Œd S r)   r  )rm  r-  rK  s     €r   rŒ  z+TritonSemantic.reduction.<locals>.<genexpr>‹  s,   øè è € Ð9Ð9¨Q�1”6”< 5Ò(Ð9Ð9Ð9Ð9Ð9Ð9r   z-all reduction inputs must have the same shapec                ó   — g | ]	}|j         ‘Œ
S r  ©r®   ©rm  r-  s     r   rn  z,TritonSemantic.reduction.<locals>.<listcomp>�  s   € Ð/IÐ/IÐ/I¸Q°´Ð/IÐ/IÐ/Ir   c              3  óŒ   •K  — | ]>}‰                      ‰                     |¦  «        ‰|         j        j        ‰¦  «        V — Œ?d S r)   ©r'  Ú
get_resultrg   rŠ   )rm  r™  r(  Ú	reduce_opr¥  r   s     €€€€r   rŒ  z+TritonSemantic.reduction.<locals>.<genexpr>‘  se   øè è € ð uð uØ\]ˆD×Ò˜Y×1Ò1°!Ñ4Ô4°f¸Q´i´nÔ6KÈYÑWÔWðuð uð uð uð uð ur   )
Útuplerg   rK  r^  r–  Úallr'   Úcreate_reduceÚverifyrJ  )r   r(  r*   Úregion_builder_fnÚrankr6  r¥  rK  s   ```  @@@r   Ú	reductionzTritonSemantic.reduction‚  su  øøøøøø€ Øˆ<ÝÐ^Ð^Ð^Ð^ÐW]Ð^Ñ^Ô^Ñ^Ô^ˆFØˆDà�q”	”Ô$ˆÝ�5‰zŒzˆØ�dŠ{ˆ{ˆ{ÐLÀTÐLÐLÐL‰{Œ{ˆ{ØAÐAÐAÐA¥9¨UÑ#3Ô#3ÐAÑAÔAˆ	ÝÐ9Ð9Ð9Ð9°&Ð9Ñ9Ô9Ñ9Ô9ÐjÐjÐ;jÑjÔjÐjà”L×.Ò.Ð/IÐ/IÀ&Ð/IÑ/IÔ/IÈ4ÑPÔPˆ	ØÐ˜)Ñ$Ô$Ð$Ø×ÒÑ!Ô!Ð!Ð!Ð!åð uð uð uð uð uð uð uÝafÕgjÐkqÑgrÔgrÑasÔasðuñ uô uñ uô uð 	ur   Úreversec                óà  ‡ ‡‡‡— ‰d         j         j        Št          ‰¦  «        }| |cxk    r|k     sn J d|› d|› d�¦   «         ‚|dk     r||z  }‰D ]}|j         j        ‰k    s
J d¦   «         ‚Œ‰ j                             d„ ‰D ¦   «         ||¦  «        Š |‰¦  «         ‰                     ¦   «         sJ ‚t          ˆˆˆ ˆfd„t          t          ‰¦  «        ¦  «        D ¦   «         ¦  «        S )Nr   z
scan axis z must be < inputs rank (r  z(all scan inputs must have the same shapec                ó   — g | ]	}|j         ‘Œ
S r  r1  r2  s     r   rn  z3TritonSemantic.associative_scan.<locals>.<listcomp>¥  s   € Ð+EÐ+EÐ+E¸¨A¬HÐ+EÐ+EÐ+Er   c              3  óŒ   •K  — | ]>}‰                      ‰                     |¦  «        ‰|         j        j        ‰¦  «        V — Œ?d S r)   r4  )rm  r™  r(  Úscan_opr   rK  s     €€€€r   rŒ  z2TritonSemantic.associative_scan.<locals>.<genexpr>©  sN   øè è € ÐwÐwÐ_`�T×%Ò% g×&8Ò&8¸Ñ&;Ô&;¸VÀA¼Y¼^Ô=RÐTYÑZÔZÐwÐwÐwÐwÐwÐwr   )rg   rK  r^  r'   Úcreate_scanr:  r7  rJ  )	r   r(  r*   r;  r>  r<  r-  rB  rK  s	   ``     @@r   Úassociative_scanzTritonSemantic.associative_scan˜  s5  øøøø€ à�q”	”Ô$ˆÝ�5‰zŒzˆàˆu˜Ð#Ð#Ò#Ð#˜tÒ#Ð#Ð#Ð#Ð#Ð%W°$Ð%WÐ%WÐPTÐ%WÐ%WÐ%WÑ#Ô#Ð#à�!Š8ˆ8Ø�D‰LˆDàð 	Uð 	UˆAØ”6”< 5Ò(Ð(Ð(Ð*TÑ(Ô(Ð(Ð(à”,×*Ò*Ð+EÐ+E¸fÐ+EÑ+EÔ+EÀtÈWÑUÔUˆØÐ˜'Ñ"Ô"Ð"Ø�~Š~ÑÔÐÐÐåÐwÐwÐwÐwÐwÐwÐwÕdiÕjmÐntÑjuÔjuÑdvÔdvÐwÑwÔwÑwÔwÐwr   ÚsrcÚindexc                ó\  — |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 ]=}||k    rŒ	|j        j        |         |j        j        |         k    sJ d|› d�¦   «         ‚Œ>| j                             |j        |j        |¦  «        }|  	                    ||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)r>   rY   r^  rg   rK  rJ  r'   Úcreate_gatherr®   r'  rŠ   )r   rE  rF  r*   r<  r‹  Úgathers          r   rI  zTritonSemantic.gather¯  sO  € ØŒ{×!Ò!Ñ#Ô#ÐFÐFÐ%FÑFÔFÐFå�3”8”>Ñ"Ô"ˆÝ�5”:Ô#Ñ$Ô$¨Ò,Ð,Ð,Ð.`Ñ,Ô,Ð,àˆu˜Ð#Ð#Ò#Ð#˜tÒ#Ð#Ð#Ð#Ð#Ð%Y°DÐ%YÐ%YÐRVÐ%YÐ%YÐ%YÑ#Ô#Ð#Ø�!Š8ˆ8Ø�D‰LˆDå�t‘”ð 	yð 	yˆAØ�DŠyˆyØØ”:Ô# AÔ&¨#¬(¬.¸Ô*;Ò;Ð;Ð;Ð=xÈ$Ð=xÐ=xÐ=xÑ;Ô;Ð;Ð;à”×+Ò+¨C¬J¸¼ÀdÑKÔKˆØ×Ò ¨¬¬¸¼Ô9IÑJÔJÐJr   c                ó(  — |sdS |^}}t          t          |¦  «        ¦  «        D ]$}|                      |||         ¦  «        \  }||<   Œ%t          t          |¦  «        dz
  ¦  «        D ]$}|                      |||         ¦  «        \  }||<   Œ%|g|¢R S )Nr  r   )rJ  r^  r�   )r   r(  ÚheadÚtailr™  s        r   rp  z TritonSemantic.broadcast_tensorsÅ  s®   € Øð 	Ø�2ØˆˆˆtÝ•s˜4‘y”yÑ!Ô!ð 	Eð 	EˆAØ ×5Ò5°d¸DÀ¼GÑDÔD‰MˆD�$�q‘'�'Ý•s˜4‘y”y 1‘}Ñ%Ô%ð 	Eð 	EˆAØ ×5Ò5°d¸DÀ¼GÑDÔD‰MˆD�$�q‘'�'Øˆ}�tˆ}ˆ}Ðr   úSequence[tl.tensor]Úresult_typesúSequence[tl.dtype]ÚpackúTuple[tl.tensor, ...]c                ój  ‡ ‡‡—  ‰ j         ‰Ž Št          ‰¦  «        dk    s
J d¦   «         ‚ˆfd„|D ¦   «         }‰ j                             d„ ‰D ¦   «         ˆ fd„|D ¦   «         |¦  «        Š |‰¦  «         ‰                     ¦   «         sJ ‚t          ˆˆ fd„t          |¦  «        D ¦   «         ¦  «        S )Nr   z1map_elementwise must have at least 1 input tensorc                óZ   •— g | ]'}‰d          j                              |j        ¦  «        ‘Œ(S )r   )rg   r¯   rŠ   )rm  rÐ  r(  s     €r   rn  z2TritonSemantic.map_elementwise.<locals>.<listcomp>Ô  s/   ø€ ÐYÐYÐYÀb˜˜qœ	œ×6Ò6°r´yÑAÔAÐYÐYÐYr   c                ó   — g | ]	}|j         ‘Œ
S r  r1  r2  s     r   rn  z2TritonSemantic.map_elementwise.<locals>.<listcomp>Ö  s   € Ð&Ð&Ð&˜!ˆQŒXÐ&Ð&Ð&r   c                óD   •— g | ]}|                      ‰j        ¦  «        ‘ŒS r  )r°   r'   )rm  rÐ  r   s     €r   rn  z2TritonSemantic.map_elementwise.<locals>.<listcomp>×  s'   ø€ Ð;Ð;Ð;¨ˆR�XŠX�d”lÑ#Ô#Ð;Ð;Ð;r   c              3  óp   •K  — | ]0\  }}‰                      ‰                     |¦  «        |¦  «        V — Œ1d S r)   )r&   r5  )rm  r™  rÐ  Úelementwise_opr   s      €€r   rŒ  z1TritonSemantic.map_elementwise.<locals>.<genexpr>Ý  sC   øè è € ÐhÐhÁuÀqÈ"�T—[’[ ×!:Ò!:¸1Ñ!=Ô!=¸rÑBÔBÐhÐhÐhÐhÐhÐhr   )rp  r^  r'   Úcreate_map_elementwiser:  r7  r–  )r   r(  rN  rP  r;  rW  s   ``   @r   Úmap_elementwisezTritonSemantic.map_elementwiseÏ  sä   øøø€ à'�Ô'¨Ð0ˆå�6‰{Œ{˜QŠˆˆÐ S‰ŒˆØYÐYÐYÐYÈLÐYÑYÔYˆØœ×<Ò<Ø&Ð&˜vÐ&Ñ&Ô&Ø;Ð;Ð;Ð;¨lÐ;Ñ;Ô;Øñ
ô 
ˆð
 	Ð˜.Ñ)Ô)Ð)Ø×$Ò$Ñ&Ô&Ð&Ð&Ð&åÐhÐhÐhÐhÐhÕPYÐZfÑPgÔPgÐhÑhÔhÑhÔhÐhr   Únum_binsc                óÔ  — t          |j        ¦  «        dk    s
J d¦   «         ‚|j                             ¦   «         s
J d¦   «         ‚|�O|                      ||j        ¦  «        }|j        j                             ¦   «         st          d¦  «        ‚|j	        }|  
                    | j                             |j	        ||¦  «        t          j        t          j        |g¦  «        ¦  «        S )Nr   z histogram only supports 1D inputz%histogram only supports integer inputrv  )r^  rK  r>   rY   r›  rg   rŠ   rÃ  r/   r®   r&   r'   Úcreate_histogramr1   rF  r2   )r   r«   rZ  r  s       r   Ú	histogramzTritonSemantic.histogramä  sÛ   € Ý�5”;ÑÔ 1Ò$Ð$Ð$Ð&HÑ$Ô$Ð$ØŒ{×!Ò!Ñ#Ô#ÐLÐLÐ%LÑLÔLÐLØÐØ×,Ò,¨T°5´;Ñ?Ô?ˆDØ”9Ô#×+Ò+Ñ-Ô-ð GÝ Ð!EÑFÔFÐFØ”;ˆDØ�{Š{˜4œ<×8Ò8¸¼ÀxÐQUÑVÔVÝœ=­¬°H°:Ñ>Ô>ñ@ô @ð 	@r   r�  c                ó  — t          dt          |j        ¦  «        ¦  «        t          |¦  «        k    rt          d¦  «        ‚|j                             dt          j        ||j                             ¦   «         ¦  «        ¦  «         |S )Nr   zAShape of input to multiple_of does not match the length of valuesztt.divisibility)	rð   r^  rK  r/   r®   Úset_attrr   Ú	make_attrÚget_context©r   rh   r�  s      r   Úmultiple_ofzTritonSemantic.multiple_ofï  sl   € Ýˆq•#�a”g‘,”,ÑÔ¥3 v¡;¤;Ò.Ð.ÝÐ`ÑaÔaÐaØ	Œ×ÒÐ+­R¬\¸&À!Ä(×BVÒBVÑBXÔBXÑ-YÔ-YÑZÔZÐZØˆr   c                óø   — t          |j        ¦  «        t          |¦  «        k    rt          d¦  «        ‚|j                             dt          j        ||j                             ¦   «         ¦  «        ¦  «         |S )NzDShape of input to max_contiguous does not match the length of valuesztt.contiguity©r^  rK  r/   r®   r_  r   r`  ra  rb  s      r   Úmax_contiguouszTritonSemantic.max_contiguousõ  sa   € ÝˆqŒw‰<Œ<�3˜v™;œ;Ò&Ð&ÝÐcÑdÔdÐdØ	Œ×Ò˜/­2¬<¸ÀÄ×@TÒ@TÑ@VÔ@VÑ+WÔ+WÑXÔXÐXØˆr   c                óø   — t          |j        ¦  «        t          |¦  «        k    rt          d¦  «        ‚|j                             dt          j        ||j                             ¦   «         ¦  «        ¦  «         |S )NzCShape of input to max_constancy does not match the length of valuesztt.constancyre  rb  s      r   Úmax_constancyzTritonSemantic.max_constancyû  sa   € ÝˆqŒw‰<Œ<�3˜v™;œ;Ò&Ð&ÝÐbÑcÔcÐcØ	Œ×Ò˜.­"¬,°v¸q¼x×?SÒ?SÑ?UÔ?UÑ*VÔ*VÑWÔWÐWØˆr   c                óp   — |                       | j                             ¦   «         t          j        ¦  «        S r)   )r&   r'   Úcreate_barrierr1   r7  )r   s    r   Údebug_barrierzTritonSemantic.debug_barrier  s&   € Ø�{Š{˜4œ<×6Ò6Ñ8Ô8½"¼'ÑBÔBÐBr   ÚprefixÚargsúList[TensorTy]Úhexc                ó‚  — |                      d¦  «        s|r|dz  }|                      d¦  «        s|r|d d…         dz   }t          |¦  «        dk    r|                     d¦  «        sd|z   }d„ |D ¦   «         }d„ |D ¦   «         }|                      | j                             ||||¦  «        t          j        ¦  «        S )Nú r”  r{  r   c                ó   — g | ]	}|j         ‘Œ
S r  r1  ©rm  Úargs     r   rn  z/TritonSemantic.device_print.<locals>.<listcomp>  s   € Ð/Ð/Ð/ 3�C”JÐ/Ð/Ð/r   c                ó@   — g | ]}|j                              ¦   «         ‘ŒS r  )r>   rË   rs  s     r   rn  z/TritonSemantic.device_print.<locals>.<listcomp>  s&   € Ð?Ð?Ð?°3�S”Y×,Ò,Ñ.Ô.Ð?Ð?Ð?r   )Úendswithr^  Ú
startswithr&   r'   Úcreate_printr1   r7  )r   rl  rm  ro  Únew_argsÚ	is_signeds         r   Údevice_printzTritonSemantic.device_print  sÓ   € ð �Š˜sÑ#Ô#ð 	¨ð 	Ø�c‰MˆFØ�Š˜tÑ$Ô$ð 	(¨ð 	(Ø˜C˜R˜C”[ 4Ñ'ˆFÝˆv‰;Œ;˜Š?ˆ? 6×#4Ò#4°SÑ#9Ô#9ˆ?Ø˜6‘\ˆFà/Ð/¨$Ð/Ñ/Ô/ˆØ?Ð?¸$Ð?Ñ?Ô?ˆ	Ø�{Š{˜4œ<×4Ò4°V¸SÀ(ÈIÑVÔVÕXZÔX_Ñ`Ô`Ð`r   r¨   r©   c                óú   — | j         j        j        sd S |�)|                      ||                      |¦  «        ¦  «        }|                      | j                              |j        |¦  «        t          j	        ¦  «        S r)   )
r'   rŸ   Údebugrþ   r  r&   Úcreate_assertr®   r1   r7  )r   r¨   r©   r  s       r   r¤   zTritonSemantic.device_assert  sd   € ØŒ|Ô#Ô)ð 	ØˆFØÐØ—8’8˜D $§)¢)¨D¡/¤/Ñ2Ô2ˆDØ�{Š{˜4œ<×5Ò5°d´kÀ3ÑGÔGÍÌÑQÔQÐQr   c                ó|   — |                       | j                             |j        ¦  «        t          j        ¦  «        S r)   )r&   r'   Úcreate_assumer®   r1   r7  )r   r¨   s     r   ÚassumezTritonSemantic.assume  s*   € Ø�{Š{˜4œ<×5Ò5°d´kÑBÔBÅBÄGÑLÔLÐLr   c                ó  — t          |t          ¦  «        rt          j        |¦  «        }t          |t          j        ¦  «        rÇt          |j        t
          ¦  «        r| j                             |j        ¦  «        S |rFd|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                             ¦   «         s
J d
¦   «         ‚|j        t          j        k    rQ|rO| j                             |j        | j                             ¦   «         |j                             ¦   «         ¦  «        S |j        t          j        k    r|s
J d¦   «         ‚|j        S J dt'          |¦  «        › �¦   «         ‚)Nrn   ro   z@Block pointers only support 64 bit `shape/strides`, got a value z which is out of the rangerl   rm   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/offsetsFzzBlock 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: )rb   r+   r1   rc   rO   rH   r'   r…  rÊ  rÞ  r&   rV  r>   rY   rv   r±   r®   Úget_int64_tyrË   rg   )r   Úelemr)  s      r   Ú_convert_elem_to_ir_valuez(TritonSemantic._convert_elem_to_ir_value  s3  € Ý�d�CÑ Ô ð 	&Ý”< Ñ%Ô%ˆDÝ�d�BœLÑ)Ô)ð 	Ý˜$œ*¥dÑ+Ô+ð 9Ø”|×,Ò,¨T¬ZÑ8Ô8Ð8Øð :Ø ¤Ð3Ð3Ò3Ð3¨eÒ3Ð3Ð3Ð3Ð3ð 6JØ#'¤:ð6Jð 6Jð 6JÑ3Ô3Ð3à”|×-Ò-¨d¬jÑ9Ô9Ð9à ¤Ð3Ð3Ò3Ð3¨eÒ3Ð3Ð3Ð3Ð3ð 6JØ#'¤:ð6Jð 6Jð 6JÑ3Ô3Ð3à”|×-Ò-¨d¬jÑ9Ô9Ð9Ý˜�bœiÑ(Ô(ð 		Ø”:Ô# qÒ(Ð(Ð(Ð*VÑ(Ô(Ð(Ø”:×$Ò$Ñ&Ô&ÐbÐbÐ(bÑbÔbÐbØŒz�RœXÒ%Ð%¨+Ð%Ø”|×3Ò3°D´KÀÄ×AZÒAZÑA\ÔA\Ø48´J×4LÒ4LÑ4NÔ4NñPô Pð Pà”�rœxÒ'Ð'°Ð'ðWð Wñ Wô Wð Wà”;ÐØXÐXÍDÐQUÉJÌJÐXÐXÑXÔXÐXr   c                ór   ‡ ‡— t          |d¦  «        rˆˆ fd„|D ¦   «         S ‰                      |‰¦  «        gS )NÚ__iter__c                ó<   •— g | ]}‰                      |‰¦  «        ‘ŒS r  )r…  )rm  r„  r)  r   s     €€r   rn  z8TritonSemantic._convert_to_ir_values.<locals>.<listcomp>8  s)   ø€ Ð\Ð\Ð\È$�D×2Ò2°4¸ÑEÔEÐ\Ð\Ð\r   )Úhasattrr…  )r   Ú	list_liker)  s   ` `r   r,  z$TritonSemantic._convert_to_ir_values6  sN   øø€ Ý�9˜jÑ)Ô)ð 	]Ø\Ð\Ð\Ð\Ð\ÐR[Ð\Ñ\Ô\Ð\Ø×.Ò.¨y¸+ÑFÔFÐGÐGr   rð  ÚbaseÚstridesr+  úList[tl.constexpr]útl.tensor_descriptorc           	     ó2  ‡ — t          |¦  «        }d|cxk    rdk    sn t          d|› d�¦  «        ‚t          |¦  «        |k    r"t          d|› dt          |¦  «        › �¦  «        ‚t          |¦  «        |k    r"t          d|› dt          |¦  «        › �¦  «        ‚t          |j        t          j        ¦  «        sJ d	|j        › �¦   «         ‚|j        j        j        d
z  }t	          j        |d         ¦  «        }||z  dk     rt          d|› d|› d||z  › d�¦  «        ‚t	          j        |d         ¦  «        }	|	dk    rt          d|	› �¦  «        ‚ˆ fd„|D ¦   «         }ˆ fd„|D ¦   «         }t	          j	        |¦  «        }t          |j
        t          j        ¦  «        sJ d	|j
        › �¦   «         ‚t	          j        |j
        j        |¦  «        }
|j        }|j
        j                             ¦   «         }‰                      |¦  «        }|j
        j                             ¦   «         r$|t           j        j        k    rt          d¦  «        ‚‰ j                             |d„ |D ¦   «         d„ |D ¦   «         |||¦  «        }t	          j        ||||
¦  «        S )Nr   é   z Expected 1 <= ndim <= 5 but got z dimensionsz	Expected z strides but got zExpected block_shape to have z dimensions but got z!base must be a pointer type, got rc  r{  r}  zRDescriptor block shape must have at least 16 bytes in the last dimension, but got z * z = z bytesz-Tensor descriptor last dim must be 1 but got c                óP   •— g | ]"}‰                      |t          j        ¦  «        ‘Œ#S r  )rW  r1   r2   ©rm  rh   r   s     €r   rn  z9TritonSemantic.make_tensor_descriptor.<locals>.<listcomp>P  s+   ø€ Ð>Ð>Ð>°1�×!Ò! !¥R¤XÑ.Ô.Ð>Ð>Ð>r   c                ót   •— g | ]4}‰                      t          j        |¦  «        t          j        ¦  «        ‘Œ5S r  )rW  r1   rl  rv   r’  s     €r   rn  z9TritonSemantic.make_tensor_descriptor.<locals>.<listcomp>Q  s6   ø€ Ð[Ð[Ð[Èa�4×#Ò#¥BÔ$;¸AÑ$>Ô$>ÅÄÑIÔIÐ[Ð[Ð[r   z8Padding option `nan` is not supported for integer blocksc                ó   — g | ]	}|j         ‘Œ
S r  r1  ©rm  rg  s     r   rn  z9TritonSemantic.make_tensor_descriptor.<locals>.<listcomp>`  s   € ÐIbÐIbÐIbÐWXÈ!Ì(ÐIbÐIbÐIbr   c                ó   — g | ]	}|j         ‘Œ
S r  r1  r•  s     r   rn  z9TritonSemantic.make_tensor_descriptor.<locals>.<listcomp>a  s   € Ð<WÐ<WÐ<WÈ!¸Q¼XÐ<WÐ<WÐ<Wr   )r^  r/   rb   r>   r1   r  r  r³  rl  Ú_unwrap_shaperg   rF  r®   rË   r÷  rY   r   rò  rô  r'   Úcreate_make_tensor_descriptorÚtensor_descriptor)r   r‹  rK  rŒ  r+  rõ  r/  Ú	elem_sizeÚcontig_dim_sizeÚlast_striderg   Úbase_handleÚis_signed_intrö  r®   s   `              r   Úmake_tensor_descriptorz%TritonSemantic.make_tensor_descriptor;  s  ø€ å�5‰zŒzˆØ�T��’�˜Q’���ÝÐQÀÐQÐQÐQÑRÔRÐRÝˆw‰<Œ<˜4ÒÐÝÐN¨ÐNÐNÅÀGÁÄÐNÐNÑOÔOÐOÝˆ{ÑÔ˜tÒ#Ð#ÝÐe¸TÐeÐeÕWZÐ[bÑWcÔWcÐeÐeÑfÔfÐfÝ˜$œ*¥b¤oÑ6Ô6ÐhÐhÐ8hÐ\`Ô\fÐ8hÐ8hÑhÔhÐhØ”JÔ)Ô<ÀÑAˆ	ÝÔ1°+¸b´/ÑBÔBˆØ˜YÑ&¨Ò+Ð+Ýð kÐetð  kð  kð  zCð  kð  kð  HWð  Zcñ  Hcð  kð  kð  kñô ð õ Ô-¨g°b¬kÑ:Ô:ˆØ˜!ÒÐÝÐZÈ[ÐZÐZÑ[Ô[Ð[à>Ð>Ð>Ð>¸Ð>Ñ>Ô>ˆØ[Ð[Ð[Ð[ÐSZÐ[Ñ[Ô[ˆõ Ô& {Ñ3Ô3ˆå˜$œ)¥R¤_Ñ5Ô5ÐfÐfÐ7fÐ[_Ô[dÐ7fÐ7fÑfÔfÐfÝŒ}˜TœYÔ1°;Ñ?Ô?ˆØ”kˆØœ	Ô,×:Ò:Ñ<Ô<ˆà×-Ò-¨nÑ=Ô=ˆàŒ9Ô×&Ò&Ñ(Ô(ð 	Y¨W½Ô8IÔ8QÒ-QÐ-QÝÐWÑXÔXÐXà”×;Ò;¸KÐIbÐIbÐ\aÐIbÑIbÔIbØ<WÐ<WÈwÐ<WÑ<WÔ<WÐYdÐfsØ<CñEô Eˆõ Ô# F¨E°7¸DÑAÔAÐAr   )r*   r+   r,   r   )r8   r9   r:   r9   r,   r9   )r8   r9   rG   rH   r:   r9   rI   rH   rJ   rH   r,   r9   )T)r   r9   r   r9   r}   rH   r,   r~   )FFTF)rƒ   r„   r…   r„   r,   r†   )rƒ   r   r…   r   r›   rœ   )r«   r„   r¬   r„   r    rH   r,   r   )r«   r„   r¬   r„   r,   r   )r«   r„   r¬   r„   rÐ   rH   r,   r   )rh   r   rÙ   r   rÚ   rÛ   )rh   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"  )rB  r+   rC  r+   rÎ   r"  r,   r   )r>   r9   r,   r   )rK  rX  r>   r9   r,   r   )rO   r   rK  rX  r,   r   )rO   r   r,   r   )r«   r   rc  rX  rd  rH   r,   r   )r«   r   r*   r+   r,   r   )rƒ   r   r…   r   rd  rH   r,   r   )rx  r   ry  r   r,   r   )rx  r   r,   r†   )r«   r   r‡  rˆ  r,   r   )r«   r   rK  rˆ  r,   r   )rƒ   r   r…   r   r,   r   )r¨  r©  )r«   r   r±  r9   r,   r   r)   )r«   r   r±  r9   rº  r©  r,   r   )r  r   r  r  r¬   r  r  r   rõ  rŸ  rÝ  rŸ  rì  rŸ  r  rH   r,   r   )r"  r#  rÝ  rŸ  rì  rŸ  r,   r   )r"  r#  rO   r   r,   r~   )r"  r#  rO   r   r,   r   )rÝ  rŸ  rì  rŸ  r,   r   )r  r   rq  r   r  r  rÝ  rŸ  rì  rŸ  r,   r   )r  r   r{  r   rq  r   r  rŸ  r  rŸ  r,   r   )
r  r   rq  r   r  r   r€  rŸ  r,   r�  )rh   r   r,   r   )r  r   rq  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Ÿ  )rq  r   rï  rŸ  )rƒ   r   r  r   r  rŸ  r…   r   r  r  r  rŸ  rÆ  r  r  rH   r  rH   r	  rH   rÈ  r9   r,   r   )r"  r   rh   r   rÙ   r   r,   r   )r(  r)  r*   r+   r,   r*  )r(  r)  r*   r+   r>  rH   r,   r*  )rE  r   rF  r   r*   r+   r,   r   )r(  rM  rN  rO  rP  r+   r,   rQ  )r«   r   rZ  r+   r  r  r,   r   )rh   r   r�  rX  r,   r   )r,   r   )rl  rŸ  rm  rn  ro  rH   r,   r   )r¨   r   r©   rŸ  r  r  r,   r   )rð  )r‹  r   rK  rn  rŒ  rn  r+  r�  rõ  rŸ  r,   rŽ  )yr    r!   r"   r1   r&   Ú__annotations__Úlangr   r4   r7   rF   r\   rj   re   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£   r4  r¢   r<  rA  rM  rf   rW  r[  rZ  rb  rh  rr  rw  r�  r†  r‘  r›  r�   r°  r  rŽ   rß  rä  rî  r÷  r  r  r!  r0  r3  r8  r@  rI  rK  rN  rR  rU  rY  r\  rk  rn  rs  rz  r  rˆ  r�  r¢  r¥  r¨  r¬  r°  r³  r·  rÅ  rî  ró  rþ  r  r  r!  r•  r'  r=  rD  rI  rp  rY  r]  rc  rf  rh  rk  r{  r¤   r�  r…  r,  rŸ  r  r   r   r%   r%      sí	  € € € € € € ØœY€FÐ&Ð&Ð&Ñ&Ø€DàÐÐÑðð ð ðOð Oð Oð Oð
Qð Qð Qð QðDð Dð Dð Dð05ð 05ð 05ð 05ðd	ð 	ð 	ð 	ðKð Kð KðF	@ð 	@ð 	@ð 	@ð aeØ05ð#ð #ð #ð #ð #ðJ,ð ,ð ,ð ,ð$>ð >ð >ð >ð>8ð 8ð 8ð 8ð"8ð 8ð 8ð 8ð]ð ]ð ]ð ]ð2>ð >ð >ð >ð
,ð 
,ð 
,ð 
,ð8ð 8ð 8ð 8ð.9ð 9ð 9ð 9ð"9ð 9ð 9ð 9ð"	að 	að 	að 	aðð ð ð ð\ð \ð \ð \ð[ð [ð [ð [ð\ð \ð \ð \ð'ð 'ð 'ð 'ð&ð &ð &ð &ð"ð "ð "ð "ð
]ð ]ð ]ð ]ð]ð ]ð ]ð ]ð\ð \ð \ð \ðð ð ð ð)ð )ð )ð )ð$ð $ð $ð $ð/ð /ð /ð /ð8ð 8ð 8ð 8ð8ð 8ð 8ð 8ð8ð 8ð 8ð 8ð8ð 8ð 8ð 8ð	8ð 	8ð 	8ð 	8ð	8ð 	8ð 	8ð 	8ð GKð Zð Zð Zð Zð Zð Zð$)ð )ð )ð )ð2ð 2ð 2ð 2ðAð Að Að Aðhð hð hð hðSð Sð Sð Sðfð fð fð fðXð Xð Xð XñVð Vð Vð Vñð ð ð ñ0
ð 
ð 
ð 
ñTð Tð Tð TñWð Wð Wð Wñ 3ð 3ð 3ð 3ñrtð tð tð tñjð jð jð jñ$i9ñ i9ð i9ð i9ð i9ð^ð ð ðð ð ð	ð 	ð 	ð	ð 	ð 	ðð ð ðð ð ñ=ð =ð =ð =ñ~
/ð 
/ð 
/ð 
/ñJð Jð Jð Jñnð nð nð nñuð uð uð uð@ð @ð @ð[ð [ð [ñ
uð uð uð uñuð uð uð uñuð uð uð uñuð uð uð uñuð uð uð uñ$ð $ð $ð $ð6*ð *ð *ð *ð0	ð 	ð 	ñ.$ð .$ð .$ð .$ñhuð uð uð uñð ð ð ñ6+ð +ð +ð +ñ$)ð $)ð $)ð $)ñL$)ð $)ð $)ð $)ñL%ð %ð %ð %ñyð yð yð yñxð xð xð xñyð yð yð yñð ð ð ð
<ð 
<ð 
<ñbyð byð byð byñHð ð ð ñ0ð 0ð 0ð 0ð$	Tð 	Tð 	Tð
mð 
mð 
mñ8tð 8tð 8tð 8tñ|eð eð eð eð(&ð &ð &ñuð uð uð uñ,xð xð xð xñ.Kð Kð Kð Kð,ð ð ñið ið ið iñ*	@ð 	@ð 	@ð 	@ñð ð ð ñð ð ð ñð ð ð ñCð Cð Cð Cñað að að añRð Rð Rð RñMð Mð Mð MðYð Yð Yð4Hð Hð Hð Hð W]ñ(Bñ (Bð (Bð (Bð (Bð (Bð (Br   r%   )Ú
__future__r   rx   rÙ  Útypingr   r   r   r   r   r	   r
   rˆ   Útriton.runtimer   Ú_C.libtritonr   r^  r   r1   r   r   Ú	Exceptionr   r%   r  r   r   ú<module>r§     s9  ðØ "Ð "Ð "Ð "Ð "Ð "Ø €€€Ø €€€à JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JØ €€€à !Ð !Ð !Ð !Ð !Ð !à Ð Ð Ð Ð Ð Ø Ð Ð Ð Ð Ð à€GˆC�L„L€Øˆ7�:ÑÔ€ðFð Fð Fð Fð F 	ñ Fô Fð FðIBð IBð IBð IBð IB�W˜XÔ&ñ IBô IBð IBð IBð IBr   