Ë
    1²j­Œ ã                  óÂ   — d dl mZ d dlZd dlZd dlmZmZmZmZm	Z	m
Z
mZ d dlZd dlmZ ddlmZ ddlmZ  e	d	«      Z e	d
«      Z G d„ de«      Z G d„ de
e   «      Zy)é    )Ú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                   j                  «       z   dz   | j                  j                  «       z   | _        t        t
        | �  | j                  «       y )Nzinvalid operands of type ú and )Útype_aÚtype_bÚ__repr__ÚmessageÚsuperr   Ú__init__)Úselfr   r   Ú	__class__s      €úW/var/www/html/ia-enhanced/venv/lib/python3.12/site-packages/triton/language/semantic.pyr   z"IncompatibleTypeErrorImpl.__init__   sX   ø€ ØˆŒØˆŒØ2°T·[±[×5IÑ5IÓ5KÑKÈgÑUÐX\×XcÑXc×XlÑXlÓXnÑnˆŒÜÔ'¨Ñ7¸¿¹ÕEó    )Ú__name__Ú
__module__Ú__qualname__r   Ú__classcell__)r   s   @r   r   r      s   ø„ ÷Fð Fr   r   c                  óp  — 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‡d„Zd‰d„Zd‰d„ZdŠd„Zd‹d„ZdŒd„ZdŒd„ZdŒd„ZdŒd„ZdŒd„Z d�d„Z!dŒd „Z"dŒd!„Z#dŒd"„Z$dŽd#„Z%dŽd$„Z&dŽd%„Z'd�d&„Z(dŒd'„Z)dŒd(„Z*dŒd)„Z+dŒd*„Z,dŒd+„Z-dŒd,„Z.d-d.œd�d/„Z/d‘d0„Z0d‘d1„Z1d’d2„Z2d“d3„Z3d”d4„Z4d•d5„Z5d–d6„Z6d—d7„Z7d˜d8„Z8d™d9„Z9dšd:„Z:d›d;„Z;dœd<„Z<d�d=„Z=džd>„Z>dŸd d?„Z?d@„ Z@dA„ ZAdB„ ZBdC„ ZCdD„ ZDdE„ ZEdF„ ZFdG„ ZGdH„ ZH	 	 	 	 	 	 	 	 	 	 d¡dI„ZI	 	 	 	 d¢dJ„ZJd£dK„ZKd¤dL„ZLd¤dM„ZMdN„ ZNdO„ ZOd¤dP„ZPd¤dQ„ZQd¤dR„ZRd¤dS„ZSd¤dT„ZTd¥dU„ZUd”dV„ZVdW„ ZWdX„ ZX	 	 	 	 d¦dY„ZYd§dZ„ZZ	 	 	 	 d¨d[„Z[d©d\„Z\dªd]„Z]dªd^„Z^dªd_„Z_dªd`„Z`dªda„Zadªdb„Zbdªdc„Zcdd„ Zd	 	 	 	 	 	 d«de„Zed¬df„Zfd­dg„Zgdh„ Zh	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 d®di„Zid¯dj„Zjdk„ Zkd°dl„Zl	 	 	 	 d±dm„Zmd²dn„Zndo„ Zo	 	 d³dp„Zpd´dq„Zqdµdr„Zrdµds„Zsdµdt„Ztd¶du„Zud·dv„Zvd¸dw„Zwd¶dx„Zxdy„ Zyd�dz„Zzd¹d{„Z{d¹d|„Z|	 dº	 	 	 	 	 d»d}„Z}y-)¼ÚTritonSemanticzType[TensorTy]Útensorz
ir.builderÚbuilderc                ó   — || _         y ©N)r'   )r   r'   s     r   r   zTritonSemantic.__init__    s	   € Øˆ�r   c                ó˜   — |dvrt        d|› �«      ‚| j                  | j                  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   Úaxiss     r   Ú
program_idzTritonSemantic.program_id'   sA   € Ø�yÑ ÜÐJÈ4È&ÐQÓRÐRØ�{‰{˜4Ÿ<™<×=Ñ=¸dÓCÄRÇXÁXÓNÐNr   c                ó˜   — |dvrt        d|› �«      ‚| j                  | j                  j                  |«      t        j
                  «      S )Nr+   z-num_programs axis must be 0, 1, or 2 but got )r,   r&   r'   Úcreate_get_num_programsr.   r/   r0   s     r   Únum_programszTritonSemantic.num_programs,   sA   € Ø�yÑ ÜÐLÈTÈFÐSÓTÐTØ�{‰{˜4Ÿ<™<×?Ñ?ÀÓEÄrÇxÁxÓPÐPr   c                ó`  — |j                   }|j                   }|j                  }|j                  }||k(  r	||kD  r|S |S |t        j                  j                  j
                  k(  r	||k\  r|S |S |t        j                  j                  j
                  k(  r	||k\  r|S |S t        d|› d|› �«      ‚)Nzunexpected signedness r   )Úint_bitwidthÚint_signednessr.   ÚdtypeÚ
SIGNEDNESSÚUNSIGNEDÚ	TypeError)r   Úa_tyÚb_tyÚa_rankÚb_rankÚa_snÚb_sns          r   Úinteger_promote_implz#TritonSemantic.integer_promote_impl5   sµ   € Ø×"Ñ"ˆØ×"Ñ"ˆØ×"Ñ"ˆØ×"Ñ"ˆð �4Š<Ø! Fš?�4Ð4°Ð4Ø”R—X‘X×(Ñ(×1Ñ1Ò1Ø! VÒ+�4Ð5°Ð5Ø”R—X‘X×(Ñ(×1Ñ1Ò1Ø! VÒ+�4Ð5°Ð5ÜÐ0°°°e¸D¸6ÐBÓCÐCr   c                ó´  — ||k7  rx|r||fn||f\  }}|j                  «       j                  |j                  «       j                  k  r6|r2|t        j                  t        j                  fv rt        j
                  S |S |j                  «       s|j                  «       rt        j                  S |j                  «       s|j                  «       rt        j
                  S |j                  «       s|j                  «       r"|rt        j
                  S t        j                  S |j                  «       r2|j                  «       r"|rt        j
                  S t        j                  S |j                  «       s|j                  «       rt        j
                  S |j                  «       r'|j                  «       r||k(  r|S t        j                  S |j                  «       r|j                  «       st        d|› d|› �«      ‚|rL|j                  |j                  k7  r3t        d|j                  «       z   dz   |j                  «       z   dz   «      ‚| j!                  ||«      S )Núunexpected type r   zCannot use /, #, or % with úx because they have different signedness;this is unlikely to result in a useful answer. Cast them to the same signedness.)ÚkindÚvaluer.   Úfloat16Úbfloat16Úfloat32Úis_fp64Úfloat64Úis_fp32Úis_fp16Úis_bf16Úis_fp8Úis_intr<   r8   r   rC   )r   r=   Úa_is_scalarr>   Úb_is_scalarÚ
div_or_modÚ	scalar_tyÚ	tensor_tys           r   Úcomputation_type_implz$TritonSemantic.computation_type_implD   sÌ  € ð
 ˜+Ò%Ù3> D¨$¡<ÀTÈ4ÀLÑ ˆI�yØ�~‰~Ó×%Ñ%¨¯©Ó)9×)?Ñ)?Ò?á 9´·±¼R¿[¹[Ð0IÑ#IÜŸ:™:Ð%Ø Ð ð �<‰<Œ>˜TŸ\™\œ^Ü—:‘:Ðð �<‰<Œ>˜TŸ\™\œ^Ü—:‘:Ðð �<‰<Œ>˜TŸ\™\œ^ÙÜ—z‘zÐ!ä—z‘zÐ!à�<‰<Œ>˜dŸl™lœnÙÜ—z‘zÐ!ä—{‘{Ð"Ø�<‰<Œ>˜TŸ\™\œ^Ü—:‘:Ðà�;‰;Œ=˜TŸ[™[œ]Ø 4š<�4Ð7¬R¯Z©ZÐ7Ø�{‰{Œ} D§K¡K¤MÜÐ.¨t¨f°E¸$¸Ð@ÓAÐAñ ˜$×-Ñ-°×1DÑ1DÒDÜÐ9¸D¿M¹M»OÑKÈgÑUÐX\×XeÑXeÓXgÑgðoñoó pð pð ×(Ñ(¨¨tÓ4Ð4r   c                ó>  — t        || j                  «      r|S t        |t        j                  «      r|j                  n|}t        |t
        t        t        f«      r$| j                  |«      }| j                  ||¬«      S |rt        d|› dt        |«      › d�«      ‚|S )N©r9   úcannot convert ú	 of type ú
 to tensor)Ú
isinstancer&   r.   Ú	constexprrH   ÚintÚfloatÚboolÚto_tensor_typeÚscalar_constantr<   Útype)r   ÚxÚ
check_typer9   s       r   Ú	to_tensorzTritonSemantic.to_tensorv   sŠ   € Ü�a˜Ÿ™Ô%ØˆHÜ! !¤R§\¡\Ô2ˆA�GŠG¸ˆÜ�aœ#œu¤dÐ+Ô,Ø×'Ñ'¨Ó*ˆEØ×'Ñ'¨°Ð'Ó7Ð7ÙÜ˜o¨a¨S°	¼$¸q»'¸À*ÐMÓ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t        j                  S  d|cxk  rdk  rt        j                  S  d|cxk  rdk  rt        j                  S  d|cxk  rdk  rt        j                  S  t        d|› d�«      ‚t        |t        «      rƒd	}d
ddz  z  }t        j                  |«      }|t        d«      k(  s6|dk(  s1||k7  s,||cxk  r|k  r1t        j                   S  t        j"                  S t        j                   S t        j"                  S t%        d|› dt'        |«      › d�«      ‚)Né   €ì        l        ìûÿÿÿ         ì            l            zNonrepresentable integer ú.g      8g   àÿÿÿ?r   é   Úinfg        r[   r\   r]   )r^   r.   r9   Úconstexpr_typerH   rb   Úint1r`   r/   Úuint32Úint64Úuint64r,   ra   ÚbuiltinsÚabsrK   rM   r<   re   )r   rf   Úmin_float32Úmax_float32Úabs_xs        r   rc   zTritonSemantic.to_tensor_type�   sn  € Ü�aœŸ™Ô"ØˆHÜ˜œ2×,Ñ,Ô-Ø—‘ˆAä�aœÔÜ—7‘7ˆNÜ˜œ3ÔØ˜Ô"˜UÒ"Ü—x‘x�ð #à˜!Ô#˜eÒ#Ü—y‘yÐ ð $à˜1Ô$˜uÒ$Ü—x‘x�ð %à˜!Ô#˜eÒ#Ü—y‘yÐ ð $äÐ8¸¸¸1Ð=Ó>Ð>Ü˜œ5Ô!Ø!ˆKØ%¨¨C©Ñ/ˆKÜ—L‘L “OˆEØœ˜e›Ò$Ø˜Š|Ø�AŠvØ˜eÔ2 {Ò2Ü—z‘zÐ!ð 3ô —z‘zÐ!ô —z‘zÐ!ä—z‘zÐ!Ü˜/¨!¨¨I´d¸1³g°Y¸jÐIÓJÐJr   c                ó¼   — |j                  «       rL|st        ||«      ‚|j                  «       r||k7  rt        ||«      ‚|j                  «       rt        ||«      ‚y y r)   )Úis_ptrr   Úis_floating)r   r   r   Úallow_ptr_as       r   Úcheck_ptr_type_implz"TritonSemantic.check_ptr_type_impl¤   s[   € Ø�=‰=Œ?ÙÜ/°¸Ó?Ð?à�}‰}Œ F¨fÒ$4Ü/°¸Ó?Ð?à×!Ñ!Ô#Ü/°¸Ó?Ð?ð $ð r   c                ó&  — t        |t        j                  «      }t        |t        j                  «      }|r|}	| j                  |«      }|r|}
| j                  |«      }|j                  j
                  }|j                  j
                  }| j                  |||«       | j                  |||«       |�rF|j                  «       �s5|j                  «       �s$| j                  |||||«      }|r	dk  r|j                  «       s|r 
dk  r|j                  «       rt        d«      ‚|j                  «       rx|r:|j                  «       	cxk  r|j                  «       k  sn t        d|	› d|› �«      ‚|r:|j                  «       
cxk  r|j                  «       k  sn t        d|
› d|› �«      ‚|r| j                  	|¬«      n| j                  ||«      }|r| j                  
|¬«      n| j                  ||«      }| j!                  ||«      \  }}||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 rZ   )r^   ÚnumbersÚNumberrh   re   Úscalarr   r|   rX   Úis_int_unsignedr,   rR   Úget_int_min_valueÚget_int_max_valuerd   ÚcastÚbroadcast_impl_value)r   ÚlhsÚrhsÚallow_lhs_ptrÚallow_rhs_ptrÚarithmetic_checkrU   Ú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ˆÜ" 3¬¯©Ó7ˆÙØˆJØ—.‘. Ó%ˆCÙØˆJØ—.‘. Ó%ˆCð —X‘X—_‘_ˆ
Ø—X‘X—_‘_ˆ
Ø× Ñ  ¨Z¸ÔGØ× Ñ  ¨Z¸ÔGÚ J×$5Ñ$5Õ$7À
×@QÑ@QÕ@SØ×3Ñ3°JÀÈzÐ[hÐjtÓuˆJÙ *¨q¢.°Z×5OÑ5OÔ5QÙ$¨°aª¸J×<VÑ<VÔ<XÜ ð "Kó Lð Là× Ñ Ô"Ù ¨*×*FÑ*FÓ*HÈJô +IØ*4×*FÑ*FÓ*Hô+Iä$ w¨z¨lÐ:TÐU_ÐT`Ð%aÓbÐbÙ ¨*×*FÑ*FÓ*HÈJô +IØ*4×*FÑ*FÓ*Hô+Iä$ w¨z¨lÐ:TÐU_ÐT`Ð%aÓbÐbÙHU�$×&Ñ& z¸Ð&ÔDÐ[_×[dÑ[dÐehÐjtÓ[uˆCÙHU�$×&Ñ& z¸Ð&ÔDÐ[_×[dÑ[dÐehÐjtÓ[uˆCð ×,Ñ,¨S°#Ó6‰ˆˆSØ�Cˆxˆr   c                ó,  — |j                   j                  j                  dk\  s | j                  j                  j
                  sy |j                   j                  }|j                   j                  }||k(  sJ ‚|j                  «       sJ ‚| j                  |t        j                  «      }| j                  |t        j                  «      } |||d«      }|j                  «       }| j                  |t        j                  «      }|j                  «       }| j                  |t        j                  «      }| j                  | j                  ||«      | j                  ||«      «      }	d|j                  › d|j                   › �}
| j#                  |	|
d «       y )Né@   Fr`   z! overflow detected for operation )re   rƒ   r7   r'   ÚoptionsÚsanitize_overflowrR   r‡   r.   rt   r†   rd   r…   Úand_Ú
less_equalÚgreater_equalr    Údevice_assert)r   r‰   rŠ   Ú	binary_opr’   r“   ÚretÚ	max_valueÚ	min_valueÚcondÚmsgs              r   Ú binary_op_sanitize_overflow_implz/TritonSemantic.binary_op_sanitize_overflow_implÔ   sA  € Ø�8‰8�?‰?×'Ñ'¨2Ò-°T·\±\×5IÑ5I×5[Ò5[ØØ—X‘X—_‘_ˆ
Ø—X‘X—_‘_ˆ
Ø˜ZÒ'Ð'Ð'Ø× Ñ Ô"Ð"Ð"Ø�i‰i˜œRŸX™XÓ&ˆØ�i‰i˜œRŸX™XÓ&ˆÙ˜˜S %Ó(ˆØ×0Ñ0Ó2ˆ	Ø×(Ñ(¨´B·H±HÓ=ˆ	Ø×0Ñ0Ó2ˆ	Ø×(Ñ(¨´B·H±HÓ=ˆ	Ø�y‰y˜Ÿ™¨¨iÓ8¸$×:LÑ:LÈSÐR[Ó:\Ó]ˆØ�J×+Ñ+Ð,Ð,MÈi×N`ÑN`ÐMaÐbˆØ×Ñ˜4  dÕ+r   c                ó0  — | j                  ||dd«      \  }}|j                  j                  }|j                  j                  }|j                  «       r|j                  «       rt	        d«      ‚|j                  «       r@|j                  «       s0||}}|j                  j                  }|j                  j                  }|j                  «       rè|j
                  }|j                  j                  «       r‚|j                  j                  dk  ri|j                  j                  t        j                  «      j                  | j                  «      }| j                  j                  |j
                  |d«      }| j                  | j                  j!                  |j
                  |«      |j                  «      S |j#                  «       rJ| j                  | j                  j%                  |j
                  |j
                  «      |j                  «      S |j'                  «       ri|r| j)                  ||| j*                  «       | j                  | j                  j-                  |j
                  |j
                  «      |j                  «      S t	        d|› �«      ‚)NTzcannot add pointers togetherr—   FrE   )r•   re   rƒ   r|   r<   Úhandler9   r„   r7   Úwith_element_tyr.   rt   Úto_irr'   Úcreate_int_castr&   Úcreate_addptrr}   Úcreate_faddrR   r¤   ÚaddÚ
create_add)r   ÚinputÚotherr™   Úinput_scalar_tyÚother_scalar_tyÚother_handleÚi64_tys           r   r¬   zTritonSemantic.addæ   së  € à×8Ñ8¸ÀÀtÈTÓR‰ˆˆuØŸ*™*×+Ñ+ˆØŸ*™*×+Ñ+ˆØ×!Ñ!Ô#¨×(>Ñ(>Ô(@ÜÐ:Ó;Ð;ð ×!Ñ!Ô#¨O×,BÑ,BÔ,DØ  %�5ˆEØ#Ÿj™j×/Ñ/ˆOØ#Ÿj™j×/Ñ/ˆOØ×!Ñ!Ô#Ø Ÿ<™<ˆLØ�{‰{×*Ñ*Ô,°·±×1IÑ1IÈBÒ1NàŸ™×3Ñ3´B·H±HÓ=×CÑCÀDÇLÁLÓQ�Ø#Ÿ|™|×;Ñ;¸E¿L¹LÈ&ÐRWÓX�Ø—;‘;˜tŸ|™|×9Ñ9¸%¿,¹,ÈÓUÐW\×WaÑWaÓbÐbà×(Ñ(Ô*Ø—;‘;˜tŸ|™|×7Ñ7¸¿¹ÀeÇlÁlÓSÐUZ×U_ÑU_Ó`Ð`à×#Ñ#Ô%Ù Ø×5Ñ5°e¸UÀDÇHÁHÔMØ—;‘;˜tŸ|™|×6Ñ6°u·|±|ÀUÇ\Á\ÓRÐTY×T^ÑT^Ó_Ð_ÜÐ*¨?Ð*;Ð<Ó=Ð=r   c                ó„  — | j                  ||dd«      \  }}|j                  j                  }|j                  «       r#| j	                  || j                  |«      d¬«      S |j                  «       rJ| j                  | j                  j                  |j                  |j                  «      |j                  «      S |j                  «       ri|r| j                  ||| j                  «       | j                  | j                  j                  |j                  |j                  «      |j                  «      S t        d|› �«      ‚)NTF)r™   rE   )r•   re   rƒ   r|   r¬   Úminusr}   r&   r'   Úcreate_fsubr¦   rR   r¤   ÚsubÚ
create_subr<   ©r   r®   r¯   r™   rV   s        r   r·   zTritonSemantic.sub  sû   € à×8Ñ8¸ÀÀtÈUÓS‰ˆˆuØ—J‘J×%Ñ%ˆ	à×ÑÔØ—8‘8˜E 4§:¡:¨eÓ#4È�8ÓNÐNà× Ñ Ô"Ø—;‘;˜tŸ|™|×7Ñ7¸¿¹ÀeÇlÁlÓSÐUZ×U_ÑU_Ó`Ð`à×ÑÔÙ Ø×5Ñ5°e¸UÀDÇHÁHÔMØ—;‘;˜tŸ|™|×6Ñ6°u·|±|ÀUÇ\Á\ÓRÐTY×T^ÑT^Ó_Ð_ÜÐ*¨9¨+Ð6Ó7Ð7r   c                ó  — | j                  ||«      \  }}|j                  j                  }|j                  «       rJ| j	                  | j
                  j                  |j                  |j                  «      |j                  «      S |j                  «       ri|r| j                  ||| j                  «       | j	                  | j
                  j                  |j                  |j                  «      |j                  «      S t        d|› �«      ‚©NrE   )r•   re   rƒ   r}   r&   r'   Úcreate_fmulr¦   rR   r¤   ÚmulÚ
create_mulr<   r¹   s        r   r½   zTritonSemantic.mul  sÍ   € à×8Ñ8¸ÀÓF‰ˆˆuØ—J‘J×%Ñ%ˆ	à× Ñ Ô"Ø—;‘;˜tŸ|™|×7Ñ7¸¿¹ÀeÇlÁlÓSÐUZ×U_ÑU_Ó`Ð`à×ÑÔÙ Ø×5Ñ5°e¸UÀDÇHÁHÔMØ—;‘;˜tŸ|™|×6Ñ6°u·|±|ÀUÇ\Á\ÓRÐTY×T^ÑT^Ó_Ð_ÜÐ*¨9¨+Ð6Ó7Ð7r   c                óŠ  — | j                  ||dddd«      \  }}|j                  j                  }|j                  j                  }|j                  «       r$|j	                  «       r| j                  ||«      }�n|j	                  «       r#|j                  «       r| j                  ||«      }nÎ|j	                  «       rQ|j	                  «       rA| j                  |t        j                  «      }| j                  |t        j                  «      }nm|j                  «       rO|j                  «       r?|j                  |j                  kD  r| j                  ||«      }n!| j                  ||«      }nt        d|› �«      ‚| j                  | j                  j                  |j                  |j                  «      |j                  «      S ©NFTrE   )r•   re   rƒ   r}   rR   r‡   r.   rK   Úfp_mantissa_widthr<   r&   r'   Úcreate_fdivr¦   )r   r®   r¯   r°   r±   s        r   ÚtruedivzTritonSemantic.truediv$  s_  € Ø×8Ñ8¸ÀÀuÈeÐUYÐ[_Ó`‰ˆˆuØŸ*™*×+Ñ+ˆØŸ*™*×+Ñ+ˆà×&Ñ&Ô(¨_×-CÑ-CÔ-EØ—I‘I˜e _Ó5ŠEà×#Ñ#Ô%¨/×*EÑ*EÔ*GØ—I‘I˜e _Ó5‰Eà×#Ñ#Ô%¨/×*@Ñ*@Ô*BØ—I‘I˜e¤R§Z¡ZÓ0ˆEØ—I‘I˜e¤R§Z¡ZÓ0‰Eà×(Ñ(Ô*¨×/JÑ/JÔ/LØ×0Ñ0°?×3TÑ3TÒTØŸ	™	 %¨Ó9‘àŸ	™	 %¨Ó9‘ô Ð.¨Ð.?Ð@ÓAÐAØ�{‰{˜4Ÿ<™<×3Ñ3°E·L±LÀ%Ç,Á,ÓOÐQV×Q[ÑQ[Ó\Ð\r   c                óœ  — | j                  ||dddd«      \  }}|j                  j                  }|j                  j                  }|j                  «       rê|j                  «       rÚ| j	                  ||«      }| j                  ||«      }| j                  ||«      }|j                  «       rJ| j                  | j                  j                  |j                  |j                  «      |j                  «      S | j                  | j                  j                  |j                  |j                  «      |j                  «      S t        d|› �«      ‚rÀ   )r•   re   rƒ   rR   rC   r‡   Úis_int_signedr&   r'   Úcreate_sdivr¦   Úcreate_udivr<   )r   r®   r¯   r°   r±   Úret_tys         r   ÚfloordivzTritonSemantic.floordiv=  s  € Ø×8Ñ8¸ÀÀuÈeÐUYÐ[_Ó`‰ˆˆuØŸ*™*×+Ñ+ˆØŸ*™*×+Ñ+ˆØ×!Ñ!Ô#¨×(>Ñ(>Ô(@Ø×.Ñ.¨ÀÓPˆFØ—I‘I˜e VÓ,ˆEØ—I‘I˜e VÓ,ˆEØ×#Ñ#Ô%Ø—{‘{ 4§<¡<×#;Ñ#;¸E¿L¹LÈ%Ï,É,Ó#WÐY^×YcÑYcÓdÐdà—{‘{ 4§<¡<×#;Ñ#;¸E¿L¹LÈ%Ï,É,Ó#WÐY^×YcÑYcÓdÐdÜÐ*¨?Ð*;Ð<Ó=Ð=r   c                óz  — |j                   j                  }|j                   j                  }|j                  «       r|j                  «       st        d«      ‚| j	                  ||dddd«      \  }}| j
                  j                  |j                  |j                  «      }| j                  ||j                   «      S )Nz4both operands of fdiv must have floating scalar typeFT)	re   rƒ   r}   r<   r•   r'   rÂ   r¦   r&   )r   r®   r¯   Úieee_roundingr°   r±   rŸ   s          r   ÚfdivzTritonSemantic.fdivK  s–   € ØŸ*™*×+Ñ+ˆØŸ*™*×+Ñ+ˆØ×*Ñ*Ô,°O×4OÑ4OÔ4QÜÐRÓSÐSØ×8Ñ8¸ÀÀuÈeÐUZÐ\`Óa‰ˆˆuØ�l‰l×&Ñ& u§|¡|°U·\±\ÓBˆØ�{‰{˜3 §
