§
    9˜ºh¶z ã                  óÚ   — d dl m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      €ú`/var/www/html/services/rag-service/venv/lib/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å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„Z/dõdX„Z0död[„Z1d÷d]„Z2død`„Z3dùda„Z4dúdb„Z5dûde„Z6düdf„Z7dýdi„Z8dþdj„Z9dÿdk„Z:�d dn„Z;�ddp„Z<�d�ddr„Z=ds„ Z>dt„ Z?du„ Z@dv„ ZAdw„ ZBdx„ ZCdy„ ZDdz„ ZEd{„ ZF�dd†„ZG�dd‰„ZH�ddŠ„ZI�dd‹„ZJ�ddŒ„ZKd�„ ZLdŽ„ ZM�dd�„ZN�dd�„ZO�dd‘„ZP�dd’„ZQ�dd“„ZR�dd”„ZS�d	d•„ZTd–„ ZUd—„ 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d§„Z`�dd¨„Zad©„ Zb�dd®„Zc�dd°„Zd�dd±„Ze�ddº„Zf�dd¼„Zgd½„ Zh�ddÁ„Zi�ddÃ„Zj�ddÆ„Zk�ddÈ„Zl�ddÊ„Zm�ddË„Zn�ddÌ„Zo�ddÍ„Zp�ddÒ„Zq�ddÕ„Zr�ddÖ„Zsd×„ ZtdädØ„Zu�ddÚ„Zv�ddÛ„Zw�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_impl4   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_implC   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   TÚ
check_typec                ó  — t          |t          ¦  «        r8|                      | j                             |¦  «        t
          j        ¦  «        S t          |t          ¦  «        ržd|cxk    rdk     rn nt
          j        }njd|cxk    rdk     rn nt
          j	        }nMd|cxk    rdk     rn nt
          j
        }n0d|cxk    rdk     rn nt
          j        }nt          d|› d�¦  «        ‚|                      ||¬	¦  «        S t          |t          ¦  «        r~d
}dddz  z  }t          d         |¦  «        }|t          d¦  «        k    s|dk    s||k    s||cxk    r|k    rn nt
          j        }nt
          j        }|                      ||¬	¦  «        S t          |t
          j        ¦  «        r|                      |j        ¦  «        S t          || j        ¦  «        r|S |r#t+          d|› dt-          |¦  «        › d�¦  «        ‚|S )Né   €ì        l        ìûÿÿÿ         ì            l            zNonrepresentable integer ú.©r>   g      8g   àÿÿÿ?r   é   ÚabsÚinfg        zcannot convert z	 of type z
 to tensor)Ú
isinstancerH   r&   r'   Úget_int1r1   Úint1r+   r2   Úuint32Úint64Úuint64r/   Úscalar_constantÚfloatÚ__builtins__rR   rT   Ú	constexprÚ	to_tensorrO   rA   Útype)r   Úxr]   r>   Úmin_float32Úmax_float32Úabs_xs          r   rr   zTritonSemantic.to_tensoru   sG  € Ý�a�ÑÔð 	Ø—;’;˜tœ|×4Ò4°QÑ7Ô7½¼ÑAÔAÐAå˜�3ÑÔð 	Ø˜Ð"Ð"Ò"Ð"˜UÒ"Ð"Ð"Ð"Ð"Ýœ��Ø˜!Ð#Ð#Ò#Ð#˜eÒ#Ð#Ð#Ð#Ð#Ýœ	��Ø˜1Ð$Ð$Ò$Ð$˜uÒ$Ð$Ð$Ð$Ð$Ýœ��Ø˜!Ð#Ð#Ò#Ð#˜eÒ#Ð#Ð#Ð#Ð#Ýœ	��å Ð!A¸QÐ!AÐ!AÐ!AÑBÔBÐBØ×'Ò'¨°Ð'Ñ7Ô7Ð7Ý˜�5Ñ!Ô!ð 	Ø!ˆKØ%¨¨C©Ñ/ˆKÝ  Ô'¨Ñ*Ô*ˆEØ�˜e™œÒ$Ð$Ø˜Š|ˆ|Ø�AŠvˆvØ˜eÐ2Ð2Ò2Ð2 {Ò2Ð2Ð2Ð2Ð2Ýœ
��åœ
�Ø×'Ò'¨°Ð'Ñ7Ô7Ð7å˜�2œ<Ñ(Ô(ð 	Ø—>’> !¤'Ñ*Ô*Ð*Ý˜˜4œ;Ñ'Ô'ð 	ØˆHØð 	OÝÐM¨aÐMÐM½$¸q¹'¼'ÐMÐMÐMÑNÔNÐNØˆr   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   rx   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 rd   )rh   ÚnumbersÚNumberrr   rs   Úscalarr}   r{   r\   Úis_int_unsignedr/   rY   Úget_int_min_valueÚget_int_max_valuern   Ú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 ‚|                     ¦   «         sJ ‚|                      |t          j	        ¦  «        }|                      |t          j	        ¦  «        } |||d¦  «        }| 
                    ¦   «         }|                      |t          j	        ¦  «        }|                     ¦   «         }|                      |t          j	        ¦  «        }|                      |                      ||¦  «        |                      ||¦  «        ¦  «        }	d|j        › d|j        › �}
|                      |	|
¦  «         d S )Né@   Fr+   z! overflow detected for operation )rs   r…   r<   r'   ÚoptionsÚsanitize_overflowrY   r‰   r1   rl   rˆ   rn   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Î   sW  € ØŒ8Œ?Ô'¨2Ò-Ð-°T´\Ô5IÔ5[Ð-ØˆFØ”X”_ˆ
Ø”X”_ˆ
Ø˜ZÒ'Ð'Ð'Ð'Ø× Ò Ñ"Ô"Ð"Ð"Ð"Ø�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 Ñ%Ô%Ð%Ð%Ð%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•   rs   r…   r{   rA   Úhandler>   r†   r<   Úwith_element_tyr1   rl   Ú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•   rs   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•   rs   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•   rs   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•   rs   r…   rY   rF   r‰   Úis_int_signedr&   r'   Úcreate_sdivr©   Úcreate_udivrA   )r   r¦   r§   r±   r²   Úret_tys         r   ÚfloordivzTritonSemantic.floordiv7  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                óZ  — |j         j        }|j         j        }|                     ¦   «         r|                     ¦   «         st          d¦  «        ‚|                      ||dddd¦  «        \  }}| j                             |j        |j        ¦  «        }|                      ||j         ¦  «        S )Nz4both operands of fdiv must have floating scalar typeFT)	rs   r…   r|   rA   r•   r'   rÃ   r©   r&   )r   r¦   r§   rË   r±   r²   r    s          r   ÚfdivzTritonSemantic.fdivE  s¢   € Øœ*Ô+ˆØœ*Ô+ˆØ×*Ò*Ñ,Ô,ð 	T°O×4OÒ4OÑ4QÔ4Qð 	TÝÐRÑSÔSÐSØ×8Ò8¸ÀÀuÈeÐUZÐ\`ÑaÔa‰ˆˆuØŒl×&Ò& u¤|°U´\ÑBÔBˆØ�{Š{˜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•   rs   r…   r|   r&   r'   Úcreate_fremr©   rY   r=   rA   r   rÆ   Úcreate_sremÚcreate_urem)r   r¦   r§   rZ   r²   s        r   ÚmodzTritonSemantic.modN  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   rt   Ú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©   rs   ÚNONEÚcreate_minnumfr/   rÆ   Úcreate_minsir†   Úcreate_minuirA   ©r   rt   rÓ   rÔ   r>   s        r   ÚminimumzTritonSemantic.minimume  ó`  € Ø×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©   rs   rÜ   Úcreate_maxnumfr/   rÆ   Úcreate_maxsir†   Úcreate_maxuirA   rà   s        r   ÚmaximumzTritonSemantic.maximumv  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©   rs   rA   )r   rt   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•   rs   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©   rs   ©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©   rs   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©   rs   rô   s      r   Úxor_zTritonSemantic.xor_«  rõ   r   c                ó  — |j                              ¦   «         s |                      |t          j        ¦  «        }|j                              ¦   «         s |                      |t          j        ¦  «        }|                      ||¦  «        S r)   )rs   Úis_int1Úbitcastr1   rj   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)   )rs   rý   rþ   r1   rj   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)   )rs   rý   rþ   r1   rj   Ú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©   rs   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©   rs   rô   s      r   ÚashrzTritonSemantic.ashrÆ  r	  r   c                ó°   — |                       ||¦  «        \  }}|                      | j                             |j        |j        ¦  «        |j        ¦  «        S r)   )rñ   r&   r'   Ú
create_shlr©   rs   rô   s      r   ÚshlzTritonSemantic.shlÊ  rõ   r   c                ó   — |S r)   © r  s     r   ÚpluszTritonSemantic.plusÒ  s   € Øˆr   c                óJ  — |j         j        }|                     ¦   «         r't          d|                     ¦   «         z   dz   ¦  «        ‚|                      | j                             |                     | j        ¦  «        ¦  «        |¦  «        }|  	                    ||d¦  «        S )Nz$wrong type argument to unary minus (ú)T)
rs   r…   r{   r/   r   r&   r'   Ú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Ø�[Š[˜œ×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  )rs   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)   )rs   rª   r1   rj   )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•   rs   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•   rs   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•   rs   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•   rs   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•   rs   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•   rs   r…   r|   r&   r'   Úcreate_fcmpUNEr©   r  rY   Úcreate_icmpNErA   r"  s       r   Ú	not_equalzTritonSemantic.not_equal-  r6  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)rh   r+   r/   rH   r1   Ú
block_typer2   r«   r'   r&   Úcreate_make_range)	r   r;  r<  rÉ   Úis_start_int64Úis_end_int64ÚrangeÚshapeÚ	ret_ty_irs	            r   ÚarangezTritonSemantic.arange<  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(t	          | j        d|j        › �¦  «        } ||¦  «        }|                      ||¦  «        S )Nz2dtype must be specified when value is not a tensorr   Úget_)r/   r'   r  r«   ÚgetattrÚnamer&   )r   rO   r>   Úget_value_fns       r   rn   zTritonSemantic.scalar_constantN  sƒ   € àˆ=ÝÐQÑRÔRÐRØ�AŠ:ˆ:Ø”L×/Ò/°·²¸D¼LÑ0IÔ0IÑJÔJˆ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)rh   r1   r&   ÚnumelrO   r‰   rn   )r   rO   r>   s      r   Úmake_scalarzTritonSemantic.make_scalarY  s_   € Ý�e�RœYÑ'Ô'ð 	+Ø”;Ô$¨Ò)Ð)Ð)Ð+GÑ)Ô)Ð)Ø—9’9˜U EÑ*Ô*Ð*à×#Ò# E¨5Ñ1Ô1Ð1r   rD  ú	List[int]c                óV   — |                       |                      ||¦  «        |¦  «        S r)   )ÚsplatrN  )r   rD  rO   r>   s       r   ÚfullzTritonSemantic.full`  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   )rs   Úis_blockÚlenr1   r?  r>   r&   r'   Úcreate_splatr«   r©   )r   rO   rD  rÉ   s       r   rQ  zTritonSemantic.splatg  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   Ú	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)
rs   rM  r/   r1   r?  r…   r&   r'   Úcreate_reshaper©   )r   r¦   rW  rX  rM  ÚsrÉ   s          r   ÚreshapezTritonSemantic.reshapen  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)Ú.0rt   s     r   ú
<listcomp>z.TritonSemantic.expand_dims.<locals>.<listcomp>x  s#   € ÐEÐEÐE°A•RÔ,¨QÑ/Ô/ÐEÐEÐEr   r   ©rD  )rD  Úinsertrs   rT  rQ  r1   r?  r…   r&   r'   Úcreate_expand_dimsr©   )r   r¦   r*   rW  rÉ   s        r   Úexpand_dimszTritonSemantic.expand_dimsw  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                ó2  — |s
J d¦   «         ‚t          |j        ¦  «        dk    sJ ‚t          j        |j        j        |j        d         |j        d         z   g¦  «        }|                      | j                             |j	        |j	        ¦  «        |¦  «        S )Nz;current implementation of `cat` always may reorder elementsr   r   )
rU  rD  r1   r?  rs   r…   r&   r'   Ú
create_catr©   )r   r~   r€   rX  Úret_types        r   ÚcatzTritonSemantic.cat�  sƒ   € ØÐYÐYÐYÑYÔYÐYÝ�3”9‰~Œ~ Ò"Ð"Ð"Ð"Ý”= ¤¤°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©rX  )rŠ   rD  rf  rh   r1   rq   r?  rs   r…   r&   r'   Úcreate_joinr©   r\  )r   rk  rl  Ú
was_rank_1ÚtwoÚ	new_shaperi  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                óx  — t          |j        ¦  «        dk    sJ ‚t          j        |j        d         ¦  «        dk    sJ ‚|j        d d…         }t          j        |j        j        |¦  «        }| j                             |j	        ¦  «        \  }}|  
                    ||¦  «        |  
                    ||¦  «        fS )Nr   rn  r   )rU  rD  r1   r`  r?  rs   r…   r'   Úcreate_splitr©   r&   )r   rk  rs  ri  ÚoutLHSÚoutRHSs         r   ÚsplitzTritonSemantic.splitŸ  s¨   € Ý�A”G‘”˜qÒ Ð Ð Ð ÝÔ'¨¬°¬Ñ4Ô4¸Ò9Ð9Ð9Ð9à”G˜C˜R˜C”Lˆ	Ý”= ¤¤°	Ñ:Ô:ˆØœ×2Ò2°1´8Ñ<Ô<‰ˆ�à�KŠK˜ Ñ)Ô)Ø�KŠK˜ Ñ)Ô)ð
ð 	
r   Údimsú
Tuple[int]c                óÖ  ‡— t          ‰j        ¦  «        t          |¦  «        k    rt          d¦  «        ‚t          d„ |D ¦   «         ¦  «        t	          t          t          |¦  «        ¦  «        ¦  «        k    rt          d|› �¦  «        ‚t          j        ‰j        j	        ˆfd„|D ¦   «         ¦  «        }|  
                    | j                             ‰j        |¦  «        |¦  «        S )Nz5permute dims must have the same length as input shapec              3  ó>   K  — | ]}t          j        |¦  «        V — Œd S r)   r_  )ra  Ú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  rc  )ra  r~  r¦   s     €r   rb  z*TritonSemantic.permute.<locals>.<listcomp>±  s   ø€ Ð4RÐ4RÐ4RÈ°U´[À´^Ð4RÐ4RÐ4Rr   )rU  rD  r/   ÚsortedÚlistrC  r1   r?  rs   r…   r&   r'   Úcreate_transr©   )r   r¦   rz  ri  s    `  r   ÚpermutezTritonSemantic.permute«  sÐ   ø€ ÝˆuŒ{ÑÔ�s 4™yœyÒ(Ð(ÝÐTÑUÔUÐUÝÐ;Ð;°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: z, r   z3Cannot broadcast, the expanded size of the tensor (z ) must match the existing size (z) at non-singleton dimension ú: )rs   rT  rQ  Úget_block_shapesrU  r/   Ú	enumerater1   r?  r…   r&   r'   Úcreate_broadcastr©   )r   r¦   rD  Ú	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        ¦  «        |¦  «        }�n8|                     ¦   «         �r#|                     ¦   «         �r| 	                    ¦   «         }| 	                    ¦   «         }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 ‚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   z?Cannot make_shape_compatible: incompatible dimensions at index r†  r   )rs   rT  rª   r…   r&   r'   rV  r«   r©   r‡  rU  rC  re  r1   r?  Ú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Ð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 )rs   rT  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 )1rs   r…   rT  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   Úcreate_fp_truncÚcreate_fp_extrY   r<   r=   rÆ   Úis_boolr>   r  r:  r¬   Úis_standard_floatingÚcreate_fp_to_siÚcreate_fp_to_uiÚcreate_ui_to_fpÚcreate_si_to_fpr{   Úcreate_ptr_to_intrl   Ú	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        }|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 )Nz.caú.cgz.cvúCache modifier ú not supported)r   ÚCACHE_MODIFIERrÜ   ÚCAÚCGÚCVr/   ©r   Úcache_modifierÚcaches      r   Ú_str_to_load_cache_modifierz*TritonSemantic._str_to_load_cache_modifier†  s}   € ÝÔ!Ô&ˆØð 	SØ Ò&Ð&ÝÔ)Ô,��Ø 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Æ  z.csz.wtrÇ  rÈ  )r   rÉ  rÜ   ÚWBrË  ÚCSÚWTr/   rÍ  s      r   Ú_str_to_store_cache_modifierz+TritonSemantic._str_to_store_cache_modifier“  s•   € ÝÔ!Ô&ˆØð 
	SØ Ò&Ð&ÝÔ)Ô,��Ø 5Ò(Ð(ÝÔ)Ô,��Ø 5Ò(Ð(ÝÔ)Ô,��Ø 5Ò(Ð(ÝÔ)Ô,��å Ð!Q°>Ð!QÐ!QÐ!QÑRÔRÐ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   c                ón  — |r²t          |d¦  «        s|g}d„ |D ¦   «         }|D ]5}t          |t          ¦  «        rd|cxk    rt          |¦  «        k     sn J ‚Œ6t          |¦  «        dk    sJ ‚t          |¦  «        t          t	          |¦  «        ¦  «        k    s
