§
    OŠtj­Œ ã                  óâ   — d dl mZ d dlZd dlZd dlmZmZmZmZm	Z	m
Z
mZ d dlZd dlmZ ddlmZ ddlmZ  e	d	¦  «        Z e	d
¦  «        Z G d„ de¦  «        Z G d„ de
e         ¦  «        ZdS )é    )ÚannotationsN)ÚListÚOptionalÚSequenceÚTupleÚTypeVarÚGenericÚType)Údriveré   )Úiré   )ÚcoreÚTÚTensorTyc                  ó   ‡ — e Zd Zˆ fd„Zˆ xZS )ÚIncompatibleTypeErrorImplc                óú   •— || _         || _        d| j                              ¦   «         z   dz   | j                             ¦   «         z   | _        t	          t
          | ¦  «                             | j        ¦  «         d S )Nzinvalid operands of type ú and )Útype_aÚtype_bÚ__repr__ÚmessageÚsuperr   Ú__init__)Úselfr   r   Ú	__class__s      €úV/var/www/html/CA-Chatbot/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                  óL  — e Zd ZU ej        Zded<   eZded<   d„ Zdëd
„Zdëd„Z	dìd„Z
díd„Zdîd„Zd„ Zdïd„Z	 	 dðdñd"„Zdòd%„Zdód)„Zdód*„Zdód+„Zdôd,„Zdôd-„Zdõd/„Zdôd0„Zdöd5„Zdöd6„Zd÷d9„Zdød:„Zdùd;„Zdùd<„Zdùd=„Zdùd>„Zdùd?„Z dúd@„Z!dùdA„Z"dùdB„Z#dùdC„Z$dûdD„Z%dûdE„Z&dûdF„Z'düdI„Z(dùdJ„Z)dùdK„Z*dùdL„Z+dùdM„Z,dùdN„Z-dùdO„Z.dPdQœdýdU„Z/dþdW„Z0dþdX„Z1dÿd[„Z2�d d]„Z3�dd^„Z4�dda„Z5�ddb„Z6�ddc„Z7�ddf„Z8�ddg„Z9�ddj„Z:�ddk„Z;�d	dl„Z<�d
do„Z=�ddq„Z>�d�dds„Z?dt„ Z@du„ ZAdv„ ZBdw„ ZCdx„ ZDdy„ ZEdz„ ZFd{„ ZGd|„ ZH�dd‡„ZI�ddŠ„ZJ�dd‹„ZK�ddŒ„ZL�dd�„ZMdŽ„ ZNd�„ ZO�dd�„ZP�dd‘„ZQ�dd’„ZR�dd“„ZS�dd”„ZT�dd•„ZU�dd–„ZVd—„ ZWd˜„ 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§„Za�dd¨„Zb�dd©„Zcdª„ Zd�dd¯„Ze�dd±„Zf�dd²„Zgd³„ Zh�dd¼„Zi�dd¾„Zjd¿„ Zk�ddÃ„Zl�ddÅ„Zm�ddÈ„ZndÉ„ Zo�d dÏ„Zp�d!dÑ„Zq�d"dÓ„Zr�d"dÔ„Zs�d"dÕ„Zt�d#dÖ„Zu�d$dÛ„Zv�d%dÞ„Zw�d#dß„Zxdà„ Zydîdá„Zz�d&dã„Z{�d&dä„Z|	 �d'�d(dê„Z}dPS ()  ÚTritonSemanticzType[TensorTy]Útensorz
ir.builderÚbuilderc                ó   — || _         d S ©N)r'   )r   r'   s     r   r   zTritonSemantic.__init__    s   € ØˆŒˆˆr   ÚaxisÚintÚreturnr   c                óž   — |dvrt          d|› �¦  «        ‚|                      | j                             |¦  «        t          j        ¦  «        S )N©r   r   r   z+program_id axis must be 0, 1, or 2 but got )Ú
ValueErrorr&   r'   Úcreate_get_program_idÚtlÚint32©r   r*   s     r   Ú
program_idzTritonSemantic.program_id'   sJ   € Ø�yÐ Ð ÝÐQÈ4ÐQÐQÑRÔRÐRØ�{Š{˜4œ<×=Ò=¸dÑCÔCÅRÄXÑNÔNÐNr   c                óž   — |dvrt          d|› �¦  «        ‚|                      | j                             |¦  «        t          j        ¦  «        S )Nr.   z-num_programs axis must be 0, 1, or 2 but got )r/   r&   r'   Úcreate_get_num_programsr1   r2   r3   s     r   Únum_programszTritonSemantic.num_programs,   sJ   € Ø�yÐ Ð ÝÐSÈTÐSÐSÑTÔTÐTØ�{Š{˜4œ<×?Ò?ÀÑEÔEÅrÄxÑPÔPÐPr   Úa_tyútl.dtypeÚb_tyc                ó  — |j         }|j         }|j        }|j        }||k    r
||k    r|n|S |t          j        j        j        k    r
||k    r|n|S |t          j        j        j        k    r
||k    r|n|S t          d|› d|› �¦  «        ‚)Nzunexpected signedness r   )Úint_bitwidthÚint_signednessr1   ÚdtypeÚ
SIGNEDNESSÚUNSIGNEDÚ	TypeError)r   r8   r:   Úa_rankÚb_rankÚa_snÚb_sns          r   Úinteger_promote_implz#TritonSemantic.integer_promote_impl5   s¬   € ØÔ"ˆØÔ"ˆØÔ"ˆØÔ"ˆð �4Š<ˆ<Ø! Fš?˜?�4�4°Ð4Ø•R”XÔ(Ô1Ò1Ð1Ø! VÒ+Ð+�4�4°Ð5Ø•R”XÔ(Ô1Ò1Ð1Ø! VÒ+Ð+�4�4°Ð5ÝÐB°ÐBÐB¸DÐBÐBÑCÔCÐCr   Úa_is_scalarÚboolÚb_is_scalarÚ
div_or_modc                óä  — ||k    rk|r||fn||f\  }}|                      ¦   «         j        |                      ¦   «         j        k    r*|r&|t          j        t          j        fv rt          j        S |S |                     ¦   «         s|                     ¦   «         rt          j        S |                     ¦   «         s|                     ¦   «         rt          j        S | 	                    ¦   «         s| 	                    ¦   «         r|rt          j        S t          j        S | 
                    ¦   «         r.| 
                    ¦   «         r|rt          j        S t          j        S | 
                    ¦   «         s| 
                    ¦   «         rt          j        S |                     ¦   «         r(|                     ¦   «         r||k    r|nt          j        S |                     ¦   «         r|                     ¦   «         st          d|› d|› �¦  «        ‚|rO|j        |j        k    r?t          d|                     ¦   «         z   dz   |                     ¦   «         z   dz   ¦  «        ‚|                      ||¦  «        S )Núunexpected type r   zCannot use /, #, or % with úx because they have different signedness;this is unlikely to result in a useful answer. Cast them to the same signedness.)ÚkindÚvaluer1   Úfloat16Úbfloat16Úfloat32Úis_fp64Úfloat64Úis_fp32Úis_fp16Úis_bf16Úis_fp8Úis_intrA   r=   r   rF   )r   r8   rG   r:   rI   rJ   Ú	scalar_tyÚ	tensor_tys           r   Úcomputation_type_implz$TritonSemantic.computation_type_implD   sO  € ð
 ˜+Ò%Ð%Ø3>Ð#P D¨$ < <ÀTÈ4ÀLÑ ˆI�yØ�~Š~ÑÔÔ%¨¯ªÑ)9Ô)9Ô)?Ò?Ð?àð & 9µ´½R¼[Ð0IÐ#IÐ#IÝœ:Ð%Ø Ð ð �<Š<‰>Œ>ð 	˜TŸ\š\™^œ^ð 	Ý”:Ðð �<Š<‰>Œ>ð 	˜TŸ\š\™^œ^ð 	Ý”:Ðð �<Š<‰>Œ>ð 	"˜TŸ\š\™^œ^ð 	"Øð "Ý”zÐ!å”zÐ!à�<Š<‰>Œ>ð 	#˜dŸlšl™nœnð 	#Øð #Ý”zÐ!å”{Ð"Ø�<Š<‰>Œ>ð 	˜TŸ\š\™^œ^ð 	Ý”:Ðà�;Š;‰=Œ=ð 	8˜TŸ[š[™]œ]ð 	8Ø 4š<˜<�4�4­R¬ZÐ7Ø�{Š{‰}Œ}ð 	B D§K¢K¡M¤Mð 	BÝÐ@¨tÐ@Ð@¸$Ð@Ð@ÑAÔAÐAð ð 	p˜$Ô-°Ô1DÒDÐDÝÐ9¸D¿MºM¹O¼OÑKÈgÑUÐX\×XeÒXeÑXgÔXgÑgðoñoñ pô pð pð ×(Ò(¨¨tÑ4Ô4Ð4r   Tc                ó`  — t          || j        ¦  «        r|S t          |t          j        ¦  «        r|j        n|}t          |t
          t          t          f¦  «        r,|                      |¦  «        }|  	                    ||¬¦  «        S |r#t          d|› dt          |¦  «        › d�¦  «        ‚|S )N©r>   úcannot convert ú	 of type ú
 to tensor)Ú
isinstancer&   r1   Ú	constexprrO   r+   ÚfloatrH   Úto_tensor_typeÚscalar_constantrA   Útype)r   ÚxÚ
check_typer>   s       r   Ú	to_tensorzTritonSemantic.to_tensorv   s³   € Ý�a˜œÑ%Ô%ð 	ØˆHÝ! !¥R¤\Ñ2Ô2Ð9ˆAŒGˆG¸ˆÝ�a�#�u¥dÐ+Ñ,Ô,ð 	OØ×'Ò'¨Ñ*Ô*ˆEØ×'Ò'¨°Ð'Ñ7Ô7Ð7Øð 	OÝÐM¨aÐMÐM½$¸q¹'¼'ÐMÐMÐMÑNÔNÐNØˆr   c                ó(  — t          |t          j        ¦  «        r|S t          |t          j        ¦  «        r|j        }t          |t
          ¦  «        rt          j        S t          |t          ¦  «        rƒd|cxk    rdk     rn nt          j        S d|cxk    rdk     rn nt          j	        S d|cxk    rdk     rn nt          j
        S d|cxk    rdk     rn nt          j        S t          d|› d�¦  «        ‚t          |t          ¦  «        red	}d
ddz  z  }t          j        |¦  «        }|t          d¦  «        k    s|dk    s||k    s||cxk    r|k    rn nt          j        S t          j        S t%          d|› dt'          |¦  «        › d�¦  «        ‚)Né   €ì        l        ìûÿÿÿ         ì            l            zNonrepresentable integer ú.g      8g   àÿÿÿ?r   é   Úinfg        r_   r`   ra   )rb   r1   r>   Úconstexpr_typerO   rH   Úint1r+   r2   Úuint32Úint64Úuint64r/   rd   ÚbuiltinsÚabsrR   rT   rA   rg   )r   rh   Úmin_float32Úmax_float32Úabs_xs        r   re   zTritonSemantic.to_tensor_type�   sÑ  € Ý�a�œÑ"Ô"ð 	ØˆHÝ˜�2Ô,Ñ-Ô-ð 	Ø”ˆAå�a�ÑÔð 	"Ý”7ˆNÝ˜�3ÑÔð 	"Ø˜Ð"Ð"Ò"Ð"˜UÒ"Ð"Ð"Ð"Ð"Ý”x�Ø˜!Ð#Ð#Ò#Ð#˜eÒ#Ð#Ð#Ð#Ð#Ý”yÐ Ø˜1Ð$Ð$Ò$Ð$˜uÒ$Ð$Ð$Ð$Ð$Ý”x�Ø˜!Ð#Ð#Ò#Ð#˜eÒ#Ð#Ð#Ð#Ð#Ý”yÐ ÝÐ=¸Ð=Ð=Ð=Ñ>Ô>Ð>Ý˜�5Ñ!Ô!ð 
	"Ø!ˆKØ%¨¨C©Ñ/ˆKÝ”L ‘O”OˆEØ�˜e™œÒ$Ð$Ø˜Š|ˆ|Ø�AŠvˆvØ˜eÐ2Ð2Ò2Ð2 {Ò2Ð2Ð2Ð2Ð2Ý”zÐ!å”zÐ!ÝÐI¨!ÐIÐIµd¸1±g´gÐIÐIÐIÑJÔJÐJr   r   r   Úallow_ptr_aÚNonec                óò   — |                      ¦   «         r`|st          ||¦  «        ‚|                      ¦   «         r||k    rt          ||¦  «        ‚|                     ¦   «         rt          ||¦  «        ‚d S d S r)   )Úis_ptrr   Úis_floating)r   r   r   r}   s       r   Úcheck_ptr_type_implz"TritonSemantic.check_ptr_type_impl¤   s”   € Ø�=Š=‰?Œ?ð 	@Øð @Ý/°¸Ñ?Ô?Ð?à�}Š}‰Œð @ F¨fÒ$4Ð$4Ý/°¸Ñ?Ô?Ð?à×!Ò!Ñ#Ô#ð @Ý/°¸Ñ?Ô?Ð?ð	@ð 	@ð@ð @r   FÚlhsúTensorTy | numbers.NumberÚrhsúTuple[TensorTy, TensorTy]c                óº  — t          |t          j        ¦  «        }t          |t          j        ¦  «        }|r|}	|                      |¦  «        }|r|}