¡
Ó+Ð+r   c                ó\  — | j                  ||dddd«      \  }}|j                  j                  }|j                  j                  }|j                  «       rJ| j	                  | j
                  j                  |j                  |j                  «      |j                  «      S |j                  «       rð|j                  |j                  k7  r3t        d|j                  «       z   dz   |j                  «       z   dz   «      ‚|j                  «       rJ| j	                  | j
                  j                  |j                  |j                  «      |j                  «      S | j	                  | j
                  j                  |j                  |j                  «      |j                  «      S t        d|› �«      ‚)NFTzCannot mod z by rF   rE   )r•   re   rƒ   r}   r&   r'   Úcreate_fremr¦   rR   r8   r<   r   rÅ   Úcreate_sremÚcreate_urem)r   r®   r¯   rV   r±   s        r   ÚmodzTritonSemantic.modT  s`  € Ø×8Ñ8¸ÀÀuÈeÐUYÐ[_Ó`‰ˆˆuØ—J‘J×%Ñ%ˆ	ØŸ*™*×+Ñ+ˆà× Ñ Ô"Ø—;‘;˜tŸ|™|×7Ñ7¸¿¹ÀeÇlÁlÓSÐUZ×U_ÑU_Ó`Ð`à×ÑÔØ×'Ñ'¨?×+IÑ+IÒIÜ °	×0BÑ0BÓ0DÑ DÀvÑ MÐP_×PhÑPhÓPjÑ jð nsñ !só tð tð ×&Ñ&Ô(Ø—{‘{ 4§<¡<×#;Ñ#;¸E¿L¹LÈ%Ï,É,Ó#WÐY^×YcÑYcÓdÐdà—{‘{ 4§<¡<×#;Ñ#;¸E¿L¹LÈ%Ï,É,Ó#WÐY^×YcÑYcÓdÐdÜÐ*¨9¨+Ð6Ó7Ð7r   c                ó   — | j                  ||«      \  }}|j                  }|j                  «       rÜ|t        j                  j
                  k(  rJ| j                  | j                  j                  |j                  |j                  «      |j                  «      S |t        j                  j                  k(  rJ| j                  | j                  j                  |j                  |j                  «      |j                  «      S t        d|› �«      ‚|j                  «       rJ| j                  | j                  j                  |j                  |j                  «      |j                  «      S |j!                  «       rJ| j                  | j                  j#                  |j                  |j                  «      |j                  «      S t%        d|› �«      ‚©NzUnexpected propagate_nan úUnexpected dtype )r•   r9   r}   r.   ÚPropagateNanÚALLr&   r'   Úcreate_minimumfr¦   re   ÚNONEÚcreate_minnumfr,   rÅ   Úcreate_minsir„   Úcreate_minuir<   ©r   rf   ÚyÚpropagate_nanr9   s        r   ÚminimumzTritonSemantic.minimumk  óM  € Ø×0Ñ0°°AÓ6‰ˆˆ1Ø—‘ˆØ×ÑÔØ¤§¡× 3Ñ 3Ò3Ø—{‘{ 4§<¡<×#?Ñ#?ÀÇÁÈ!Ï(É(Ó#SÐUV×U[ÑU[Ó\Ð\Ø¤"§/¡/×"6Ñ"6Ò6Ø—{‘{ 4§<¡<×#>Ñ#>¸q¿x¹xÈÏÉÓ#RÐTU×TZÑTZÓ[Ð[ä Ð#<¸]¸OÐ!LÓMÐMØ× Ñ Ô"Ø—;‘;˜tŸ|™|×8Ñ8¸¿¹À1Ç8Á8ÓLÈaÏfÉfÓUÐUØ×"Ñ"Ô$Ø—;‘;˜tŸ|™|×8Ñ8¸¿¹À1Ç8Á8ÓLÈaÏfÉfÓUÐUäÐ/°¨wÐ7Ó8Ð8r   c                ó   — | j                  ||«      \  }}|j                  }|j                  «       rÜ|t        j                  j
                  k(  rJ| j                  | j                  j                  |j                  |j                  «      |j                  «      S |t        j                  j                  k(  rJ| j                  | j                  j                  |j                  |j                  «      |j                  «      S t        d|› �«      ‚|j                  «       rJ| j                  | j                  j                  |j                  |j                  «      |j                  «      S |j!                  «       rJ| j                  | j                  j#                  |j                  |j                  «      |j                  «      S t%        d|› �«      ‚rÓ   )r•   r9   r}   r.   rÕ   rÖ   r&   r'   Úcreate_maximumfr¦   re   rØ   Úcreate_maxnumfr,   rÅ   Úcreate_maxsir„   Úcreate_maxuir<   rÜ   s        r   ÚmaximumzTritonSemantic.maximum|  rà   r   c                ó‚  — | j                  ||«      \  }}| j                  ||«      \  }}| j                  ||«      \  }}|j                  }|j                  «       rV| j                  | j                  j                  |j                  |j                  |j                  |«      |j                  «      S t        d|› d�«      ‚)NrÔ   z(. Only floating point clamp is supported)	r•   r9   r}   r&   r'   Úcreate_clampfr¦   re   r<   )r   rf   ÚminÚmaxrÞ   r9   s         r   ÚclampzTritonSemantic.clamp�  s¬   € Ø×4Ñ4°S¸#Ó>‰ˆˆSØ×2Ñ2°1°cÓ:‰ˆˆ3Ø×2Ñ2°1°cÓ:‰ˆˆ3à—‘ˆØ×ÑÔØ—;‘;˜tŸ|™|×9Ñ9¸!¿(¹(ÀCÇJÁJÐPS×PZÑPZÐ\iÓjÐlm×lrÑlrÓsÐsäÐ/°¨wÐ6^Ð_Ó`Ð`r   c                ód  — | j                  ||«      \  }}|j                  j                  }|j                  j                  }|j                  «       r|j                  «       st	        ||«      ‚| j                  ||«      }||k7  r| j                  ||«      }||k7  r| j                  ||«      }||fS r)   )r•   re   rƒ   rR   r   rC   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‰ˆˆuØ—z‘z×(Ñ(ˆØ—z‘z×(Ñ(ˆØ×"Ñ"Ô$¨L×,?Ñ,?Ô,AÜ+¨L¸,ÓGÐGØ×.Ñ.¨|¸\ÓJˆ
Ø˜Ò%Ø—I‘I˜e ZÓ0ˆEØ˜Ò%Ø—I‘I˜e ZÓ0ˆEØ�eˆ|Ðr   c                óÀ   — | j                  ||«      \  }}| j                  | j                  j                  |j                  |j                  «      |j
                  «      S r)   )rï   r&   r'   Ú
create_andr¦   re   ©r   r®   r¯   s      r   rš   zTritonSemantic.and_©  óI   € Ø×9Ñ9¸%ÀÓG‰ˆˆuØ�{‰{˜4Ÿ<™<×2Ñ2°5·<±<ÀÇÁÓNÐPU×PZÑPZÓ[Ð[r   c                óÀ   — | j                  ||«      \  }}| j                  | j                  j                  |j                  |j                  «      |j
                  «      S r)   )rï   r&   r'   Ú	create_orr¦   re   rò   s      r   Úor_zTritonSemantic.or_­  sF   € Ø×9Ñ9¸%ÀÓG‰ˆˆuØ�{‰{˜4Ÿ<™<×1Ñ1°%·,±,ÀÇÁÓMÈuÏzÉzÓZÐZr   c                óÀ   — | j                  ||«      \  }}| j                  | j                  j                  |j                  |j                  «      |j
                  «      S r)   )rï   r&   r'   Ú
create_xorr¦   re   rò   s      r   Úxor_zTritonSemantic.xor_±  ró   r   c                ó  — |j                   j                  «       s | j                  |t        j                  «      }|j                   j                  «       s | j                  |t        j                  «      }| j                  ||«      S r)   )re   Úis_int1Úbitcastr.   rr   rš   rò   s      r   Úlogical_andzTritonSemantic.logical_andµ  s[   € Ø�z‰z×!Ñ!Ô#Ø—L‘L ¬¯©Ó0ˆEØ�z‰z×!Ñ!Ô#Ø—L‘L ¬¯©Ó0ˆEØ�y‰y˜ Ó&Ð&r   c                ó  — |j                   j                  «       s | j                  |t        j                  «      }|j                   j                  «       s | j                  |t        j                  «      }| j                  ||«      S r)   )re   rû   rü   r.   rr   rö   rò   s      r   Ú
logical_orzTritonSemantic.logical_or¼  s[   € Ø�z‰z×!Ñ!Ô#Ø—L‘L ¬¯©Ó0ˆEØ�z‰z×!Ñ!Ô#Ø—L‘L ¬¯©Ó0ˆEØ�x‰x˜˜uÓ%Ð%r   c                ó˜   — |j                   j                  «       s | j                  |t        j                  «      }| j                  |«      S r)   )re   rû   rü   r.   rr   Úinvert©r   r®   s     r   Únot_zTritonSemantic.not_Ã  s5   € Ø�z‰z×!Ñ!Ô#Ø—L‘L ¬¯©Ó0ˆEØ�{‰{˜5Ó!Ð!r   c                óÀ   — | j                  ||«      \  }}| j                  | j                  j                  |j                  |j                  «      |j
                  «      S r)   )rï   r&   r'   Úcreate_lshrr¦   re   rò   s      r   ÚlshrzTritonSemantic.lshrÈ  óI   € Ø×9Ñ9¸%ÀÓG‰ˆˆuØ�{‰{˜4Ÿ<™<×3Ñ3°E·L±LÀ%Ç,Á,ÓOÐQV×Q[ÑQ[Ó\Ð\r   c                óÀ   — | j                  ||«      \  }}| j                  | j                  j                  |j                  |j                  «      |j
                  «      S r)   )rï   r&   r'   Úcreate_ashrr¦   re   rò   s      r   ÚashrzTritonSemantic.ashrÌ  r  r   c                óÀ   — | j                  ||«      \  }}| j                  | j                  j                  |j                  |j                  «      |j
                  «      S r)   )rï   r&   r'   Ú