J d¦   «         ‚t          |¦  «        S dS )NÚ__iter__c                óT   — g | ]%}t          |t          j        ¦  «        r|j        n|‘Œ&S r  ©rh   r1   rq   rO   ©ra  Úelems     r   rb  z?TritonSemantic._canonicalize_boundary_check.<locals>.<listcomp>Ø  s1   € ÐpÐpÐpÐY]­J°t½R¼\Ñ,JÔ,JÐT˜dœj˜jÐPTÐpÐpÐpr   r   z'Duplicate dimension in `boundary_check`r  )Úhasattrrh   r+   rU  Úsetr�  )r   Úboundary_checkÚblock_shapeÚdims       r   Ú_canonicalize_boundary_checkz+TritonSemantic._canonicalize_boundary_checkÔ  sæ   € Øð 	*Ý˜>¨:Ñ6Ô6ð 2Ø"0Ð!1�ØpÐpÐaoÐpÑpÔpˆNØ%ð Lð L�Ý! #¥sÑ+Ô+ÐK°°SÐ0KÐ0KÒ0KÐ0K½3¸{Ñ;KÔ;KÒ0KÐ0KÐ0KÐ0KÐ0KÐ0KÐ0KÝ�~Ñ&Ô&¨Ò*Ð*Ð*Ð*Ý�~Ñ&Ô&­#­c°.Ñ.AÔ.AÑ*BÔ*BÒBÐBÐBÐDmÑBÔBÐBÝ˜.Ñ)Ô)Ð)Øˆrr   c	           
     óÆ  — |€|�t          d¦  «        ‚|j        j        j        }	|	t          j        k    s