|                      |¦  «        }|j        j        }|j        j        }|                      |||¦  «         |                      |||¦  «         |�r�|                     ¦   «         �s{|                     ¦   «         �sf|                      |||||¦  «        }|r|	dk     r| 	                    ¦   «         s|r)|
dk     r#| 	                    ¦   «         rt          d¦  «        ‚|                     ¦   «         r”|rH|                     ¦   «         |	cxk    r|                     ¦   «         k    sn t          d|	› d|› �¦  «        ‚|rH|                     ¦   «         |
cxk    r|                     ¦   «         k    sn t          d|
› d|› �¦  «        ‚|r|                      |	|¬¦  «        n|                      ||¦  «        }|r|                      |
|¬¦  «        n|                      ||¦  «        }|                      ||¦  «        \  }}||fS )Nr   z{Cannot perform a binary operation between an unsigned tensor and a negative scalar. Perform a explicit cast on one of them.zScalar z is out of range for type r^   )rb   ÚnumbersÚNumberrj   rg   Úscalarr‚   r€   r\   Úis_int_unsignedr/   rY   Úget_int_min_valueÚget_int_max_valuerf   ÚcastÚbroadcast_impl_value)r   rƒ   r…   Úallow_lhs_ptrÚallow_rhs_ptrÚarithmetic_checkrJ   Úlhs_is_scalarÚrhs_is_scalarÚ
lhs_scalarÚ
rhs_scalarÚ
lhs_sca_tyÚ
rhs_sca_tyÚ
ret_sca_tys                 r   Úbinary_op_type_checking_implz+TritonSemantic.binary_op_type_checking_impl¯   sä  € õ # 3­¬Ñ7Ô7ˆÝ" 3­¬Ñ7Ô7ˆØð 	&ØˆJØ—.’. Ñ%Ô%ˆCØð 	&ØˆJØ—.’. Ñ%Ô%ˆCð ”X”_ˆ
Ø”X”_ˆ
Ø× Ò  ¨Z¸ÑGÔGÐGØ× Ò  ¨Z¸ÑGÔGÐGØñ 	v J×$5Ò$5Ñ$7Ô$7ñ 	vÀ
×@QÒ@QÑ@SÔ@Sñ 	vØ×3Ò3°JÀÈzÐ[hÐjtÑuÔuˆJØð L *¨q¢. .°Z×5OÒ5OÑ5QÔ5Q .Ø$ð #1Ø)3°aª¨¸J×<VÒ<VÑ<XÔ<X¨Ý ð "Kñ Lô Lð Là× Ò Ñ"Ô"ð cØ ð c¨*×*FÒ*FÑ*HÔ*HÈJð +Ið +Iò +Ið +IØ*4×*FÒ*FÑ*HÔ*Hò+Ið +Ið +Ið +Iå$Ð%a¨zÐ%aÐ%aÐU_Ð%aÐ%aÑbÔbÐbØ ð c¨*×*FÒ*FÑ*HÔ*HÈJð +Ið +Iò +Ið +IØ*4×*FÒ*FÑ*HÔ*Hò+Ið +Ið +Ið +Iå$Ð%a¨zÐ%aÐ%aÐU_Ð%aÐ%aÑbÔbÐbØHUÐu�$×&Ò& z¸Ð&ÑDÔDÐDÐ[_×[dÒ[dÐehÐjtÑ[uÔ[uˆCØHUÐu�$×&Ò& z¸Ð&ÑDÔDÐDÐ[_×[dÒ[dÐehÐjtÑ[uÔ[uˆCð ×,Ò,¨S°#Ñ6Ô6‰ˆˆSØ�Cˆxˆr   Ú	binary_opÚcallablec                óú  — |j         j        j        dk    s| j        j        j        sd S |j         j        }|j         j        }||k    sJ ‚|                     ¦   «         sJ ‚|                      |t          j	        ¦  «        }|                      |t          j	        ¦  «        } |||d¦  «        }| 
                    ¦   «         }|                      |t          j	        ¦  «        }|                     ¦   «         }|                      |t          j	        ¦  «        }|                      |                      ||¦  «        |                      ||¦  «        ¦  «        }	d|j        › d|j        › �}
|                      |	|
d ¦  «         d S )Né@   Fr+   z! overflow detected for operation )rg   rŠ   r<   r'   ÚoptionsÚsanitize_overflowrY   rŽ   r1   rv   r�   rf   rŒ   Úand_Ú
less_equalÚgreater_equalr    Údevice_assert)r   rƒ   r…   r›   r—   r˜   ÚretÚ	max_valueÚ	min_valueÚcondÚmsgs              r   Ú binary_op_sanitize_overflow_implz/TritonSemantic.binary_op_sanitize_overflow_implÔ   sY  € ØŒ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  dÑ+Ô+Ð+Ð+Ð+r   ÚinputÚotherr    c                óº  — |                       ||dd¦  «        \  }}|j        j        }|j        j        }|                     ¦   «         r#|                     ¦   «         rt	          d¦  «        ‚|                     ¦   «         r0|                     ¦   «         s||}}|j        j        }|j        j        }|                     ¦   «         rÆ|j        }|j                             ¦   «         rm|j        j        dk     r]|j         	                    t          j        ¦  «                             | j        ¦  «        }| j                             |j        |d¦  «        }|                      | j                             |j        |¦  «        |j        ¦  «        S |                     ¦   «         r>|                      | j                             |j        |j        ¦  «        |j        ¦  «        S |                     ¦   «         r\|r|                      ||| j        ¦  «         |                      | j                             |j        |j        ¦  «        |j        ¦  «        S t	          d|› �¦  «        ‚)NTzcannot add pointers togetherrž   FrL   )rš   rg   rŠ   r€   rA   Úhandler>   r‹   r<   Úwith_element_tyr1   rv   Úto_irr'   Úcreate_int_castr&   Úcreate_addptrr�   Úcreate_faddrY   rª   ÚaddÚ
create_add)r   r«   r¬   r    Úinput_scalar_tyÚother_scalar_tyÚother_handleÚi64_tys           r   r´   zTritonSemantic.addæ   s   € à×8Ò8¸ÀÀtÈTÑRÔR‰ˆˆuØœ*Ô+ˆØœ*Ô+ˆØ×!Ò!Ñ#Ô#ð 	<¨×(>Ò(>Ñ(@Ô(@ð 	<ÝÐ:Ñ;Ô;Ð;ð ×!Ò!Ñ#Ô#ð 	0¨O×,BÒ,BÑ,DÔ,Dð 	0Ø  %�5ˆEØ#œjÔ/ˆOØ#œjÔ/ˆOØ×!Ò!Ñ#Ô#ð 	`Ø œ<ˆLØŒ{×*Ò*Ñ,Ô,ð Y°´Ô1IÈBÒ1NÐ1Nàœ×3Ò3µB´HÑ=Ô=×CÒCÀDÄLÑQÔQ�Ø#œ|×;Ò;¸E¼LÈ&ÐRWÑXÔX�Ø—;’;˜tœ|×9Ò9¸%¼,ÈÑUÔUÐW\ÔWaÑbÔbÐbà×(Ò(Ñ*Ô*ð 	`Ø—;’;˜tœ|×7Ò7¸¼ÀeÄlÑSÔSÐUZÔU_Ñ`Ô`Ð`à×#Ò#Ñ%Ô%ð 	`Ø ð NØ×5Ò5°e¸UÀDÄHÑMÔMÐMØ—;’;˜tœ|×6Ò6°u´|ÀUÄ\ÑRÔRÐTYÔT^Ñ_Ô_Ð_ÝÐ<¨?Ð<Ð<Ñ=Ô=Ð=r   c                óv  — |                       ||dd¦  «        \  }}|j        j        }|                     ¦   «         r+|                      ||                      |¦  «        d¬¦  «        S |                     ¦   «         r>|                      | j         	                    |j
        |j
        ¦  «        |j        ¦  «        S |                     ¦   «         r\|r|                      ||| j        ¦  «         |                      | j                             |j
        |j
        ¦  «        |j        ¦  «        S t          d|› �¦  «        ‚)NTF)r    rL   )rš   rg   rŠ   r€   r´   Úminusr�   r&   r'   Úcreate_fsubr®   rY   rª   ÚsubÚ
create_subrA   ©r   r«   r¬   r    rZ   s        r   r½   zTritonSemantic.sub  s"  € à×8Ò8¸ÀÀtÈUÑSÔS‰ˆˆuØ”JÔ%ˆ	à×ÒÑÔð 	OØ—8’8˜E 4§:¢:¨eÑ#4Ô#4È�8ÑNÔNÐNà× Ò Ñ"Ô"ð 	`Ø—;’;˜tœ|×7Ò7¸¼ÀeÄlÑSÔSÐUZÔU_Ñ`Ô`Ð`à×ÒÑÔð 	`Ø ð NØ×5Ò5°e¸UÀDÄHÑMÔMÐMØ—;’;˜tœ|×6Ò6°u´|ÀUÄ\ÑRÔRÐTYÔT^Ñ_Ô_Ð_ÝÐ6¨9Ð6Ð6Ñ7Ô7Ð7r   c                óô  — |                       ||¦  «        \  }}|j        j        }|                     ¦   «         r>|                      | j                             |j        |j        ¦  «        |j        ¦  «        S |                     ¦   «         r\|r|  	                    ||| j
        ¦  «         |                      | j                             |j        |j        ¦  «        |j        ¦  «        S t          d|› �¦  «        ‚©NrL   )rš   rg   rŠ   r�   r&   r'   Úcreate_fmulr®   rY   rª   ÚmulÚ
create_mulrA   r¿   s        r   rÃ   zTritonSemantic.mul  så   € à×8Ò8¸ÀÑFÔF‰ˆˆuØ”JÔ%ˆ	à× Ò Ñ"Ô"ð 	`Ø—;’;˜tœ|×7Ò7¸¼ÀeÄlÑSÔSÐUZÔU_Ñ`Ô`Ð`à×ÒÑÔð 	`Ø ð NØ×5Ò5°e¸UÀDÄHÑMÔMÐMØ—;’;˜tœ|×6Ò6°u´|ÀUÄ\ÑRÔRÐTYÔT^Ñ_Ô_Ð_ÝÐ6¨9Ð6Ð6Ñ7Ô7Ð7r   c                ó¨  — |                       ||dddd¦  «        \  }}|j        j        }|j        j        }|                     ¦   «         r,|                     ¦   «         r|                      ||¦  «        }�n |                     ¦   «         r+|                     ¦   «         r|                      ||¦  «        }ná|                     ¦   «         rU|                     ¦   «         rA|                      |t          j        ¦  «        }|                      |t          j        ¦  «        }nx|                     ¦   «         rR|                     ¦   «         r>|j        |j        k    r|                      ||¦  «        }n)|                      ||¦  «        }nt          d|› �¦  «        ‚|  
                    | j                             |j        |j        ¦  «        |j        ¦  «        S ©NFTrL   )rš   rg   rŠ   r�   rY   rŽ   r1   rR   Úfp_mantissa_widthrA   r&   r'   Úcreate_fdivr®   )r   r«   r¬   r¶   r·   s        r   ÚtruedivzTritonSemantic.truediv$  sµ  € Ø×8Ò8¸ÀÀuÈeÐUYÐ[_Ñ`Ô`‰ˆˆuØœ*Ô+ˆØœ*Ô+ˆà×&Ò&Ñ(Ô(ð 	B¨_×-CÒ-CÑ-EÔ-Eð 	BØ—I’I˜e _Ñ5Ô5ˆE‰Eà×#Ò#Ñ%Ô%ð 	B¨/×*EÒ*EÑ*GÔ*Gð 	BØ—I’I˜e _Ñ5Ô5ˆEˆEà×#Ò#Ñ%Ô%ð 	B¨/×*@Ò*@Ñ*BÔ*Bð 	BØ—I’I˜e¥R¤ZÑ0Ô0ˆEØ—I’I˜e¥R¤ZÑ0Ô0ˆEˆEà×(Ò(Ñ*Ô*ð 	B¨×/JÒ/JÑ/LÔ/Lð 	BØÔ0°?Ô3TÒTÐTØŸ	š	 %¨Ñ9Ô9��àŸ	š	 %¨Ñ9Ô9��õ Ð@¨Ð@Ð@ÑAÔAÐAØ�{Š{˜4œ<×3Ò3°E´LÀ%Ä,ÑOÔOÐQVÔQ[Ñ\Ô\Ð\r   c                ó„  — |                       ||dddd¦  «        \  }}|j        j        }|j        j        }|                     ¦   «         ræ|                     ¦   «         rÒ|                      ||¦  «        }|                      ||¦  «        }|                      ||¦  «        }|                     ¦   «         r>|                      | j         	                    |j
        |j
        ¦  «        |j        ¦  «        S |                      | j                             |j
        |j
        ¦  «        |j        ¦  «        S t          d|› �¦  «        ‚rÆ   )rš   rg   rŠ   rY   rF   rŽ   Úis_int_signedr&   r'   Úcreate_sdivr®   Úcreate_udivrA   )r   r«   r¬   r¶   r·   Úret_tys         r   ÚfloordivzTritonSemantic.floordiv=  s%  € Ø×8Ò8¸ÀÀuÈeÐUYÐ[_Ñ`Ô`‰ˆˆuØœ*Ô+ˆØœ*Ô+ˆØ×!Ò!Ñ#Ô#ð 	e¨×(>Ò(>Ñ(@Ô(@ð 	eØ×.Ò.¨ÀÑPÔPˆFØ—I’I˜e VÑ,Ô,ˆEØ—I’I˜e VÑ,Ô,ˆEØ×#Ò#Ñ%Ô%ð eØ—{’{ 4¤<×#;Ò#;¸E¼LÈ%Ì,Ñ#WÔ#WÐY^ÔYcÑdÔdÐdà—{’{ 4¤<×#;Ò#;¸E¼LÈ%Ì,Ñ#WÔ#WÐY^ÔYcÑdÔdÐdÝÐ<¨?Ð<Ð<Ñ=Ô=Ð=r   Úieee_roundingc                ó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)	rg   rŠ   r�   rA   rš   r'   rÈ   r®   r&   )r   r«   r¬   rÐ   r¶   r·   r¥   s          r   ÚfdivzTritonSemantic.fdivK  s¢   € Øœ*Ô+ˆØœ*Ô+ˆØ×*Ò*Ñ,Ô,ð 	T°O×4OÒ4OÑ4QÔ4Qð 	TÝÐRÑSÔSÐSØ×8Ò8¸ÀÀuÈeÐUZÐ\`ÑaÔa‰ˆˆuØŒ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š   rg   rŠ   r�   r&   r'   Úcreate_fremr®   rY   r=   rA   r   rË   Úcreate_sremÚcreate_urem)r   r«   r¬   rZ   r·   s        r   ÚmodzTritonSemantic.modT  sx  € Ø×8Ò8¸ÀÀuÈeÐUYÐ[_Ñ`Ô`‰ˆˆuØ”JÔ%ˆ	Øœ*Ô+ˆà× Ò Ñ"Ô"ð 	eØ—;’;˜tœ|×7Ò7¸¼ÀeÄlÑSÔSÐUZÔU_Ñ`Ô`Ð`à×ÒÑÔð 	eØÔ'¨?Ô+IÒIÐIÝ °	×0BÒ0BÑ0DÔ0DÑ DÀvÑ MÐP_×PhÒPhÑPjÔPjÑ jð nsñ !sñ tô tð tð ×&Ò&Ñ(Ô(ð eØ—{’{ 4¤<×#;Ò#;¸E¼LÈ%Ì,Ñ#WÔ#WÐY^ÔYcÑdÔdÐdà—{’{ 4¤<×#;Ò#;¸E¼LÈ%Ì,Ñ#WÔ#WÐY^ÔYcÑdÔdÐdÝÐ6¨9Ð6Ð6Ñ7Ô7Ð7r   rh   ÚyÚpropagate_nanútl.PropagateNanc                óF  — |                       ||¦  «        \  }}|j        }|                     ¦   «         r¸|t          j        j        k    r>|                      | j                             |j	        |j	        ¦  «        |j
        ¦  «        S |t          j        j        k    r>|                      | j                             |j	        |j	        ¦  «        |j
        ¦  «        S t          d|› �¦  «        ‚|                     ¦   «         r>|                      | j                             |j	        |j	        ¦  «        |j
        ¦  «        S |                     ¦   «         r>|                      | j                             |j	        |j	        ¦  «        |j
        ¦  «        S t%          d|› �¦  «        ‚©NzUnexpected propagate_nan úUnexpected dtype )rš   r>   r�   r1   ÚPropagateNanÚALLr&   r'   Úcreate_minimumfr®   rg   ÚNONEÚcreate_minnumfr/   rË   Úcreate_minsir‹   Úcreate_minuirA   ©r   rh   rØ   rÙ   r>   s        r   ÚminimumzTritonSemantic.minimumk  ó`  € Ø×0Ò0°°AÑ6Ô6‰ˆˆ1Ø”ˆØ×ÒÑÔð 	9Ø¥¤Ô 3Ò3Ð3Ø—{’{ 4¤<×#?Ò#?ÀÄÈ!Ì(Ñ#SÔ#SÐUVÔU[Ñ\Ô\Ð\Ø¥"¤/Ô"6Ò6Ð6Ø—{’{ 4¤<×#>Ò#>¸q¼xÈÌÑ#RÔ#RÐTUÔTZÑ[Ô[Ð[å Ð!L¸]Ð!LÐ!LÑMÔMÐMØ× Ò Ñ"Ô"ð 	9Ø—;’;˜tœ|×8Ò8¸¼À1Ä8ÑLÔLÈaÌfÑUÔUÐUØ×"Ò"Ñ$Ô$ð 	9Ø—;’;˜tœ|×8Ò8¸¼À1Ä8ÑLÔLÈaÌfÑUÔUÐUåÐ7°Ð7Ð7Ñ8Ô8Ð8r   c                óF  — |                       ||¦  «        \  }}|j        }|                     ¦   «         r¸|t          j        j        k    r>|                      | j                             |j	        |j	        ¦  «        |j
        ¦  «        S |t          j        j        k    r>|                      | j                             |j	        |j	        ¦  «        |j
        ¦  «        S t          d|› �¦  «        ‚|                     ¦   «         r>|                      | j                             |j	        |j	        ¦  «        |j
        ¦  «        S |                     ¦   «         r>|                      | j                             |j	        |j	        ¦  «        |j
        ¦  «        S t%          d|› �¦  «        ‚rÜ   )rš   r>   r�   r1   rÞ   rß   r&   r'   Úcreate_maximumfr®   rg   rá   Úcreate_maxnumfr/   rË   Úcreate_maxsir‹   Úcreate_maxuirA   rå   s        r   ÚmaximumzTritonSemantic.maximum|  rç   r   ÚminÚmaxc                ó~  — |                       ||¦  «        \  }}|                       ||¦  «        \  }}|                       ||¦  «        \  }}|j        }|                     ¦   «         rE|                      | j                             |j        |j        |j        |¦  «        |j        ¦  «        S t          d|› d�¦  «        ‚)NrÝ   z(. Only floating point clamp is supported)	rš   r>   r�   r&   r'   Úcreate_clampfr®   rg   rA   )r   rh   rî   rï   rÙ   r>   s         r   ÚclampzTritonSemantic.clamp�  sº   € Ø×4Ò4°S¸#Ñ>Ô>‰ˆˆSØ×2Ò2°1°cÑ:Ô:‰ˆˆ3Ø×2Ò2°1°cÑ:Ô:‰ˆˆ3à”ˆØ×ÒÑÔð 	aØ—;’;˜tœ|×9Ò9¸!¼(ÀCÄJÐPSÔPZÐ\iÑjÔjÐlmÔlrÑsÔsÐsåÐ_°Ð_Ð_Ð_Ñ`Ô`Ð`r   c                óx  — |                       ||¦  «        \  }}|j        j        }|j        j        }|                     ¦   «         r|                     ¦   «         st	          ||¦  «        ‚|                      ||¦  «        }||k    r|                      ||¦  «        }||k    r|                      ||¦  «        }||fS r)   )rš   rg   rŠ   rY   r   rF   rŽ   )r   r«   r¬   Úinput_sca_tyÚother_sca_tyr™   s         r   Úbitwise_op_type_checking_implz,TritonSemantic.bitwise_op_type_checking_implœ  sÃ   € Ø×8Ò8¸ÀÑFÔF‰ˆˆuØ”zÔ(ˆØ”zÔ(ˆØ×"Ò"Ñ$Ô$ð 	H¨L×,?Ò,?Ñ,AÔ,Að 	HÝ+¨L¸,ÑGÔGÐGØ×.Ò.¨|¸\ÑJÔJˆ