create_shlr¦   re   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|j	                  «       z   dz   «      ‚| j                  | j                  j                  |j                  | j                  «      «      |«      }| j                  ||d«      S )Nz$wrong type argument to unary minus (ú)T)
re   rƒ   r|   r,   r   r&   r'   Úget_null_valuer¨   r·   )r   r®   rí   Ú_0s       r   rµ   zTritonSemantic.minusÛ  s   € Ø—z‘z×(Ñ(ˆØ×ÑÔ ÜÐCÀl×F[ÑF[ÓF]Ñ]Ð`cÑcÓdÐdØ�[‰[˜Ÿ™×4Ñ4°\×5GÑ5GÈÏÉÓ5UÓVÐXdÓeˆØ�x‰x˜˜E 4Ó(Ð(r   c                óX  — |j                   j                  }|j                  «       s|j                  «       rt	        d|j                  «       z   dz   «      ‚| j                  | j                  j                  |j                  | j                  «      «      |«      }| j                  ||«      S )Nz%wrong type argument to unary invert (r  )re   rƒ   r|   r}   r,   r   r&   r'   Úget_all_ones_valuer¨   rù   )r   r®   rí   Ú_1s       r   r  zTritonSemantic.invertâ  sˆ   € Ø—z‘z×(Ñ(ˆØ×ÑÔ  L×$<Ñ$<Ô$>ÜÐDÀ|×G\ÑG\ÓG^Ñ^ÐadÑdÓeÐeØ�[‰[˜Ÿ™×8Ñ8¸×9KÑ9KÈDÏLÉLÓ9YÓZÐ\hÓiˆØ�y‰y˜ Ó#Ð#r   c                óT   — |j                   j                  t        j                  «      S r)   )re   r§   r.   rr   )r   Úvs     r   Ú
_bool_likezTritonSemantic._bool_likeí  s   € Ø�v‰v×%Ñ%¤b§g¡gÓ.Ð.r   c                ó®  — | j                  ||«      \  }}|j                  j                  }|j                  «       rO| j	                  | j
                  j                  |j                  |j                  «      | j                  |«      «      S |j                  «       r®|j                  «       rO| j	                  | j
                  j                  |j                  |j                  «      | j                  |«      «      S | j	                  | j
                  j                  |j                  |j                  «      | j                  |«      «      S t        d|› �«      ‚r»   )r•   re   rƒ   r}   r&   r'   Úcreate_fcmpOGTr¦   r  rR   rÅ   Úcreate_icmpSGTÚcreate_icmpUGTr<   ©r   r®   r¯   rV   s       r   Úgreater_thanzTritonSemantic.greater_thanð  ó  € Ø×8Ñ8¸ÀÓF‰ˆˆuØ—J‘J×%Ñ%ˆ	à× Ñ Ô"Ø—;‘;˜tŸ|™|×:Ñ:¸5¿<¹<ÈÏÉÓVÐX\×XgÑXgÐhmÓXnÓoÐoà×ÑÔØ×&Ñ&Ô(Ø—{‘{ 4§<¡<×#>Ñ#>¸u¿|¹|ÈUÏ\É\Ó#ZÐ\`×\kÑ\kÐlqÓ\rÓsÐsà—{‘{ 4§<¡<×#>Ñ#>¸u¿|¹|ÈUÏ\É\Ó#ZÐ\`×\kÑ\kÐlqÓ\rÓsÐsÜÐ*¨9¨+Ð6Ó7Ð7r   c                ó®  — | j                  ||«      \  }}|j                  j                  }|j                  «       rO| j	                  | j
                  j                  |j                  |j                  «      | j                  |«      «      S |j                  «       r®|j                  «       rO| j	                  | j
                  j                  |j                  |j                  «      | j                  |«      «      S | j	                  | j
                  j                  |j                  |j                  «      | j                  |«      «      S t        d|› �«      ‚r»   )r•   re   rƒ   r}   r&   r'   Úcreate_fcmpOGEr¦   r  rR   rÅ   Úcreate_icmpSGEÚcreate_icmpUGEr<   r  s       r   rœ   zTritonSemantic.greater_equalþ  r!  r   c                ó®  — | j                  ||«      \  }}|j                  j                  }|j                  «       rO| j	                  | j
                  j                  |j                  |j                  «      | j                  |«      «      S |j                  «       r®|j                  «       rO| j	                  | j
                  j                  |j                  |j                  «      | j                  |«      «      S | j	                  | j
                  j                  |j                  |j                  «      | j                  |«      «      S t        d|› �«      ‚r»   )r•   re   rƒ   r}   r&   r'   Úcreate_fcmpOLTr¦   r  rR   rÅ   Úcreate_icmpSLTÚcreate_icmpULTr<   r  s       r   Ú	less_thanzTritonSemantic.less_than  r!  r   c                ó®  — | j                  ||«      \  }}|j                  j                  }|j                  «       rO| j	                  | j
                  j                  |j                  |j                  «      | j                  |«      «      S |j                  «       r®|j                  «       rO| j	                  | j
                  j                  |j                  |j                  «      | j                  |«      «      S | j	                  | j
                  j                  |j                  |j                  «      | j                  |«      «      S t        d|› �«      ‚r»   )r•   re   rƒ   r}   r&   r'   Úcreate_fcmpOLEr¦   r  rR   rÅ   Úcreate_icmpSLEÚcreate_icmpULEr<   r  s       r   r›   zTritonSemantic.less_equal  r!  r   c                óð  — | j                  ||«      \  }}|j                  j                  }|j                  «       rO| j	                  | j
                  j                  |j                  |j                  «      | j                  |«      «      S |j                  «       rO| j	                  | j
                  j                  |j                  |j                  «      | j                  |«      «      S t        d|› �«      ‚r»   )r•   re   rƒ   r}   r&   r'   Úcreate_fcmpOEQr¦   r  rR   Úcreate_icmpEQr<   r  s       r   ÚequalzTritonSemantic.equal(  óÀ   € Ø×8Ñ8¸ÀÓF‰ˆˆuØ—J‘J×%Ñ%ˆ	à× Ñ Ô"Ø—;‘;˜tŸ|™|×:Ñ:¸5¿<¹<ÈÏÉÓVÐX\×XgÑXgÐhmÓXnÓoÐoà×ÑÔØ—;‘;˜tŸ|™|×9Ñ9¸%¿,¹,ÈÏÉÓUÐW[×WfÑWfÐglÓWmÓnÐnÜÐ*¨9¨+Ð6Ó7Ð7r   c                óð  — | j                  ||«      \  }}|j                  j                  }|j                  «       rO| j	                  | j
                  j                  |j                  |j                  «      | j                  |«      «      S |j                  «       rO| j	                  | j
                  j                  |j                  |j                  «      | j                  |«      «      S t        d|› �«      ‚r»   )r•   re   rƒ   r}   r&   r'   Úcreate_fcmpUNEr¦   r  rR   Úcreate_icmpNEr<   r  s       r   Ú	not_equalzTritonSemantic.not_equal3  r3  r   N)rÈ   c               óæ  — t        |t        «      rt        |t        «      st        d«      ‚t        |dz	  «      }t        |dz	  «      }|s|rt        d«      ‚||k  rt        d«      ‚||z
  }||dz
  z  dk7  rt        d«      ‚|g}|€$t	        j
                  t        j                  |«      }|j                  | j                  «      }| 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)r^   r`   r,   rb   r.   Ú
block_typer/   r¨   r'   r&   Úcreate_make_range)	r   ÚstartÚendrÈ   Úis_start_int64Úis_end_int64ÚrangeÚshapeÚ	ret_ty_irs	            r   ÚarangezTritonSemantic.arangeB  sÝ   € Ü˜%¤Ô%¬Z¸¼SÔ-AÜÐNÓOÐOÜ˜e r™kÓ*ˆÜ˜C 2™I“ˆÙ™\ÜÐ7Ó8Ð8Ø�%Š<ÜÐ\Ó]Ð]Ø�e‘ˆØ�U˜Q‘YÑ AÒ%ÜÐBÓCÐCØ�ˆØˆ>Ü—]‘]¤2§8¡8¨UÓ3ˆFØ—L‘L §¡Ó.ˆ	Ø�{‰{˜4Ÿ<™<×9Ñ9¸)ÀUÈCÓPÐRXÓYÐYr   c                óÌ  — |€t        d«      ‚|dk(  r5| j                  j                  |j                  | j                  «      «      }nŒ|j	                  «       rQ| j                  j                  |«      }| j                  j                  ||j                  | j                  «      «      }n+t        | j                  d|j                  › �«      } ||«      }| j                  ||«      S )Nz2dtype must be specified when value is not a tensorr   Úget_)
r,   r'   r  r¨   rQ   Úget_fp32Úcreate_fp_truncÚgetattrÚnamer&   )r   rH   r9   Úget_value_fns       r   rd   zTritonSemantic.scalar_constantT  s±   € àˆ=ÜÐQÓRÐRØ�AŠ:Ø—L‘L×/Ñ/°·±¸D¿L¹LÓ0IÓJ‰EØ�\‰\Œ^Ø—L‘L×)Ñ)¨%Ó0ˆEØ—L‘L×0Ñ0°¸¿¹ÀDÇLÁLÓ8QÓR‰Eä" 4§<¡<°4¸¿
¹
°|Ð1DÓEˆLÙ  Ó'ˆEØ�{‰{˜5 %Ó(Ð(r   c                ó¾   — t        |t        j                  «      r2|j                  j                  dk(  sJ d«       ‚| j                  ||«      S | j                  ||«      S )Nr   zonly accepts size-1 tensor)r^   r.   r&   ÚnumelrH   r‡   rd   )r   rH   r9   s      r   Úmake_scalarzTritonSemantic.make_scalarb  sR   € Ü�eœRŸY™YÔ'Ø—;‘;×$Ñ$¨Ò)ÐGÐ+GÓGÐ)Ø—9‘9˜U EÓ*Ð*à×#Ñ# E¨5Ó1Ð1r   c                óF   — | j                  | j                  ||«      |«      S r)   )ÚsplatrM  )r   rA  rH   r9   s       r   ÚfullzTritonSemantic.fulli  s    € Ø�z‰z˜$×*Ñ*¨5°%Ó8¸%Ó@Ð@r   c                óB  — |j                   j                  «       rJ d«       ‚t        |«      dk(  r|S t        j                  |j
                  |«      }| j                  | j                  j                  |j                  | j                  «      |j                  «      |«      S )NzCannot splat a block tensorr   )re   Úis_blockÚlenr.   r:  r9   r&   r'   Úcreate_splatr¨   r¦   )r   rH   rA  rÈ   s       r   rO  zTritonSemantic.splatp  sy   € Ø—:‘:×&Ñ&Ô(ÐGÐ*GÓGÐ(Üˆu‹:˜Š?ØˆLÜ—‘˜uŸ{™{¨EÓ2ˆØ�{‰{˜4Ÿ<™<×4Ñ4°V·\±\À$Ç,Á,Ó5OÐQV×Q]ÑQ]Ó^Ð`fÓgÐgr   c                ó€   — | j                  | j                  j                  |j                  «      |j                  «      S r)   )r&   r'   Úcreate_unsplatr¦   r9   )r   rH   s     r   ÚunsplatzTritonSemantic.unsplatw  s*   € Ø�{‰{˜4Ÿ<™<×6Ñ6°u·|±|ÓDÀeÇkÁkÓRÐRr   c                ó(  — d}|D ]  }||z  }Œ	 |j                   j                  |k7  rt        d«      ‚t        j                  |j                   j
                  |«      }| j                  | j                  j                  |j                  ||«      |«      S )Nr   z:reshape() cannot change total number of elements in tensor)
re   rL  r,   r.   r:  rƒ   r&   r'   Úcreate_reshaper¦   )r   r®   Ú	dst_shapeÚcan_reorderrL  ÚsrÈ   s          r   ÚreshapezTritonSemantic.reshapez  s�   € ØˆØò 	ˆAØ�Q‰J‰Eð	à�:‰:×Ñ˜uÒ$ÜÐYÓZÐZÜ—‘˜uŸz™z×0Ñ0°)Ó<ˆØ�{‰{˜4Ÿ<™<×6Ñ6°u·|±|ÀYÐP[Ó\Ð^dÓeÐer   c                ó¢  — |j                   D �cg c]  }t        j                  |«      ‘Œ }}|j                  |d«       |j                  j                  «       s| j                  ||¬«      S t        j                  |j                  j                  |«      }| j                  | j                  j                  |j                  |«      |«      S c c}w )Nr   )rA  )rA  r.   Ú_unwrap_if_constexprÚinsertre   rR  rO  r:  rƒ   r&   r'   Úcreate_expand_dimsr¦   )r   r®   r1   rf   rZ  rÈ   s         r   Úexpand_dimszTritonSemantic.expand_dimsƒ  sž   € Ø9>¿¹ÖE°A”R×,Ñ,¨QÕ/ÐEˆ	ÐEØ×Ñ˜˜qÔ!à�z‰z×"Ñ"Ô$Ø—:‘:˜e¨9�:Ó5Ð5ä—‘˜uŸz™z×0Ñ0°)Ó<ˆØ�{‰{˜4Ÿ<™<×:Ñ:¸5¿<¹<ÈÓNÐPVÓWÐWùò Fs   �Cc                óX  — |sJ d«       ‚t        |j                  «      dk(  sJ ‚t        j                  |j                  j
                  |j                  d   |j                  d   z   g«      }| j                  | j                  j                  |j                  |j                  «      |«      S )Nz;current implementation of `cat` always may reorder elementsr   r   )
rS  rA  r.   r:  re   rƒ   r&   r'   Ú
create_catr¦   )r   r‰   rŠ   r[  Úret_types        r   ÚcatzTritonSemantic.cat�  s€   € ÙÐYÐYÓYˆ{Ü�3—9‘9‹~ Ò"Ð"Ð"Ü—=‘= §¡§¡°3·9±9¸Q±<À#Ç)Á)ÈAÁ,Ñ3NÐ2OÓPˆØ�{‰{˜4Ÿ<™<×2Ñ2°3·:±:¸s¿z¹zÓJÈHÓUÐUr   c                ó:  — | j                  ||«      \  }}|j                  g k(  }|r$| j                  |d«      }| j                  |d«      }t        |j                  d   t        j
                  «      rt	        j
                  d«      }nd}|j                  |gz   }t	        j                  |j                  j                  |«      }| j                  | j                  j                  |j                  |j                  «      |«      }|r| j                  |dgd¬«      }|S )Nr   éÿÿÿÿr   F©r[  )rˆ   rA  rb  r^   r.   r_   r:  re   rƒ   r&   r'   Úcreate_joinr¦   r]  )r   ÚaÚbÚ
was_rank_1ÚtwoÚ	new_shapere  rŸ   s           r   ÚjoinzTritonSemantic.join“  sä   € Ø×(Ñ(¨¨AÓ.‰ˆˆ1ð —W‘W ‘]ˆ
ÙØ× Ñ   AÓ&ˆAØ× Ñ   AÓ&ˆAä�a—g‘g˜b‘k¤2§<¡<Ô0Ü—,‘,˜q“/‰CàˆCØ—G‘G˜s˜e‘Oˆ	ä—=‘= §¡§¡°	Ó:ˆØ�k‰k˜$Ÿ,™,×2Ñ2°1·8±8¸Q¿X¹XÓFÈÓQˆáØ—,‘,˜s Q C°U�,Ó;ˆCàˆ
r   c                óŽ  — t        |j                  «      dkD  sJ ‚t        j                  |j                  d   «      dk(  sJ ‚|j                  d d }t        j                  |j
                  j                  |«      }| j                  j                  |j                  «      \  }}| j                  ||«      | j                  ||«      fS )Nr   rh  r   )rS  rA  r.   r_  r:  re   rƒ   r'   Úcreate_splitr¦   r&   )r   rk  ro  re  ÚoutLHSÚoutRHSs         r   ÚsplitzTritonSemantic.split«  s¢   € Ü�A—G‘G“˜qÒ Ð!Ð Ü×'Ñ'¨¯©°©Ó4¸Ò9Ð:Ð9à—G‘G˜C˜R�Lˆ	Ü—=‘= §¡§¡°	Ó:ˆØŸ™×2Ñ2°1·8±8Ó<‰ˆ�à�K‰K˜ Ó)Ø�K‰K˜ Ó)ð
ð 	
r   c                ó  — t        |j                  «      t        |«      k7  r-t        dt        |j                  «      › dt        |«      › �«      ‚t        d„ |D «       «      t	        t        t        |«      «      «      k7  rt        d|› �«      ‚t        j                  |j                  j                  |D �cg c]  }|j                  |   ‘Œ c}«      }| j                  | j                  j                  |j                  |«      |«      S c c}w )Nz;permute dims must have the same length as input shape, got r   c              3  óF   K  — | ]  }t        j                  |«      –— Œ y ­wr)   )r.   r_  )Ú.0Úds     r   ú	<genexpr>z)TritonSemantic.permute.<locals>.<genexpr>»  s   è ø€ Ò;°”"×)Ñ)¨!×,Ñ;ùs   ‚!z?permute dims must be a permutation of 0, 1, ..., n-1, but were )rS  rA  r,   ÚsortedÚlistr@  r.   r:  re   rƒ   r&   r'   Úcreate_transr¦   )r   r®   Údimsry  re  s        r   ÚpermutezTritonSemantic.permute·  sÛ   € Üˆu�{‰{Óœs 4›yÒ(ÜØMÌcÐRW×R]ÑR]ÓN^ÐM_Ð_dÔehÐimÓenÐdoÐpórð räÑ;°dÔ;Ó;¼tÄEÌ#ÈdË)ÓDTÓ?UÒUÜÐ^Ð_cÐ^dÐeÓfÐfä—=‘= §¡×!2Ñ!2ÈTÖ4RÈ°U·[±[À³^Ò4RÓSˆØ�{‰{˜4Ÿ<™<×4Ñ4°U·\±\À4ÓHÈ(ÓSÐSùò 5Ss   Â5D
c                ó(  — |j                   j                  «       s| j                  ||«      S |j                   j                  «       }t	        |«      t	        |«      k7  rt        d|› d|› �«      ‚||k(  r|S t        |«      D ]0  \  }}||   |k7  sŒ|dk7  sŒt        d||   › d|› d|› d|› d|› �
«      ‚ t        j                  |j                   j                  |«      }| 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 ú: )re   rR  rO  Úget_block_shapesrS  r,   Ú	enumerater.   r:  rƒ   r&   r'   Úcreate_broadcastr¦   )r   r®   rA  Ú	src_shapeÚiÚitemrÈ   s          r   Úbroadcast_impl_shapez#TritonSemantic.broadcast_impl_shapeÁ  s  € Ø�z‰z×"Ñ"Ô$Ø—:‘:˜e UÓ+Ð+Ø—J‘J×/Ñ/Ó1ˆ	Üˆy‹>œS ›ZÒ'ÜÐ@ÀÀÈ2ÈeÈWÐUÓVÐVØ�IÒØˆLÜ  Ó+ò 	@‰GˆAˆtØ�Q‰x˜4Ó D¨A£IÜ Ð#VÐW\Ð]^ÑW_ÐV`ð aCØCGÀ&ð I%Ø%& C r¨)¨°B°u°gð"?ó @ð @ð	@ô
 —‘˜uŸz™z×0Ñ0°%Ó8ˆØ�{‰{˜4Ÿ<™<×8Ñ8¸¿¹ÀuÓMÈvÓVÐVr   c           	     óx  — |j                   }|j                   }|j                  «       r~|j                  «       sn|j                  |j                  «      }| j	                  | j
                  j                  |j                  | j
                  «      |j                  «      |«      }||fS |j                  «       s~|j                  «       rn|j                  |j                  «      }| j	                  | j
                  j                  |j                  | j
                  «      |j                  «      |«      }||fS |j                  «       �rò|j                  «       �rá|j                  «       }|j                  «       }t        |«      t        |«      k  r¢t        t        |«      t        |«      «      D ]€  }| j	                  | j
                  j                  |j                  d«      t        j                  |j                  dg|j                  z   «      «      }|j                   }|j                  «       }Œ‚ n¸t        |«      t        |«      k  r¡t        t        |«      t        |«      «      D ]€  }| j	                  | j
                  j                  |j                  d«      t        j                  |j                  dg|j                  z   «      «      }|j                   }|j                  «       }Œ‚ t        |«      t        |«      k(  sJ ‚g }t!        |«      D ]q  \  }	}
||	   }|
dk(  r|j#                  |«       Œ"|dk(  s||
k(  r|j#                  |
«       Œ>t%        dt'        |	«      z   dz   t'        |
«      z   dz   t'        |«      z   «      ‚ ||k7  rVt        j                  |j                  |«      }| j	                  | j
                  j)                  |j                  |«      |«      }||k7  rVt        j                  |j                  |«      }| j	                  | j
                  j)                  |j                  |«      |«      }||fS )Nr   r   z?Cannot make_shape_compatible: incompatible dimensions at index r‚  r   )re   rR  r§   rƒ   r&   r'   rT  r¨   r¦   rƒ  rS  r@  ra  r.   r:  Úvaluesr„  Úappendr,   Ústrr…  )r   r‰   rŠ   Úlhs_tyÚrhs_tyÚ	lhs_shapeÚ	rhs_shapeÚ_Ú	ret_shaper‡  ÚleftÚrightrÈ   s                r   rˆ   z#TritonSemantic.broadcast_impl_valueÑ  s‚  € Ø—‘ˆØ—‘ˆð �?‰?Ô V§_¡_Ô%6Ø×+Ñ+¨F¯M©MÓ:ˆFØ—+‘+˜dŸl™l×7Ñ7¸¿¹ÀTÇ\Á\Ó8RÐTW×T^ÑT^Ó_ÐagÓhˆCðV �CˆxˆðS —‘Ô" v§¡Ô'8Ø×+Ñ+¨F¯M©MÓ:ˆFØ—+‘+˜dŸl™l×7Ñ7¸¿¹ÀTÇ\Á\Ó8RÐTW×T^ÑT^Ó_ÐagÓhˆCðN �CˆxˆðK �_‰_Õ 6§?¡?Õ#4Ø×/Ñ/Ó1ˆIØ×/Ñ/Ó1ˆIä�9‹~¤ I£Ò.äœs 9›~¬s°9«~Ó>ò :�AØŸ+™+ d§l¡l×&EÑ&EÀcÇjÁjÐRSÓ&TÜ&(§m¡m°F·M±MÀAÀ3È×IYÑIYÑCYÓ&Zó\�Cà ŸX™X�FØ &× 7Ñ 7Ó 9‘Iñ	:ô
 �Y“¤# i£.Ò0äœs 9›~¬s°9«~Ó>ò :�AØŸ+™+ d§l¡l×&EÑ&EÀcÇjÁjÐRSÓ&TÜ&(§m¡m°F·M±MÀAÀ3È×IYÑIYÑCYÓ&Zó\�Cà ŸX™X�FØ &× 7Ñ 7Ó 9‘Ið	:ô
 �y“>¤S¨£^Ò3Ð3Ð3àˆIÜ$ YÓ/ò e‘��4Ø! !™�Ø˜1’9Ø×$Ñ$ UÕ+Ø˜q’j e¨t¢mØ×$Ñ$ TÕ*ä$ð &1Ü36°q³6ñ&:Ø<@ñ&AÜCFÀtÃ9ñ&MØOVñ&WÜY\Ð]bÓYcñ&dó eð eðeð ˜IÒ%ÜŸ™ v§}¡}°iÓ@�Ø—k‘k $§,¡,×"?Ñ"?ÀÇ