J d¦   «         ‚|	                     ¦   «         r$|t          j        j        k    rt          d¦  «        ‚|j        j        }
|  	                    ||
 
                    ¦   «         ¦  «        }|                      | j                             |j        |||||¦  «        |
¦  «        S )NúK`mask` and `other` arguments cannot be specified for loading block pointersú4`tl.int1` should be rewritten in `tl.make_block_ptr`z@Padding option `nan` is not supported for integer block pointers)r/   rs   Ú
element_tyr1   rj   rY   r   rã  rå  r  r‡  r&   r'   Úcreate_tensor_pointer_loadr©   )r   ÚptrÚmaskr§   r
  rç  rÏ  rÞ  Úis_volatileÚelt_tyr£  s              r   Ú_load_block_pointerz"TritonSemantic._load_block_pointerà  sä   € ð Ð˜uÐ0ÝÐjÑkÔkÐkà”Ô$Ô/ˆØ�œÒ Ð Ð Ð"XÑ Ô Ð Ø�=Š=‰?Œ?ð 	a˜w­"Ô*;Ô*CÒCÐCÝÐ_Ñ`Ô`Ð`ð ”Ô$ˆð ×:Ò:¸>È6×KbÒKbÑKdÔKdÑeÔeˆð �{Š{ØŒL×3Ò3°C´JÀÐPWÐY^Ð`hÐjuÑvÔvØñô ð 	r   c	           
     óJ  — |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                              ¦   «         r^|�-|                      ||j                              ¦   «         ¦  «        }|�-|                      ||j                              ¦   «         ¦  «        }|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)rs   r…   r{   r/   r   rT  r�  r‡  r  r1   rj   Úint8Úpointer_typeÚaddress_spacer‰   rª   r&   r'   Úcreate_loadr©   Úcreate_masked_load)r   r  r  r§   r
  rç  rÏ  rÞ  r  Úptr_tyr  r¶  r£  r    s                 r   Ú_load_legacyzTritonSemantic._load_legacyö  sº  € àŒ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×ÒÑÔð 	VØÐØ×0Ò0°°s´x×7PÒ7PÑ7RÔ7RÑSÔS�ØÐ Ø×1Ò1°%¸¼×9RÒ9RÑ9TÔ9TÑUÔ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   r  r  úOptional[TensorTy]r
  r   ræ  r‘  rÎ  rÝ  r  c	           
     ó^  — |                       |¦  «        }	|                      |¦  «        }