Ø˜Ò%Ð%Ø—I’I˜e ZÑ0Ô0ˆEØ˜Ò%Ð%Ø—I’I˜e ZÑ0Ô0ˆEØ�eˆ|Ðr   c                ó°   — |                       ||¦  «        \  }}|                      | j                             |j        |j        ¦  «        |j        ¦  «        S r)   )rö   r&   r'   Ú
create_andr®   rg   ©r   r«   r¬   s      r   r¡   zTritonSemantic.and_©  óI   € Ø×9Ò9¸%ÀÑGÔG‰ˆˆuØ�{Š{˜4œ<×2Ò2°5´<ÀÄÑNÔNÐPUÔPZÑ[Ô[Ð[r   c                ó°   — |                       ||¦  «        \  }}|                      | j                             |j        |j        ¦  «        |j        ¦  «        S r)   )rö   r&   r'   Ú	create_orr®   rg   rù   s      r   Úor_zTritonSemantic.or_­  sG   € Ø×9Ò9¸%ÀÑGÔG‰ˆˆuØ�{Š{˜4œ<×1Ò1°%´,ÀÄÑMÔMÈuÌzÑZÔZÐZr   c                ó°   — |                       ||¦  «        \  }}|                      | j                             |j        |j        ¦  «        |j        ¦  «        S r)   )rö   r&   r'   Ú
create_xorr®   rg   rù   s      r   Úxor_zTritonSemantic.xor_±  rú   r   c                ó  — |j                              ¦   «         s |                      |t          j        ¦  «        }|j                              ¦   «         s |                      |t          j        ¦  «        }|                      ||¦  «        S r)   )rg   Úis_int1Úbitcastr1   rt   r¡   rù   s      r   Úlogical_andzTritonSemantic.logical_andµ  sl   € ØŒz×!Ò!Ñ#Ô#ð 	1Ø—L’L ­¬Ñ0Ô0ˆEØŒz×!Ò!Ñ#Ô#ð 	1Ø—L’L ­¬Ñ0Ô0ˆEØ�yŠy˜ Ñ&Ô&Ð&r   c                ó  — |j                              ¦   «         s |                      |t          j        ¦  «        }|j                              ¦   «         s |                      |t          j        ¦  «        }|                      ||¦  «        S r)   )rg   r  r  r1   rt   rý   rù   s      r   Ú
logical_orzTritonSemantic.logical_or¼  sl   € ØŒz×!Ò!Ñ#Ô#ð 	1Ø—L’L ­¬Ñ0Ô0ˆEØŒz×!Ò!Ñ#Ô#ð 	1Ø—L’L ­¬Ñ0Ô0ˆEØ�xŠx˜˜uÑ%Ô%Ð%r   c                óž   — |j                              ¦   «         s |                      |t          j        ¦  «        }|                      |¦  «        S r)   )rg   r  r  r1   rt   Úinvert©r   r«   s     r   Únot_zTritonSemantic.not_Ã  s?   € ØŒz×!Ò!Ñ#Ô#ð 	1Ø—L’L ­¬Ñ0Ô0ˆEØ�{Š{˜5Ñ!Ô!Ð!r   c                ó°   — |                       ||¦  «        \  }}|                      | j                             |j        |j        ¦  «        |j        ¦  «        S r)   )rö   r&   r'   Úcreate_lshrr®   rg   rù   s      r   ÚlshrzTritonSemantic.lshrÈ  óI   € Ø×9Ò9¸%ÀÑGÔG‰ˆˆuØ�{Š{˜4œ<×3Ò3°E´LÀ%Ä,ÑOÔOÐQVÔQ[Ñ\Ô\Ð\r   c                ó°   — |                       ||¦  «        \  }}|                      | j                             |j        |j        ¦  «        |j        ¦  «        S r)   )rö   r&   r'   Úcreate_ashrr®   rg   rù   s      r   ÚashrzTritonSemantic.ashrÌ  r  r   c                ó°   — |                       ||¦  «        \  }}|                      | j                             |j        |j        ¦  «        |j        ¦  «        S r)   )rö   r&   r'   Ú
create_shlr®   rg   rù   s      r   ÚshlzTritonSemantic.shlÐ  rú   r   c                ó   — |S r)   © r	  s     r   ÚpluszTritonSemantic.plusØ  s   € Øˆr   c                óJ  — |j         j        }|                     ¦   «         r't          d|                     ¦   «         z   dz   ¦  «        ‚|                      | j                             |                     | j        ¦  «        ¦  «        |¦  «        }|  	                    ||d¦  «        S )Nz$wrong type argument to unary minus (ú)T)
rg   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  )rg   rŠ   r€   r�   r/   r   r&   r'   Úget_all_ones_valuer°   r   )r   r«   rô   Ú_1s       r   r  zTritonSemantic.invertâ  s£   € Ø”zÔ(ˆØ×ÒÑ Ô ð 	f L×$<Ò$<Ñ$>Ô$>ð 	fÝÐDÀ|×G\ÒG\ÑG^ÔG^Ñ^ÐadÑdÑeÔeÐeØ�[Š[˜œ×8Ò8¸×9KÒ9KÈDÌLÑ9YÔ9YÑZÔZÐ\hÑiÔiˆØ�yŠy˜ Ñ#Ô#Ð#r   Úvútl.block_typec                óJ   — |j                              t          j        ¦  «        S r)   )rg   r¯   r1   rt   )r   r  s     r   Ú
_bool_likezTritonSemantic._bool_likeí  s   € ØŒv×%Ò%¥b¤gÑ.Ô.Ð.r   c                ó°  — |                       ||¦  «        \  }}|j        j        }|                     ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S | 	                    ¦   «         r¬| 
                    ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S |                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S t          d|› �¦  «        ‚rÁ   )rš   rg   rŠ   r�   r&   r'   Úcreate_fcmpOGTr®   r"  rY   rË   Úcreate_icmpSGTÚcreate_icmpUGTrA   ©r   r«   r¬   rZ   s       r   Úgreater_thanzTritonSemantic.greater_thanð  ó)  € Ø×8Ò8¸ÀÑFÔF‰ˆˆuØ”JÔ%ˆ	à× Ò Ñ"Ô"ð 	tØ—;’;˜tœ|×:Ò:¸5¼<ÈÌÑVÔVÐX\×XgÒXgÐhmÑXnÔXnÑoÔoÐoà×ÒÑÔð 	tØ×&Ò&Ñ(Ô(ð tØ—{’{ 4¤<×#>Ò#>¸u¼|ÈUÌ\Ñ#ZÔ#ZÐ\`×\kÒ\kÐlqÑ\rÔ\rÑsÔsÐsà—{’{ 4¤<×#>Ò#>¸u¼|ÈUÌ\Ñ#ZÔ#ZÐ\`×\kÒ\kÐlqÑ\rÔ\rÑsÔsÐsÝÐ6¨9Ð6Ð6Ñ7Ô7Ð7r   c                ó°  — |                       ||¦  «        \  }}|j        j        }|                     ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S | 	                    ¦   «         r¬| 
                    ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S |                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S t          d|› �¦  «        ‚rÁ   )rš   rg   rŠ   r�   r&   r'   Úcreate_fcmpOGEr®   r"  rY   rË   Úcreate_icmpSGEÚcreate_icmpUGErA   r'  s       r   r£   zTritonSemantic.greater_equalþ  r)  r   c                ó°  — |                       ||¦  «        \  }}|j        j        }|                     ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S | 	                    ¦   «         r¬| 
                    ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S |                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S t          d|› �¦  «        ‚rÁ   )rš   rg   rŠ   r�   r&   r'   Úcreate_fcmpOLTr®   r"  rY   rË   Úcreate_icmpSLTÚcreate_icmpULTrA   r'  s       r   Ú	less_thanzTritonSemantic.less_than  r)  r   c                ó°  — |                       ||¦  «        \  }}|j        j        }|                     ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S | 	                    ¦   «         r¬| 
                    ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S |                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S t          d|› �¦  «        ‚rÁ   )rš   rg   rŠ   r�   r&   r'   Úcreate_fcmpOLEr®   r"  rY   rË   Úcreate_icmpSLEÚcreate_icmpULErA   r'  s       r   r¢   zTritonSemantic.less_equal  r)  r   c                óð  — |                       ||¦  «        \  }}|j        j        }|                     ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S | 	                    ¦   «         rL|                      | j         
                    |j        |j        ¦  «        |                      |¦  «        ¦  «        S t          d|› �¦  «        ‚rÁ   )rš   rg   rŠ   r�   r&   r'   Úcreate_fcmpOEQr®   r"  rY   Úcreate_icmpEQrA   r'  s       r   ÚequalzTritonSemantic.equal(  óÚ   € Ø×8Ò8¸ÀÑFÔF‰ˆˆuØ”JÔ%ˆ	à× Ò Ñ"Ô"ð 	oØ—;’;˜tœ|×:Ò:¸5¼<ÈÌÑVÔVÐX\×XgÒXgÐhmÑXnÔXnÑoÔoÐoà×ÒÑÔð 	oØ—;’;˜tœ|×9Ò9¸%¼,ÈÌÑUÔUÐW[×WfÒWfÐglÑWmÔWmÑnÔnÐnÝÐ6¨9Ð6Ð6Ñ7Ô7Ð7r   c                óð  — |                       ||¦  «        \  }}|j        j        }|                     ¦   «         rL|                      | j                             |j        |j        ¦  «        |                      |¦  «        ¦  «        S | 	                    ¦   «         rL|                      | j         
                    |j        |j        ¦  «        |                      |¦  «        ¦  «        S t          d|› �¦  «        ‚rÁ   )rš   rg   rŠ   r�   r&   r'   Úcreate_fcmpUNEr®   r"  rY   Úcreate_icmpNErA   r'  s       r   Ú	not_equalzTritonSemantic.not_equal3  r;  r   N)rÎ   ÚstartÚendrÎ   c               ó(  — t          |t          ¦  «        rt          |t          ¦  «        st          d¦  «        ‚t          |dz	  ¦  «        }t          |dz	  ¦  «        }|s|rt          d¦  «        ‚||k    rt          d¦  «        ‚||z
  }||dz
  z  dk    rt          d¦  «        ‚|g}|€t	          j        t          j        |¦  «        }|                     | j        ¦  «        }|  	                    | j         
                    |||¦  «        |¦  «        S )Nz/arange's arguments must be of type tl.constexpré    zarange must fit in int32z=arange's end argument must be greater than the start argumentr   r   z#arange's range must be a power of 2)rb   r+   r/   rH   r1   Ú