Á
ÈIÓ"VÐX^Ó_�Ø˜IÒ%ÜŸ™ v§}¡}°iÓ@�Ø—k‘k $§,¡,×"?Ñ"?ÀÇ
Á
ÈIÓ"VÐX^Ó_�à�Cˆxˆr   c                ó¢   — |€y |dk(  rt         j                  j                  S |dk(  rt         j                  j                  S t	        d|› d�«      ‚)NÚrtneÚrtzzInvalid rounding mode: z0. Supported rounding modes are 'rtne' and 'rtz'.)r   ÚROUNDING_MODEÚRTNEÚRTZr,   )r   Úrounding_modes     r   Ú_str_to_rounding_modez$TritonSemantic._str_to_rounding_mode	  sU   € ØÐ ØØ˜FÒ"Ü×#Ñ#×(Ñ(Ð(Ø˜EÒ!Ü×#Ñ#×'Ñ'Ð'ÜÐ2°=°/ÐAqÐrÓsÐsr   c                ó6  — |j                   }|j                  «       r|j                  |j                  «      }||k(  r|S |j                  }|j                  }|j	                  «       s|j	                  «       r| j                  ||«      S |j                  }|j                  }||k7  r&t        dt        |«      z   dz   t        |«      z   «      ‚| j                  | j                  j                  |j                  |j                  | j                  «      «      |«      S )Nz!Cannot bitcast data-type of size z to data-type of size )re   rR  r§   rƒ   r|   r‡   Úprimitive_bitwidthr,   r�  r&   r'   Úcreate_bitcastr¦   r¨   )r   r®   Údst_tyÚsrc_tyÚ
src_sca_tyÚ
dst_sca_tyÚsrc_bitsÚdst_bitss           r   rü   zTritonSemantic.bitcast  sø   € Ø—‘ˆØ�?‰?ÔØ×+Ñ+¨F¯M©MÓ:ˆFØ�VÒØˆLØ—]‘]ˆ
Ø—]‘]ˆ
Ø×ÑÔ *×"3Ñ"3Ô"5Ø—9‘9˜U FÓ+Ð+à×0Ñ0ˆØ×0Ñ0ˆØ�xÒÜÐ@Ä3ÀxÃ=ÑPð T2ñ 2Ü47¸³MñBó Cð Cà�{‰{˜4Ÿ<™<×6Ñ6°u·|±|ÀVÇ\Á\ÐRV×R^ÑR^ÓE_Ó`ÐbhÓiÐir   c                ó   — |j                   }|j                  }|j                  }||k(  r|S |j                  «       r|j                  |«      }| j	                  |«      }d}|j                  «       rf|j                  «       rV|j                  |j                  k  r=|€t        j                  j                  }nH|t        j                  j                  k7  r+d}n(|�&t        dt        |«      z   dz   t        |«      z   «      ‚|j                  «       s|j                  «       rP| j                  j                  j                  d«      	 €J d«       ‚ | j                  j                  d   |||| ¬«      S |j!                  «       r|j                  «       s"|j                  «       r|j!                  «       s|rP| j#                  | j                  j%                  |j&                  |j)                  | j                  «      |«      |«      S |j+                  «       r|j-                  «       r |j/                  «       r@|j-                  «       s0| j1                  | j1                  |t2        j4                  «      |«      S |j                  «       xr+ |j                  «       xr |j                  |j                  kD  }|rO| j#                  | j                  j7                  |j&                  |j)                  | j                  «      «      |«      S |j                  «       xr+ |j                  «       xr |j                  |j                  k  }	|	rO| j#                  | j                  j9                  |j&                  |j)                  | j                  «      «      |«      S |j;                  «       �r2|j;                  «       �r!|j<                  |j<                  k7  s|j>                  |j>                  k7  rï|jA                  «       xr |jC                  «        }
|jC                  «       rl|jD                  j)                  | j                  «      }| j#                  | j                  jG                  |«      |jD                  «      }| jI                  ||«      S | j#                  | j                  jK                  |j&                  |j)                  | j                  «      |
«      |«      S |jM                  «       �r;|j;                  «       �r*|jC                  «       rl|jD                  j)                  | j                  «      }| j#                  | j                  jG                  |«      |jD                  «      }| jI                  ||«      S |jA                  «       rO| j#                  | j                  jO                  |j&                  |j)                  | j                  «      «      |«      S | j#                  | j                  jQ                  |j&                  |j)                  | j                  «      «      |«      S |j;                  «       rÎ|jM                  «       r¾|jC                  «       s|jA                  «       sO| j#                  | j                  jS                  |j&                  |j)                  | j                  «      «      |«      S | j#                  | j                  jU                  |j&                  |j)                  | j                  «      «      |«      S |jW                  «       rÜ|j;                  «       rÌ|j<                  }|dk(  rO| j#                  | j                  jY                  |j&                  |j)                  | j                  «      «      |«      S |d	k(  rg| jI                  | j1                  |t2        jZ                  «      | j#                  | j                  j]                  d
«      t2        jZ                  «      «      S |j;                  «       r_|jW                  «       rO| j#                  | j                  j_                  |j&                  |j)                  | j                  «      «      |«      S |jW                  «       r_|jW                  «       rO| j#                  | j                  ja                  |j&                  |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 )1re   rƒ   rR  r§   r�  r}   rŸ  r   r™  rš  r,   r�  Úis_fp8e4b15r'   Úcodegen_fnsÚgetrQ   r&   Úcreate_fp_to_fpr¦   r¨   rO   rN   rP   r‡   r.   rK   rG  Úcreate_fp_extrR   r7   r8   rÅ   Úis_boolr9   r  r7  r©   Úis_standard_floatingÚcreate_fp_to_siÚcreate_fp_to_uiÚcreate_ui_to_fpÚcreate_si_to_fpr|   Úcreate_ptr_to_intrt   Ú	get_int64Úcreate_int_to_ptrr   )r   r®   r¡  Úfp_downcast_roundingr¢  r£  r¤  Úuse_custom_roundingÚtruncate_fpÚext_fpÚsign_extendÚtyr  Úbitwidths                 r   r‡   zTritonSemantic.cast$  s  € Ø—‘ˆØ—]‘]ˆ
Ø—]‘]ˆ
Ø˜Ò#ØˆLØ�?‰?ÔØ×+Ñ+¨JÓ7ˆFð  $×9Ñ9Ð:NÓOÐØ#ÐØ×!Ñ!Ô#¨
×(>Ñ(>ô )
Ø×+Ñ+¨j×.KÑ.KÒKØ#Ð+ÄB×DTÑDT×DYÑDYÑ-AØ%¬×)9Ñ)9×)>Ñ)>Ò>ÐVZÑ@Sà#Ð/Ü ð ":Ü<?À
»Oñ"LØNiñ"jä!$ Z£ñ"1ó 2ð 2ð ×"Ñ"Ô$¨
×(>Ñ(>Ô(@Ø—<‘<×+Ñ+×/Ñ/Ø&ó(Ø/3ð4ð hØ5góhð 4àC�4—<‘<×+Ñ+Ð,BÑCÀEÈ6ÐSgÐswÔxÐxð ×ÑÔ J×$:Ñ$:Ô$<Ø×"Ñ"Ô$¨×):Ñ):Ô)<ÙØ—;‘;Ø—‘×,Ñ,¨U¯\©\¸6¿<¹<ÈÏÉÓ;UÐWkÓlÐntóvð vð ×ÑÔ ¨×);Ñ);Ô)=Ø×ÑÔ ¨×);Ñ);Ô)=Ø—9‘9˜TŸY™Y u¬b¯j©jÓ9¸:ÓFÐFð
 !×,Ñ,Ó.ò JØ×"Ñ"Ó$òJà×)Ñ)¨J×,IÑ,IÑIð 	ñ Ø—;‘;˜tŸ|™|×;Ñ;¸E¿L¹LÈ&Ï,É,ÐW[×WcÑWcÓJdÓeÐgmÓnÐnð ×'Ñ'Ó)ò JØ×"Ñ"Ó$òJà×)Ñ)¨J×,IÑ,IÑIð 	ñ Ø—;‘;˜tŸ|™|×9Ñ9¸%¿,¹,ÈÏÉÐUY×UaÑUaÓHbÓcÐekÓlÐlð ×ÑÕ :×#4Ñ#4Õ#6Ø×#Ñ# z×'>Ñ'>Ò>À*×B[ÑB[Ð_i×_xÑ_xÒBxØ$×2Ñ2Ó4ÒQ¸Z×=OÑ=OÓ=QÐ9QˆKØ×!Ñ!Ô#Ø—[‘[×&Ñ& t§|¡|Ó4�Ø—[‘[ §¡×!<Ñ!<¸RÓ!@À%Ç+Á+ÓN�Ø—~‘~ e¨RÓ0Ð0à—{‘{ 4§<¡<×#?Ñ#?ÀÇÁÈfÏlÉlÐ[_×[gÑ[gÓNhÐjuÓ#vØ#)ó+ð +ð ×*Ñ*Õ,°×1BÑ1BÕ1DØ×!Ñ!Ô#Ø—[‘[×&Ñ& t§|¡|Ó4�Ø—[‘[ §¡×!<Ñ!<¸RÓ!@À%Ç+Á+ÓN�Ø—~‘~ e¨RÓ0Ð0Ø×)Ñ)Ô+Ø—{‘{ 4§<¡<×#?Ñ#?ÀÇÁÈfÏlÉlÐ[_×[gÑ[gÓNhÓ#iÐkqÓrÐrà—{‘{ 4§<¡<×#?Ñ#?ÀÇÁÈfÏlÉlÐ[_×[gÑ[gÓNhÓ#iÐkqÓrÐrð ×ÑÔ :×#BÑ#BÔ#DØ×!Ñ!Ô#¨:×+CÑ+CÔ+EØ—{‘{ 4§<¡<×#?Ñ#?ÀÇÁÈfÏlÉlÐ[_×[gÑ[gÓNhÓ#iÐkqÓrÐrà—{‘{ 4§<¡<×#?Ñ#?ÀÇÁÈfÏlÉlÐ[_×[gÑ[gÓNhÓ#iÐkqÓrÐrð ×ÑÔ :×#4Ñ#4Ô#6Ø!×.Ñ.ˆHØ˜2Š~Ø—{‘{ 4§<¡<×#AÑ#AÀ%Ç,Á,ÐPV×P\ÑP\Ð]a×]iÑ]iÓPjÓ#kÐmsÓtÐtØ˜1Š}Ø—~‘~ d§i¡i°´r·x±xÓ&@À$Ç+Á+ÈdÏlÉl×NdÑNdÐefÓNgÔik×iqÑiqÓBrÓsÐsð ×ÑÔ :×#4Ñ#4Ô#6Ø—;‘;˜tŸ|™|×=Ñ=¸e¿l¹lÈFÏLÉLÐY]×YeÑYeÓLfÓgÐioÓpÐpð ×ÑÔ :×#4Ñ#4Ô#6Ø—;‘;˜tŸ|™|×:Ñ:¸5¿<¹<ÈÏÉÐVZ×VbÑVbÓIcÓdÐflÓmÐmà8˜ U G¨4°¨xÐ8Ó8ˆur   c                ó"  — t         j                  j                  }|rr|dk(  rt         j                  j                  }|S |dk(  rt         j                  j                  }|S |dk(  rt         j                  j
                  }|S t        d|› d�«      ‚|S )Nz.caú.cgz.cvúCache modifier ú not supported)r   ÚCACHE_MODIFIERrØ   ÚCAÚCGÚCVr,   ©r   Úcache_modifierÚcaches      r   Ú_str_to_load_cache_modifierz*TritonSemantic._str_to_load_cache_modifier“  s–   € Ü×!Ñ!×&Ñ&ˆÙØ Ò&Ü×)Ñ)×,Ñ,�ð ˆð   5Ò(Ü×)Ñ)×,Ñ,�ð
 ˆð	   5Ò(Ü×)Ñ)×,Ñ,�ð ˆô ! ?°>Ð2BÀ.Ð!QÓRÐRØˆr   c                ód  — t         j                  j                  }|r“|dk(  rt         j                  j                  }|S |dk(  rt         j                  j                  }|S |dk(  rt         j                  j
                  }|S |dk(  rt         j                  j                  }|S t        d|› d�«      ‚|S )Nz.wbrÀ  z.csz.wtrÁ  rÂ  )r   rÃ  rØ   ÚWBrÅ  ÚCSÚWTr,   rÇ  s      r   Ú_str_to_store_cache_modifierz+TritonSemantic._str_to_store_cache_modifier   s¸   € Ü×!Ñ!×&Ñ&ˆÙØ Ò&Ü×)Ñ)×,Ñ,�ð ˆð   5Ò(Ü×)Ñ)×,Ñ,�ð ˆð   5Ò(Ü×)Ñ)×,Ñ,�ð
 ˆð	   5Ò(Ü×)Ñ)×,Ñ,�ð ˆô ! ?°>Ð2BÀ.Ð!QÓRÐRØˆr   c                óà   — t         j                  j                  }|rQ|dk(  rt         j                  j                  }|S |dk(  rt         j                  j                  }|S t        d|› d�«      ‚|S )NÚ
evict_lastÚevict_firstzEviction policy rÂ  )r   ÚEVICTION_POLICYÚNORMALÚ
EVICT_LASTÚEVICT_FIRSTr,   )r   Úeviction_policyÚevictions      r   Ú_str_to_eviction_policyz&TritonSemantic._str_to_eviction_policy¯  su   € Ü×%Ñ%×,Ñ,ˆÙØ ,Ò.Ü×-Ñ-×8Ñ8�ð
 ˆð	 ! MÒ1Ü×-Ñ-×9Ñ9�ð ˆô !Ð#3°OÐ3DÀNÐ!SÓTÐTØˆr   c                ó°   — d }|rQ|dk(  rt         j                  j                  }|S |dk(  rt         j                  j                  }|S t	        d|› d�«      ‚|S )NÚzeroÚnanzPadding option rÂ  )r   ÚPADDING_OPTIONÚPAD_ZEROÚPAD_NANr,   )r   Úpadding_optionÚpaddings      r   Ú_str_to_padding_optionz%TritonSemantic._str_to_padding_optionº  sh   € ØˆÙØ Ò'Ü×+Ñ+×4Ñ4�ð
 ˆð	   5Ò(Ü×+Ñ+×3Ñ3�ð ˆô ! ?°>Ð2BÀ.Ð!QÓRÐRØˆr   c                ód  — t         j                  j                  }|r“|dk(  rt         j                  j                  }|S |dk(  rt         j                  j                  }|S |dk(  rt         j                  j                  }|S |dk(  rt         j                  j
                  }|S t        d|› d�«      ‚|S )NÚacquireÚreleaseÚacq_relÚrelaxedúMemory semantic rÂ  )r   ÚMEM_SEMANTICÚACQUIRE_RELEASEÚACQUIREÚRELEASEÚRELAXEDr,   )r   Ú
sem_optionÚsems      r   Ú_str_to_semzTritonSemantic._str_to_semÅ  s®   € Ü�o‰o×-Ñ-ˆÙØ˜YÒ&Ü—o‘o×-Ñ-�ð ˆ
ð ˜yÒ(Ü—o‘o×-Ñ-�ð ˆ
ð ˜yÒ(Ü—o‘o×5Ñ5�ð
 ˆ
ð	 ˜yÒ(Ü—o‘o×-Ñ-�ð ˆ
ô !Ð#3°J°<¸~Ð!NÓOÐOØˆ
r   c                ó"  — t         j                  j                  }|rr|dk(  rt         j                  j                  }|S |dk(  rt         j                  j                  }|S |dk(  rt         j                  j                  }|S t        d|› d�«      ‚|S )NÚgpuÚctaÚsysrè  rÂ  )r   ÚMEM_SYNC_SCOPEÚGPUÚCTAÚSYSTEMr,   )r   Úscope_optionÚscopes      r   Ú_str_to_scopezTritonSemantic._str_to_scopeÔ  s–   € Ü×!Ñ!×%Ñ%ˆÙØ˜uÒ$Ü×)Ñ)×-Ñ-�ð ˆð  Ò&Ü×)Ñ)×-Ñ-�ð
 ˆð	  Ò&Ü×)Ñ)×0Ñ0�ð ˆô !Ð#3°L°>ÀÐ!PÓQÐQØˆr   c                ó~  — |r¶t        |d«      s|g}|D �cg c]*  }t        |t        j                  «      r|j                  n|‘Œ, }}|D ]+  }t        |t
        «      rd|cxk  rt        |«      k  rŒ(J ‚ J ‚ t        |«      dkD  sJ ‚t        |«      t        t        |«      «      k(  sJ d«       ‚t        |«      S yc c}w )NÚ__iter__r   z'Duplicate dimension in `boundary_check`r  )	Úhasattrr^   r.   r_   rH   r`   rS  Úsetr{  )r   Úboundary_checkÚblock_shapeÚelemÚdims        r   Ú_canonicalize_boundary_checkz+TritonSemantic._canonicalize_boundary_checká  sÇ   € ÙÜ˜>¨:Ô6Ø"0Ð!1�ØaoÖpÐY]¬J°t¼R¿\¹\Ô,J˜dŸjšjÐPTÑTÐpˆNÐpØ%ò L�Ü! #¤sÔ+°°SÔ0K¼3¸{Ó;KÓ0KÐKÐKÐ0KÐKÐKðLä�~Ó&¨Ò*Ð*Ð*Ü�~Ó&¬#¬c°.Ó.AÓ*BÒBÐmÐDmÓmÐBÜ˜.Ó)Ð)Øùò qs   –/B:c	           
     óä  — |€|�t        d«      ‚|j                  j                  j                  }	|	t        j                  k7  sJ d«       ‚|	j                  «       r(|t        j                  j                  k(  rt        d«      ‚|j                  j                  }
| j                  ||
j                  «       «      }| 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,   re   Ú
element_tyr.   rr   rR   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à—‘×$Ñ$×/Ñ/ˆØœŸ™Ò ÐXÐ"XÓXÐ Ø�=‰=Œ?˜w¬"×*;Ñ*;×*CÑ*CÒCÜÐ_Ó`Ð`ð —‘×$Ñ$ˆð ×:Ñ:¸>È6×KbÑKbÓKdÓeˆð �{‰{Ø�L‰L×3Ñ3°C·J±JÀÐPWÐY^Ð`hÐjuÓvØóð 	r   c	           
     ó   — |j                   j                  j                  «       s't        d|j                   j	                  «       › d�«      ‚|€|�t        d«      ‚|s|rt        d«      ‚|j                   j                  «       sN|r%|j                   j                  «       rt        d«      ‚|r%|j                   j                  «       rt        d«      ‚|j                   j                  «       r.|�| j                  ||«      \  }}|�| j                  ||«      \  }}|j                   j                  }	|	j                  }
|
t        j                  k(  }|rBt        j                  }
t        j                  |
|	j                  «      }	| j                  ||	«      }|�| j                  ||
«      }|j                   j                  «       r|j                   j                  |
«      }n|
}|€9| j                  | j                   j#                  |j$                  |||«      |«      }nR| j                  | j                   j'                  |j$                  |j$                  |r|j$                  nd |||«      |«      }|r | j                  |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)re   rƒ   r|   r,   r   rR  rˆ   r  r.   rr   Ú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  s9  € à�x‰x�‰×%Ñ%Ô'ÜÐ4°S·X±X×5FÑ5FÓ5HÐ4IÈÐWÓXÐXð ˆ<˜EÐ-ÜÐHÓIÐIÙ‘nÜð Xó Yð Yð
 �x‰x× Ñ Ô"Ù˜Ÿ	™	×*Ñ*Ô,Ü Ð!hÓiÐiÙ˜Ÿ™×,Ñ,Ô.Ü Ð!iÓjÐjð �8‰8×ÑÔØÐØ ×5Ñ5°c¸4Ó@‘	��TØÐ Ø!×6Ñ6°s¸EÓB‘
��Uð —‘—‘ˆØ×"Ñ"ˆð œBŸG™GÑ#ˆÙÜ—W‘WˆFÜ—_‘_ V¨V×-AÑ-AÓBˆFØ—)‘)˜C Ó(ˆCð ÐØ—I‘I˜e VÓ,ˆEð �8‰8×ÑÔØ—X‘X×-Ñ-¨fÓ5‰Fð ˆFð ˆ<Ø—+‘+˜dŸl™l×6Ñ6°s·z±zÀ5È(ÐT_Ó`ÐbhÓi‰Cà—+‘+Ø—‘×/Ñ/°·
±
¸D¿K¹KÑY^ÈÏÊÐdhÐjoØ08¸+óGØHNóPˆCñ Ø—)‘)˜C¤§¡Ó)ˆCØˆ
r   c	           
     óD  — | j                  |«      }	| j                  |«      }
| j                  |«      }|j                  j	                  «       r<|j                  j
                  j                  «       r| j                  ||||||	|