|                      |¦  «        }|j                             ¦   «         r:|j        j                             ¦   «         r|                      ||||||	|
|¦  «        S |                      ||||||	|
|¦  «        S r)   )	rÐ  rß  rè  rs   r{   r  rT  r  r!  )r   r  r  r§   r
  ræ  rÎ  rÝ  r  rÏ  rÞ  rç  s               r   ÚloadzTritonSemantic.load2  s·   € ð ×0Ò0°Ñ@Ô@ˆØ×/Ò/°Ñ@Ô@ˆØ×-Ò-¨nÑ=Ô=ˆàŒ8�?Š?ÑÔð 	n ¤Ô!4×!=Ò!=Ñ!?Ô!?ð 	nà×+Ò+¨C°°u¸nÈgÐW\Ð^fÐhsÑtÔtÐtð ×$Ò$ S¨$°°~ÀwÐPUÐW_ÐalÑmÔmÐmr   Údescútl.tensor_descriptor_basec                ó¶  — t          |t          j        ¦  «        sJ ‚t          |j        ¦  «        }t          |¦  «        |k    sJ d|› dt          |¦  «        › �¦   «         ‚|                      |d¬¦  «        }| j                             |j        ||  	                    |¦  «        |  
                    |¦  «        ¦  «        }|                      ||j        ¦  «        S )Nú	expected ú offsets, but got F©Úrequire_i64)rh   r1   Útensor_descriptor_baserU  r  Ú_convert_to_ir_valuesr'   Úcreate_descriptor_loadr©   rÐ  rß  r&   r?  )r   r%  ÚoffsetsrÎ  rÝ  Úndimrt   s          r   Údescriptor_loadzTritonSemantic.descriptor_load@  sÐ   € å˜$¥Ô 9Ñ:Ô:Ð:Ð:Ð:Ý�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          j        ¦  «        sJ ‚t          |j        ¦  «        }t          |¦  «        |k    sJ d|› dt          |¦  «        › �¦   «         ‚|j        |j        k    sJ ‚d S )Nr(  r)  )rh   r1   r,  rU  r  rD  )r   r%  rO   r/  r0  s        r   Úvalidate_store_likez"TritonSemantic.validate_store_likeK  s}   € Ý˜$¥Ô 9Ñ:Ô:Ð:Ð:Ð:Ý�4Ô#Ñ$Ô$ˆÝ�7‰|Œ|˜tÒ#Ð#Ð#Ð%W°Ð%WÐ%WÍÈWÉÌÐ%WÐ%WÑ#Ô#Ð#ØŒ{˜dÔ.Ò.Ð.Ð.Ð.Ð.Ð.r   c                óæ   — |                       |||¦  «         |                      |d¬¦  «        }|                      | j                             |j        |j        |¦  «        t          j        ¦  «        S ©NFr*  )r3  r-  r&   r'   Úcreate_descriptor_storer©   r1   Úvoid)r   r%  rO   r/  s       r   Údescriptor_storezTritonSemantic.descriptor_storeQ  sc   € Ø× Ò   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   rk   r2   rm   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_addV  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_tma]  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   rk   r2   rm   rl   rP   rQ   rI  )r   r>   s     r   Ú$_descriptor_atomic_min_max_supportedz3TritonSemantic._descriptor_atomic_min_max_supporteda  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_minf  óˆ   € Ø× Ò   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_maxm  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   rk   r2   rm   rl   r-  r   r<  ÚANDr&   r'   r>  r©   r7  r?  s        r   Údescriptor_atomic_andz$TritonSemantic.descriptor_atomic_andt  óŸ   € Ø× Ò   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   rk   r2   rm   rl   r-  r   r<  ÚORr&   r'   r>  r©   r7  r?  s        r   Údescriptor_atomic_orz#TritonSemantic.descriptor_atomic_or{  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   rk   r2   rm   rl   r-  r   r<  ÚXORr&   r'   r>  r©   r7  r?  s        r   Údescriptor_atomic_xorz$TritonSemantic.descriptor_atomic_xor‚  rV  r   c                ó~  — t          |t          j        ¦  «        sJ ‚|dk    s
J d¦   «         ‚|dk    s
J d¦   «         ‚t          |j        ¦  «        dk    sJ d|j        › �¦   «         ‚|j        d         dk    sJ d|j        › �¦   «         ‚t          |j        ¦  «        dk    sJ d	|j        › �¦   «         ‚|j        d         d
k    sJ d|j        › �¦   «         ‚|j        }d|j        z  d
z  }|j        d         |k    sJ d|› d|› d|j        d         › �¦   «         ‚t          j        |j        |j        d         |j        d         g¦  «        }|  	                    |fd¬¦  «        d         }| j
                             |j        |j        ||                     | j
        ¦  «        ¦  «        }	|                      |	|¦  «        S )NÚ z#cache modifier is not supported yetz$eviction policy is not supported yetr   údescriptor must be 2D, but got r   r   ú*descriptor block must have 1 row, but got úx offsets must be 1D, but got é   z5descriptor gather must have at least 8 rows, but got r>  zdescriptor gather of ú must have at least ú columns, but got Fr*  )rh   r1   r,  rU  r  rD  r>   r¥  r?  r-  r'   Úcreate_descriptor_gatherr©   r«   r&   )
r   r%  Ú	x_offsetsÚy_offsetrÎ  rÝ  r>   Úmin_colsrs   rt   s
             r   Údescriptor_gatherz TritonSemantic.descriptor_gather‰  s   € Ý˜$¥Ô 9Ñ:Ô:Ð:Ð:Ð:Ø Ò#Ð#Ð#Ð%JÑ#Ô#Ð#Ø "Ò$Ð$Ð$Ð&LÑ$Ô$Ð$õ �4Ô#Ñ$Ô$¨Ò)Ð)Ð)Ð+_ÈTÔM]Ð+_Ð+_Ñ)Ô)Ð)ØÔ Ô" aÒ'Ð'Ð'Ð)hÐVZÔVfÐ)hÐ)hÑ'Ô'Ð'õ �9”?Ñ#Ô# qÒ(Ð(Ð(Ð*\È9Ì?Ð*\Ð*\Ñ(Ô(Ð(ð Œ˜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                óÈ  — t          |t          j        ¦  «        sJ ‚t          |j        ¦  «        dk    sJ d|j        › �¦   «         ‚|j        d         dk    sJ d|j        › �¦   «         ‚t          |j        ¦  «        dk    sJ d|j        › �¦   «         ‚|j        d         dk    sJ d|j        › �¦   «         ‚|j        }d	|j        z  dz  }|j        d         |k    sJ d