block_typer2   r°   r'   r&   Úcreate_make_range)	r   r@  rA  rÎ   Úis_start_int64Úis_end_int64ÚrangeÚshapeÚ	ret_ty_irs	            r   ÚarangezTritonSemantic.arangeB  s  € Ý˜%¥Ñ%Ô%ð 	P­Z¸½SÑ-AÔ-Að 	PÝÐNÑOÔOÐOÝ˜e r™kÑ*Ô*ˆÝ˜C 2™I‘”ˆØð 	9˜\ð 	9ÝÐ7Ñ8Ô8Ð8Ø�%Š<ˆ<ÝÐ\Ñ]Ô]Ð]Ø�e‘ˆØ�U˜Q‘YÑ AÒ%Ð%ÝÐBÑCÔCÐCØ�ˆØˆ>Ý”]¥2¤8¨UÑ3Ô3ˆFØ—L’L ¤Ñ.Ô.ˆ	Ø�{Š{˜4œ<×9Ò9¸)ÀUÈCÑPÔPÐRXÑYÔYÐYr   r>   c                óÖ  — |€t          d¦  «        ‚|dk    r3| j                             |                     | j        ¦  «        ¦  «        }nŠ|                     ¦   «         rN| j                             |¦  «        }| j                             ||                     | j        ¦  «        ¦  «        }n(t          | j        d|j        › �¦  «        } ||¦  «        }|  	                    ||¦  «        S )Nz2dtype must be specified when value is not a tensorr   Úget_)
r/   r'   r  r°   rX   Úget_fp32Úcreate_fp_truncÚgetattrÚnamer&   )r   rO   r>   Úget_value_fns       r   rf   zTritonSemantic.scalar_constantT  sÏ   € àˆ=ÝÐQÑRÔRÐRØ�AŠ:ˆ:Ø”L×/Ò/°·²¸D¼LÑ0IÔ0IÑJÔJˆEˆEØ�\Š\‰^Œ^ð 	(Ø”L×)Ò)¨%Ñ0Ô0ˆEØ”L×0Ò0°¸¿ºÀDÄLÑ8QÔ8QÑRÔRˆEˆEå" 4¤<Ð1D¸¼
Ð1DÐ1DÑEÔEˆLØ �L Ñ'Ô'ˆEØ�{Š{˜5 %Ñ(Ô(Ð(r   c                óÂ   — t          |t          j        ¦  «        r0|j        j        dk    s
J d¦   «         ‚|                      ||¦  «        S |                      ||¦  «        S )Nr   zonly accepts size-1 tensor)rb   r1   r&   ÚnumelrO   rŽ   rf   )r   rO   r>   s      r   Úmake_scalarzTritonSemantic.make_scalarb  s_   € Ý�e�RœYÑ'Ô'ð 	+Ø”;Ô$¨Ò)Ð)Ð)Ð+GÑ)Ô)Ð)Ø—9’9˜U EÑ*Ô*Ð*à×#Ò# E¨5Ñ1Ô1Ð1r   rI  ú	List[int]c                óV   — |                       |                      ||¦  «        |¦  «        S r)   )ÚsplatrU  )r   rI  rO   r>   s       r   ÚfullzTritonSemantic.fulli  s&   € Ø�zŠz˜$×*Ò*¨5°%Ñ8Ô8¸%Ñ@Ô@Ð@r   rO   c                ó>  — |j                              ¦   «         r
J d¦   «         ‚t          |¦  «        dk    r|S t          j        |j        |¦  «        }|                      | j                             | 	                    | j        ¦  «        |j
        ¦  «        |¦  «        S )NzCannot splat a block tensorr   )rg   Úis_blockÚlenr1   rD  r>   r&   r'   Úcreate_splatr°   r®   )r   rO   rI  rÎ   s       r   rX  zTritonSemantic.splatp  s…   € Ø”:×&Ò&Ñ(Ô(ÐGÐGÐ*GÑGÔGÐ(Ýˆu‰:Œ:˜Š?ˆ?ØˆLÝ”˜uœ{¨EÑ2Ô2ˆØ�{Š{˜4œ<×4Ò4°V·\²\À$Ä,Ñ5OÔ5OÐQVÔQ]Ñ^Ô^Ð`fÑgÔgÐgr   c                ór   — |                       | j                             |j        ¦  «        |j        ¦  «        S r)   )r&   r'   Úcreate_unsplatr®   r>   )r   rO   s     r   ÚunsplatzTritonSemantic.unsplatw  s*   € Ø�{Š{˜4œ<×6Ò6°u´|ÑDÔDÀeÄkÑRÔRÐRr   Ú	dst_shapeÚcan_reorderc                ó   — d}|D ]}||z  }Œ|j         j        |k    rt          d¦  «        ‚t          j        |j         j        |¦  «        }|                      | j                             |j	        ||¦  «        |¦  «        S )Nr   z:reshape() cannot change total number of elements in tensor)
rg   rT  r/   r1   rD  rŠ   r&   r'   Úcreate_reshaper®   )r   r«   ra  rb  rT  ÚsrÎ   s          r   ÚreshapezTritonSemantic.reshapez  s‚   € ØˆØð 	ð 	ˆAØ�Q‰JˆEˆEØŒ:Ô˜uÒ$Ð$ÝÐYÑZÔZÐZÝ”˜uœzÔ0°)Ñ<Ô<ˆØ�{Š{˜4œ<×6Ò6°u´|ÀYÐP[Ñ\Ô\Ð^dÑeÔeÐer   c                óV  — d„ |j         D ¦   «         }|                     |d¦  «         |j                             ¦   «         s|                      ||¬¦  «        S t          j        |j        j        |¦  «        }|                      | j	         
                    |j        |¦  «        |¦  «        S )Nc                ó6   — g | ]}t          j        |¦  «        ‘ŒS r  ©r1   Ú_unwrap_if_constexpr)Ú.0rh   s     r   ú
<listcomp>z.TritonSemantic.expand_dims.<locals>.<listcomp>„  s#   € ÐEÐEÐE°A•RÔ,¨QÑ/Ô/ÐEÐEÐEr   r   ©rI  )rI  Úinsertrg   r[  rX  r1   rD  rŠ   r&   r'   Úcreate_expand_dimsr®   )r   r«   r*   ra  rÎ   s        r   Úexpand_dimszTritonSemantic.expand_dimsƒ  s˜   € ØEÐE¸¼ÐEÑEÔEˆ	Ø×Ò˜˜qÑ!Ô!Ð!àŒz×"Ò"Ñ$Ô$ð 	6Ø—:’:˜e¨9�:Ñ5Ô5Ð5å”˜uœzÔ0°)Ñ<Ô<ˆØ�{Š{˜4œ<×:Ò:¸5¼<ÈÑNÔNÐPVÑWÔWÐWr   c                ó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   )
r\  rI  r1   rD  rg   rŠ   r&   r'   Ú
create_catr®   )r   rƒ   r…   rb  Úret_types        r   ÚcatzTritonSemantic.cat�  s‚   € ØÐ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©rb  )r�   rI  rp  rb   r1   rc   rD  rg   rŠ   r&   r'   Úcreate_joinr®   rf  )r   ru  rv  Ú
was_rank_1ÚtwoÚ	new_shapers  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   rx  r   )r\  rI  r1   rj  rD  rg   rŠ   r'   Úcreate_splitr®   r&   )r   ru  r}  rs  Ú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    r4t          dt          ‰j        ¦  «        › dt          |¦  «        › �¦  «        ‚t          d„ |D ¦   «         ¦  «        t	          t          t          |¦  «        ¦  «        ¦  «        k    rt          d|› �¦  «        ‚t          j        ‰j        j	        ˆfd„|D ¦   «         ¦  «        }|  
                    | j                             ‰j        |¦  «        |¦  «        S )Nz;permute dims must have the same length as input shape, got r   c              3  ó>   K  — | ]}t          j        |¦  «        V — Œd S r)   ri  )rk  Ú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  rm  )rk  rˆ  r«   s     €r   rl  z*TritonSemantic.permute.<locals>.<listcomp>¾  s   ø€ Ð4RÐ4RÐ4RÈ°U´[À´^Ð4RÐ4RÐ4Rr   )r\  rI  r/   ÚsortedÚlistrH  r1   rD  rg   rŠ   r&   r'   Úcreate_transr®   )r   r«   r„  rs  s    `  r   ÚpermutezTritonSemantic.permute·  sÿ   ø€ ÝˆuŒ{ÑÔ�s 4™yœyÒ(Ð(ÝØpÍcÐRWÔR]ÑN^ÔN^ÐpÐpÕehÐimÑenÔenÐpÐpñrô rð råÐ;Ð;°dÐ;Ñ;Ô;Ñ;Ô;½tÅEÍ#ÈdÉ)Ì)ÑDTÔDTÑ?UÔ?UÒUÐUÝÐeÐ_cÐeÐeÑfÔfÐfå”= ¤Ô!2Ð4RÐ4RÐ4RÐ4RÈTÐ4RÑ4RÔ4RÑSÔSˆØ�{Š{˜4œ<×4Ò4°U´\À4ÑHÔHÈ(ÑSÔSÐSr   c                óH  — |j                              ¦   «         s|                      ||¦  «        S |j                              ¦   «         }t	          |¦  «        t	          |¦  «        k    rt          d|› d|› �¦  «        ‚||k    r|S t          |¦  «        D ];\  }}||         |k    r*|dk    r$t          d||         › d|› d|› d|› d|› �
¦  «        ‚Œ<t          j        |j         j	        |¦  «        }|  
                    | j                             |j        |¦  «        |¦  «        S )Nz!Cannot broadcast, rank mismatch: ú, r   z3Cannot broadcast, the expanded size of the tensor (z ) must match the existing size (z) at non-singleton dimension ú: )rg   r[  rX  Úget_block_shapesr\  r/   Ú	enumerater1   rD  rŠ   r&   r'   Úcreate_broadcastr®   )r   r«   rI  Ú	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   )rg   r[  r¯   rŠ   r&   r'   r]  r°   r®   r’  r\  rH  ro  r1   rD  Ú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 )rg   r[  r¯   rŠ   r€   rŽ   Úprimitive_bitwidthr/   rœ  r&   r'   Úcreate_bitcastr®   r°   )r   r«   r®  Úsrc_tyÚ
src_sca_tyÚ
dst_sca_tyÚsrc_bitsÚdst_bitss           r   r  zTritonSemantic.bitcast  s  € Ø”ˆØ�?Š?ÑÔð 	;Ø×+Ò+¨F¬MÑ:Ô:ˆFØ�VÒÐØˆLØ”]ˆ
Ø”]ˆ
Ø×ÒÑÔð 	, *×"3Ò"3Ñ"5Ô"5ð 	,Ø—9’9˜U FÑ+Ô+Ð+àÔ0ˆØÔ0ˆØ�xÒÐÝÐ@Å3ÀxÁ=Ä=ÑPð T2ñ 2Ý47¸±M´MñBñ Cô Cð Cà�{Š{˜4œ<×6Ò6°u´|ÀVÇ\Â\ÐRVÔR^ÑE_ÔE_Ñ`Ô`ÐbhÑiÔiÐir   Úfp_downcast_roundingc                óê  — |j         }|j        }|j        }||k    r|S |                     ¦   «         r|                     |¦  «        }|                      |¦  «        }d}|                     ¦   «         rP|                     ¦   «         r<|j        |j        k     r,|€t          j        j	        }nL|t          j        j	        k    rd}n4|�2t          dt          |¦  «        z   dz   t          |¦  «        z   ¦  «        ‚|                     ¦   «         s|                     ¦   «         rI| j        j                             d¦  «        	 €
J d¦   «         ‚ | j        j        d         |||| ¬¦  «        S |                     ¦   «         r|                     ¦   «         s*|                     ¦   «         r|                     ¦   «         s|rM|                      | j                             |j        |                     | j        ¦  «        |¦  «        |¦  «        S |                     ¦   «         r|                     ¦   «         r(|                     ¦   «         rH|                     ¦   «         s4|                      |                      |t2          j        ¦  «        |¦  «        S |                     ¦   «         o#|                     ¦   «         o|j        |j        k    }|rL|                      | j                             |j        |                     | j        ¦  «        ¦  «        |¦  «        S |                     ¦   «         o#|                     ¦   «         o|j        |j        k     }	|	rL|                      | j                             |j        |                     | j        ¦  «        ¦  «        |¦  «        S |                     ¦   «         �r'|                     ¦   «         �r|j        |j        k    s|j        |j        k    rò|                      ¦   «         o| !                    ¦   «          }
| !                    ¦   «         rh|j"                             | j        ¦  «        }|                      | j         #                    |¦  «        |j"        ¦  «        }|  $                    ||¦  «        S |                      | j         %                    |j        |                     | j        ¦  «        |
¦  «        |¦  «        S | &                    ¦   «         �r=|                     ¦   «         �r(| !                    ¦   «         rh|j"                             | j        ¦  «        }|                      | j         #                    |¦  «        |j"        ¦  «        }|  $                    ||¦  «        S |                      ¦   «         rL|                      | j         '                    |j        |                     | j        ¦  «        ¦  «        |¦  «        S |                      | j         (                    |j        |                     | j        ¦  «        ¦  «        |¦  «        S |                     ¦   «         rÔ| &                    ¦   «         rÀ| !                    ¦   «         s|                      ¦   «         sL|                      | j         )                    |j        |                     | j        ¦  «        ¦  «        |¦  «        S |                      | j         *                    |j        |                     | j        ¦  «        ¦  «        |¦  «        S | +                    ¦   «         rÝ|                     ¦   «         rÉ|j        }|dk    rL|                      | j         ,                    |j        |                     | j        ¦  «        ¦  «        |¦  «        S |d	k    rj|  $                    |                      |t2          j-        ¦  «        |                      | j         .                    d
¦  «        t2          j-        ¦  «        ¦  «        S |                     ¦   «         r`| +                    ¦   «         rL|                      | j         /                    |j        |                     | j        ¦  «        ¦  «        |¦  «        S | +                    ¦   «         r`| +                    ¦   «         rL|                      | j         0                    |j        |                     | j        ¦  «        ¦  «        |¦  «        S J d|› d|› �¦   «         ‚)NFTz]fp_downcast_rounding should be set only for truncating fp conversions. Source scalar type is z and destination type is Úconvert_custom_typesz0target doesn't provide conversion for this type.)Ú	_semanticrž   r   r   zcannot cast z to )1rg   rŠ   r[  r¯   r­  r�   r°  r   rª  r«  r/   rœ  Úis_fp8e4b15r'   Úcodegen_fnsÚgetrX   r&   Úcreate_fp_to_fpr®   r°   rV   rU   rW   rŽ   r1   rR   rO  Ú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_intrv   Ú	get_int64Úcreate_int_to_ptrr±  )r   r«   r®  r·  r²  r³  r´  Úuse_custom_roundingÚtruncate_fpÚext_fpÚsign_extendÚtyr  Úbitwidths                 r   rŽ   zTritonSemantic.cast$  sc  € Ø”ˆØ”]ˆ
Ø”]ˆ
Ø˜Ò#Ð#Øˆ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ˆur   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  ©rb   r1   rc   rO   ©rk  Úelems     r   rl  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  )Úhasattrrb   r+   r\  Ú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ÐKÐ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/   rg   Ú
element_tyr1   rt   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                             ¦   «         s*t          d|j                              ¦   «         › d�¦  «        ‚|€|�t          d¦  «        ‚|s|rt          d¦  «        ‚|j                              ¦   «         sT|r(|j                              ¦   «         rt          d¦  «        ‚|r(|j                              ¦   «         rt          d¦  «        ‚|j                              ¦   «         r6|�|                      ||¦  «        \  }}|�|                      ||¦  «        \  }}|j         j        }	|	j        }
|
t          j	        k    }|r<t          j
        }