|«      S | j                  ||||||	|
|«      S r)   )	rÊ  rÙ  râ  re   r|   r  rR  r  r  )r   r
  r  r¯   r   rà  rÈ  r×  r  rÉ  rØ  rá  s               r   ÚloadzTritonSemantic.load?  s�   € ð ×0Ñ0°Ó@ˆØ×/Ñ/°Ó@ˆØ×-Ñ-¨nÓ=ˆà�8‰8�?‰?Ô §¡×!4Ñ!4×!=Ñ!=Ô!?à×+Ñ+¨C°°u¸nÈgÐW\Ð^fÐhsÓtÐtð ×$Ñ$ S¨$°°~ÀwÐPUÐW_ÐalÓmÐmr   c                ó–  — t        |t        j                  «      sJ ‚t        |j                  «      }t        |«      |k(  sJ d|› dt        |«      › �«       ‚| j                  |d¬«      }| j                  j                  |j                  || j                  |«      | j                  |«      «      }| j                  ||j                  «      S )Nú	expected ú offsets, but got F©Úrequire_i64)r^   r.   Útensor_descriptor_baserS  r  Ú_convert_to_ir_valuesr'   Úcreate_descriptor_loadr¦   rÊ  rÙ  r&   r:  )r   ÚdescÚoffsetsrÈ  r×  Úndimrf   s          r   Údescriptor_loadzTritonSemantic.descriptor_loadM  s´   € ä˜$¤× 9Ñ 9Ô:Ð:Ð:Ü�4×#Ñ#Ó$ˆÜ�7‹|˜tÒ#ÐW y°°Ð6HÌÈWËÈÐ%WÓWÐ#à×,Ñ,¨WÀ%Ð,ÓHˆØ�L‰L×/Ñ/°·±¸WÀd×FfÑFfÐguÓFvØ04×0LÑ0LÈ_Ó0]ó_ˆà�{‰{˜1˜dŸo™oÓ.Ð.r   c                óä   — t        |t        j                  «      sJ ‚t        |j                  «      }t        |«      |k(  sJ d|› dt        |«      › �«       ‚|j
                  |j                  k(  sJ ‚y )Nr  r  )r^   r.   r   rS  r  rA  )r   r#  rH   r$  r%  s        r   Úvalidate_store_likez"TritonSemantic.validate_store_likeX  si   € Ü˜$¤× 9Ñ 9Ô:Ð:Ð:Ü�4×#Ñ#Ó$ˆÜ�7‹|˜tÒ#ÐW y°°Ð6HÌÈWËÈÐ%WÓWÐ#Ø�{‰{˜d×.Ñ.Ò.Ð.Ñ.r   c                ó$  — | j                  |||«       | j                  ||j                  «      }| j                  |d¬«      }| j	                  | j
                  j                  |j                  |j                  |«      t        j                  «      S ©NFr  )
r(  r‡   r9   r!  r&   r'   Úcreate_descriptor_storer¦   r.   Úvoid)r   r#  rH   r$  s       r   Údescriptor_storezTritonSemantic.descriptor_store^  sq   € Ø× Ñ   u¨gÔ6à—	‘	˜% §¡Ó,ˆØ×,Ñ,¨WÀ%Ð,ÓHˆØ�{‰{˜4Ÿ<™<×?Ñ?ÀÇÁÈUÏ\É\Ð[bÓcÔeg×elÑelÓmÐmr   c                ó   — | j                  |||«       |j                  t        j                  t        j                  t        j
                  t        j                  t        j                  t        j                  hv sJ d«       ‚| j                  |d¬«      }t        j                  j                  }| j                  | j                  j                  ||j                   |j                   |«      t        j"                  «      S ©NúUnsupported dtypeFr  )r(  r9   r.   rs   r/   ru   rK   rI   rJ   r!  r   ÚDESCRIPTOR_REDUCE_KINDÚADDr&   r'   Úcreate_descriptor_reducer¦   r,  ©r   r#  rH   r$  rG   s        r   Údescriptor_atomic_addz$TritonSemantic.descriptor_atomic_adde  s³   € Ø× Ñ   u¨gÔ6Ø�z‰zœbŸi™i¬¯©´2·9±9¼b¿j¹jÌ"Ï*É*ÔVX×VaÑVaÐbÑbÐwÐdwÓwÐbØ×,Ñ,¨WÀ%Ð,ÓHˆÜ×(Ñ(×,Ñ,ˆØ�{‰{˜4Ÿ<™<×@Ñ@ÀÀtÇ{Á{ÐTY×T`ÑT`ÐbiÓjÔln×lsÑlsÓtÐtr   c                ó~   — t         j                  j                  «       }|j                  dk(  xr |j                  dk\  S )NÚcudaéZ   )r   ÚactiveÚget_current_targetÚbackendÚarch)r   Útargets     r   Ú_has_native_tmazTritonSemantic._has_native_tmal  s1   € Ü—‘×1Ñ1Ó3ˆØ—‘ &Ñ(Ò>¨V¯[©[¸BÑ->Ð?r   c                óB  — |t         j                  t         j                  t         j                  t         j                  t         j
                  t         j                  hv sJ d«       ‚|t         j
                  t         j                  hv r| j                  «       sJ d«       ‚y y )Nr0  z-16-bit float types require native tma support)r.   rs   r/   ru   rt   rI   rJ   r>  )r   r9   s     r   Ú$_descriptor_atomic_min_max_supportedz3TritonSemantic._descriptor_atomic_min_max_supportedp  sm   € ØœŸ™¤B§H¡H¬b¯i©i¼¿¹Ä2Ç:Á:ÌrÏ{É{Ð[Ñ[ÐpÐ]pÓpÐ[Ø”R—Z‘Z¤§¡Ð-Ñ-Ø×'Ñ'Ô)ÐZÐ+ZÓZÑ)ð .r   c                óX  — | j                  |||«       | j                  |j                  «       | j                  |d¬«      }t        j
                  j                  }| j                  | j                  j                  ||j                  |j                  |«      t        j                  «      S r*  )r(  r@  r9   r!  r   r1  ÚMINr&   r'   r3  r¦   r.   r,  r4  s        r   Údescriptor_atomic_minz$TritonSemantic.descriptor_atomic_minu  ó…   € Ø× Ñ   u¨gÔ6Ø×1Ñ1°$·*±*Ô=Ø×,Ñ,¨WÀ%Ð,ÓHˆÜ×(Ñ(×,Ñ,ˆØ�{‰{˜4Ÿ<™<×@Ñ@ÀÀtÇ{Á{ÐTY×T`ÑT`ÐbiÓjÔln×lsÑlsÓtÐtr   c                óX  — | j                  |||«       | j                  |j                  «       | j                  |d¬«      }t        j
                  j                  }| j                  | j                  j                  ||j                  |j                  |«      t        j                  «      S r*  )r(  r@  r9   r!  r   r1  ÚMAXr&   r'   r3  r¦   r.   r,  r4  s        r   Údescriptor_atomic_maxz$TritonSemantic.descriptor_atomic_max|  rD  r   c                óÄ  — | j                  |||«       |j                  t        j                  t        j                  t        j
                  t        j                  hv sJ d«       ‚| j                  |d¬«      }t        j                  j                  }| j                  | j                  j                  ||j                  |j                  |«      t        j                  «      S r/  )r(  r9   r.   rs   r/   ru   rt   r!  r   r1  ÚANDr&   r'   r3  r¦   r,  r4  s        r   Údescriptor_atomic_andz$TritonSemantic.descriptor_atomic_andƒ  ó¤   € Ø× Ñ   u¨gÔ6Ø�z‰zœbŸi™i¬¯©´2·9±9¼b¿h¹hÐGÑGÐ\ÐI\Ó\ÐGØ×,Ñ,¨WÀ%Ð,ÓHˆÜ×(Ñ(×,Ñ,ˆØ�{‰{˜4Ÿ<™<×@Ñ@ÀÀtÇ{Á{ÐTY×T`ÑT`ÐbiÓjÔln×lsÑlsÓtÐtr   c                óÄ  — | j                  |||«       |j                  t        j                  t        j                  t        j
                  t        j                  hv sJ d«       ‚| j                  |d¬«      }t        j                  j                  }| j                  | j                  j                  ||j                  |j                  |«      t        j                  «      S r/  )r(  r9   r.   rs   r/   ru   rt   r!  r   r1  ÚORr&   r'   r3  r¦   r,  r4  s        r   Údescriptor_atomic_orz#TritonSemantic.descriptor_atomic_orŠ  s¤   € Ø× Ñ   u¨gÔ6Ø�z‰zœbŸi™i¬¯©´2·9±9¼b¿h¹hÐGÑGÐ\ÐI\Ó\ÐGØ×,Ñ,¨WÀ%Ð,ÓHˆÜ×(Ñ(×+Ñ+ˆØ�{‰{˜4Ÿ<™<×@Ñ@ÀÀtÇ{Á{ÐTY×T`ÑT`ÐbiÓjÔln×lsÑlsÓtÐtr   c                óÄ  — | j                  |||«       |j                  t        j                  t        j                  t        j
                  t        j                  hv sJ d«       ‚| j                  |d¬«      }t        j                  j                  }| j                  | j                  j                  ||j                  |j                  |«      t        j                  «      S r/  )r(  r9   r.   rs   r/   ru   rt   r!  r   r1  ÚXORr&   r'   r3  r¦   r,  r4  s        r   Údescriptor_atomic_xorz$TritonSemantic.descriptor_atomic_xor‘  rK  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«      }| j                  |fd¬«      d   }| j                  j                  |j                  |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 r9  zdescriptor gather of ú must have at least ú columns, but got Fr  )r^   r.   r   rS  r  rA  r9   rŸ  r:  r!  r'   Úcreate_descriptor_gatherr¦   r¨   r&   )
r   r#  Ú	x_offsetsÚy_offsetrÈ  r×  r9   Úmin_colsre   rf   s
             r   Údescriptor_gatherz TritonSemantic.descriptor_gather˜  sû  € Ü˜$¤× 9Ñ 9Ô:Ð:Ð:Ø Ò#ÐJÐ%JÓJÐ#Ø "Ò$ÐLÐ&LÓLÐ$ô �4×#Ñ#Ó$¨Ò)Ð_Ð-LÈT×M]ÑM]ÐL^Ð+_Ó_Ð)Ø×Ñ Ñ" aÒ'ÐhÐ+UÐVZ×VfÑVfÐUgÐ)hÓhÐ'ô �9—?‘?Ó# qÒ(Ð\Ð,JÈ9Ï?É?ÐJ[Ð*\Ó\Ð(ð �‰˜qÑ! QÒ&ÐqÐ*_Ð`i×`oÑ`oÐ_pÐ(qÓqÐ&Ø—
‘
ˆØ˜×1Ñ1Ñ1°AÑ5ˆØ×ÑØñØòð 	AØ3°E°7Ð:NÈxÈjÐXjÐko×k{Ñk{Ð|}Ñk~Ðjð  Aó	Að ô �}‰}˜TŸZ™Z¨)¯/©/¸!Ñ*<¸d×>NÑ>NÈqÑ>QÐ)RÓSˆØ×-Ñ-¨x¨lÈÐ-ÓNÈqÑQˆØ�L‰L×1Ñ1°$·+±+¸y×?OÑ?OÐQYÐ[_×[eÑ[eÐfj×frÑfrÓ[sÓ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   › �«       ‚| j                  |fd¬«      d   }| j                  j                  |j                  |j                  |j                  |«       | j                  d t        j                  «      S )Nr   rT  r   r   rU  rV  rW  z6descriptor scatter must have at least 8 rows, but got r9  zdescriptor scatter of rX  rY  Fr  )r^   r.   r   rS  r  rA  Úshapaer9   rŸ  r!  r'   Úcreate_descriptor_scatterr¦   r&   r,  )r   r#  rH   r[  r\  r9   r]  s          r   Údescriptor_scatterz!TritonSemantic.descriptor_scatter°  s   € Ü˜$¤× 9Ñ 9Ô:Ð:Ð:ô �4×#Ñ#Ó$¨Ò)Ð_Ð-LÈT×M]ÑM]ÐL^Ð+_Ó_Ð)Ø×Ñ Ñ" aÒ'ÐhÐ+UÐVZ×VfÑVfÐUgÐ)hÓhÐ'ô �9—?‘?Ó# qÒ(Ð]Ð,JÈ9×K[ÑK[ÐJ\Ð*]Ó]Ð(ð �‰˜qÑ! QÒ&ÐrÐ*`Ðaj×apÑapÐ`qÐ(rÓrÐ&Ø—
‘
ˆØ˜×1Ñ1Ñ1°AÑ5ˆØ×ÑØñØòð 	BØ4°U°GÐ;OÐPXÈzÐYkÐlp×l|Ñl|Ð}~Ñlð  lAð  Bó	Bð ð ×-Ñ-¨x¨lÈÐ-ÓNÈqÑQˆØ�‰×.Ñ.¨t¯{©{¸E¿L¹LÈ)×JZÑJZÐ\dÔeØ�{‰{˜4¤§¡Ó)Ð)r   c           	     óÐ  — |�t        d«      ‚|j                  j                  j                  «       }|j                  j	                  «       s| j                  ||«      }|j                  j	                  «       sJ d«       ‚||j                  j                  «       k(  s&J d|› d|j                  j                  «       › d�«       ‚|j                  j                  j                  |j                  j                  k(  s@J d|j                  j                  j                  › d|j                  j                  › d�«       ‚|j                  j                  j                  }|t        j                  k7  sJ d«       ‚| j                  ||«      }| j                  ||«      }| j                  | j                  j                  |j                  |j                  |||«      t        j                  «      S )	Nr  z-Value argument must be block type or a scalarzBlock shape(z) and value shape(z
) mismatchzBlock element type(z) and value element type(r  )r,   re   r  rƒ  rR  r‰  r.   rr   r  r‡   r&   r'   Úcreate_tensor_pointer_storer¦   r,  )	r   r
  Úvalr  r   rÉ  rØ  r  r  s	            r   Ú_store_block_pointerz#TritonSemantic._store_block_pointerÅ  sÚ  € ð ÐÜÐjÓkÐkð —h‘h×)Ñ)×:Ñ:Ó<ˆØ�x‰x× Ñ Ô"Ø×+Ñ+¨C°Ó=ˆCØ�x‰x× Ñ Ô"ÐSÐ$SÓSÐ"Ø˜cŸh™h×7Ñ7ó 
ò 
ð 	aØ˜+˜Ð&8¸¿¹×9RÑ9RÓ9TÐ8UÐU_Ð`ó	að 
à�x‰x×"Ñ"×-Ñ-°·±×1DÑ1DÒDð  	uÐH[Ð\_×\dÑ\d×\oÑ\o×\zÑ\zÐ[{ð  |Uð  VY÷  V^ñ  V^÷  Viñ  Við  Ujð  jtð  Guó  	uÐDà—‘×$Ñ$×/Ñ/ˆØœŸ™Ò ÐXÐ"XÓXÐ ð ×:Ñ:¸>È;ÓWˆð �i‰i˜˜VÓ$ˆð �{‰{Ø�L‰L×4Ñ4°S·Z±ZÀÇÁÈ^Ð]bÐdlÓmÔoq×ovÑovóxð 	xr   c           	     óF  — |j                   j                  j                  «       s't        d|j                   j	                  «       › d�«      ‚|rt        d«      ‚|j                   j                  «       sL|j                   j                  «       rt        d«      ‚|r%|j                   j                  «       rt        d«      ‚|j                   j                  «       re|j                  }|€| j                  ||«      \  }}n| j                  |||«      \  }}}||j                  k7  rt        d|j                  › d|› �«      ‚|j                   j                  }|j                  }	|	t        j                  k(  rBt        j                  }	t        j                  |	|j                  «      }| j                  ||«      }| j                  ||	«      }|€P| j                  | j                   j#                  |j$                  |j$                  ||«      t        j&                  «      S |j                   j                  j)                  «       st        d«      ‚| j                  | j                   j+                  |j$                  |j$                  |j$                  ||«      t        j&                  «      S )	Nr  z in `tl.store`zÑ`boundary_check` argument is not supported for storing a tensor of pointers or storing a scalar. Because the compiler does not know the boundary; please use block pointers (defined by `make_block_ptr`) insteadzFValue argument cannot be block type if pointer argument is not a blockr  z(Expected pointer argument to have shape z	 but got ú"Mask must have boolean scalar type)re   rƒ   r|   r,   r   rR  rA  Úbroadcast_tensorsr  r.   rr   r  r  r  r‡   r&   r'   Úcreate_storer¦   r,  r¯  Úcreate_masked_store)
r   r
  re  r  r   rÉ  rØ  Ú	ptr_shaper  r  s
             r   Ú_store_legacyzTritonSemantic._store_legacyá  s$  € à�x‰x�‰×%Ñ%Ô'ÜÐ4°S·X±X×5FÑ5FÓ5HÐ4IÈÐXÓYÐYñ Üð Eó Fð Fð
 �x‰x× Ñ Ô"Ø�x‰x× Ñ Ô"Ü Ð!iÓjÐjÙ˜Ÿ	™	×*Ñ*Ô,Ü Ð!hÓiÐið �8‰8×ÑÔØŸ	™	ˆIØˆ|Ø×1Ñ1°#°sÓ;‘�‘Sà!%×!7Ñ!7¸¸SÀ$Ó!G‘��S˜$Ø˜CŸI™IÒ%Ü Ð#KÈCÏIÉIÈ;ÐV_Ð`iÐ_jÐ!kÓlÐlà—‘—‘ˆØ×"Ñ"ˆð ”R—W‘WÒÜ—W‘WˆFÜ—_‘_ V¨V×-AÑ-AÓBˆFØ—)‘)˜C Ó(ˆCð �i‰i˜˜VÓ$ˆð ˆ<Ø—;‘;˜tŸ|™|×8Ñ8¸¿¹ÀSÇZÁZÐQVÐX`ÓaÔce×cjÑcjÓkÐkØ�y‰y×Ñ×'Ñ'Ô)ÜÐAÓBÐBØ�{‰{˜4Ÿ<™<×;Ñ;¸C¿J¹JÈÏ
É
ÐTX×T_ÑT_ÐafÐhpÓqÜŸ7™7ó$ð 	$r   c                ó¬  — | j                  |«      }| j                  |«      }|j                  j                  «       s$|j                  j                  j                  «       rt        d«      ‚|j                  j                  «       r:|j                  j                  j                  «       r| j                  ||||||«      S | j                  ||||||«      S )Nú"Cannot store to a constant pointer)rÏ  rÙ  re   Úis_constrƒ   r,   r|   r  rR  rf  rm  )	r   r
  re  r  r   rÈ  r×  rÉ  rØ  s	            r   ÚstorezTritonSemantic.store  s¯   € ð ×1Ñ1°.ÓAˆØ×/Ñ/°Ó@ˆà�8‰8×ÑÔ #§(¡(§/¡/×":Ñ":Ô"<ÜÐAÓBÐBà�8‰8�?‰?Ô §¡×!4Ñ!4×!=Ñ!=Ô!?à×,Ñ,¨S°#°t¸^ÈUÐT\Ó]Ð]ð ×%Ñ% c¨3°°nÀeÈXÓVÐVr   c           	     óf  — | j                  |«      }| j                  |«      }|j                  j                  j                  }|j
                  dvrt        d«      ‚| j                  | j                  j                  |j                  |j                  |j                  ||«      |j                  «      S )N)é   r9  r—   z9atomic_cas only supports elements with width {16, 32, 64})rð  rû  re   rƒ   r  rŸ  r,   r&   r'   Úcreate_atomic_casr¦   )r   r
  Úcmpre  rï  rú  r  s          r   Ú
atomic_caszTritonSemantic.atomic_cas%  s�   € Ø×Ñ˜sÓ#ˆØ×"Ñ" 5Ó)ˆØ—X‘X—_‘_×/Ñ/ˆ
Ø×(Ñ(°Ñ<ÜÐXÓYÐYØ�{‰{˜4Ÿ<™<×9Ñ9¸#¿*¹*ÀcÇjÁjÐRU×R\ÑR\Ð^aÐchÓiÐkn×ksÑksÓtÐtr   c                ó`  — |j                   j                  j                  «       s&t        d|j                   j	                  «       z   «      ‚|j                   j                  «       s$|j                   j                  j                  «       rt        d«      ‚|j                   j                  j                  }|t        j                  u r|dk7  rt        d|z   dz   «      ‚|t        j                  u r|dk7  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                   j                  «       rX|�*| j                  ||j                   j!                  «       «      }|�*| j                  ||j                   j!                  «       «      }| j#                  ||j                   j                  j                  «      }|€µ| j$                  j'                  d	«      }t        j(                  }|j                   j                  «       r^|j                   j+                  t        j(                  «      }| j$                  j-                  |j/                  | j$                  «      |«      }| j1                  ||«      }|||fS )
Nz)Pointer argument of store instruction is ro  r¬   Úatomic_z does not support fp16z does not support bf16rs  z does not support T)re   rƒ   r|   r,   r   rp  r  r.   rI   rJ   Úint16Úuint16rŸ  r�  rR  r‰  rƒ  r‡   r'   Úget_int1rr   r§   rT  r¨   r&   )r   r
  re  r  Úopr  Úmask_irÚmask_tys           r   Úatom_red_typechecking_implz)TritonSemantic.atom_red_typechecking_impl-  s  € à�x‰x�‰×%Ñ%Ô'ÜÐHÈ3Ï8É8×K\ÑK\ÓK^Ñ^Ó_Ð_Ø�8‰8×ÑÔ #§(¡(×"5Ñ"5×">Ñ">Ô"@ÜÐAÓBÐBØ—X‘X—_‘_×/Ñ/ˆ