|› d|› d|j        d         › �¦   «         ‚|  	                    |fd¬¦  «        d         }| j
                             |j        |j        |j        |¦  «         |                      d t          j        ¦  «        S )Nr   r_  r   r   r`  ra  rb  z6descriptor scatter must have at least 8 rows, but got r>  zdescriptor scatter of rc  rd  Fr*  )rh   r1   r,  rU  r  rD  Úshapaer>   r¥  r-  r'   Úcreate_descriptor_scatterr©   r&   r7  )r   r%  rO   rf  rg  r>   rh  s          r   Údescriptor_scatterz!TritonSemantic.descriptor_scatter¡  s¹  € Ý˜$¥Ô 9Ñ:Ô:Ð:Ð:Ð:õ �4Ô#Ñ$Ô$¨Ò)Ð)Ð)Ð+_ÈTÔM]Ð+_Ð+_Ñ)Ô)Ð)ØÔ Ô" aÒ'Ð'Ð'Ð)hÐVZÔVfÐ)hÐ)hÑ'Ô'Ð'õ �9”?Ñ#Ô# qÒ(Ð(Ð(Ð*]È9ÔK[Ð*]Ð*]Ñ(Ô(Ð(ð Œ˜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           	     óZ  — |�t          d¦  «        ‚|j        j                             ¦   «         }|j                             ¦   «         s|                      ||¦  «        }|j                             ¦   «         s
J d¦   «         ‚||j                             ¦   «         k    s(J d|› d|j                             ¦   «         › d�¦   «         ‚|j        j        j        |j        j        k    s*J d|j        j        j        › d|j        j        › d�¦   «         ‚|j        j        j        }|t          j        k    s
J d¦   «         ‚|                      ||¦  «        }|  	                    ||¦  «        }|  
                    | j                             |j        |j        |||¦  «        t          j        ¦  «        S )	Nr  z-Value argument must be block type or a scalarzBlock shape(z) and value shape(z
) mismatchzBlock element type(z) and value element type(r  )r/   rs   r  r‡  rT  r�  r1   rj   r  r‰   r&   r'   Úcreate_tensor_pointer_storer©   r7  )	r   r  Úvalr  r
  rÏ  rÞ  r  r  s	            r   Ú_store_block_pointerz#TritonSemantic._store_block_pointer¶  sÞ  € ð ÐÝÐjÑkÔkÐkð ”hÔ)×:Ò:Ñ<Ô<ˆØŒx× Ò Ñ"Ô"ð 	>Ø×+Ò+¨C°Ñ=Ô=ˆCØŒx× Ò Ñ"Ô"ÐSÐSÐ$SÑSÔSÐSØ˜cœh×7Ò7ñ 
ô 
ò 
ð 
ð 
Ø`˜+Ð`Ð`¸¼×9RÒ9RÑ9TÔ9TÐ`Ð`Ð`ñ
ô 
ð 
àŒxÔ"Ô-°´Ô1DÒDÐDÐDð  GuÐ\_Ô\dÔ\oÔ\zð  Guð  Guð  VYô  V^ô  Við  Guð  Guð  GuÑDÔDÐDà”Ô$Ô/ˆØ�œÒ Ð Ð Ð"XÑ Ô Ð ð ×:Ò:¸>È;ÑWÔWˆð �iŠi˜˜VÑ$Ô$ˆð �{Š{ØŒL×4Ò4°S´ZÀÄÈ^Ð]bÐdlÑmÔmÕoqÔovñxô xð 	xr   c           	     óÜ  — |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         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 )Nr  z in `tl.store`zÑ`boundary_check` argument is not supported for storing a tensor of pointers or storing a scalar. Because the compiler does not know the boundary; please use block pointers (defined by `make_block_ptr`) insteadzFValue argument cannot be block type if pointer argument is not a blockr  ú"Mask must have boolean scalar type)rs   r…   r{   r/   r   rT  r�  r‡  r  r1   rj   r  r  r  r‰   r&   r'   Úcreate_storer©   r7  r¶  Úcreate_masked_store)	r   r  rp  r  r
  rÏ  rÞ  r   r  s	            r   Ú_store_legacyzTritonSemantic._store_legacyÒ  sA  € àŒ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×ÒÑÔð 	TØ×+Ò+¨C°´×1JÒ1JÑ1LÔ1LÑMÔMˆCØÐØ×0Ò0°°s´x×7PÒ7PÑ7RÔ7RÑSÔS�à””ˆØÔ"ˆð •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   rp  c                ó¸  — |                       |¦  «        }|                      |¦  «        }|j                             ¦   «         s|j        j                             ¦   «         rt          d¦  «        ‚|j                             ¦   «         r8|j        j                             ¦   «         r|  	                    ||||||¦  «        S |  
                    ||||||¦  «        S )Nú"Cannot store to a constant pointer)rÕ  rß  rs   Úis_constr…   r/   r{   r  rT  rq  rv  )	r   r  rp  r  r
  rÎ  rÝ  rÏ  rÞ  s	            r   ÚstorezTritonSemantic.storeþ  sÕ   € ð ×1Ò1°.ÑAÔAˆØ×/Ò/°Ñ@Ô@ˆàŒ8×ÒÑÔð 	C #¤(¤/×":Ò":Ñ"<Ô"<ð 	CÝÐAÑBÔBÐBàŒ8�?Š?ÑÔð 	W ¤Ô!4×!=Ò!=Ñ!?Ô!?ð 	Wà×,Ò,¨S°#°t¸^ÈUÐT\Ñ]Ô]Ð]ð ×%Ò% c¨3°°nÀeÈXÑVÔVÐVr   Úcmprõ  r   c           	     ó4  — |                       |¦  «        }|                      |¦  «        }|j        j        j        }|j        dvrt          d¦  «        ‚|                      | j         	                    |j
        |j
        |j
        ||¦  «        |j        ¦  «        S )N)é   r>  r™   z9atomic_cas only supports elements with width {16, 32, 64})rö  r  rs   r…   r  r¥  r/   r&   r'   Úcreate_atomic_casr©   )r   r  r{  rp  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        j        ¦  «        }|€¬| j                             d	¦  «        }t          j        }|j                              ¦   «         rW|j                              t          j        ¦  «        }| j                             |                     | j        ¦  «        |¦  «        }|                      ||¦  «        }|||fS )
Nz)Pointer argument of store instruction is rx  r¯   Úatomic_z does not support fp16z does not support bf16r}  z does not support T)rs   r…   r{   r/   r   ry  r  r1   rP   rQ   Úint16Úuint16r¥  r‘  rT  r�  r‡  r‰   r'   ri   rj   rª   rV  r«   r&   )r   r  rp  r  r€  r  Úmask_irÚmask_tys           r   Úatom_red_typechecking_implz)TritonSemantic.atom_red_typechecking_impl  s=  € àŒ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×ÒÑÔð 	RØÐØ×0Ò0°°s´x×7PÒ7PÑ7RÔ7RÑSÔS�ØˆØ×/Ò/°°S´X×5NÒ5NÑ5PÔ5PÑQÔQ�Ø�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Ø�C˜ˆ~Ðr   c                óä   — |j         j        }t          j        |d¬¦  «        }|                      ||¦  «        }|                      ||dz
  ¦  «        }|                      |t          j        ¦  «        S )NF)rÄ  Úsignedr   )r>   r¥  r1   Úget_int_dtyperþ   r  r‰   rj   )r   rt   rÄ  ÚidtypeÚixÚsignbits         r   Ú_signbitzTritonSemantic._signbit6  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  rs   r…   rY   rÆ   r&   r'   Úcreate_atomic_rmwr   Ú	ATOMIC_OPrQ  r©   ÚUMAXr1   rR   rT   rA   r2   rl   rþ   r  rk   rm   r�  r  rœ   ÚUMINÚwhere©r   r  rp  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  rs   r…   rY   rÆ   r&   r'   r‘  r   r’  rM  r©   r”  r1   rR   rT   rA   r2   rl   rþ   r  rk   rm   r�  r  rœ   r“  r•  r–  s                     r   Ú
atomic_minzTritonSemantic.atomic_minc  r£  r   c           
     ó   — |                       |||d¦  «        \  }}}|                      |¦  «        }|                      |¦  «        }|j        j        }|                     ¦   «         rt          j        j        nt          j        j	        }|  
                    | j                             ||j        |j        |j        ||¦  «        |j        ¦  «        S )Nr¯   )rˆ  rö  r  rs   r…   r|   r   r’  ÚFADDr=  r&   r'   r‘  r©   )r   r  rp  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©   rs   ©r   r  rp  r  rõ  r   s         r   Ú
atomic_andzTritonSemantic.atomic_and’  ó’   € Ø×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©   rs   r«  s         r   Ú	atomic_orzTritonSemantic.atomic_or™  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©   rs   r«  s         r   Ú
atomic_xorzTritonSemantic.atomic_xor   r­  r   c           
     ó:  — |                       |||d¦  «        \  }}}|                      |¦  «        }|                      |¦  «        }|                      | j                             t          j        j        |j	        |j	        |j	        ||¦  «        |j
        ¦  «        S )NÚxchg)rˆ  rö  r  r&   r'   r‘  r   r’  ÚXCHGr©   rs   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}t          t          j        |¦  «        S )Nzinput_precision must be one of ú. Got ÚTF32X3ÚTF32x3)Úlowerr'   rš   Úallowed_dot_input_precisionsÚupperrI  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Ý•rÔ)¨?Ñ;Ô;Ð;r   ÚaccrÀ  Úmax_num_imprecise_accÚ	out_dtypec           
     ó˜  — |j                              ¦   «         r|j                              ¦   «         sJ ‚|j                             ¦   «         r|j                             ¦   «         rnÎ|j        t          j        t          j        t          j        t          j        t          j	        fv sJ d|j        › �¦   «         ‚|j        t          j        t          j        t          j        t          j        t          j	        fv sJ d|j        › �¦   «         ‚|j        |j        k    sJ d|j        › d|j        › �¦   «         ‚|j         
                    ¦   «         s|j         
                    ¦   «         rgd| j        j        j        v rt          j        d¦  «         |                      |t          j        ¦  «        }|                      |t          j        ¦  «        }|€| j        j        j        }|                      |¦  «        }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�¦   «         ‚|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                             ¦   «         rK|j         j        t          j        k    s
J d¦   «         ‚| j                             d¦  «        }
t          j        }nÐ|                     ¦   «         rt;          d¦  «        ‚|j         j                             ¦   «         s|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        ¦  «        |
¦  «        }n|j%        }|j         |k    sJ ‚|€G|j                             ¦   «         r+|j                             ¦   «         r| j        j        j&        }nQd}nN|j                             ¦   «         r5|j                             ¦   «         r||k    rt;          d|› d|› d�¦  «        ‚|  '                    | j         (                    |j%        |j%        |||¦  «        |¦  «        S )Nz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 releaser   é   ú+Both inputs must be either 2D or 3D; (lhs: ú	 vs rhs: r  rn  éþÿÿÿzFirst input shape (z) and second input shape z= are not compatible for matmul (second index of first shape (z0) must be equal to first index of second shape (Úmin_dot_sizez2target doesn't provide lower shape bounds for dot.r   r   zInput shapes should have M >= z, N >= z
 and K >= zonly int8 supported!zhout_dtype=bfloat16 is unsupported. Please use out_dtype=float32/float16 and cast with `.to(tl.bfloat16)`zmax_num_imprecise_acc (z) must be <= K ())rs   rT  r>   rX   r1   r  Úuint8rP   rQ   rR   r°  r'   rš   Ú!deprecated_fp8_dot_operand_dtypesÚwarningsÚwarnr‰   Údefault_dot_input_precisionrÁ  rU  rD  rO   r±  r²  r…   rY   Ú	get_int32r2   rW   r/   rU   Úget_fp32rV   Úget_fp16r?  rV  r«   r©   Úmax_num_imprecise_acc_defaultr&   Ú
create_dot)r   r~   r€   rÂ  rÀ  rÃ  rÄ  Úlhs_rankÚrhs_rankrË  r  Úret_scalar_tyÚMÚNÚKÚBrÉ   Ú
acc_handles                     r   ÚdotzTritonSemantic.dot»  s¹  € àŒx× Ò Ñ"Ô"Ð: s¤x×'8Ò'8Ñ':Ô':Ð:Ð:Ð:àŒ9×ÒÑÔð 	p #¤)×"2Ò"2Ñ"4Ô"4ð 	pàà”9¥¤­"¬(µB´JÅÄÝ!#¤ð!-ð -ð -ð -Ø.RÀsÄyÐ.RÐ.Rñ-ô -ð -à”9¥¤­"¬(µB´JÅÄÝ!#¤ð!-ð -ð -ð -Ø.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àÐ"Ø"œlÔ2ÔNˆOà×:Ò:¸?Ñ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à-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Ò%Ð%Ð%Ð%ð !Ð(ØŒ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: rc   )rI  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   rp  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   Ú	lhs_scaleÚ
lhs_formatÚ	rhs_scaleÚ
rhs_formatúTensorTy | NoneÚ	fast_mathÚ
lhs_k_packÚ
rhs_k_packc                ó  — |j                              ¦   «         r|j                              ¦   «         sJ ‚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�¦   «         ‚|j        }|j        }|                      |¦  «        }|                      |¦  «        }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  }|
s||z  }t          j
        ||r|||gn||g¦  «        }| j                             d¦  «        }|€4| j                             |                     | j        ¦  «        |¦  «        }n|j        }|j         |k    sJ ‚|rd n|j        }|rd n|j        } |                      | j                             |j        | ||j        ||||	|
|¦
  «
        |¦  «        S )Nr   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   )rs   rT  rU  rD  rO   rã  rh   r1   rq   rî  r?  r'   rÒ  rV  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Ý  Úrhs_scale_handleÚlhs_scale_handles!                                    r   Ú
dot_scaledzTritonSemantic.dot_scaled!  s  € ð Œx× Ò Ñ"Ô"Ð: s¤x×'8Ò'8Ñ':Ô':Ð:Ð:Ð:å�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Øð 	Ø�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Ò%Ð%Ð%Ð%Ø#4ÐJ˜4˜4¸)Ô:JÐØ#4ÐJ˜4˜4¸)Ô:JÐØ�{Š{ØŒ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   rj   rÎ  rÏ  r‰   r•   rs   rT  rŠ   r&   r'   Úcreate_selectr©   )r   r  rt   rÓ   r–  rÉ   s         r   r•  zTritonSemantic.whereU  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   r?  r&   )r   rt   rZ   r—  Úres_tys        r   Úwrap_tensorzTritonSemantic.wrap_tensori  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 )Tro  N)r\  rM  rO   )ra  Útr   s     €r   r  z+TritonSemantic.reduction.<locals>.<genexpr>s  s<   øè è € Ð^Ð^ÐRS˜4Ÿ<š<¨¨A¬G¬M¨?È˜<ÑMÔMÐ^Ð^Ð^Ð^Ð^Ð^r   r   z&reduction axis must be < inputs rank (r  c                ó&   •— g | ]\  }}|‰k    ¯|‘ŒS r  r  )ra  r‹  r[  r*   s      €r   rb  z,TritonSemantic.reduction.<locals>.<listcomp>y  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)   )rs   rD  )ra  r  rD  s     €r   r  z+TritonSemantic.reduction.<locals>.<genexpr>z  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©   ©ra  r  s     r   rb  z,TritonSemantic.reduction.<locals>.<listcomp>|  s   € Ð/IÐ/IÐ/I¸Q°´Ð/IÐ/IÐ/Ir   c              3  óŒ   •K  — | ]>}‰                      ‰                     |¦  «        ‰|         j        j        ‰¦  «        V — Œ?d S r)   ©r  Ú
get_resultrs   r…   )ra  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   )
Útuplers   rD  rU  rˆ  Úallr'   Úcreate_reduceÚverifyrC  )r   r  r*   Úregion_builder_fnÚrankr  r—  rD  s   ```  @@@r   Ú	reductionzTritonSemantic.reductionq  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  r  r  s     r   rb  z3TritonSemantic.associative_scan.<locals>.<listcomp>”  s   € Ð+EÐ+EÐ+E¸¨A¬HÐ+EÐ+EÐ+Er   c              3  óŒ   •K  — | ]>}‰                      ‰                     |¦  «        ‰|         j        j        ‰¦  «        V — Œ?d S r)   r  )ra  r‹  r  Úscan_opr   rD  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   )rs   rD  rU  r'   Úcreate_scanr  r  rC  )	r   r  r*   r   r#  r!  r  r'  rD  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   rU  rs   rD  rC  r'   Úcreate_gatherr©   r  r…   )r   r*  r+  r*   r!  r~  Úgathers          r   r.  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   Ú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 inputrs  )rU  rD  r>   rY   r�  rs   r…   r¶  r/   r©   r&   r'   Úcreate_histogramr1   r?  r2   )r   r¦   r/  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ê   rU  rD  r/   r©   Úset_attrr   Ú	make_attrÚget_context©r   rt   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©rU  rD  r/   r©   r4  r   r5  r6  r7  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.constancyr:  r7  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†  rn  r   c                ó   — g | ]	}|j         ‘Œ
S r  r  ©ra  Úargs     r   rb  z/TritonSemantic.device_print.<locals>.<listcomp>ß  s   € Ð/Ð/Ð/ 3�C”JÐ/Ð/Ð/r   c                ó@   — g | ]}|j                              ¦   «         ‘ŒS r  )r>   rÆ   rH  s     r   rb  z/TritonSemantic.device_print.<locals>.<listcomp>à  s&   € Ð?Ð?Ð?°3�S”Y×,Ò,Ñ.Ô.Ð?Ð?Ð?r   )ÚendswithrU  Ú
startswithr&   r'   Úcreate_printr1   r7  )r   rA  rB  rD  Ú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&   Úcreate_assertr©   r1   r7  )r   r£   r¤   s      r   rŸ   zTritonSemantic.device_assertã  sA   € ØŒ|Ô#Ô)ð 	ØˆFØ�{Š{˜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)          |¦  «        › �¦   «         ‚)Nra   rb   z@Block pointers only support 64 bit `shape/strides`, got a value z which is out of the ranger_   r`   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: )rh   r+   r1   rq   rO   rH   r'   ri   r½  rÑ  r&   rM  r>   rY   rl   r¬   r©   Úget_int64_tyrÆ   r2   rs   )r   r  r+  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 )Nr  c                ó<   •— g | ]}‰                      |‰¦  «        ‘ŒS r  )rY  )ra  r  r+  r   s     €€r   rb  z8TritonSemantic._convert_to_ir_values.<locals>.<listcomp>  s)   ø€ Ð\Ð\Ð\È$�D×2Ò2°4¸ÑEÔEÐ\Ð\Ð\r   )r  rY  )r   Ú	list_liker+  s   ` `r   r-  z$TritonSemantic._convert_to_ir_values  sN   øø€ Ý�9˜jÑ)Ô)ð 	]Ø\Ð\Ð\Ð\Ð\ÐR[Ð\Ñ\Ô\Ð\Ø×.Ò.¨y¸+ÑFÔFÐGÐGr   Úbasec           	     ó.  ‡— |                       |¦  «        }|                       |¦  «        }|                       |d¬¦  «        }|j                             ¦   «         r|j        j                             ¦   «         rt          d¦  «        ‚|j        j        t          j        k    r=|                      |t          j	        t          j
        |j        j        ¦  «        ¦  «        }t          ‰d¦  «        s‰gŠd„ ‰D ¦   «         Št          d„ ‰D ¦   «         ¦  «        s