t          j        |
|	j        ¦  «        }	|                      ||	¦  «        }|�|                      ||
¦  «        }|j                              ¦   «         r|j                              |
¦  «        }n|
}|€7|                      | j                             |j        |||¦  «        |¦  «        }nF|                      | j                             |j        |j        |r|j        nd |||¦  «        |¦  «        }|r |                      |t          j	        ¦  «        }|S )NúUnsupported ptr type z in `tl.load`z)`other` cannot be provided without `mask`zä`padding_option` or `boundary_check` argument is not supported for loading a tensor ofpointers or loading a scalar. Because the compiler does not know the boundary; please use block pointers (defined by `make_block_ptr`) insteadúEMask argument cannot be block type if pointer argument is not a blockzFOther argument cannot be block type if pointer argument is not a block)rg   rŠ   r€   r/   r   r[  r�   r  r1   rt   Ú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×ÒÑÔð 	CØÐØ ×5Ò5°c¸4Ñ@Ô@‘	��TØÐ Ø!×6Ò6°s¸EÑBÔB‘
��Uð ””ˆØÔ"ˆð �BœGÒ#ˆØð 	)Ý”WˆFÝ”_ V¨VÔ-AÑBÔBˆFØ—)’)˜C Ñ(Ô(ˆCð ÐØ—I’I˜e VÑ,Ô,ˆEð Œ8×ÒÑÔð 	Ø”X×-Ò-¨fÑ5Ô5ˆFˆFð ˆFð ˆ<Ø—+’+˜dœl×6Ò6°s´zÀ5È(ÐT_Ñ`Ô`ÐbhÑiÔiˆCˆCà—+’+Ø”×/Ò/°´
¸D¼KÐY^ÐIhÈÌÈÐdhÐjoØ08¸+ñGô GØHNñPô PˆCð ð 	*Ø—)’)˜C¥¤Ñ)Ô)ˆCØˆ
r   r  r  úOptional[TensorTy]r  r   rð  rœ  rØ  rç  r  c	           
     ó^  — |                       |¦  «        }	|                      |¦  «        }
|                      |¦  «        }|j                             ¦   «         r:|j        j                             ¦   «         r|                      ||||||	|
|¦  «        S |                      ||||||	|
|¦  «        S r)   )	rÚ  ré  rò  rg   r€   r  r[  r!  r+  )r   r  r  r¬   r  rð  rØ  rç  r  rÙ  rè  rñ  s               r   ÚloadzTritonSemantic.load?  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)rb   r1   Útensor_descriptor_baser\  r  Ú_convert_to_ir_valuesr'   Úcreate_descriptor_loadr®   rÚ  ré  r&   rD  )r   r/  ÚoffsetsrØ  rç  Úndimrh   s          r   Údescriptor_loadzTritonSemantic.descriptor_loadM  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 )Nr2  r3  )rb   r1   r6  r\  r  rI  )r   r/  rO   r9  r:  s        r   Úvalidate_store_likez"TritonSemantic.validate_store_likeX  s}   € Ý˜$¥Ô 9Ñ:Ô:Ð:Ð:Ð:Ý�4Ô#Ñ$Ô$ˆÝ�7‰|Œ|˜tÒ#Ð#Ð#Ð%W°Ð%WÐ%WÍÈWÉÌÐ%WÐ%WÑ#Ô#Ð#ØŒ{˜dÔ.Ò.Ð.Ð.Ð.Ð.Ð.r   c                ó  — |                       |||¦  «         |                      ||j        ¦  «        }|                      |d¬¦  «        }|                      | j                             |j        |j        |¦  «        t          j	        ¦  «        S ©NFr4  )
r=  rŽ   r>   r7  r&   r'   Úcreate_descriptor_storer®   r1   Úvoid)r   r/  rO   r9  s       r   Údescriptor_storezTritonSemantic.descriptor_store^  sx   € Ø× Ò   u¨gÑ6Ô6Ð6à—	’	˜% ¤Ñ,Ô,ˆØ×,Ò,¨WÀ%Ð,ÑHÔHˆØ�{Š{˜4œ<×?Ò?ÀÄÈUÌ\Ð[bÑcÔcÕegÔelÑmÔmÐmr   c                ó´  — |                       |||¦  «         |j        t          j        t          j        t          j        t          j        t          j        t          j        hv s
J d¦   «         ‚|  	                    |d¬¦  «        }t          j        j        }|                      | j                             ||j        |j        |¦  «        t          j        ¦  «        S ©NúUnsupported dtypeFr4  )r=  r>   r1   ru   r2   rw   rR   rP   rQ   r7  r   ÚDESCRIPTOR_REDUCE_KINDÚADDr&   r'   Úcreate_descriptor_reducer®   rA  ©r   r/  rO   r9  rN   s        r   Údescriptor_atomic_addz$TritonSemantic.descriptor_atomic_adde  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_tmal  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 )NrE  z-16-bit float types require native tma support)r1   ru   r2   rw   rv   rP   rQ   rS  )r   r>   s     r   Ú$_descriptor_atomic_min_max_supportedz3TritonSemantic._descriptor_atomic_min_max_supportedp  sx   € Ø�œ¥B¤H­b¬i½¼Å2Ä:ÍrÌ{Ð[Ð[Ð[Ð[Ð]pÑ[Ô[Ð[Ø•R”Z¥¤Ð-Ð-Ð-Ø×'Ò'Ñ)Ô)ÐZÐZÐ+ZÑZÔZÐ)ð .Ð-ØZÐZr   c                ó>  — |                       |||¦  «         |                      |j        ¦  «         |                      |d¬¦  «        }t          j        j        }|                      | j         	                    ||j
        |j
        |¦  «        t          j        ¦  «        S r?  )r=  rU  r>   r7  r   rF  ÚMINr&   r'   rH  r®   r1   rA  rI  s        r   Údescriptor_atomic_minz$TritonSemantic.descriptor_atomic_minu  óˆ   € Ø× Ò   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 r?  )r=  rU  r>   r7  r   rF  ÚMAXr&   r'   rH  r®   r1   rA  rI  s        r   Údescriptor_atomic_maxz$TritonSemantic.descriptor_atomic_max|  rY  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 rD  )r=  r>   r1   ru   r2   rw   rv   r7  r   rF  ÚANDr&   r'   rH  r®   rA  rI  s        r   Údescriptor_atomic_andz$TritonSemantic.descriptor_atomic_andƒ  óŸ   € Ø× Ò   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 rD  )r=  r>   r1   ru   r2   rw   rv   r7  r   rF  ÚORr&   r'   rH  r®   rA  rI  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 rD  )r=  r>   r1   ru   r2   rw   rv   r7  r   rF  ÚXORr&   r'   rH  r®   rA  rI  s        r   Údescriptor_atomic_xorz$TritonSemantic.descriptor_atomic_xor‘  r`  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 rC  zdescriptor gather of ú must have at least ú columns, but got Fr4  )rb   r1   r6  r\  r  rI  r>   r°  rD  r7  r'   Úcreate_descriptor_gatherr®   r°   r&   )
r   r/  Ú	x_offsetsÚy_offsetrØ  rç  r>   Úmin_colsrg   rh   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   ri  r   r   rj  rk  rl  z6descriptor scatter must have at least 8 rows, but got rC  zdescriptor scatter of rm  rn  Fr4  )rb   r1   r6  r\  r  rI  Úshapaer>   r°  r7  r'   Úcreate_descriptor_scatterr®   r&   rA  )r   r/  rO   rp  rq  r>   rr  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/   rg   r  r’  r[  r˜  r1   rt   r  rŽ   r&   r'   Úcreate_tensor_pointer_storer®   rA  )	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Ð"Ø˜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                              ¦   «         rc|j        }|€|                      ||¦  «        \  }}n|                      |||¦  «        \  }}}||j        k    rt          d|j        › d|› �¦  «        ‚|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$  z(Expected pointer argument to have shape z	 but got ú"Mask must have boolean scalar type)rg   rŠ   r€   r/   r   r[  rI  Úbroadcast_tensorsr  r1   rt   r%  r&  r'  rŽ   r&   r'   Úcreate_storer®   rA  rÀ  Úcreate_masked_store)
r   r  rz  r  r  rÙ  rè  Ú	ptr_shaper*  r   s
             r   Ú_store_legacyzTritonSemantic._store_legacyá  sh  € àŒ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×ÒÑÔð 	mØœ	ˆIØˆ|Ø×1Ò1°#°sÑ;Ô;‘��S�Sà!%×!7Ò!7¸¸SÀ$Ñ!GÔ!G‘��S˜$Ø˜CœIÒ%Ð%Ý Ð!kÈCÌIÐ!kÐ!kÐ`iÐ!kÐ!kÑlÔlÐlà””ˆØÔ"ˆð •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   rz  c                ó¸  — |                       |¦  «        }|                      |¦  «        }|j                             ¦   «         s|j        j                             ¦   «         rt          d¦  «        ‚|j                             ¦   «         r8|j        j                             ¦   «         r|  	                    ||||||¦  «        S |  
                    ||||||¦  «        S )Nú"Cannot store to a constant pointer)rß  ré  rg   Úis_constrŠ   r/   r€   r  r[  r{  r‚  )	r   r  rz  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)é   rC  rž   z9atomic_cas only supports elements with width {16, 32, 64})r   r  rg   rŠ   r  r°  r/   r&   r'   Úcreate_atomic_casr®   )r   r  r‡  rz  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 r„  r´   Úatomic_z does not support fp16z does not support bf16r‰  z does not support T)rg   rŠ   r€   r/   r   r…  r  r1   rP   rQ   Úint16Úuint16r°  rœ  r[  r˜  r’  rŽ   r'   Úget_int1rt   r¯   r]  r°   r&   )r   r  rz  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Ž   rt   )r   rh   rÎ  ÚidtypeÚixÚsignbits         r   Ú_signbitzTritonSemantic._signbitI  sa   € Ø”7Ô-ˆÝÔ!¨8¸EÐBÑBÔBˆØ�\Š\˜!˜VÑ$Ô$ˆØ—)’)˜B ¨1¡Ñ-Ô-ˆØ�yŠy˜¥"¤'Ñ*Ô*Ð*r   c                ó˜  — |                       |||d¦  «        \  }}}|                      |¦  «        }|                      |¦  «        }|j        j        }|                     ¦   «         rÀ|                     ¦   «         rV|                      | j         	                    t          j        j        |j        |j        |j        ||¦  «        |j        ¦  «        S |                      | j         	                    t          j        j        |j        |j        |j        ||¦  «        |j        ¦  «        S |t          j        t          j        hvrt%          d|› �¦  «        ‚|t          j        k    rt          j        nt          j        }|                      ||¦  «        }|                      |t          j        |d¦  «        ¦  «        }	|t          j        k    rt          j        nt          j        }
|                      ||
¦  «        }|                      |t          j        |
d¦  «        ¦  «        }|                      |¦  «        }|                      |¦  «        }|                      | j         	                    t          j        j        |	j        |j        |                      ||¦  «        j        ||¦  «        |j        ¦  «        }|                      | j         	                    t          j        j        |j        |j        |                      ||¦  «        j        ||¦  «        |j        ¦  «        }|                      |||¦  «        }|                      ||¦  «        S )Nrï   z#atomic_max not supported for dtype r   )r•  r   r  rg   rŠ   rY   rË   r&   r'   Úcreate_atomic_rmwr   Ú	ATOMIC_OPr[  r®   ÚUMAXr1   rR   rT   rA   r2   rv   r  r&  ru   rw   rœ  r
  r¡   ÚUMINÚwhere©r   r  rz  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_maxP  óÊ  € Ø×8Ò8¸¸cÀ4ÈÑOÔO‰ˆˆS�$Ø×Ò˜sÑ#Ô#ˆØ×"Ò" 5Ñ)Ô)ˆØ””ˆà�=Š=‰?Œ?ð 	Ø×#Ò#Ñ%Ô%ð Ø—{’{Ø”L×2Ò2µ2´<Ô3CÀSÄZÐQTÔQ[Ð]aÔ]hÐjmÐotÑuÔuØ”Hñô ð ð —{’{Ø”L×2Ò2µ2´<Ô3DÀcÄjÐRUÔR\Ð^bÔ^iÐknÐpuÑvÔvØ”Hñô ð ð �"œ*¥b¤jÐ1Ð1Ð1ÝÐJÀ&ÐJÐJÑKÔKÐKà#¥r¤zÒ1Ð1•”�µr´xˆØ—’˜S &Ñ)Ô)ˆØ—’˜S¥"¤/°&¸!Ñ"<Ô"<Ñ=Ô=ˆØ%­¬Ò3Ð3•"”)�)½¼ˆØ—’˜c 7Ñ+Ô+ˆØ—’˜c¥2¤?°7¸AÑ#>Ô#>Ñ?Ô?ˆØ�mŠm˜CÑ Ô ˆØ�iŠi˜‰nŒnˆØ—+’+ØŒL×*Ò*­2¬<Ô+;¸U¼\È5Ì<Ø+/¯9ª9°T¸3Ñ+?Ô+?Ô+FÈÈUñTô TØUZÔU_ñaô aˆð —+’+ØŒL×*Ò*­2¬<Ô+<¸f¼mÈVÌ]Ø+/¯9ª9°T¸3Ñ+?Ô+?Ô+FÈÈUñTô TØU[ÔU`ñbô bˆð �jŠj˜˜g wÑ/Ô/ˆØ�|Š|˜C Ñ(Ô(Ð(r   c                ó˜  — |                       |||d¦  «        \  }}}|                      |¦  «        }|                      |¦  «        }|j        j        }|                     ¦   «         rÀ|                     ¦   «         rV|                      | j         	                    t          j        j        |j        |j        |j        ||¦  «        |j        ¦  «        S |                      | j         	                    t          j        j        |j        |j        |j        ||¦  «        |j        ¦  «        S |t          j        t          j        hvrt%          d|› �¦  «        ‚|t          j        k    rt          j        nt          j        }|                      ||¦  «        }|                      |t          j        |d¦  «        ¦  «        }	|t          j        k    rt          j        nt          j        }