ØœŸ™Ñ#¨¨eªÜ˜Y¨™^Ð.FÑFÓGÐGØœŸ™Ñ$¨¨uªÜ˜Y¨™^Ð.FÑFÓGÐGØœ"Ÿ(™(¤B§I¡IÐ.Ñ.°*×2OÑ2OÐRTÒ2TÜ˜Y¨™^Ð.BÑBÄSÈÃ_ÑTÓUÐUØ�8‰8×ÑÔØÐØ×0Ñ0°°s·x±x×7PÑ7PÓ7RÓS�ØˆØ×/Ñ/°°S·X±X×5NÑ5NÓ5PÓQ�Ø�i‰i˜˜SŸX™XŸ_™_×7Ñ7Ó8ˆØˆ<Ø—l‘l×+Ñ+¨DÓ1ˆGÜ—g‘gˆGØ�x‰x× Ñ Ô"ØŸ(™(×2Ñ2´2·7±7Ó;�ØŸ,™,×3Ñ3°G·M±MÀ$Ç,Á,Ó4OÐQXÓY�Ø—;‘;˜w¨Ó0ˆDØ�C˜ˆ~Ðr   c                óê   — |j                   j                  }t        j                  |d¬«      }| j	                  ||«      }| j                  ||dz
  «      }| j                  |t        j                  «      S )NF)r¾  Úsignedr   )r9   rŸ  r.   Úget_int_dtyperü   r  r‡   rr   )r   rf   r¾  ÚidtypeÚixÚsignbits         r   Ú_signbitzTritonSemantic._signbitI  s\   € Ø—7‘7×-Ñ-ˆÜ×!Ñ!¨8¸EÔBˆØ�\‰\˜!˜VÓ$ˆØ—)‘)˜B ¨1¡Ó-ˆØ�y‰y˜¤"§'¡'Ó*Ð*r   c                ó<  — | j                  |||d«      \  }}}| j                  |«      }| j                  |«      }|j                  j                  }|j                  «       rð|j                  «       rp| j                  | j                  j                  t        j                  j                  |j                  |j                  |j                  ||«      |j                  «      S | j                  | j                  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(                  }| j+                  ||«      }| j+                  |t        j,                  |d«      «      }	|t        j                   k(  rt        j.                  nt        j0                  }
| j+                  ||
«      }| j+                  |t        j,                  |
d«      «      }| j3                  |«      }| j5                  |«      }| j                  | j                  j                  t        j                  j                  |	j                  |j                  | j7                  ||«      j                  ||«      |j                  «      }| j                  | j                  j                  t        j                  j8                  |j                  |j                  | j7                  ||«      j                  ||«      |j                  «      }| j;                  |||«      }| j+                  ||«      S )Nrê   z#atomic_max not supported for dtype r   )r  rð  rû  re   rƒ   rR   rÅ   r&   r'   Úcreate_atomic_rmwr   Ú	ATOMIC_OPrF  r¦   ÚUMAXr.   rK   rM   r<   r/   rt   rü   r  rs   ru   r†  r  rš   ÚUMINÚwhere©r   r
  re  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‰ˆˆS�$Ø×Ñ˜sÓ#ˆØ×"Ñ" 5Ó)ˆØ—‘—‘ˆà�=‰=Œ?Ø×#Ñ#Ô%Ø—{‘{Ø—L‘L×2Ñ2´2·<±<×3CÑ3CÀSÇZÁZÐQT×Q[ÑQ[Ð]a×]hÑ]hÐjmÐotÓuØ—H‘Hóð ð —{‘{Ø—L‘L×2Ñ2´2·<±<×3DÑ3DÀcÇjÁjÐRU×R\ÑR\Ð^b×^iÑ^iÐknÐpuÓvØ—H‘Hóð ð œ"Ÿ*™*¤b§j¡jÐ1Ñ1ÜÐAÀ&ÀÐJÓKÐKà#¤r§z¡zÒ1”—’´r·x±xˆØ—‘˜S &Ó)ˆØ—‘˜S¤"§/¡/°&¸!Ó"<Ó=ˆØ%¬¯©Ò3”"—)’)¼¿¹ˆØ—‘˜c 7Ó+ˆØ—‘˜c¤2§?¡?°7¸AÓ#>Ó?ˆØ�m‰m˜CÓ ˆØ�i‰i˜‹nˆØ—+‘+Ø�L‰L×*Ñ*¬2¯<©<×+;Ñ+;¸U¿\¹\È5Ï<É<Ø+/¯9©9°T¸3Ó+?×+FÑ+FÈÈUóTØUZ×U_ÑU_óaˆð —+‘+Ø�L‰L×*Ñ*¬2¯<©<×+<Ñ+<¸f¿m¹mÈVÏ]É]Ø+/¯9©9°T¸3Ó+?×+FÑ+FÈÈUóTØU[×U`ÑU`óbˆð �j‰j˜˜g wÓ/ˆØ�|‰|˜C Ó(Ð(r   c                ó<  — | j                  |||d«      \  }}}| j                  |«      }| j                  |«      }|j                  j                  }|j                  «       rð|j                  «       rp| j                  | j                  j                  t        j                  j                  |j                  |j                  |j                  ||«      |j                  «      S | j                  | j                  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(                  }| j+                  ||«      }| j+                  |t        j,                  |d«      «      }	|t        j                   k(  rt        j.                  nt        j0                  }
| j+                  ||
«      }| j+                  |t        j,                  |
d«      «      }| j3                  |«      }| j5                  |«      }| j                  | j                  j                  t        j                  j                  |	j                  |j                  | j7                  ||«      j                  ||«      |j                  «      }| j                  | j                  j                  t        j                  j8                  |j                  |j                  | j7                  ||«      j                  ||«      |j                  «      }| j;                  |||«      }| j+                  ||«      S )Nré   z#atomic_min not supported for dtype r   )r  rð  rû  re   rƒ   rR   rÅ   r&   r'   rˆ  r   r‰  rB  r¦   r‹  r.   rK   rM   r<   r/   rt   rü   r  rs   ru   r†  r  rš   rŠ  rŒ  r�  s                     r   Ú
atomic_minzTritonSemantic.atomic_minv  rš  r   c           
     óÚ  — | j                  |||d«      \  }}}| j                  |«      }| j                  |«      }|j                  j                  }|j                  «       rt        j                  j                  nt        j                  j                  }| j                  | j                  j                  ||j                  |j                  |j                  ||«      |j                  «      S )Nr¬   )r  rð  rû  re   rƒ   r}   r   r‰  ÚFADDr2  r&   r'   rˆ  r¦   )r   r
  re  r  rï  rú  rŽ  r|  s           r   Ú
atomic_addzTritonSemantic.atomic_addœ  sµ   € Ø×8Ñ8¸¸cÀ4ÈÓO‰ˆˆS�$Ø×Ñ˜sÓ#ˆØ×"Ñ" 5Ó)ˆØ—‘—‘ˆØ"(×"4Ñ"4Ô"6ŒR�\‰\×Ò¼B¿L¹L×<LÑ<LˆØ�{‰{˜4Ÿ<™<×9Ñ9¸"¸c¿j¹jÈ#Ï*É*ÐVZ×VaÑVaÐcfÐhmÓnØŸ8™8ó%ð 	%r   c           
     óV  — | j                  |||d«      \  }}}| j                  |«      }| j                  |«      }| j                  | j                  j                  t        j                  j                  |j                  |j                  |j                  ||«      |j                  «      S )NÚand)r  rð  rû  r&   r'   rˆ  r   r‰  rI  r¦   re   ©r   r
  re  r  rï  rú  s         r   Ú
atomic_andzTritonSemantic.atomic_and¥  ó‘   € Ø×8Ñ8¸¸cÀ4ÈÓO‰ˆˆS�$Ø×Ñ˜sÓ#ˆØ×"Ñ" 5Ó)ˆØ�{‰{Ø�L‰L×*Ñ*¬2¯<©<×+;Ñ+;¸S¿Z¹ZÈÏÉÐUY×U`ÑU`ÐbeÐglÓmÐor×owÑowóyð 	yr   c           
     óV  — | j                  |||d«      \  }}}| j                  |«      }| j                  |«      }| j                  | j                  j                  t        j                  j                  |j                  |j                  |j                  ||«      |j                  «      S )NÚor)r  rð  rû  r&   r'   rˆ  r   r‰  rM  r¦   re   r¢  s         r   Ú	atomic_orzTritonSemantic.atomic_or¬  s�   € Ø×8Ñ8¸¸cÀ4ÈÓN‰ˆˆS�$Ø×Ñ˜sÓ#ˆØ×"Ñ" 5Ó)ˆØ�{‰{Ø�L‰L×*Ñ*¬2¯<©<¯?©?¸C¿J¹JÈÏ
É
ÐTX×T_ÑT_ÐadÐfkÓlÐnq×nvÑnvóxð 	xr   c           
     óV  — | j                  |||d«      \  }}}| j                  |«      }| j                  |«      }| j                  | j                  j                  t        j                  j                  |j                  |j                  |j                  ||«      |j                  «      S )NÚxor)r  rð  rû  r&   r'   rˆ  r   r‰  rP  r¦   re   r¢  s         r   Ú
atomic_xorzTritonSemantic.atomic_xor³  r¤  r   c           
     óV  — | j                  |||d«      \  }}}| j                  |«      }| j                  |«      }| j                  | j                  j                  t        j                  j                  |j                  |j                  |j                  ||«      |j                  «      S )NÚxchg)r  rð  rû  r&   r'   rˆ  r   r‰  ÚXCHGr¦   re   r¢  s         r   Úatomic_xchgzTritonSemantic.atomic_xchgº  s�   € Ø×8Ñ8¸¸cÀ4ÈÓP‰ˆˆS�$Ø×Ñ˜sÓ#ˆØ×"Ñ" 5Ó)ˆØ�{‰{Ø�L‰L×*Ñ*¬2¯<©<×+<Ñ+<¸c¿j¹jÈ#Ï*É*ÐVZ×VaÑVaÐcfÐhmÓnØ�H‰Hóð 	r   c                ó6  — |j                  «       | j                  j                  j                  v s+J d| j                  j                  j                  › d|› �«       ‚|j	                  «       }|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ÚupperrH  r   ÚINPUT_PRECISION)r   Úinput_precisions     r   Ú_str_to_dot_input_precisionz*TritonSemantic._str_to_dot_input_precisionÆ  sž   € Ø×$Ñ$Ó&¨$¯,©,×*>Ñ*>×*[Ñ*[Ñ[ð 	yØ-¨d¯l©l×.BÑ.B×._Ñ._Ð-`Ð`fÐgvÐfwÐxó	yÐ[à)×/Ñ/Ó1ˆØ˜hÒ&Ø&ˆOØ˜hÒ&Ø&ˆOØ˜hÒ&Ø&ˆOÜ”r×)Ñ)¨?Ó;Ð;r   c           
     óæ  — |j                   j                  «       r|j                   j                  «       sJ ‚|j                  j                  «       r|j                  j                  «       r�n2|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                  j                  «       s|j                  j                  «       rwd| j                  j                  j                  v rt        j                   d«       | j#                  |t        j                  «      }| j#                  |t        j                  «      }|j                  j%                  «       xs |j                  j%                  «       }|j                  j'                  «       xs |j                  j'                  «       }|s|r¦|rdnd}	|	| j                  j                  j                  v r~| j                  j                  j(                  }
t        j                   |	› d	|
› d
|
› d�«       | j#                  |t        j                  «      }| j#                  |t        j                  «      }|€ | j                  j                  j*                  }| j-                  |«      }t/        |j0                  «      }t/        |j0                  «      }||cxk(  rdk(  s1n ||cxk(  rdk(  s$n J d|j0                  › d|j0                  › d�«       ‚|j0                  d   j2                  |j0                  d   j2                  k(  sVJ d|j0                  › d|j0                  › d|j0                  d   j2                  › d|j0                  d   j2                  › d�	«       ‚| j                  j4                  j7                  d«      	 €J d«       ‚ | j                  j4                  d   |j                   |j                   «      }|j0                  d   j2                  |d   k\  r>|j0                  d   j2                  |d   k\  r|j0                  d   j2                  |d   k\  sJ d|d   › d|d   › d|d   › �«       ‚|j                   j8                  j;                  «       r[|j                   j8                  t        j
                  k(  sJ d«       ‚| j                  j=                  d«      }t        j>                  }�n'|jA                  «       rtC        d«      ‚|j                   j8                  jE                  «       s$|j                   j8                  jA                  «       r,| j                  jG                  d«      }t        j                  }n˜|j                   j8                  jI                  «       r,| j                  jK                  d«      }t        j                  }nH|jM                  «       r| j                  jO                  d«      n| j                  jG                  d«      }|}|j                   j0                  d   }|j                   j0                  d   }|j                   j0                  d   }|dk(  r|j                   j0                  d   nd }t	        jP                  ||r|||gn||g«      }|€6| j                  jS                  |jU                  | j                  «      |«      }nJ|jV                  }|j                   j0                  |j0                  k(  r|j                   jX                  |k(  sJ ‚|€X|j                  j                  «       r;|j                  j                  «       r!| j                  j                  jZ                  }nNd}nK|j                  j                  «       r1|j                  j                  «       r||kD  rtC        d |› d!|› d�«      ‚| j]                  | j                  j_                  |jV                  |jV                  |||«      |«      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  rh  éþÿÿÿ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 ()0re   rR  r9   rQ   r.   r  Úuint8rI   rJ   rK   rM   rª  r'   r˜   Ú!deprecated_fp8_dot_operand_dtypesÚwarningsÚwarnr‡   Ú
is_fp8e4b8Úis_fp8e5b16r<  Údefault_dot_input_precisionr¼  rS  rA  rH   r«  r¬  rƒ   rR   Ú	get_int32r/   rP   r,   rN   rF  rL   Úget_fp64rO   Úget_fp16r:  rT  r¨   r¦   r  Úmax_num_imprecise_acc_defaultr&   Ú
create_dot)r   r‰   rŠ   Úaccr»  Úmax_num_imprecise_accÚ	out_dtypeÚuses_fp8e4b8Úuses_fp8e5b16Ú	type_namer<  Úlhs_rankÚrhs_rankrÅ  r  Úret_scalar_tyÚMÚNÚKÚBrÈ   Ú
acc_handles                         r   ÚdotzTritonSemantic.dotÒ  s˜  € à�x‰x× Ñ Ô" s§x¡x×'8Ñ'8Ô':Ð:Ð:à�9‰9×ÑÔ #§)¡)×"2Ñ"2Ô"4áà—9‘9¤§¡¬"¯(©(´B·J±JÄÇÁÌRÏZÉZÜ!#§¡ð!-ñ -ð SØ0FÀsÇyÁyÀkÐ.RóSð -à—9‘9¤§¡¬"¯(©(´B·J±JÄÇÁÌRÏZÉZÜ!#§¡ð!-ñ -ð SØ0FÀsÇyÁyÀkÐ.RóSð -à—9‘9 §	¡	Ò)ÐoÐ-SÐTW×T]ÑT]ÐS^Ð^cÐdg×dmÑdmÐcnÐ+oÓoÐ)à�9‰9× Ñ Ô" c§i¡i×&;Ñ&;Ô&=Ø˜TŸ\™\×1Ñ1×SÑSÑSÜ—‘ð môð —)‘)˜C¤§¡Ó,ˆCØ—)‘)˜C¤§¡Ó,ˆCà—y‘y×+Ñ+Ó-ÒG°·±×1EÑ1EÓ1GˆØŸ	™	×-Ñ-Ó/ÒJ°3·9±9×3HÑ3HÓ3JˆÙ™=Ù%1™	°zˆIØ˜DŸL™L×0Ñ0×RÑRÑRØ—|‘|×+Ñ+×0Ñ0�Ü—‘Ø �kÐ!NÈtÈfð U6Ø6:°VÐ;KðMôNð —i‘i ¤R§Z¡ZÓ0�Ø—i‘i ¤R§Z¡ZÓ0�àÐ"Ø"Ÿl™l×2Ñ2×NÑNˆOà×:Ñ:¸?ÓKˆä�s—y‘y“>ˆÜ�s—y‘y“>ˆØ˜8Ô( qÔ(¨H¸Ô,EÀAÔ,Eð  	VÐItÐux×u~Ñu~Ðtð  @Ið  JM÷  JSñ  JSð  ITð  TUð  HVó  	VÐEØ�y‰y˜‰}×"Ñ" c§i¡iØñ'ß‘òð 	uØ,¨S¯Y©Y¨KÐ7PÐQT×QZÑQZÐP[ð  \Yð  Z]÷  Zcñ  Zcð  dfñ  Zg÷  Zmñ  Zmð  Ynð  n^ð  _b÷  _hñ  _hð  ikñ  _l÷  _rñ  _rð  ^sð  stð  uó	uð à�|‰|×'Ñ'×+Ñ+ØóØ#'ð(ð 	^Ø)]ó	^ð (à?�t—|‘|×/Ñ/°Ñ?ÀÇÁÈ#Ï(É(ÓSˆØ�y‰y˜‰}×"Ñ" l°1¡oÒ5¸#¿)¹)ÀB¹-×:MÑ:MÐQ]Ð^_ÑQ`Ò:`Ø—	‘	˜"‘×#Ñ# |°A¡Ò6ð	và0°¸a±Ð0AÀÈÐVWÉÐHYÐYcÐdpÐqrÑdsÐctÐuó	vð 7ð �8‰8�?‰?×!Ñ!Ô#Ø—8‘8—?‘?¤b§g¡gÒ-ÐEÐ/EÓEÐ-Ø—‘×'Ñ'¨Ó*ˆBÜŸH™HŠMØ×ÑÔ ÜØzóð ð �X‰X�_‰_×$Ñ$Ô&¨#¯(©(¯/©/×*AÑ*AÔ*CØ—‘×&Ñ& qÓ)ˆBÜŸJ™J‰MØ�X‰X�_‰_×$Ñ$Ô&Ø—‘×&Ñ& qÓ)ˆBÜŸJ™J‰Mà-6×->Ñ->Ô-@�—‘×&Ñ& qÔ)ÀdÇlÁl×F[ÑF[Ð\]ÓF^ˆBØ%ˆMà�H‰H�N‰N˜2ÑˆØ�H‰H�N‰N˜2ÑˆØ�H‰H�N‰N˜2ÑˆØ!)¨Q¢ˆC�H‰H�N‰N˜1Ò°DˆÜ—‘˜}¹1¨q°!°Q©iÀ1ÀaÀ&ÓIˆØˆ;ØŸ™×2Ñ2°6·<±<ÀÇÁÓ3MÈrÓR‰JàŸ™ˆJØ—8‘8—>‘> V§\¡\Ò1°c·h±h×6IÑ6IÈYÒ6VÐVÐVð !Ð(Ø�y‰y×ÑÔ! c§i¡i×&6Ñ&6Ô&8Ø(,¯©×(<Ñ(<×(ZÑ(ZÑ%à()Ñ%à�y‰y×ÑÔ! c§i¡i×&6Ñ&6Ô&8Ð=RÐUVÒ=VÜ Ð#:Ð;PÐ:QÐQaÐbcÐadÐdeÐ!fÓgÐgà�{‰{Ø�L‰L×#Ñ# C§J¡J°·
±
¸JÈÐYnÓoÐqwóyð 	yr   c                óz   — t        t        j                  |j                  «       d «      }|€t	        d|› d�«      ‚|S )NzInvalid float format: rn   )rH  r   ÚScaleDotElemTypeTYr¹  r,   )r   Úfloat_formatÚty_enums      r   Ú_str_to_fp_typezTritonSemantic._str_to_fp_type/  s>   € Üœ"×/Ñ/°×1CÑ1CÓ1EÀtÓLˆØˆ?ÜÐ5°l°^À1ÐEÓFÐFØˆr   c                ó<  — t         j                  t         j                  t         j                  t         j                  dœj                  |«      }|€B|dk(  s
J d|› �«       ‚|j                  t         j                  k(  sJ d|j                  › �«       ‚|S |j                  |k(  r|S t         j                  t         j                  t         j                  t         j                  dœ|   }|j                  |k(  sJ d|› d|j                  › �«       ‚| 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Úfp16Úe2m1z)Internal Error: Unexpected float format: z)e2m1 format must be packed as uint8. Got zUnexpected dtype for r°  )
r.   Úfloat8e5Ú
float8e4nvrJ   rI   r¬  r9   rÆ  rz  rü   )r   re  rã  Ú	triton_tyÚunsigned_tys        r   Ú_bitcast_to_fp_typez"TritonSemantic._bitcast_to_fp_type5  sü   € ô
  Ÿ[™[´"·-±-ÌÏÉÜ—Z‘Zñ!ß!$¡ \Ó!2ð 	àÐØ 6Ò)ÐeÐ-VÐWcÐVdÐ+eÓeÐ)Ø—9‘9¤§¡Ò(ÐaÐ,UÐVY×V_ÑV_ÐU`Ð*aÓaÐ(ØˆJØ�9‰9˜	Ò!ØˆJä#%§8¡8´R·X±XÄrÇyÁyÔZ\×ZcÑZcÑdÐeqÑrˆKØ—9‘9 Ò+ÐdÐ/DÀ\ÀNÐRXÐY\×YbÑYbÐXcÐ-dÓdÐ+Ø—<‘<  YÓ/Ð/r   c                ó^  — |�T|j                   j                  «       rdnd}|j                  j                  }|dd  |||z  gk(  sJ d|› d||z  › d|› �«       ‚|�U|j                   j                  «       rdnd}|j                  j                  }|dd  |||z  gk(  sJ d|› d||z  › d|› �«       ‚y y )Nrs  r9  rÄ  z*lhs_scale must be a tensor of shape [..., r�  z]. Got z*rhs_scale must be a tensor of shape [..., )r9   Ú
is_fp8e4nvre   rA  )	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û   € ØÐ Ø!*§¡×!;Ñ!;Ô!=™2À2ˆLØ'Ÿn™n×2Ñ2ˆOØ" 2 3Ð'Ø�1˜Ñ$ð,ò ð mà;¸A¸3¸bÀÀlÑARÐ@SÐSZÐ[jÐZkÐlómð ð Ð Ø!*§¡×!;Ñ!;Ô!=™2À2ˆLØ'Ÿn™n×2Ñ2ˆOØ" 2 3Ð'Ø�1˜Ñ$ð,ò ð mà;¸A¸3¸bÀÀlÑARÐ@SÐSZÐ[jÐZkÐlómñ ð !r   c                ó"  — |j                   j                  «       r|j                   j                  «       sJ ‚t        |j                  «      }t        |j                  «      }||cxk(  rdk(  s1n ||cxk(  rdk(  s$n J d|j                  › d|j                  › d�«       ‚|j                  }|j                  }| j                  |«      }| j                  |«      }h d£}||v s
J d|› �«       ‚||v s
J d|› �«       ‚|d u xs* t        |t        j                  «      xr |j                  d u }|d u xs* t        |t        j                  «      xr |j                  d u }| j                  ||«      }| j                  ||«      }|	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                  j                  d«      }|€6| j                  j                  |j                  | j                  «      |«      }nJ|j                  }|j                   j                  |j                  k(  r|j                   j                   |k(  sJ ‚|rd n|j                  } |rd n|j                  }!| j#                  ||||rd n||rd n|«       | j%                  | j                  j'                  |j                  |!||j                  | |||	|
|«
      |«      S )Nr   rÁ  rÂ  rÃ  r  >   ré  rë  rè  rç  rê  zNYI: lhs_format zNYI: rhs_format rë  zBonly mxfp4 inputs can be packed along a dimension different than KrÄ  r   zCReduction dimension should pack the same number of elements; (lhs: r   )re   rR  rS  rA  rH   rå  r^   r.   r_   rð  r:  r'   rF  rT  r¨   r¦   r  rø  r&   Úcreate_dot_scaled)"r   r‰   ró  Ú