J d¦   «         ‚t          |d¦  «        s|g}d„ |D ¦   «         }t          |¦  «        t          t!          t#          |¦  «        ¦  «        ¦  «        k    s
J d	¦   «         ‚t          ˆfd
„||||fD ¦   «         ¦  «        s
J d¦   «         ‚| j                             |j        |||‰|¦  «        }|                      |t          j	        t          j        |j        j        ‰¦  «        ¦  «        ¦  «        S )NFr*  zMExpected `base` to be a pointer type (but not a block pointer type or others)r  c                óT   — g | ]%}t          |t          j        ¦  «        r|j        n|‘Œ&S r  r  r  s     r   rb  z1TritonSemantic.make_block_ptr.<locals>.<listcomp>  s0   € ÐfÐfÐfÐRV¥Z°µb´lÑ%CÔ%CÐM�t”z�zÈÐfÐfÐfr   c              3  ó`   K  — | ])}t          |t          ¦  «        od |cxk    odk     nc V — Œ*dS )r_   r`   N)rh   r+   r  s     r   r  z0TritonSemantic.make_block_ptr.<locals>.<genexpr>  sM   è è € Ð\Ð\È•:˜d¥CÑ(Ô(ÐC¨V°tÐ-CÐ-CÒ-CÐ-C¸eÒ-CÐ-CÐ-CÐ-CÐ\Ð\Ð\Ð\Ð\Ð\r   zGExpected a list of constant integers (`int32_t` range) in `block_shape`c                óT   — g | ]%}t          |t          j        ¦  «        r|j        n|‘Œ&S r  r  r  s     r   rb  z1TritonSemantic.make_block_ptr.<locals>.<listcomp>#  s/   € ÐZÐZÐZÈD�z¨$µ´Ñ=Ô=ÐG�”�À4ÐZÐZÐZr   z<Expected a permutation of (0, 1, ..., len(order)-1) in orderc              3  óX   •K  — | ]$}t          ‰¦  «        t          |¦  «        k    V — Œ%d S r)   )rU  )ra  r\  r  s     €r   r  z0TritonSemantic.make_block_ptr.<locals>.<genexpr>'  s6   øè è € ÐhÐh¸)•3�{Ñ#Ô#¥s¨9¡~¤~Ò5ÐhÐhÐhÐhÐhÐhr   zBExpected shape/strides/offsets/block_shape to have the same length)r-  rs   r{   r  rT  r/   r1   rj   r‰   r  r  r  r  r  r�  r‚  rC  rU  r'   Úcreate_make_block_ptrr©   r&   r?  )r   r]  rD  Ústridesr/  r  Úorderr©   s        `  r   Úmake_block_ptrzTritonSemantic.make_block_ptr
  s2  ø€ ð ×*Ò*¨5Ñ1Ô1ˆØ×,Ò,¨WÑ5Ô5ˆØ×,Ò,¨WÀ%Ð,ÑHÔHˆð Œy×ÒÑ!Ô!ð 	n T¤YÔ%9×%BÒ%BÑ%DÔ%Dð 	nÝÐlÑmÔmÐmð Œ9Ô¥2¤7Ò*Ð*Ø—9’9˜T¥2¤?µ2´7¸D¼IÔ<SÑ#TÔ#TÑUÔUˆDõ �{ JÑ/Ô/ð 	(Ø&˜-ˆKØfÐfÐZeÐfÑfÔfˆÝÐ\Ð\ÐP[Ð\Ñ\Ô\Ñ\Ô\ð 	Vð 	VØUñ	Vô 	Vð 	Võ �u˜jÑ)Ô)ð 	Ø�GˆEØZÐZÐTYÐZÑZÔZˆÝ�e‰}Œ}¥¥U­3¨u©:¬:Ñ%6Ô%6Ñ 7Ô 7Ò7Ð7Ð7Ð9wÑ7Ô7Ð7õ ÐhÐhÐhÐhÈÈwÐX_ÐafÐGgÐhÑhÔhÑhÔhð 	Qð 	QØPñ	Qô 	Qð 	Qð ”×3Ò3°D´KÀÈÐQXÐZeÐglÑmÔmˆØ�{Š{˜6¥2¤?µ2´=ÀÄÔAUÐWbÑ3cÔ3cÑ#dÔ#dÑeÔeÐer   c                ó¢   — |                       |d¬¦  «        }|                      | j                             |j        |¦  «        |j        ¦  «        S r5  )r-  r&   r'   Úcreate_advancer©   rs   )r   r]  r/  s      r   ÚadvancezTritonSemantic.advance0  sF   € à×,Ò,¨WÀ%Ð,ÑHÔHˆð �{Š{˜4œ<×6Ò6°t´{ÀGÑLÔLÈdÌiÑXÔXÐXr   rd  r  úList[tl.constexpr]útl.tensor_descriptorc           	     ó`  ‡ — 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 ‚|j        j        j        d	z  }t	          j        |d
         ¦  «        }||z  dk     rt          d|› d|› d||z  › d�¦  «        ‚t	          j        |d
         ¦  «        |d