|                      ||
¦  «        }|                      |t          j        |
d¦  «        ¦  «        }|                      |¦  «        }|                      |¦  «        }|                      | j         	                    t          j        j        |	j        |j        |                      ||¦  «        j        ||¦  «        |j        ¦  «        }|                      | j         	                    t          j        j        |j        |j        |                      ||¦  «        j        ||¦  «        |j        ¦  «        }|                      |||¦  «        }|                      ||¦  «        S )Nrî   z#atomic_min not supported for dtype r   )r•  r   r  rg   rŠ   rY   rË   r&   r'   rž  r   rŸ  rW  r®   r¡  r1   rR   rT   rA   r2   rv   r  r&  ru   rw   rœ  r
  r¡   r   r¢  r£  s                     r   Ú
atomic_minzTritonSemantic.atomic_minv  r°  r   c           
     ó   — |                       |||d¦  «        \  }}}|                      |¦  «        }|                      |¦  «        }|j        j        }|                     ¦   «         rt          j        j        nt          j        j	        }|  
                    | j                             ||j        |j        |j        ||¦  «        |j        ¦  «        S )Nr´   )r•  r   r  rg   rŠ   r�   r   rŸ  ÚFADDrG  r&   r'   rž  r®   )r   r  rz  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Ÿ  r^  r®   rg   ©r   r  rz  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Ÿ  rb  r®   rg   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Ÿ  re  r®   rg   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®   rg   r¸  s         r   Úatomic_xchgzTritonSemantic.atomic_xchgº  sŽ   € Ø×8Ò8¸¸cÀ4ÈÑPÔP‰ˆˆS�$Ø×Ò˜sÑ#Ô#ˆØ×"Ò" 5Ñ)Ô)ˆØ�{Š{ØŒL×*Ò*­2¬<Ô+<¸c¼jÈ#Ì*ÐVZÔVaÐcfÐhmÑnÔnØŒHñô ð 	r   c                ó  — |                      ¦   «         | j        j        j        v sJ d| j        j        j        › d|› �¦   «         ‚|                     ¦   «         }|dk    rd}|dk    rd}|dk    rd}t          t          j        |¦  «        S )	Nzinput_precision must be one of ú. Got ÚTF32X3ÚTF32x3ÚBF16X3ÚBF16x3ÚBF16X6ÚBF16x6)Úlowerr'   rŸ   Úallowed_dot_input_precisionsÚupperrP  r   ÚINPUT_PRECISION)r   Úinput_precisions     r   Ú_str_to_dot_input_precisionz*TritonSemantic._str_to_dot_input_precisionÆ  s    € Ø×$Ò$Ñ&Ô&¨$¬,Ô*>Ô*[Ð[Ð[Ð[Øx¨d¬lÔ.BÔ._ÐxÐxÐgvÐxÐxñ \Ô[Ð[à)×/Ò/Ñ1Ô1ˆØ˜hÒ&Ð&Ø&ˆOØ˜hÒ&Ð&Ø&ˆOØ˜hÒ&Ð&Ø&ˆOÝ•rÔ)¨?Ñ;Ô;Ð;r   ÚaccrÑ  Úmax_num_imprecise_accÚ	out_dtypec           
     ób  — |j                              ¦   «         r|j                              ¦   «         sJ ‚|j                             ¦   «         r|j                             ¦   «         rnä|j        t          j        t          j        t          j        t          j        t          j	        t          j
        fv sJ d|j        › �¦   «         ‚|j        t          j        t          j        t          j        t          j        t          j	        t          j
        fv sJ d|j        › �¦   «         ‚|j        |j        k    sJ d|j        › d|j        › �¦   «         ‚|j                             ¦   «         s|j                             ¦   «         rgd| j        j        j        v rt          j        d¦  «         |                      |t          j        ¦  «        }|                      |t          j        ¦  «        }|j                             ¦   «         p|j                             ¦   «         }|j                             ¦   «         p|j                             ¦   «         }|s|r‡|rdnd}	|	| j        j        j        v rn| j        j        j        }
t          j        |	› d	|
› d
|
› d�¦  «         |                      |t          j        ¦  «        }|                      |t          j        ¦  «        }|€| j        j        j        }|                      |¦  «        }t/          |j        ¦  «        }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                             ¦   «         rL|j         j        t          j        k    s
J d¦   «         ‚| j                             d¦  «        }t          j        }�n|                      ¦   «         rtC          d¦  «        ‚|j         j         "                    ¦   «         s|j         j                              ¦   «         r'| j         #                    d¦  «        }t          j	        }n�|j         j         $                    ¦   «         r'| j         %                    d¦  «        }t          j
        }nJ| &                    ¦   «         r| j         '                    d¦  «        n| j         #                    d¦  «        }|}|j         j        d         }|j         j        d         }|j         j        d         }|dk    r|j         j        d         nd }t	          j(        ||r|||gn||g¦  «        }|€4| j         )                    | *                    | j        ¦  «        |¦  «        }n.|j+        }|j         j        |j        k    r|j         j,        |k    sJ ‚|€G|j                             ¦   «         r+|j                             ¦   «         r| j        j        j-        }nQd}nN|j                             ¦   «         r5|j                             ¦   «         r||k    rtC          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 releaseÚfp8e4b8Úfp8e5b16z- is AMD gfx942 specific and not supported on z^ so it's upcasted to fp16 and can cause significant slow down. Please use OCP fp8 variants on z for performancer   é   ú+Both inputs must be either 2D or 3D; (lhs: ú	 vs rhs: r  rx  éþÿÿÿ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 ()0rg   r[  r>   rX   r1   r%  Úuint8rP   rQ   rR   rT   r»  r'   rŸ   Ú!deprecated_fp8_dot_operand_dtypesÚwarningsÚwarnrŽ   Ú
is_fp8e4b8Úis_fp8e5b16rQ  Údefault_dot_input_precisionrÒ  r\  rI  rO   r¼  r½  rŠ   rY   Ú	get_int32r2   rW   r/   rU   rN  rS   Úget_fp64rV   Úget_fp16rD  r]  r°   r®   r  Úmax_num_imprecise_acc_defaultr&   Ú
create_dot)r   rƒ   r…   rÓ  rÑ  rÔ  rÕ  Úuses_fp8e4b8Úuses_fp8e5b16Ú	type_namerQ  Ú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ÌZÝ!#¤ð!-ð -ð -ð -Ø.RÀsÄyÐ.RÐ.Rñ-ô -ð -à”9¥¤­"¬(µB´JÅÄÍRÌZÝ!#¤ð!-ð -ð -ð -Ø.RÀsÄyÐ.RÐ.Rñ-ô -ð -à”9 ¤	Ò)Ð)Ð)Ð+oÐTWÔT]Ð+oÐ+oÐdgÔdmÐ+oÐ+oÑ)Ô)Ð)àŒ9× Ò Ñ"Ô"ð 	- c¤i×&;Ò&;Ñ&=Ô&=ð 	-Ø˜Tœ\Ô1ÔSÐSÐSÝ”ð mñô ð ð —)’)˜C¥¤Ñ,Ô,ˆCØ—)’)˜C¥¤Ñ,Ô,ˆCà”y×+Ò+Ñ-Ô-ÐG°´×1EÒ1EÑ1GÔ1GˆØœ	×-Ò-Ñ/Ô/ÐJ°3´9×3HÒ3HÑ3JÔ3JˆØð 	1˜=ð 	1Ø%1ÐA˜	˜	°zˆIØ˜DœLÔ0ÔRÐRÐRØ”|Ô+Ô0�Ý”Ø ð Mð MÈtð Mð MØ6:ðMð Mð MñNô Nð Nð —i’i ¥R¤ZÑ0Ô0�Ø—i’i ¥R¤ZÑ0Ô0�àÐ"Ø"œlÔ2ÔNˆOà×:Ò:¸?ÑKÔ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ð7ð Œ8Œ?×!Ò!Ñ#Ô#ð 	&Ø”8”?¥b¤gÒ-Ð-Ð-Ð/EÑ-Ô-Ð-Ø”×'Ò'¨Ñ*Ô*ˆBÝœHˆM‰MØ×ÒÑ Ô ð 	&ÝØzñô ð ð ŒXŒ_×$Ò$Ñ&Ô&ð 	&¨#¬(¬/×*AÒ*AÑ*CÔ*Cð 	&Ø”×&Ò& qÑ)Ô)ˆBÝœJˆMˆMØŒXŒ_×$Ò$Ñ&Ô&ð 	&Ø”×&Ò& qÑ)Ô)ˆBÝœJˆMˆMà-6×->Ò->Ñ-@Ô-@Ð^�”×&Ò& qÑ)Ô)Ð)ÀdÄl×F[ÒF[Ð\]ÑF^ÔF^ˆBØ%ˆMàŒHŒN˜2ÔˆØŒHŒN˜2ÔˆØŒHŒN˜2ÔˆØ!)¨Q¢ ˆCŒHŒN˜1ÔÐ°DˆÝ”˜}¸1Ð.H¨q°!°Q¨i¨iÀ1ÀaÀ&ÑIÔIˆØˆ;Øœ×2Ò2°6·<²<ÀÄÑ3MÔ3MÈrÑRÔRˆJˆJàœˆJØ”8”> V¤\Ò1Ð1°c´hÔ6IÈYÒ6VÐ6VÐ6VÐ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: rp   )rP  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   rz  r÷  Ú	triton_tyÚunsigned_tys        r   Ú_bitcast_to_fp_typez"TritonSemantic._bitcast_to_fp_type5  s  € õ
  œ[µ"´-ÍÌÝ”Zð!ð !ß!$¢ \Ñ!2Ô!2ð 	àÐØ 6Ò)Ð)Ð)Ð+eÐWcÐ+eÐ+eÑ)Ô)Ð)Ø”9¥¤Ò(Ð(Ð(Ð*aÐVYÔV_Ð*aÐ*aÑ(Ô(Ð(ØˆJØŒ9˜	Ò!Ð!ØˆJå#%¤8µR´XÅrÄyÕZ\ÔZcÐdÐdÐeqÔrˆKØ”9 Ò+Ð+Ð+Ð-dÀ\Ð-dÐ-dÐY\ÔYbÐ-dÐ-dÑ+Ô+Ð+Ø—<’<  YÑ/Ô/Ð/r   c                óZ  — |�R|j                              ¦   «         rdnd}|j        j        }|dd …         |||z  gk    sJ d|› d||z  › d|› �¦   «         ‚|�R|j                              ¦   «         rdnd}|j        j        }|dd …         |||z  gk    sJ d|› d||z  › d|› �¦   «         ‚d S d S )Nr‰  rC  rÝ  z*lhs_scale must be a tensor of shape [..., r�  z]. Got z*rhs_scale must be a tensor of shape [..., )r>   Ú
is_fp8e4nvrg   rI  )	r   rñ  rò  ró  Ú	lhs_scaleÚ	rhs_scaleÚscale_factorÚlhs_scale_shapeÚrhs_scale_shapes	            r   Úverify_scaled_shapez"TritonSemantic.verify_scaled_shapeG  s"  € ØÐ Ø!*¤×!;Ò!;Ñ!=Ô!=ÐE˜2˜2À2ˆLØ'œnÔ2ˆOØ" 2 3 3Ô'Ø�1˜Ñ$ð,ò ð ð àl¸AÐlÐlÀÀlÑARÐlÐlÐ[jÐlÐlñô ð ð Ð Ø!*¤×!;Ò!;Ñ!=Ô!=ÐE˜2˜2À2ˆLØ'œnÔ2ˆOØ" 2 3 3Ô'Ø�1˜Ñ$ð,ò ð ð àl¸AÐlÐlÀÀlÑARÐlÐlÐ[jÐlÐlñô ð ð !Ð ðð r   r	  Ú
lhs_formatr
  Ú
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  }n||z  }|
s||z  }t          j
        ||r|||gn||g¦  «        }| j                             d¦  «        }|€4| j                             |                     | j        ¦  «        |¦  «        }n.|j        }|j         j        |j        k    r|j         j        |k    sJ ‚|rd n|j        } |rd n|j        }!|                      ||||rd n||rd n|¦  «         |                      | 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   )rg   r[  r\  rI  rO   rû  rb   r1   rc   r  rD  r'   rN  r]  r°   r®   r  r  r&   Úcreate_dot_scaled)"r   rƒ   r	  r  r…   r
  r  rÓ  r  r  r  rÕ  rî  rï  Úlhs_format_enumÚrhs_format_enumÚallowed_formatsÚrhs_scale_is_noneÚlhs_scale_is_nonerñ  ÚK_LHSÚK_RHSrò  ÚPACKED_AÚPACKED_BÚPACKED_A_DIMÚPACKED_B_DIMrô  ró  rÎ   r  rõ  Úrhs_scale_handleÚlhs_scale_handles"                                     r   Ú
dot_scaledzTritonSemantic.dot_scaledU  sl  € ð Œ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ˆAà�H‘ˆAØð 	Ø�H‘ˆAÝ”˜y°qÐ*D¨1¨a°¨)¨)¸qÀ!¸fÑEÔEˆØŒ\×"Ò" 1Ñ%Ô%ˆØˆ;Øœ×2Ò2°6·<²<ÀÄÑ3MÔ3MÈrÑRÔRˆJˆJàœˆJØ”8”> V¤\Ò1Ð1°c´hÔ6IÈYÒ6VÐ6VÐ6VÐVØ#4ÐJ˜4˜4¸)Ô:JÐØ#4ÐJ˜4˜4¸)Ô:JÐØ× Ò   A qÐ2CÐ*R¨$¨$ÈØ):Ð!I  À	ñ	Kô 	Kð 	Kà�{Š{ØŒL×*Ò*¨3¬:Ð7GÈÐZ]ÔZdÐfvØ+:¸IÀzÐS]Ð_iñkô kØlrñtô tð 	tr   Ú	conditionc                ó0  — |j         t          j        k    rt          j        d|j         › �¦  «         |                      |t          j        ¦  «        }|                      ||dd¦  «        \  }}|j                             ¦   «         r3|  	                    ||¦  «        \  }}|  	                    ||¦  «        \  }}n|  	                    ||¦  «        \  }}|j        }|  
                    | j                             |j        |j        |j        ¦  «        |¦  «        S )Nzgtl.where with a non-boolean condition is deprecated and will error out in a future triton release. Got T)r>   r1   rt   rá  râ  rŽ   rš   rg   r[  r�   r&   r'   Úcreate_selectr®   )r   r%  rh   rØ   r¡  rÎ   s         r   r¢  zTritonSemantic.whereŽ  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   rD  r&   )r   rh   rZ   r¢  Úres_tys        r   Úwrap_tensorzTritonSemantic.wrap_tensor¢  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 )Try  N)rf  rT  rO   )rk  Útr   s     €r   r‰  z+TritonSemantic.reduction.<locals>.<genexpr>¬  s<   øè è € Ð^Ð^ÐRS˜4Ÿ<š<¨¨A¬G¬M¨?È˜<ÑMÔMÐ^Ð^Ð^Ð^Ð^Ð^r   r   z&reduction axis must be < inputs rank (r  c                ó&   •— g | ]\  }}|‰k    ¯|‘ŒS r  r  )rk  r–  re  r*   s      €r   rl  z,TritonSemantic.reduction.<locals>.<listcomp>²  s"   ø€ ÐAÐAÐA™4˜1˜a°q¸D²y°y�Q°y°y°yr   c              3  ó8   •K  — | ]}|j         j        ‰k    V — Œd S r)   )rg   rI  )rk  r0  rI  s     €r   r‰  z+TritonSemantic.reduction.<locals>.<genexpr>³  s,   øè è € Ð9Ð9¨Q�1”6”< 5Ò(Ð9Ð9Ð9Ð9Ð9Ð9r   z-all reduction inputs must have the same shapec                ó   — g | ]	}|j         ‘Œ
S r  ©r®   ©rk  r0  s     r   rl  z,TritonSemantic.reduction.<locals>.<listcomp>µ  s   € Ð/IÐ/IÐ/I¸Q°´Ð/IÐ/IÐ/Ir   c              3  óŒ   •K  — | ]>}‰                      ‰                     |¦  «        ‰|         j        j        ‰¦  «        V — Œ?d S r)   ©r*  Ú
get_resultrg   rŠ   )rk  r–  r+  Ú	reduce_opr¢  r   s     €€€€r   r‰  z+TritonSemantic.reduction.<locals>.<genexpr>¹  se   øè è € ð uð uØ\]ˆD×Ò˜Y×1Ò1°!Ñ4Ô4°f¸Q´i´nÔ6KÈYÑWÔWðuð uð uð uð uð ur   )
Útuplerg   rI  r\  r“  Úallr'   Úcreate_reduceÚverifyrH  )r   r+  r*   Úregion_builder_fnÚrankr9  r¢  rI  s   ```  @@@r   Ú	reductionzTritonSemantic.reductionª  su  øøøøøø€ Øˆ<ÝÐ^Ð^Ð^Ð^ÐW]Ð^Ñ^Ô^Ñ^Ô^ˆFØˆDà�q”	”Ô$ˆÝ�5‰zŒzˆØ�dŠ{ˆ{ˆ{ÐLÀTÐLÐLÐL‰{Œ{ˆ{ØAÐAÐAÐA¥9¨UÑ#3Ô#3ÐAÑAÔAˆ	ÝÐ9Ð9Ð9Ð9°&Ð9Ñ9Ô9Ñ9Ô9ÐjÐjÐ;jÑjÔjÐ9à”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  r4  r5  s     r   rl  z3TritonSemantic.associative_scan.<locals>.<listcomp>Í  s   € Ð+EÐ+EÐ+E¸¨A¬HÐ+EÐ+EÐ+Er   c              3  óŒ   •K  — | ]>}‰                      ‰                     |¦  «        ‰|         j        j        ‰¦  «        V — Œ?d S r)   r7  )rk  r–  r+  Úscan_opr   rI  s     €€€€r   r‰  z2TritonSemantic.associative_scan.<locals>.<genexpr>Ñ  sN   øè è € ÐwÐwÐ_`�T×%Ò% g×&8Ò&8¸Ñ&;Ô&;¸VÀA¼Y¼^Ô=RÐTYÑZÔZÐwÐwÐwÐwÐwÐwr   )rg   rI  r\  r'   Úcreate_scanr=  r:  rH  )	r   r+  r*   r>  rA  r?  r0  rE  rI  s	   ``     @@r   Úassociative_scanzTritonSemantic.associative_scanÀ  s5  øøøø€ à�q”	”Ô$ˆÝ�5‰zŒzˆàˆu˜Ð#Ð#Ò#Ð#˜tÒ#Ð#Ð#Ð#Ð#Ð%W°$Ð%WÐ%WÐPTÐ%WÐ%WÐ%WÑ#Ô#Ð#à�!Š8ˆ8Ø�D‰LˆDàð 	Uð 	UˆAØ”6”< 5Ò(Ð(Ð(Ð*TÑ(Ô(Ð(Ð(à”,×*Ò*Ð+EÐ+E¸fÐ+EÑ+EÔ+EÀtÈWÑUÔUˆØÐ˜'Ñ"Ô"Ð"Ø�~Š~ÑÔÐÐÐåÐwÐwÐwÐwÐwÐwÐwÕdiÕjmÐntÑjuÔjuÑdvÔdvÐwÑwÔwÑwÔwÐwr   ÚsrcÚindexc                ó\  — |j                              ¦   «         s
J d¦   «         ‚t          |j        j        ¦  «        }t          |j        j        ¦  «        |k    s
J d¦   «         ‚| |cxk    r|k     sn J d|› d|› d�¦   «         ‚|dk     r||z  }t          |¦  «        D ]=}||k    rŒ	|j        j        |         |j        j        |         k    sJ d|› d�¦   «         ‚Œ>| j                             |j        |j        |¦  «        }|  	                    ||j        j
        |j        j        ¦  «        S )	Nzindex must be an integer tensorz0source and index tensors must have the same rankzgather axis z must be < source rank (r  r   z
index dim z( must match the corresponding source dim)r>   rY   r\  rg   rI  rH  r'   Úcreate_gatherr®   r*  rŠ   )r   rH  rI  r*   r?  rˆ  Úgathers          r   rL  zTritonSemantic.gather×  sO  € ØŒ{×!Ò!Ñ#Ô#ÐFÐFÐ%FÑFÔFÐ#å�3”8”>Ñ"Ô"ˆÝ�5”:Ô#Ñ$Ô$¨Ò,Ð,Ð,Ð.`Ñ,Ô,Ð,àˆu˜Ð#Ð#Ò#Ð#˜tÒ#Ð#Ð#Ð#Ð#Ð%Y°DÐ%YÐ%YÐRVÐ%YÐ%YÐ%YÑ#Ô#Ð#Ø�!Š8ˆ8Ø�D‰LˆDå�t‘”ð 	yð 	yˆAØ�DŠyˆyØØ”:Ô# AÔ&¨#¬(¬.¸Ô*;Ò;Ð;Ð;Ð=xÈ$Ð=xÐ=xÐ=xÑ;Ô;Ð;Ð;à”×+Ò+¨C¬J¸¼ÀdÑKÔKˆØ×Ò ¨¬¬¸¼Ô9IÑJÔJÐJr   c                ó(  — |sdS |^}}t          t          |¦  «        ¦  «        D ]$}|                      |||         ¦  «        \  }||<   Œ%t          t          |¦  «        dz
  ¦  «        D ]$}|                      |||         ¦  «        \  }||<   Œ%|g|¢R S )Nr  r   )rH  r\  r�   )r   r+  ÚheadÚtailr–  s        r   r~  z TritonSemantic.broadcast_tensorsí  s®   € Øð 	Ø�2ØˆˆˆtÝ•s˜4‘y”yÑ!Ô!ð 	Eð 	EˆAØ ×5Ò5°d¸DÀ¼GÑDÔD‰MˆD�$�q‘'�'Ý•s˜4‘y”y 1‘}Ñ%Ô%ð 	Eð 	EˆAØ ×5Ò5°d¸DÀ¼GÑDÔD‰MˆD�$�q‘'�'Øˆ}�tˆ}ˆ}Ðr   úSequence[tl.tensor]Úresult_typesúSequence[tl.dtype]ÚpackúTuple[tl.tensor, ...]c                ój  ‡ ‡‡—  ‰ j         ‰Ž Št          ‰¦  «        dk    s