lhs_formatrŠ   rô  Ú
rhs_formatrÒ  Ú	fast_mathÚ
lhs_k_packÚ
rhs_k_packrÔ  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  sâ  € ð �x‰x× Ñ Ô" s§x¡x×'8Ñ'8Ô':Ð:Ð:ä�s—y‘y“>ˆÜ�s—y‘y“>ˆØ˜8Ô( qÔ(¨H¸Ô,EÀAÔ,Eð  	VÐItÐux×u~Ñu~Ðtð  @Ið  JM÷  JSñ  JSð  ITð  TUð  HVó  	VÐEØ$×*Ñ*ˆ
Ø$×*Ñ*ˆ
Ø×.Ñ.¨zÓ:ˆØ×.Ñ.¨zÓ:ˆÚBˆØ˜_Ñ,ÐMÐ0@ÀÀÐ.MÓMÐ,Ø˜_Ñ,ÐMÐ0@ÀÀÐ.MÓMÐ,Ø%¨Ð-Òr´*¸YÌÏÉÓ2UÒ2qÐZc×ZiÑZiÐmqÐZqÐØ%¨Ð-Òr´*¸YÌÏÉÓ2UÒ2qÐZc×ZiÑZiÐmqÐZqÐØ×&Ñ& s¨JÓ7ˆØ×&Ñ& s¨JÓ7ˆá˜Z¨6Ò1ÐwÐ3wÓwÐ1Ù˜Z¨6Ò1ÐwÐ3wÓwÐ1Ø—8‘8—>‘> " #Ð&‰ˆˆ5Ø—8‘8—>‘> " #Ð&‰ˆˆqØ" fÒ,‘1°!ˆØ" fÒ,‘1°!ˆÙ+5�x %Ò'¸5ˆÙ+5�x %Ò'¸5ˆØ˜|Ò+ð  	TÐ/rÐsv×s|Ñs|Ðr}ð  ~Gð  HK÷  HQñ  HQð  GRð  RSð  .Tó  	TÐ+à!)¨Q¢ˆC�H‰H�N‰N˜1Ò°DˆØˆÙØ�H‘‰Aà�H‘ˆAÙØ�H‘ˆAÜ—‘˜y±q¨1¨a°©)¸qÀ!¸fÓEˆØ�\‰\×"Ñ" 1Ó%ˆØˆ;ØŸ™×2Ñ2°6·<±<ÀÇÁÓ3MÈrÓR‰JàŸ™ˆJØ—8‘8—>‘> V§\¡\Ò1°c·h±h×6IÑ6IÈYÒ6VÐVÐVÙ#4™4¸)×:JÑ:JÐÙ#4™4¸)×:JÑ:JÐØ× Ñ   A qÑ2C©$ÈÙ):¡À	ô	Kà�{‰{Ø�L‰L×*Ñ*¨3¯:©:Ð7GÈÐZ]×ZdÑZdÐfvØ+:¸IÀzÐS]Ð_iókØlrótð 	tr   c                óP  — |j                   t        j                  k7  r"t        j                  d|j                   › �«       | j                  |t        j                  «      }| j                  ||dd«      \  }}|j                  j                  «       r+| j                  ||«      \  }}| j                  ||«      \  }}n| j                  ||«      \  }}|j                  }| j                  | 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)r9   r.   rr   rÈ  rÉ  r‡   r•   re   rR  rˆ   r&   r'   Úcreate_selectr¦   )r   Ú	conditionrf   rÝ   r’  rÈ   s         r   rŒ  zTritonSemantic.whereŽ  s  € Ø�?‰?œbŸg™gÒ%Ü�M‰MØyð  {D÷  {Jñ  {Jð  zKð  Lôð —I‘I˜i¬¯©Ó1ˆ	Ø×0Ñ0°°A°t¸TÓB‰ˆˆ1à�>‰>×"Ñ"Ô$Ø×4Ñ4°YÀÓB‰LˆI�qØ×,Ñ,¨Q°Ó2‰DˆA‰qà×4Ñ4°YÀÓB‰LˆI�qØ—‘ˆØ�{‰{˜4Ÿ<™<×5Ñ5°i×6FÑ6FÈÏÉÐRS×RZÑRZÓ[Ð]cÓdÐdr   c                ó\   — |rt        j                  ||«      }n|}| j                  ||«      S r)   )r.   r:  r&   )r   rf   rV   r“  Úres_tys        r   Úwrap_tensorzTritonSemantic.wrap_tensor¢  s-   € ÙÜ—]‘] 9¨iÓ8‰Fð ˆFØ�{‰{˜1˜fÓ%Ð%r   c                ó&  ‡ ‡‡‡	‡
— |€t        ˆ fd„‰D «       «      Šd}‰d   j                  j                  Š
t        ‰
«      }||k  sJ d|› d�«       ‚t	        ‰
«      D ��cg c]  \  }}||k7  sŒ|‘Œ c}}Š	t        ˆ
fd„‰D «       «      sJ d«       ‚‰ j                  j                  ‰D �cg c]  }|j                  ‘Œ c}|«      Š |‰«       ‰j                  «       sJ ‚t        ˆˆˆ	ˆ fd„t        t        ‰«      «      D «       «      S c c}}w c c}w )Nc              3  óp   •K  — | ]-  }‰j                  ||j                  j                  gd ¬«      –— Œ/ y­w)Tri  N)r]  rL  rH   )rx  Útr   s     €r   rz  z+TritonSemantic.reduction.<locals>.<genexpr>¬  s+   øè ø€ Ò^ÐRS˜4Ÿ<™<¨¨A¯G©G¯M©M¨?È˜<×MÑ^ùs   ƒ36r   z&reduction axis must be < inputs rank (r  c              3  óP   •K  — | ]  }|j                   j                  ‰k(  –— Œ y ­wr)   )re   rA  )rx  r  rA  s     €r   rz  z+TritonSemantic.reduction.<locals>.<genexpr>³  s   øè ø€ Ò9¨Q�1—6‘6—<‘< 5Õ(Ñ9ùs   ƒ#&z-all reduction inputs must have the same shapec              3  ó�   •K  — | ]=  }‰j                  ‰j                  |«      ‰|   j                  j                  ‰«      –— Œ? y ­wr)   ©r  Ú
get_resultre   rƒ   )rx  r‡  ÚinputsÚ	reduce_opr“  r   s     €€€€r   rz  z+TritonSemantic.reduction.<locals>.<genexpr>¹  s@   øè ø€ ò uØ\]ˆD×Ñ˜Y×1Ñ1°!Ó4°f¸Q±i·n±n×6KÑ6KÈY×Wñuùó   ƒAA)Útuplere   rA  rS  r„  Úallr'   Úcreate_reducer¦   Úverifyr@  )r   r  r1   Úregion_builder_fnÚrankr‡  r\  r  r  r“  rA  s   ``      @@@r   Ú	reductionzTritonSemantic.reductionª  s  ü€ Øˆ<ÜÓ^ÐW]Ô^Ó^ˆFØˆDà�q‘	—‘×$Ñ$ˆÜ�5‹zˆØ�dŠ{ÐLÐDÀTÀFÈ!ÐLÓLˆ{Ü#,¨UÓ#3×A™4˜1˜a°q¸D³y’QÓAˆ	ÜÓ9°&Ô9Ô9ÐjÐ;jÓjÐ9à—L‘L×.Ñ.À&Ö/I¸Q°·³Ò/IÈ4ÓPˆ	Ù˜)Ô$Ø×ÑÔ!Ð!Ð!äö uÜafÔgjÐkqÓgrÓasôuó uð 	uùó Bùò 0Js   Á DÁ.DÂ*Dc                óæ  ‡ ‡‡‡— ‰d   j                   j                  Št        ‰«      }| |cxk  r|k  sn J d|› d|› d�«       ‚|dk  r||z  }‰D ]"  }|j                   j                  ‰k(  rŒJ d«       ‚ ‰ j                  j	                  ‰D �cg c]  }|j
                  ‘Œ c}||«      Š |‰«       ‰j                  «       sJ ‚t        ˆˆˆ ˆfd„t        t        ‰«      «      D «       «      S c c}w )Nr   z
scan axis z must be < inputs rank (r  z(all scan inputs must have the same shapec              3  ó�   •K  — | ]=  }‰j                  ‰j                  |«      ‰|   j                  j                  ‰«      –— Œ? y ­wr)   r  )rx  r‡  r  Úscan_opr   rA  s     €€€€r   rz  z2TritonSemantic.associative_scan.<locals>.<genexpr>Ñ  s;   øè ø€ ÒwÐ_`�T×%Ñ% g×&8Ñ&8¸Ó&;¸VÀA¹Y¿^¹^×=RÑ=RÐTY×ZÑwùr  )	re   rA  rS  r'   Úcreate_scanr¦   r!  r  r@  )	r   r  r1   r"  Úreverser#  r  r'  rA  s	   ``     @@r   Úassociative_scanzTritonSemantic.associative_scanÀ  sí   û€ à�q‘	—‘×$Ñ$ˆÜ�5‹zˆàˆu˜Ô#˜tÔ#ÐW z°$°Ð7OÐPTÈvÐUVÐ%WÓWÐ#à�!Š8Ø�D‰LˆDàò 	UˆAØ—6‘6—<‘< 5Ó(ÐTÐ*TÓTÐ(ð	Uð —,‘,×*Ñ*¸fÖ+E¸¨A¯H«HÒ+EÀtÈWÓUˆÙ˜'Ô"Ø�~‰~ÔÐÐäÖwÔdiÔjmÐntÓjuÓdvÔwÓwÐwùò	 ,Fs   ÂC.c                ó   — |j                   j                  «       sJ d«       ‚t        |j                  j                  «      }t        |j                  j                  «      |k(  sJ d«       ‚| |cxk  r|k  sn J d|› d|› d�«       ‚|dk  r||z  }t        |«      D ]F  }||k(  rŒ	|j                  j                  |   |j                  j                  |   k(  rŒ=J d|› d�«       ‚ | j                  j                  |j                  |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)r9   rR   rS  re   rA  r@  r'   Úcreate_gatherr¦   r  rƒ   )r   ÚsrcÚindexr1   r#  ry  Úgathers          r   r/  zTritonSemantic.gather×  s1  € Ø�{‰{×!Ñ!Ô#ÐFÐ%FÓFÐ#ä�3—8‘8—>‘>Ó"ˆÜ�5—:‘:×#Ñ#Ó$¨Ò,Ð`Ð.`Ó`Ð,àˆu˜Ô#˜tÔ#ÐY |°D°6Ð9QÐRVÐQWÐWXÐ%YÓYÐ#Ø�!Š8Ø�D‰LˆDä�t“ò 	yˆAØ�DŠyØØ—:‘:×#Ñ# AÑ&¨#¯(©(¯.©.¸Ñ*;Ó;Ðx¸zÈ$ÈÐOwÐ=xÓxÐ;ð	yð
 —‘×+Ñ+¨C¯J©J¸¿¹ÀdÓKˆØ×Ñ ¨¯©¯©¸¿¹×9IÑ9IÓJÐJr   c                óò   — |sy|^}}t        t        |«      «      D ]  }| j                  |||   «      \  }||<   Œ t        t        |«      dz
  «      D ]  }| j                  |||   «      \  }||<   Œ |g|¢­S )Nr  r   )r@  rS  rˆ   )r   r  ÚheadÚtailr‡  s        r   ri  z TritonSemantic.broadcast_tensorsí  s‘   € ÙØØˆˆˆtÜ”s˜4“yÓ!ò 	EˆAØ ×5Ñ5°d¸DÀ¹GÓD‰MˆD�$�q’'ð	Eä”s˜4“y 1‘}Ó%ò 	EˆAØ ×5Ñ5°d¸DÀ¹GÓD‰MˆD�$�q’'ð	Eàˆ}�t‰}Ðr   c           	     óþ  ‡ ‡—  ‰ j                   |Ž }t        |«      dkD  sJ d«       ‚|D �cg c]*  }|d   j                  j                  |j                  «      ‘Œ, }}‰ j
                  j                  |D �cg c]  }|j                  ‘Œ c}|D �cg c]  }|j                  ‰ j
                  «      ‘Œ c}|«      Š |‰«       ‰j                  «       sJ ‚t        ˆˆ fd„t        |«      D «       «      S c c}w c c}w c c}w )Nr   z1map_elementwise must have at least 1 input tensorc              3  óf   •K  — | ](  \  }}‰j                  ‰j                  |«      |«      –— Œ* y ­wr)   )r&   r  )rx  r‡  r½  Úelementwise_opr   s      €€r   rz  z1TritonSemantic.map_elementwise.<locals>.<genexpr>  s+   øè ø€ ÒhÁuÀqÈ"�T—[‘[ ×!:Ñ!:¸1Ó!=¸r×BÑhùs   ƒ.1)ri  rS  re   r§   rƒ   r'   Úcreate_map_elementwiser¦   r¨   r!  r  r„  )r   r  Úresult_typesÚpackr"  r½  r  r5  s   `      @r   Úmap_elementwisezTritonSemantic.map_elementwise÷  sÙ   ù€ à'�×'Ñ'¨Ð0ˆä�6‹{˜QŠÐSÐ SÓSˆØLXÖYÀb˜˜q™	Ÿ™×6Ñ6°r·y±yÕAÐYˆÐYØŸ™×<Ñ<Ø%Ö&˜!ˆQ�X‹XÒ&Ø.:Ö;¨ˆR�X‰X�d—l‘lÕ#Ò;Øó
ˆñ
 	˜.Ô)Ø×$Ñ$Ô&Ð&Ð&äÔhÔPYÐZfÓPgÔhÓhÐhùò Zùâ&ùÚ;s   «/C0Á5C5Â"C:
c                óè  — t        |j                  «      dk(  sJ d«       ‚|j                  j                  «       sJ d«       ‚|�W| j	                  ||j                  «      }|j
                  j                  j                  «       st        d«      ‚|j                  }| j                  | j                  j                  |j                  ||«      t        j                  t        j                  |g«      «      S )Nr   z histogram only supports 1D inputz%histogram only supports integer inputrh  )rS  rA  r9   rR   r‰  re   rƒ   r¯  r,   r¦   r&   r'   Úcreate_histogramr.   r:  r/   )r   r®   Únum_binsr  s       r   Ú	histogramzTritonSemantic.histogram  sÃ   € Ü�5—;‘;Ó 1Ò$ÐHÐ&HÓHÐ$Ø�{‰{×!Ñ!Ô#ÐLÐ%LÓLÐ#ØÐØ×,Ñ,¨T°5·;±;Ó?ˆDØ—9‘9×#Ñ#×+Ñ+Ô-Ü Ð!EÓFÐFØ—;‘;ˆDØ�{‰{˜4Ÿ<™<×8Ñ8¸¿¹ÀxÐQUÓVÜŸ=™=¬¯©°H°:Ó>ó@ð 	@r   c                ó  — t        dt        |j                  «      «      t        |«      k7  rt        d«      ‚|j                  j                  dt        j                  ||j                  j                  «       «      «       |S )Nr   zAShape of input to multiple_of does not match the length of valuesztt.divisibility)	rê   rS  rA  r,   r¦   Úset_attrr   Ú	make_attrÚget_context©r   rf   r‹  s      r   Úmultiple_ofzTritonSemantic.multiple_of  s[   € Üˆq”#�a—g‘g“,Ó¤3 v£;Ò.ÜÐ`ÓaÐaØ	�‰×ÑÐ+¬R¯\©\¸&À!Ç(Á(×BVÑBVÓBXÓ-YÔZØˆr   c                óî   — t        |j                  «      t        |«      k7  rt        d«      ‚|j                  j	                  dt        j                  ||j                  j                  «       «      «       |S )NzDShape of input to max_contiguous does not match the length of valuesztt.contiguity©rS  rA  r,   r¦   r?  r   r@  rA  rB  s      r   Úmax_contiguouszTritonSemantic.max_contiguous  sS   € Üˆq�w‰w‹<œ3˜v›;Ò&ÜÐcÓdÐdØ	�‰×Ñ˜/¬2¯<©<¸ÀÇÁ×@TÑ@TÓ@VÓ+WÔXØˆr   c                óî   — t        |j                  «      t        |«      k7  rt        d«      ‚|j                  j	                  dt        j                  ||j                  j                  «       «      «       |S )NzCShape of input to max_constancy does not match the length of valuesztt.constancyrE  rB  s      r   Úmax_constancyzTritonSemantic.max_constancy#  sS   € Üˆq�w‰w‹<œ3˜v›;Ò&ÜÐbÓcÐcØ	�‰×Ñ˜.¬"¯,©,°v¸q¿x¹x×?SÑ?SÓ?UÓ*VÔWØˆr   c                ór   — | j                  | j                  j                  «       t        j                  «      S r)   )r&   r'   Úcreate_barrierr.   r,  )r   s    r   Údebug_barrierzTritonSemantic.debug_barrier)  s$   € Ø�{‰{˜4Ÿ<™<×6Ñ6Ó8¼"¿'¹'ÓBÐBr   c                ó¼  — |j                  d«      s|r|dz  }|j                  d«      s