<   |d
         dk    rt          d|d
         › �¦  «        ‚ˆ fd„|D ¦   «         }ˆ fd„|D ¦   «         }t	          j	        |¦  «        }t          |j
        t          j        ¦  «        sJ ‚t	          j        |j
        j        |¦  «        }|j        }	|j
        j                             ¦   «         }
‰ 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 rb  rn  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  )rN  r1   r2   ©ra  rt   r   s     €r   rb  z9TritonSemantic.make_tensor_descriptor.<locals>.<listcomp>Q  s+   ø€ Ð>Ð>Ð>°1�×!Ò! !¥R¤XÑ.Ô.Ð>Ð>Ð>r   c                óP   •— g | ]"}‰                      |t          j        ¦  «        ‘Œ#S r  )rN  r1   rl   ro  s     €r   rb  z9TritonSemantic.make_tensor_descriptor.<locals>.<listcomp>R  s+   ø€ ÐBÐBÐB°Q�4×#Ò# A¥r¤xÑ0Ô0ÐBÐBÐBr   c                ó   — g | ]	}|j         ‘Œ
S r  r  ©ra  r[  s     r   rb  z9TritonSemantic.make_tensor_descriptor.<locals>.<listcomp>\  s   € ÐIbÐIbÐIbÐWXÈ!Ì(ÐIbÐIbÐIbr   c                ó   — g | ]	}|j         ‘Œ
S r  r  rr  s     r   rb  z9TritonSemantic.make_tensor_descriptor.<locals>.<listcomp>]  s   € Ð<WÐ<WÐ<WÈ!¸Q¼XÐ<WÐ<WÐ<Wr   )rU  r/   rh   r>   r1   r  r  r¥  r`  Ú_unwrap_shapers   r?  r©   rÆ   r'   Úcreate_make_tensor_descriptorÚtensor_descriptor)r   r]  rD  rd  r  r0  Ú	elem_sizeÚcontig_dim_sizers   Úbase_handleÚis_signed_intr©   s   `           r   Úmake_tensor_descriptorz%TritonSemantic.make_tensor_descriptor7  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Ð6Ð6Ð6Ø”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Ñ:Ô:ˆ�‰Ø�2Œ;˜!ÒÐÝÐZÈWÐUWÌ[ÐZÐZÑ[Ô[Ð[à>Ð>Ð>Ð>¸Ð>Ñ>Ô>ˆØBÐBÐBÐB¸'ÐBÑBÔBˆõ Ô& {Ñ3Ô3ˆå˜$œ)¥R¤_Ñ5Ô5Ð5Ð5Ð5ÝŒ}˜TœYÔ1°;Ñ?Ô?ˆØ”kˆØœ	Ô,×:Ò:Ñ<Ô<ˆà”×;Ò;¸KÐIbÐIbÐ\aÐIbÑIbÔIbØ<WÐ<WÈwÐ<WÑ<WÔ<WÐYdÐfsñuô uˆåÔ# 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]   rH   )r   r9   r   r9   rx   rH   r,   ry   )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   )rt   r   rÓ   r   rÔ   rÕ   )rt   r   ré   r   rê   r   rÔ   rÕ   )r¦   r   r§   r   r,   r�   )r¦   r   r§   r   r,   r   )r¦   r   )r¦   r   r,   r   )r  r   r,   r  )r;  r+   r<  r+   rÉ   r  r,   r   )r>   r9   r,   r   )rD  rO  r>   r9   r,   r   )rO   r   rD  rO  r,   r   )r¦   r   rW  rO  rX  rH   r,   r   )r¦   r   r*   r+   r,   r   )r~   r   r€   r   rX  rH   r,   r   )rk  r   rl  r   r,   r   )rk  r   r,   r�   )r¦   r   rz  r{  r,   r   )r¦   r   rD  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,   ry   )r%  r&  rO   r   r,   r   )rÎ  r‘  rÝ  r‘  r,   r   )rO   r   r,   r   )r  r   rp  r   r  r"  rÎ  r‘  rÝ  r‘  r,   r   )r  r   r{  r   rp  r   rõ  r‘  r   r‘  r,   r   )
r  r   rp  r   r  r   r€  r‘  r,   r�  )rt   r   r,   r   )r  r   rp  r   r  r   rõ  r‘  r   r‘  r,   r   )r~   r   r€   r   rÂ  r   rÀ  r›  rÃ  r+   rÄ  r9   r,   r   )rß  r‘  )rp  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   rt   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/  r+   r  r"  r,   r   )rt   r   r�  rO  r,   r   )r,   r   )rA  r‘  rB  rC  rD  rH   r,   r   )r£   r   r¤   r‘  r,   r   )r]  r   r,   r   )
r]  r   rD  rC  rd  rC  r  rj  r,   rk  )yr    r!   r"   r1   r&   Ú__annotations__Úlangr   r4   r7   rF   r\   rr   r}   r•   r¥   r¯   r¸   r¾   rÄ   rÊ   rÍ   rÒ   rá   rè   rí   rñ   rœ   rø   rû   rÿ   r  r  r  r  r  r  r¶   r  r  r#  rž   r-  r�   r5  r:  rF  rn   rN  rR  rQ  r\  rf  rj  rt  ry  r„  r�  rŠ   r¢  rþ   r‰   rÐ  rÕ  rß  rè  rö  r  r  r  r!  r$  r1  r3  r8  r@  rI  rK  rN  rR  rU  rY  r\  ri  rm  rq  rv  rz  r  rˆ  r�  r¢  r¥  r¨  r¬  r°  r³  r·  rÁ  rÞ  rã  rî  r  r•  r  r"  r)  r.  r2  r8  r;  r=  r@  rP  rŸ   rV  rY  r-  rf  ri  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#ð #ð #ð #ð #ðR	@ð 	@ð 	@ð 	@ð 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ðfð fð fð fðXð Xð Xð XðVð Vð Vð Vðð ð ð ð0