J d¦   «         ‚ˆfd„|D ¦   «         }‰ j                             d„ ‰D ¦   «         ˆ fd„|D ¦   «         |¦  «        Š |‰¦  «         ‰                     ¦   «         sJ ‚t          ˆˆ fd„t          |¦  «        D ¦   «         ¦  «        S )Nr   z1map_elementwise must have at least 1 input tensorc                óZ   •— g | ]'}‰d          j                              |j        ¦  «        ‘Œ(S )r   )rg   r¯   rŠ   )rk  rÍ  r+  s     €r   rl  z2TritonSemantic.map_elementwise.<locals>.<listcomp>ü  s/   ø€ ÐYÐYÐYÀb˜˜qœ	œ×6Ò6°r´yÑAÔAÐYÐYÐYr   c                ó   — g | ]	}|j         ‘Œ
S r  r4  r5  s     r   rl  z2TritonSemantic.map_elementwise.<locals>.<listcomp>þ  s   € Ð&Ð&Ð&˜!ˆQŒXÐ&Ð&Ð&r   c                óD   •— g | ]}|                      ‰j        ¦  «        ‘ŒS r  )r°   r'   )rk  rÍ  r   s     €r   rl  z2TritonSemantic.map_elementwise.<locals>.<listcomp>ÿ  s'   ø€ Ð;Ð;Ð;¨ˆR�XŠX�d”lÑ#Ô#Ð;Ð;Ð;r   c              3  óp   •K  — | ]0\  }}‰                      ‰                     |¦  «        |¦  «        V — Œ1d S r)   )r&   r8  )rk  r–  rÍ  Úelementwise_opr   s      €€r   r‰  z1TritonSemantic.map_elementwise.<locals>.<genexpr>  sC   øè è € ÐhÐhÁuÀqÈ"�T—[’[ ×!:Ò!:¸1Ñ!=Ô!=¸rÑBÔBÐhÐhÐhÐhÐhÐhr   )r~  r\  r'   Úcreate_map_elementwiser=  r:  r“  )r   r+  rQ  rS  r>  rZ  s   ``   @r   Úmap_elementwisezTritonSemantic.map_elementwise÷  sä   øøø€ à'�Ô'¨Ð0ˆå�6‰{Œ{˜QŠˆˆÐ S‰ŒˆØYÐYÐYÐYÈLÐYÑYÔYˆØœ×<Ò<Ø&Ð&˜vÐ&Ñ&Ô&Ø;Ð;Ð;Ð;¨lÐ;Ñ;Ô;Øñ
ô 
ˆð
 	Ð˜.Ñ)Ô)Ð)Ø×$Ò$Ñ&Ô&Ð&Ð&Ð&åÐhÐhÐhÐhÐhÕPYÐZfÑPgÔPgÐhÑhÔhÑhÔhÐhr   Únum_binsc                óÔ  — t          |j        ¦  «        dk    s
J d¦   «         ‚|j                             ¦   «         s
J d¦   «         ‚|�O|                      ||j        ¦  «        }|j        j                             ¦   «         st          d¦  «        ‚|j	        }|  
                    | j                             |j	        ||¦  «        t          j        t          j        |g¦  «        ¦  «        S )Nr   z histogram only supports 1D inputz%histogram only supports integer inputr}  )r\  rI  r>   rY   r˜  rg   rŠ   rÀ  r/   r®   r&   r'   Úcreate_histogramr1   rD  r2   )r   r«   r]  r  s       r   Ú	histogramzTritonSemantic.histogram  sÛ   € Ý�5”;ÑÔ 1Ò$Ð$Ð$Ð&HÑ$Ô$Ð$ØŒ{×!Ò!Ñ#Ô#ÐLÐLÐ%LÑLÔLÐ#ØÐØ×,Ò,¨T°5´;Ñ?Ô?ˆDØ”9Ô#×+Ò+Ñ-Ô-ð GÝ Ð!EÑFÔFÐFØ”;ˆDØ�{Š{˜4œ<×8Ò8¸¼ÀxÐQUÑVÔVÝœ=­¬°H°:Ñ>Ô>ñ@ô @ð 	@r   rš  c                ó  — t          dt          |j        ¦  «        ¦  «        t          |¦  «        k    rt          d¦  «        ‚|j                             dt          j        ||j                             ¦   «         ¦  «        ¦  «         |S )Nr   zAShape of input to multiple_of does not match the length of valuesztt.divisibility)	rï   r\  rI  r/   r®   Úset_attrr   Ú	make_attrÚget_context©r   rh   rš  s      r   Úmultiple_ofzTritonSemantic.multiple_of  sl   € Ýˆq•#�a”g‘,”,ÑÔ¥3 v¡;¤;Ò.Ð.ÝÐ`ÑaÔaÐaØ	Œ×ÒÐ+­R¬\¸&À!Ä(×BVÒBVÑBXÔBXÑ-YÔ-YÑZÔZÐZØˆr   c                óø   — t          |j        ¦  «        t          |¦  «        k    rt          d¦  «        ‚|j                             dt          j        ||j                             ¦   «         ¦  «        ¦  «         |S )NzDShape of input to max_contiguous does not match the length of valuesztt.contiguity©r\  rI  r/   r®   rb  r   rc  rd  re  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.constancyrh  re  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   rA  )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‘  rx  r   c                ó   — g | ]	}|j         ‘Œ
S r  r4  ©rk  Úargs     r   rl  z/TritonSemantic.device_print.<locals>.<listcomp>6  s   € Ð/Ð/Ð/ 3�C”JÐ/Ð/Ð/r   c                ó@   — g | ]}|j                              ¦   «         ‘ŒS r  )r>   rË   rv  s     r   rl  z/TritonSemantic.device_print.<locals>.<listcomp>7  s&   € Ð?Ð?Ð?°3�S”Y×,Ò,Ñ.Ô.Ð?Ð?Ð?r   )Úendswithr\  Ú
startswithr&   r'   Úcreate_printr1   rA  )r   ro  rp  rr  Únew_argsÚ	is_signeds         r   Údevice_printzTritonSemantic.device_print,  sÓ   € ð �Š˜sÑ#Ô#ð 	¨ð 	Ø�c‰MˆFØ�Š˜tÑ$Ô$ð 	(¨ð 	(Ø˜C˜R˜C”[ 4Ñ'ˆFÝˆv‰;Œ;˜Š?ˆ? 6×#4Ò#4°SÑ#9Ô#9ˆ?Ø˜6‘\ˆFà/Ð/¨$Ð/Ñ/Ô/ˆØ?Ð?¸$Ð?Ñ?Ô?ˆ	Ø�{Š{˜4œ<×4Ò4°V¸SÀ(ÈIÑVÔVÕXZÔX_Ñ`Ô`Ð`r   r¨   r©   c                óú   — | j         j        j        sd S |�)|                      ||                      |¦  «        ¦  «        }|                      | j                              |j        |¦  «        t          j	        ¦  «        S r)   )
r'   rŸ   Údebugrý   r
  r&   Úcreate_assertr®   r1   rA  )r   r¨   r©   r  s       r   r¤   zTritonSemantic.device_assert:  sd   € ØŒ|Ô#Ô)ð 	ØˆFØÐØ—8’8˜D $§)¢)¨D¡/¤/Ñ2Ô2ˆDØ�{Š{˜4œ<×5Ò5°d´kÀ3ÑGÔGÍÌÑQÔQÐQr   c                ó|   — |                       | j                             |j        ¦  «        t          j        ¦  «        S r)   )r&   r'   Úcreate_assumer®   r1   rA  )r   r¨   s     r   ÚassumezTritonSemantic.assumeA  s*   € Ø�{Š{˜4œ<×5Ò5°d´kÑBÔBÅBÄGÑLÔLÐLr   c                ó  — t          |t          ¦  «        rt          j        |¦  «        }t          |t          j        ¦  «        rÇt          |j        t
          ¦  «        r| j                             |j        ¦  «        S |rFd|j        cxk    rdk     sn J d|j        › d�¦   «         ‚| j                             |j        ¦  «        S d|j        cxk    rdk     sn J d|j        › d�¦   «         ‚| j         	                    |j        ¦  «        S t          |t          j
        ¦  «        rË|j        j        dk    s