|r|d d dz   }t        |«      dkD  r|j                  d«      sd|z   }|D �cg c]  }|j                  ‘Œ }}|D �cg c]  }|j                  j                  «       ‘Œ }}| j                  | j                  j                  ||||«      t        j                  «      S c c}w c c}w )Nú r‚  rh  r   )ÚendswithrS  Ú
startswithr¦   r9   rÅ   r&   r'   Úcreate_printr.   r,  )r   ÚprefixÚargsÚhexÚargÚnew_argsÚ	is_signeds          r   Údevice_printzTritonSemantic.device_print,  sÆ   € ð �‰˜sÔ#©Ø�c‰MˆFØ�‰˜tÔ$©Ø˜C˜R�[ 4Ñ'ˆFÜˆv‹;˜Š? 6×#4Ñ#4°SÔ#9Ø˜6‘\ˆFà*.Ö/ 3�C—J“JÐ/ˆÐ/Ø:>Ö?°3�S—Y‘Y×,Ñ,Õ.Ð?ˆ	Ð?Ø�{‰{˜4Ÿ<™<×4Ñ4°V¸SÀ(ÈIÓVÔXZ×X_ÑX_Ó`Ð`ùò 0ùÚ?s   ÁCÁ5!Cc                ó  — | j                   j                  j                  sy |�!| j                  || j	                  |«      «      }| j                  | j                   j                  |j                  |«      t        j                  «      S r)   )
r'   r˜   Údebugrö   r  r&   Úcreate_assertr¦   r.   r,  )r   r¢   r£   r  s       r   r�   zTritonSemantic.device_assert:  s`   € Ø�|‰|×#Ñ#×)Ò)ØØÐØ—8‘8˜D $§)¡)¨D£/Ó2ˆDØ�{‰{˜4Ÿ<™<×5Ñ5°d·k±kÀ3ÓGÌÏÉÓQÐQr   c                óˆ   — | j                  | j                  j                  |j                  «      t        j
                  «      S r)   )r&   r'   Úcreate_assumer¦   r.   r,  )r   r¢   s     r   ÚassumezTritonSemantic.assumeA  s*   € Ø�{‰{˜4Ÿ<™<×5Ñ5°d·k±kÓBÄBÇGÁGÓLÐLr   c                ór  — t        |t        «      rt        j                  |«      }t        |t        j                  «      rãt        |j                  t
        «      r%| j                  j                  |j                  «      S |rQd|j                  cxk  rdk  sn J d|j                  › d�«       ‚| j                  j                  |j                  «      S d|j                  cxk  rdk  sn J d|j                  › d�«       ‚| j                  j                  |j                  «      S t        |t        j                  «      ré|j                  j                  dk(  sJ d	«       ‚|j                  j                  «       sJ d
«       ‚|j                  t        j                  k7  rY|rW| j                  j                  |j                   | j                  j#                  «       |j                  j%                  «       «      S |j                  t        j                  k(  r	|sJ d«       ‚|j                   S J dt'        |«      › �«       ‚)Nrl   rm   z@Block pointers only support 64 bit `shape/strides`, got a value z which is out of the rangerj   rk   zFBlock pointers only support 32 bit `offsets/block_shape`, got a value r   z*Expected a scalar in shape/strides/offsetsz8Expected an integer scalar type in shape/strides/offsetszzBlock pointers only support 32 bit `offsets/block_shape`, add a `.to(tl.int32)` or use regular indexing for 64 bit supportz3Unsupported element type in shape/strides/offsets: )r^   r`   r.   r_   rH   rb   r'   r{  r¶  rÍ  r&   rL  r9   rR   rt   r©   r¦   Úget_int64_tyrÅ   re   )r   r  r  s      r   Ú_convert_elem_to_ir_valuez(TritonSemantic._convert_elem_to_ir_valueD  sÑ  € Ü�dœCÔ Ü—<‘< Ó%ˆDÜ�dœBŸL™LÔ)Ü˜$Ÿ*™*¤dÔ+Ø—|‘|×,Ñ,¨T¯Z©ZÓ8Ð8ÙØ §¡Ô3¨eÔ3ð Jð 8#Ø#'§:¡: ,Ð.Hð6Jó JÐ3à—|‘|×-Ñ-¨d¯j©jÓ9Ð9à §¡Ô3¨eÔ3ð Jð 8#Ø#'§:¡: ,Ð.Hð6Jó JÐ3à—|‘|×-Ñ-¨d¯j©jÓ9Ð9Ü˜œbŸi™iÔ(Ø—:‘:×#Ñ# qÒ(ÐVÐ*VÓVÐ(Ø—:‘:×$Ñ$Ô&ÐbÐ(bÓbÐ&Ø�z‰zœRŸX™XÒ%©+Ø—|‘|×3Ñ3°D·K±KÀÇÁ×AZÑAZÓA\Ø48·J±J×4LÑ4LÓ4NóPð Pà—‘œrŸx™xÒ'±ðWð Wó W�uà—;‘;ÐØXÐKÌDÐQUËJÈ<ÐXÓXˆur   c                óŠ   — t        |d«      r |D �cg c]  }| j                  ||«      ‘Œ c}S | j                  ||«      gS c c}w )Nrý  )rþ  r`  )r   Ú	list_liker  r  s       r   r!  z$TritonSemantic._convert_to_ir_values^  sF   € Ü�9˜jÔ)ØR[Ö\È$�D×2Ñ2°4¸ÕEÒ\Ð\Ø×.Ñ.¨y¸+ÓFÐGÐGùò ]s   ‘A c           	     óÂ  ‡— | j                  |«      }| j                  |«      }| j                  |d¬«      }|j                  j                  «       r$|j                  j                  j	                  «       rt        d«      ‚|j                  j                  t        j                  k(  rH| j                  |t        j                  t        j                  |j                  j                  «      «      }t        ‰d«      s‰gŠ‰D �cg c]*  }t        |t        j                  «      r|j                  n|‘Œ, c}Št!        d„ ‰D «       «      sJ d«       ‚t        |d«      s|g}|D �cg c]*  }t        |t        j                  «      r|j                  n|‘Œ, }}t#        |«      t%        t'        t)        |«      «      «      k(  sJ d«       ‚t!        ˆfd„||||fD «       «      sJ d	«       ‚| j*                  j-                  |j.                  |||‰|«      }| j1                  |t        j                  t        j2                  |j                  j                  ‰«      «      «      S c c}w c c}w )
NFr  zMExpected `base` to be a pointer type (but not a block pointer type or others)rý  c              3  ó`   K  — | ]&  }t        |t        «      xr d |cxk  xr dk  nc –— Œ( y­w)rj   rk   N)r^   r`   )rx  r  s     r   rz  z0TritonSemantic.make_block_ptr.<locals>.<genexpr>v  s)   è ø€ Ò\È”:˜d¤CÓ(ÒC¨V°tÖ-C¸eÔ-CÓCÑ\ùs   ‚,.zGExpected a list of constant integers (`int32_t` range) in `block_shape`z<Expected a permutation of (0, 1, ..., len(order)-1) in orderc              3  óL   •K  — | ]  }t        ‰«      t        |«      k(  –— Œ y ­wr)   )rS  )rx  rb  r  s     €r   rz  z0TritonSemantic.make_block_ptr.<locals>.<genexpr>€  s   øè ø€ Òh¸)”3�{Ó#¤s¨9£~Õ5Ñhùs   ƒ!$zBExpected shape/strides/offsets/block_shape to have the same length)r!  re   r|   r  rR  r,   r.   rr   r‡   r  r  r  rþ  r^   r_   rH   r  r{  r|  r@  rS  r'   Úcreate_make_block_ptrr¦   r&   r:  )	r   ÚbaserA  Ústridesr$  r  Úorderr  r¦   s	        `   r   Úmake_block_ptrzTritonSemantic.make_block_ptrc  sû  ø€ ð ×*Ñ*¨5Ó1ˆØ×,Ñ,¨WÓ5ˆØ×,Ñ,¨WÀ%Ð,ÓHˆð �y‰y×ÑÔ! T§Y¡Y×%9Ñ%9×%BÑ%BÔ%DÜÐlÓmÐmð �9‰9×Ñ¤2§7¡7Ò*Ø—9‘9˜T¤2§?¡?´2·7±7¸D¿I¹I×<SÑ<SÓ#TÓUˆDô �{ JÔ/Ø&˜-ˆKØZeÖfÐRV¤Z°´b·l±lÔ%C�t—z’zÈÑMÒfˆÜÑ\ÐP[Ô\Ô\ð 	VØUó	VÐ\ô �u˜jÔ)Ø�GˆEØTYÖZÈDœz¨$´·±Ô=�—’À4ÑGÐZˆÐZÜ�e‹}¤¤U¬3¨u«:Ó%6Ó 7Ò7ÐwÐ9wÓwÐ7ô ÓhÈÈwÐX_ÐafÐGgÔhÔhð 	QØPó	QÐhð —‘×3Ñ3°D·K±KÀÈÐQXÐZeÐglÓmˆØ�{‰{˜6¤2§?¡?´2·=±=ÀÇÁ×AUÑAUÐWbÓ3cÓ#dÓeÐeùò% gùò [s   Ä/IÅ /Ic                ó¨   — | j                  |d¬«      }| j                  | j                  j                  |j                  |«      |j
                  «      S r*  )r!  r&   r'   Úcreate_advancer¦   re   )r   rg  r$  s      r   ÚadvancezTritonSemantic.advance‰  sC   € à×,Ñ,¨WÀ%Ð,ÓHˆð �{‰{˜4Ÿ<™<×6Ñ6°t·{±{ÀGÓLÈdÏiÉiÓXÐXr   c           	     óþ  — t        |«      }d|cxk  rdk  sn t        d|› d�«      ‚t        |«      |k7  rt        d|› dt        |«      › �«      ‚t        |«      |k7  rt        d|› dt        |«      › �«      ‚t        |j                  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k7  rt        d|	› �«      ‚|D �
cg c]"  }
| j                  |
t        j                  «      ‘Œ$ }}
|D �
cg c]5  }
| j                  t	        j                  |
«      t        j                  «      ‘Œ7 }}
t	        j                  |«      }t        |j                  t        j
                  «      sJ ‚t	        j                  |j                  j                  |«      }|j                  }|j                  j                  j!                  «       }| j#                  |«      }|j                  j                  j%                  «       r(|t&        j(                  j*                  k(  rt        d«      ‚| j,                  j/                  ||D �cg c]  }|j                  ‘Œ c}|D �cg c]  }|j                  ‘Œ c}|||«      }t	        j0                  ||||«      S c c}
w c c}
w c c}w c c}w )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 rW  rh  rs  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 z8Padding option `nan` is not supported for integer blocks)rS  r,   r^   r9   r.   r  r  rŸ  r_  rM  r/   rt   Ú_unwrap_shapere   r:  r¦   rÅ   râ  rR   r   rÝ  rß  r'   Úcreate_make_tensor_descriptorÚtensor_descriptor)r   rg  rA  rh  r  rà  r%  Ú	elem_sizeÚcontig_dim_sizeÚlast_striderf   re   Úbase_handleÚis_signed_intrá  r\  r¦   s                    r   Úmake_tensor_descriptorz%TritonSemantic.make_tensor_descriptor�  sË  € ä�5‹zˆØ�T”˜Q”ÜÐ?À¸vÀ[ÐQÓRÐRÜˆw‹<˜4ÒÜ˜y¨¨Ð.?ÄÀGÃ¸~ÐNÓOÐOÜˆ{Ó˜tÒ#ÜÐ<¸T¸FÐBVÔWZÐ[bÓWcÐVdÐeÓfÐfÜ˜$Ÿ*™*¤b§o¡oÔ6Ð6Ð6Ø—J‘J×)Ñ)×<Ñ<ÀÑAˆ	Ü×1Ñ1°+¸b±/ÓBˆØ˜YÑ&¨Ò+ÜØdÐetÐduÐuxð  zCð  yDð  DGð  HWð  Zcñ  Hcð  Gdð  djð  kóð ô ×-Ñ-¨g°b©kÓ:ˆØ˜!ÒÜÐLÈ[ÈMÐZÓ[Ð[à8=Ö>°1�×!Ñ! !¤R§X¡XÕ.Ð>ˆÐ>ØSZÖ[Èa�4×#Ñ#¤B×$;Ñ$;¸AÓ$>ÄÇÁÕIÐ[ˆÐ[ô ×&Ñ& {Ó3ˆä˜$Ÿ)™)¤R§_¡_Ô5Ð5Ð5Ü�}‰}˜TŸY™Y×1Ñ1°;Ó?ˆØ—k‘kˆØŸ	™	×,Ñ,×:Ñ:Ó<ˆà×-Ñ-¨nÓ=ˆà�9‰9×Ñ×&Ñ&Ô(¨W¼×8IÑ8I×8QÑ8QÒ-QÜÐWÓXÐXà—‘×;Ñ;¸KÐ\aÖIbÐWXÈ!Ï(Ë(ÒIbØOVÖ<WÈ!¸Q¿X»XÒ<WÐYdÐfsØ<CóEˆô ×#Ñ# F¨E°7¸DÓAÐAùò' ?ùÚ[ùò JcùÚ<Ws   Ä('K+Å:K0ÊK5
Ê6K:)r1   r`   Úreturnr   )r=   útl.dtyper>   rz  ry  rz  )r=   rz  rS   rb   r>   rz  rT   rb   rU   rb   ry  rz  )T)r   rz  r   rz  r~   rb   ry  ÚNone)FFTF)r‰   úTensorTy | numbers.NumberrŠ   r|  ry  úTuple[TensorTy, TensorTy])r‰   r   rŠ   r   rž   Úcallable)r®   r|  r¯   r|  r™   rb   ry  r   )r®   r|  r¯   r|  ry  r   )r®   r|  r¯   r|  rË   rb   ry  r   )rf   r   rÝ   r   rÞ   útl.PropagateNan)rf   r   ré   r   rê   r   rÞ   r  )r®   r   r¯   r   ry  r}  )r®   r   r¯   r   ry  r   )r®   r   )r®   r   ry  r   )r  r   ry  útl.block_type)r<  r`   r=  r`   rÈ   r€  ry  r   )r9   rz  ry  r   )rA  ú	List[int]r9   rz  ry  r   )rH   r   rA  r�  ry  r   )rH   r   ry  r   )r®   r   rZ  r�  r[  rb   ry  r   )r®   r   r1   r`   ry  r   )r‰   r   rŠ   r   r[  rb   ry  r   )rk  r   rl  r   ry  r   )rk  r   ry  r}  )r®   r   r~  ú
Tuple[int]ry  r   )r®   r   rA  r‚  ry  r   )r‰   r   rŠ   r   ry  r   )rœ  úOptional[str])r®   r   r¡  rz  ry  r   r)   )r®   r   r¡  rz  r¸  rƒ  ry  r   )r
  r   r  úOptional[TensorTy]r¯   r„  r   r   rà  r�  rÈ  r�  r×  r�  r  rb   ry  r   )r#  útl.tensor_descriptor_baserÈ  r�  r×  r�  ry  r   )r#  r…  rH   r   ry  r{  )r#  r…  rH   r   ry  r   )rÈ  r�  r×  r�  ry  r   )r
  r   re  r   r  r„  rÈ  r�  r×  r�  ry  r   )r
  r   ru  r   re  r   rï  r�  rú  r�  ry  r   )
r
  r   re  r   r  r   r|  r�  ry  z#Tuple[TensorTy, TensorTy, TensorTy])rf   r   ry  r   )r
  r   re  r   r  r   rï  r�  rú  r�  ry  r   )r‰   r   rŠ   r   rÒ  r   r»  rƒ  rÓ  r`   rÔ  rz  ry  r   )rã  r�  )re  r   rã  r�  )r‰   r   ró  r   rû  r�  rŠ   r   rô  r„  rü  r�  rÒ  zTensorTy | Nonerý  rb   rþ  rb   rÿ  rb   rÔ  rz  ry  r   )r  r   rf   r   rÝ   r   ry  r   )r  úSequence[TensorTy]r1   r`   ry  úTuple[TensorTy, ...])r  r†  r1   r`   r)  rb   ry  r‡  )r-  r   r.  r   r1   r`   ry  r   )r  zSequence[tl.tensor]r7  zSequence[tl.dtype]r8  r`   ry  zTuple[tl.tensor, ...])r®   r   r<  r`   r  r„  ry  r   )rf   r   r‹  r�  ry  r   )ry  r   )rQ  r�  rR  úList[TensorTy]rS  rb   ry  r   )r¢   r   r£   r�  r  r„  ry  r   )rg  r   ry  r   )rÛ  )rg  r   rA  rˆ  rh  rˆ  r  zList[tl.constexpr]rà  r�  ry  ztl.tensor_descriptor)~r    r!   r"   r.   r&   Ú__annotations__Úlangr   r2   r5   rC   rX   rh   rc   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›   r2  r7  rC  rd   rM  rP  rO  rW  r]  rb  rf  rp  ru  r  r‰  rˆ   r�  rü   r‡   rÊ  rÏ  rÙ  râ  rð  rû  r  r  r  r  r&  r(  r-  r5  r>  r@  rC  rG  rJ  rN  rQ  r^  rb  rf  rm  rq  rv  r  r†  r™  rœ  rŸ  r£  r§  rª  r®  r¼  rà  rå  rð  rø  r  rŒ  r  r$  r*  r/  ri  r9  r=  rC  rF  rH  rK  rW  r�   r]  r`  r!  rj  rm  rx  r  r   r   r%   r%      s  … ØŸY™Y€FˆNÓ&Ø€DàÓòóOó
QóDð05Ø*.ð05Ø3;ó05ód	òKóF	@ð aeØ05ð#à:Só#óJ,ð$>Ø#ð>Ø(0ó>ð>8Ø#ð8Ø(0ó8ð"8Ø#ð8Ø(0ó8ó]ó2>ó,ó8ó.9ó"9ó"	aóó\ó[ó\ó'ó&ó"ó
]ó]ó\óó)ó$ó/ó8ó8ó8ó8ó	8ó	8ð GKõ Zó$)ó2óAóhóSófóXóVóó0

óTóWó 2óptójô$i9ò^òò	ò	òòò
òò,:ðxnØ ðnØ25ðnØHKðnØZ^ðnØckónð	/Ø),ð	/Ø19ó	/ó/ónóuò@ò[ó
uóuóuóuóuó$ó0*ò*xò8.$ð`WØ"ðWØ'/óWó(uðØ'*ðØ/Róó8+ó$)óL$)óL%óyóxóyóò
<ð[yØ#&ð[yØ3;ð[yØ@Hó[yózó0ò$mð3tØ0ð3tØ>Að3tØHWð3tØdhð3tà#ð3tà15ð3tàBJð3tàOWó3tóreò(&óuð,xØ"&ðxØ+?óxó.Kò,ðiØ.Cóió*	@óóóóCóaóRóMòYó4Hó
$fóLYð W]ð(BØ,>ð(BØPSð(BØauô(Br   r%   )Ú
__future__r   rv   rÈ  Útypingr   r   r   r   r   r	   r
   r�   Útriton.runtimer   Ú_C.libtritonr   rS  r   r.   r   r   Ú	Exceptionr   r%   r  r   r   ú<module>r�     sY   ðÝ "Û Û ç J× JÑ JÛ å !å Ý áˆCƒL€Ù�:Ó€ôF 	ô Fô^B�W˜XÑ&õ ^Br   