ð 

ð 

ð 

ðTð Tð Tð TðWð Wð Wð Wð 2ð 2ð 2ð 2ñptð tð tð tñjð jð jð jñ$i9ñ i9ð i9ð i9ð i9ð^ð ð ðð ð ð	ð 	ð 	ð	ð 	ð 	ðð ð ðð ð ð
ð 
ð 
ðð ð ð,:ð :ð :ñxnð nð nð nñ	/ð 	/ð 	/ð 	/ñ/ð /ð /ð /ñ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ñ$ð $ð $ð $ñ0*ð *ð *ð *ð*xð xð xð8*$ð *$ð *$ñXWð Wð Wð Wñ(uð uð uð uñð ð ð ñ8+ð +ð +ð +ñ$)ð $)ð $)ð $)ñL$)ð $)ð $)ð $)ñL%ð %ð %ð %ñyð yð yð yñxð xð xð xñyð yð yð yñð ð ð ð<ð <ð <ñLyð Lyð Lyð Lyñ\ð ð ð ñ0ð 0ð 0ð 0ñ$.tð .tð .tð .tñheð eð eð eð(&ð &ð &ñuð uð uð uñ,xð xð xð xñ.Kð Kð Kð Kñ.	@ð 	@ð 	@ð 	@ñð ð ð ñð ð ð ñð ð ð ñCð Cð Cð Cñað að að añRð Rð Rð Rñ
Mð Mð Mð MðYð Yð Yð4Hð Hð Hð Hñ
$fð $fð $fð $fñLYð Yð Yð Yñ'Bð 'Bð 'Bð 'Bð 'Bð 'Br   r%   )Ú
__future__r   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ƒ     s0  ðØ "Ð "Ð "Ð "Ð "Ð "Ø €€€à 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ðEBð EBð EBð EBð EB�W˜XÔ&ñ EBô EBð EBð EBð EBr   