J d	¦   «         ‚|j                             ¦   «         s
J d
¦   «         ‚|j        t          j        k    rQ|rO| j                             |j        | j                             ¦   «         |j                             ¦   «         ¦  «        S |j        t          j        k    r|s
J d¦   «         ‚|j        S J dt'          |¦  «        › �¦   «         ‚)Nrn   ro   z@Block pointers only support 64 bit `shape/strides`, got a value z which is out of the rangerl   rm   zFBlock pointers only support 32 bit `offsets/block_shape`, got a value r   z*Expected a scalar in shape/strides/offsetsz8Expected an integer scalar type in shape/strides/offsetsFzzBlock pointers only support 32 bit `offsets/block_shape`, add a `.to(tl.int32)` or use regular indexing for 64 bit supportz3Unsupported element type in shape/strides/offsets: )rb   r+   r1   rc   rO   rH   r'   r’  rÇ  ræ  r&   rT  r>   rY   rv   r±   r®   Úget_int64_tyrË   rg   )r   r  r5  s      r   Ú_convert_elem_to_ir_valuez(TritonSemantic._convert_elem_to_ir_valueD  s.  € Ý�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Ð&ØŒ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�uà”;ÐØXÐXÍDÐQUÉJÌJÐXÐXÑXÔXˆur   c                ór   ‡ ‡— t          |d¦  «        rˆˆ fd„|D ¦   «         S ‰                      |‰¦  «        gS )Nr  c                ó<   •— g | ]}‰                      |‰¦  «        ‘ŒS r  )r‡  )rk  r  r5  r   s     €€r   rl  z8TritonSemantic._convert_to_ir_values.<locals>.<listcomp>`  s)   ø€ Ð\Ð\Ð\È$�D×2Ò2°4¸ÑEÔEÐ\Ð\Ð\r   )r  r‡  )r   Ú	list_liker5  s   ` `r   r7  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 )NFr4  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   rl  z1TritonSemantic.make_block_ptr.<locals>.<listcomp>u  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 )rl   rm   N)rb   r+   r  s     r   r‰  z0TritonSemantic.make_block_ptr.<locals>.<genexpr>v  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   rl  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)   )r\  )rk  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)r7  rg   r€   r  r[  r/   r1   rt   rŽ   r&  r%  r'  r  r;  r‹  rŒ  rH  r\  r'   Úcreate_make_block_ptrr®   r&   rD  )r   r‹  rI  Ústridesr9  r  Úorderr®   s        `  r   Úmake_block_ptrzTritonSemantic.make_block_ptrc  s,  ø€ ð ×*Ò*¨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Ð\õ �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Ðhð ”×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 r?  )r7  r&   r'   Úcreate_advancer®   rg   )r   r‹  r9  s      r   ÚadvancezTritonSemantic.advance‰  sF   € à×,Ò,¨WÀ%Ð,ÑHÔHˆð �{Š{˜4œ<×6Ò6°t´{ÀGÑLÔLÈdÌiÑXÔXÐXr   rë  r’  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k    rt          d|	› �¦  «        ‚ˆ fd„|D ¦   «         }ˆ fd„|D ¦   «         }t	          j	        |¦  «        }t          |j
        t          j        ¦  «        sJ ‚t	          j        |j
        j        |¦  «        }
|j        }|j
        j                             ¦   «         }‰                      |¦  «        }|j
        j                             ¦   «         r$|t           j        j        k    rt          d¦  «        ‚‰ j                             |d„ |D ¦   «         d„ |D ¦   «         |||¦  «        }t	          j        ||||
¦  «        S )Nr   é   z Expected 1 <= ndim <= 5 but got z dimensionsz	Expected z strides but got zExpected block_shape to have z dimensions but got rl  rx  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  )rU  r1   r2   ©rk  rh   r   s     €r   rl  z9TritonSemantic.make_tensor_descriptor.<locals>.<listcomp>¥  s+   ø€ Ð>Ð>Ð>°1�×!Ò! !¥R¤XÑ.Ô.Ð>Ð>Ð>r   c                ót   •— g | ]4}‰                      t          j        |¦  «        t          j        ¦  «        ‘Œ5S r  )rU  r1   rj  rv   r�  s     €r   rl  z9TritonSemantic.make_tensor_descriptor.<locals>.<listcomp>¦  s6   ø€ Ð[Ð[Ð[Èa�4×#Ò#¥BÔ$;¸AÑ$>Ô$>ÅÄÑIÔIÐ[Ð[Ð[r   z8Padding option `nan` is not supported for integer blocksc                ó   — g | ]	}|j         ‘Œ
S r  r4  ©rk  re  s     r   rl  z9TritonSemantic.make_tensor_descriptor.<locals>.<listcomp>µ  s   € ÐIbÐIbÐIbÐWXÈ!Ì(ÐIbÐIbÐIbr   c                ó   — g | ]	}|j         ‘Œ
S r  r4  r   s     r   rl  z9TritonSemantic.make_tensor_descriptor.<locals>.<listcomp>¶  s   € Ð<WÐ<WÐ<WÈ!¸Q¼XÐ<WÐ<WÐ<Wr   )r\  r/   rb   r>   r1   r&  r  r°  rj  Ú_unwrap_shaperg   rD  r®   rË   rò  rY   r   rí  rï  r'   Úcreate_make_tensor_descriptorÚtensor_descriptor)r   r‹  rI  r’  r  rð  r:  Ú	elem_sizeÚcontig_dim_sizeÚlast_striderg   Úbase_handleÚis_signed_intrñ  r®   s   `              r   Úmake_tensor_descriptorz%TritonSemantic.make_tensor_descriptor�  sØ  ø€ å�5‰zŒzˆØ�T��’�˜Q’���ÝÐQÀÐQÐQÐQÑRÔRÐRÝˆw‰<Œ<˜4ÒÐÝÐN¨ÐNÐNÅÀGÁÄÐNÐNÑOÔOÐOÝˆ{ÑÔ˜tÒ#Ð#ÝÐe¸TÐeÐeÕWZÐ[bÑWcÔWcÐeÐeÑfÔfÐfÝ˜$œ*¥b¤oÑ6Ô6Ð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Ñ:Ô:ˆØ˜!ÒÐÝÐZÈ[ÐZÐZÑ[Ô[Ð[à>Ð>Ð>Ð>¸Ð>Ñ>Ô>ˆØ[Ð[Ð[Ð[ÐSZÐ[Ñ[Ô[ˆõ Ô& {Ñ3Ô3ˆå˜$œ)¥R¤_Ñ5Ô5Ð5Ð5Ð5ÝŒ}˜TœYÔ1°;Ñ?Ô?ˆØ”kˆØœ	Ô,×:Ò:Ñ<Ô<ˆà×-Ò-¨nÑ=Ô=ˆàŒ9Ô×&Ò&Ñ(Ô(ð 	Y¨W½Ô8IÔ8QÒ-QÐ-QÝÐWÑXÔXÐXà”×;Ò;¸KÐIbÐIbÐ\aÐIbÑIbÔIbØ<WÐ<WÈwÐ<WÑ<WÔ<WÐYdÐfsØ<CñEô Eˆõ Ô# F¨E°7¸DÑAÔAÐAr   )r*   r+   r,   r   )r8   r9   r:   r9   r,   r9   )r8   r9   rG   rH   r:   r9   rI   rH   rJ   rH   r,   r9   )T)r   r9   r   r9   r}   rH   r,   r~   )FFTF)rƒ   r„   r…   r„   r,   r†   )rƒ   r   r…   r   r›   rœ   )r«   r„   r¬   r„   r    rH   r,   r   )r«   r„   r¬   r„   r,   r   )r«   r„   r¬   r„   rÐ   rH   r,   r   )rh   r   rØ   r   rÙ   rÚ   )rh   r   rî   r   rï   r   rÙ   rÚ   )r«   r   r¬   r   r,   r†   )r«   r   r¬   r   r,   r   )r«   r   )r«   r   r,   r   )r  r   r,   r   )r@  r+   rA  r+   rÎ   r   r,   r   )r>   r9   r,   r   )rI  rV  r>   r9   r,   r   )rO   r   rI  rV  r,   r   )rO   r   r,   r   )r«   r   ra  rV  rb  rH   r,   r   )r«   r   r*   r+   r,   r   )rƒ   r   r…   r   rb  rH   r,   r   )ru  r   rv  r   r,   r   )ru  r   r,   r†   )r«   r   r„  r…  r,   r   )r«   r   rI  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/  r0  rØ  rœ  rç  rœ  r,   r   )r/  r0  rO   r   r,   r~   )r/  r0  rO   r   r,   r   )rØ  rœ  rç  rœ  r,   r   )r  r   rz  r   r  r,  rØ  rœ  rç  rœ  r,   r   )r  r   r‡  r   rz  r   rÿ  rœ  r
  rœ  r,   r   )
r  r   rz  r   r  r   rŒ  rœ  r,   r�  )rh   r   r,   r   )r  r   rz  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œ  )rz  r   r÷  rœ  )rƒ   r   r	  r   r  rœ  r…   r   r
  r,  r  rœ  rÓ  r  r  rH   r  rH   r  rH   rÕ  r9   r,   r   )r%  r   rh   r   rØ   r   r,   r   )r+  r,  r*   r+   r,   r-  )r+  r,  r*   r+   rA  rH   r,   r-  )rH  r   rI  r   r*   r+   r,   r   )r+  rP  rQ  rR  rS  r+   r,   rT  )r«   r   r]  r+   r  r,  r,   r   )rh   r   rš  rV  r,   r   )r,   r   )ro  rœ  rp  rq  rr  rH   r,   r   )r¨   r   r©   rœ  r  r,  r,   r   )r‹  r   r,   r   )rë  )r‹  r   rI  rq  r’  rq  r  r˜  rð  rœ  r,   r™  )~r    r!   r"   r1   r&   Ú__annotations__Úlangr   r4   r7   rF   r\   rj   re   r‚   rš   rª   r´   r½   rÃ   rÉ   rÏ   rÒ   r×   ræ   rí   rò   rö   r¡   rý   r   r  r  r
  r  r  r  r  r»   r  r"  r(  r£   r2  r¢   r:  r?  rK  rf   rU  rY  rX  r`  rf  rp  rt  r~  rƒ  rŽ  r˜  r�   r­  r  rŽ   rÚ  rß  ré  rò  r   r  r  r!  r+  r.  r;  r=  rB  rJ  rS  rU  rX  r\  r_  rc  rf  rs  rw  r{  r‚  r†  r‹  r•  rœ  r¯  r²  rµ  r¹  r½  rÀ  rÄ  rÒ  rö  rû  r  r  r$  r¢  r*  r@  rG  rL  r~  r\  r`  rf  ri  rk  rn  r~  r¤   r„  r‡  r7  r”  r—  rª  r  r   r   r%   r%      sO
  € € € € € € ØœY€FÐ&Ð&Ð&Ñ&Ø€DàÐÐÑðð ð ðOð Oð Oð Oð
Qð Qð Qð QðDð Dð Dð Dð05ð 05ð 05ð 05ðd	ð 	ð 	ð 	ðKð Kð KðF	@ð 	@ð 	@ð 	@ð aeØ05ð#ð #ð #ð #ð #ðJ,ð ,ð ,ð ,ð$>ð >ð >ð >ð>8ð 8ð 8ð 8ð"8ð 8ð 8ð 8ð]ð ]ð ]ð ]ð2>ð >ð >ð >ð,ð ,ð ,ð ,ð8ð 8ð 8ð 8ð.9ð 9ð 9ð 9ð"9ð 9ð 9ð 9ð"	að 	að 	að 	aðð ð ð ð\ð \ð \ð \ð[ð [ð [ð [ð\ð \ð \ð \ð'ð 'ð 'ð 'ð&ð &ð &ð &ð"ð "ð "ð "ð
]ð ]ð ]ð ]ð]ð ]ð ]ð ]ð\ð \ð \ð \ðð ð ð ð)ð )ð )ð )ð$ð $ð $ð $ð/ð /ð /ð /ð8ð 8ð 8ð 8ð8ð 8ð 8ð 8ð8ð 8ð 8ð 8ð8ð 8ð 8ð 8ð	8ð 	8ð 	8ð 	8ð	8ð 	8ð 	8ð 	8ð GKð Zð Zð Zð Zð Zð Zð$)ð )ð )ð )ð2ð 2ð 2ð 2ðAð Að Að Añhð hð hð hñSð Sð Sð Sñfð fð fð fñXð Xð Xð XñVð Vð Vð Vñð ð ð ñ0

ð 

ð 

ð 

ñTð Tð Tð TñWð Wð Wð Wñ 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.$ð .$ð .$ñ`Wð Wð Wð Wñ(uð uð uð uñð ð ð ñ8+ð +ð +ð +ñ$)ð $)ð $)ð $)ñL$)ð $)ð $)ð $)ñL%ð %ð %ð %ñyð yð yð yñxð xð xð xñyð yð yð yñð ð ð ð
<ð 
<ð 
<ñ[yð [yð [yð [yñzð ð ð ñ0ð 0ð 0ð 0ð$mð mð mñ3tð 3tð 3tð 3tñreð eð eð eð(&ð &ð &ñuð uð uð uñ,xð xð xð xñ.Kð Kð Kð Kð,ð ð ñið ið ið iñ*	@ð 	@ð 	@ð 	@ñð ð ð ñð ð ð ñð ð ð ñCð Cð Cð Cñað að að añRð Rð Rð RñMð Mð Mð MðYð Yð Yð4Hð Hð Hð Hñ
$fð $fð $fð $fñLYð Yð Yð Yð W]ñ(Bñ (Bð (Bð (Bð (Bð (Bð (Br   r%   )Ú
__future__r   rx   rá  Útypingr   r   r   r   r   r	   r
   rˆ   Útriton.runtimer   Ú_C.libtritonr   rh  r   r1   r   r   Ú	Exceptionr   r%   r  r   r   ú<module>r²     s9  ðØ "Ð "Ð "Ð "Ð "Ð "Ø €€€Ø €€€à JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JÐ JØ €€€à !Ð !Ð !Ð !Ð !Ð !à Ð Ð Ð Ð Ð Ø Ð Ð Ð Ð Ð à€GˆC�L„L€Øˆ7�:ÑÔ€ðFð Fð Fð Fð F 	ñ Fô Fð Fð^Bð ^Bð ^Bð ^Bð ^B�W˜XÔ&ñ ^Bô ^Bð ^Bð ^Bð ^Br   