Ë
    ÏÍ:j¶z ã                  óº   — d dl mZ d dlZd dlmZmZmZmZmZm	Z	m
Z
 d dlZd dlmZ ddlmZ ddlmZ  ed	«      Z ed
«      Z G d„ de«      Z G d„ de	e   «      Z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      €úm/home/mcse/projects/srt_converter/srt-converter-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                  óR  — e Zd ZU ej                  Zded<   eZded<   d„ Zdyd„Zdyd„Z	dzd„Z
	 	 	 	 d{d	„Zd|d}d
„Zd~d„Z	 	 d	 d€d„Zd�d„Z	 	 	 	 d‚d„Z	 	 	 	 d‚d„Z	 	 	 	 d‚d„Zdƒd„Zdƒd„Zd„d„Zdƒ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�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›d=„Z=d>„ Z>d?„ Z?d@„ Z@dA„ ZAdB„ ZBdC„ ZCdD„ ZDdE„ ZEdF„ ZF	 	 	 	 	 	 	 	 	 	 dœdG„ZG	 	 	 	 d�dH„ZHdždI„ZIdŸdJ„ZJdŸdK„ZKdL„ ZLdM„ ZMdŸdN„ZNdŸdO„ZOdŸdP„ZPdŸdQ„ZQdŸdR„ZRd dS„ZSd¡dT„ZTdU„ ZUdV„ ZV	 	 	 	 d¢dW„ZWd£dX„ZX	 	 	 	 d¤dY„ZYd¥dZ„ZZd¦d[„Z[d¦d\„Z\d¦d]„Z]d¦d^„Z^d¦d_„Z_d¦d`„Z`d¦da„Zadb„ Zb	 	 	 	 	 	 d§dc„Zcd¨dd„Zdd©de„Ze	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 dªdf„Zfd«dg„Zgdh„ Zhd¬di„Zi	 	 	 	 d­dj„Zjd®dk„Zkd¯dl„Zld°dm„Zmd°dn„Znd°do„Zod±dp„Zpd²dq„Zqd³dr„Zrd±ds„Zsdt„ Ztd|du„Zud´dv„Zvd´dw„Zw	 	 	 	 	 	 	 	 	 	 dµdx„Zxy,)¶Ú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_impl4   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_implC   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        |t        «      r9| j                  | j                  j	                  |«      t
        j                  «      S t        |t        «      ržd|cxk  rdk  rn nt
        j                  }nld|cxk  rdk  rn nt
        j                  }nMd|cxk  rdk  rn nt
        j                  }n.d|cxk  rdk  rn nt
        j                  }nt        d|› d�«      ‚| j                  ||¬	«      S t        |t        «      rrd
}dddz  z  }t        d   |«      }|t        d«      k(  s|dk(  s||k7  s||cxk  r|k  rn nt
        j                   }nt
        j"                  }| j                  ||¬	«      S t        |t
        j$                  «      r| j'                  |j(                  «      S t        || j                  «      r|S |rt+        d|› dt-        |«      › d�«      ‚|S )Né   €ì        l        ìûÿÿÿ         ì            l            zNonrepresentable integer ú.©r9   g      8g   àÿÿÿ?r   é   ÚabsÚinfg        zcannot convert z	 of type z
 to tensor)Ú
isinstanceÚboolr&   r'   Úget_int1r.   Úint1Úintr/   Úuint32Úint64Úuint64r,   Úscalar_constantÚfloatÚ__builtins__rK   rM   Ú	constexprÚ	to_tensorrH   r<   Útype)r   ÚxÚ
check_typer9   Úmin_float32Úmax_float32Úabs_xs          r   ro   zTritonSemantic.to_tensoru   s’  € Ü�aœÔØ—;‘;˜tŸ|™|×4Ñ4°QÓ7¼¿¹ÓAÐAä˜œ3ÔØ˜Ô"˜UÕ"ÜŸ™‘Ø˜!Ô#˜eÕ#ÜŸ	™	‘Ø˜1Ô$˜uÕ$ÜŸ™‘Ø˜!Ô#˜eÕ#ÜŸ	™	‘ä Ð#<¸Q¸C¸qÐ!AÓBÐBØ×'Ñ'¨°Ð'Ó7Ð7Ü˜œ5Ô!Ø!ˆKØ%¨¨C©Ñ/ˆKÜ  Ñ'¨Ó*ˆEØœ˜e›Ò$Ø˜Š|Ø�AŠvØ˜eÔ2 {Õ2ÜŸ
™
‘äŸ
™
�Ø×'Ñ'¨°Ð'Ó7Ð7ä˜œ2Ÿ<™<Ô(Ø—>‘> !§'¡'Ó*Ð*Ü˜˜4Ÿ;™;Ô'ØˆHÙÜ˜o¨a¨S°	¼$¸q»'¸À*ÐMÓNÐNØˆr   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 r_   )rc   ÚnumbersÚNumberro   rp   Úscalarrz   rw   rX   Úis_int_unsignedr,   rR   Úget_int_min_valueÚget_int_max_valuerk   Ú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#                  |	|
«       y )Né@   Frg   z! overflow detected for operation )rp   r~   r7   r'   ÚoptionsÚsanitize_overflowrR   r‚   r.   ri   r�   rk   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Î   s?  € Ø�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 Õ%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�   rp   r~   rw   r<   Úhandler9   r   r7   Úwith_element_tyr.   ri   Úto_irr'   Úcreate_int_castr&   Úcreate_addptrrx   Ú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�   rp   r~   rw   r§   Úminusrx   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�   rp   r~   rx   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�   rp   r~   rx   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�   rp   r~   rR   rC   r‚   Úis_int_signedr&   r'   Úcreate_sdivr¡   Úcreate_udivr<   )r   r©   rª   r«   r¬   Úret_tys         r   ÚfloordivzTritonSemantic.floordiv7  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)	rp   r~   rx   r<   r�   r'   r½   r¡   r&   )r   r©   rª   Úieee_roundingr«   r¬   rš   s          r   ÚfdivzTritonSemantic.fdivE  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�   rp   r~   rx   r&   r'   Úcreate_fremr¡   rR   r8   r<   r   rÀ   Úcreate_sremÚcreate_urem)r   r©   rª   rV   r¬   s        r   ÚmodzTritonSemantic.modN  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   rx   r.   ÚPropagateNanÚALLr&   r'   Úcreate_minimumfr¡   rp   ÚNONEÚcreate_minnumfr,   rÀ   Úcreate_minsir   Úcreate_minuir<   ©r   rq   ÚyÚpropagate_nanr9   s        r   ÚminimumzTritonSemantic.minimume  ó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   rx   r.   rÐ   rÑ   r&   r'   Úcreate_maximumfr¡   rp   rÓ   Úcreate_maxnumfr,   rÀ   Úcreate_maxsir   Úcreate_maxuir<   r×   s        r   ÚmaximumzTritonSemantic.maximumv  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   rx   r&   r'   Úcreate_clampfr¡   rp   r<   )r   rq   Ú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�   rp   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¡   rp   ©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¡   rp   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¡   rp   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)   )rp   Úis_int1Úbitcastr.   rf   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)   )rp   rö   r÷   r.   rf   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)   )rp   rö   r÷   r.   rf   Ú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¡   rp   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¡   rp   rí   s      r   ÚashrzTritonSemantic.ashrÆ  r  r   c                óÀ   — | j                  ||«      \  }}| j                  | j                  j                  |j                  |j                  «      |j
                  «      S r)   )rê   r&   r'   Ú
create_shlr¡   rp   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)
rp   r~   rw   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  )rp   r~   rw   rx   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)   )rp   r¢   r.   rf   )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�   rp   r~   rx   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�   rp   r~   rx   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�   rp   r~   rx   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�   rp   r~   rx   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�   rp   r~   rx   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�   rp   r~   rx   r&   r'   Úcreate_fcmpUNEr¡   r  rR   Úcreate_icmpNEr<   r  s       r   Ú	not_equalzTritonSemantic.not_equal-  r.  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)rc   rg   r,   rd   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.arange<  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+t	        | j                  d|j
                  › �«      } ||«      }| j                  ||«      S )Nz2dtype must be specified when value is not a tensorr   Úget_)r,   r'   r  r£   ÚgetattrÚnamer&   )r   rH   r9   Úget_value_fns       r   rk   zTritonSemantic.scalar_constantN  sq   € àˆ=ÜÐQÓRÐRØ�AŠ:Ø—L‘L×/Ñ/°·±¸D¿L¹LÓ0IÓJ‰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)rc   r.   r&   ÚnumelrH   r‚   rk   )r   rH   r9   s      r   Úmake_scalarzTritonSemantic.make_scalarY  sR   € Ü�eœRŸY™YÔ'Ø—;‘;×$Ñ$¨Ò)ÐGÐ+GÓGÐ)Ø—9‘9˜U EÓ*Ð*à×#Ñ# E¨5Ó1Ð1r   c                óF   — | j                  | j                  ||«      |«      S r)   )ÚsplatrF  )r   r<  rH   r9   s       r   ÚfullzTritonSemantic.full`  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   )rp   Úis_blockÚlenr.   r5  r9   r&   r'   Úcreate_splatr£   r¡   )r   rH   r<  rÃ   s       r   rH  zTritonSemantic.splatg  sy   € Ø—:‘:×&Ñ&Ô(ÐGÐ*GÓGÐ(Üˆu‹:˜Š?ØˆLÜ—‘˜uŸ{™{¨EÓ2ˆØ�{‰{˜4Ÿ<™<×4Ñ4°V·\±\À$Ç,Á,Ó5OÐQV×Q]ÑQ]Ó^Ð`fÓgÐgr   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)
rp   rE  r,   r.   r5  r~   r&   r'   Úcreate_reshaper¡   )r   r©   Ú	dst_shapeÚcan_reorderrE  ÚsrÃ   s          r   ÚreshapezTritonSemantic.reshapen  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   )r<  )r<  r.   Ú_unwrap_if_constexprÚinsertrp   rK  rH  r5  r~   r&   r'   Úcreate_expand_dimsr¡   )r   r©   r1   rq   rP  rÃ   s         r   Úexpand_dimszTritonSemantic.expand_dimsw  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   )
rL  r<  r.   r5  rp   r~   r&   r'   Ú
create_catr¡   )r   r„   r…   rQ  Ú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©rQ  )rƒ   r<  rX  rc   r.   rn   r5  rp   r~   r&   r'   Úcreate_joinr¡   rS  )r   ÚaÚbÚ
was_rank_1ÚtwoÚ	new_shaper[  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   r^  r   )rL  r<  r.   rU  r5  rp   r~   r'   Úcreate_splitr¡   r&   )r   ra  re  r[  Ú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        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 )Nz5permute dims must have the same length as input shapec              3  óF   K  — | ]  }t        j                  |«      –— Œ y ­wr)   )r.   rU  )Ú.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 )rL  r<  r,   ÚsortedÚlistr;  r.   r5  rp   r~   r&   r'   Úcreate_transr¡   )r   r©   Údimsro  r[  s        r   ÚpermutezTritonSemantic.permute«  s²   € Üˆu�{‰{Óœs 4›yÒ(ÜÐTÓUÐUÜÑ;°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   ÂC&
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: z, r   z3Cannot broadcast, the expanded size of the tensor (z ) must match the existing size (z) at non-singleton dimension ú: )rp   rK  rH  Úget_block_shapesrL  r,   Ú	enumerater.   r5  r~   r&   r'   Úcreate_broadcastr¡   )r   r©   r<  Ú	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 rw  r   )rp   rK  r¢   r~   r&   r'   rM  r£   r¡   rx  rL  r;  rW  r.   r5  Úvaluesry  Úappendr,   Ústrrz  )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 )rp   rK  r¢   r~   rw   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 )1rp   r~   rK  r¢   r’  rx   r”  r   rŽ  r�  r,   r‚  Úis_fp8e4b15r'   Úcodegen_fnsÚgetrQ   r&   Úcreate_fp_to_fpr¡   r£   rO   rN   rP   r‚   r.   rK   Úcreate_fp_truncÚcreate_fp_extrR   r7   r8   rÀ   Úis_boolr9   r  r2  r¤   Úis_standard_floatingÚcreate_fp_to_siÚcreate_fp_to_uiÚcreate_ui_to_fpÚcreate_si_to_fprw   Úcreate_ptr_to_intri   Ú	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
  )	Úhasattrrc   r.   rn   rH   rg   rL  Úsetrq  )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,   rp   Ú
element_tyr.   rf   rR   r   rÓ  rÕ  rú  rx  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	           
     ót  — |j                   j                  j                  «       s't        d|j                   j	                  «       › d�«      ‚|€|�t        d«      ‚|s|rt        d«      ‚|j                   j                  «       sN|r%|j                   j                  «       rt        d«      ‚|r%|j                   j                  «       rt        d«      ‚|j                   j                  «       rX|�*| j                  ||j                   j                  «       «      }|�*| j                  ||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)rp   r~   rw   r,   r   rK  r~  rx  rþ  r.   rf   Ú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ö  sK  € à�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×ÑÔØÐØ×0Ñ0°°s·x±x×7PÑ7PÓ7RÓS�ØÐ Ø×1Ñ1°%¸¿¹×9RÑ9RÓ9TÓ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Ø  rp   rw   rþ  rK  r  r  )r   r   r  rª   rö  rÖ  r¾  rÍ  r  r¿  rÎ  r×  s               r   ÚloadzTritonSemantic.load2  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)rc   r.   Útensor_descriptor_baserL  r÷  Ú_convert_to_ir_valuesr'   Úcreate_descriptor_loadr¡   rÀ  rÏ  r&   r5  )r   ÚdescÚoffsetsr¾  rÍ  Úndimrq   s          r   Údescriptor_loadzTritonSemantic.descriptor_load@  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  )rc   r.   r  rL  r÷  r<  )r   r  rH   r  r  s        r   Úvalidate_store_likez"TritonSemantic.validate_store_likeK  si   € Ü˜$¤× 9Ñ 9Ô:Ð:Ð:Ü�4×#Ñ#Ó$ˆÜ�7‹|˜tÒ#ÐW y°°Ð6HÌÈWËÈÐ%WÓWÐ#Ø�{‰{˜d×.Ñ.Ò.Ð.Ñ.r   c                óì   — | j                  |||«       | j                  |d¬«      }| j                  | j                  j	                  |j
                  |j
                  |«      t        j                  «      S ©NFr  )r  r  r&   r'   Úcreate_descriptor_storer¡   r.   Úvoid)r   r  rH   r  s       r   Údescriptor_storezTritonSemantic.descriptor_storeQ  s]   € Ø× Ñ   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.   rh   r/   rj   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_addV  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_tma]  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 )Nr&  z-16-bit float types require native tma support)r.   rh   r/   rj   ri   rI   rJ   r4  )r   r9   s     r   Ú$_descriptor_atomic_min_max_supportedz3TritonSemantic._descriptor_atomic_min_max_supporteda  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  r6  r9   r  r   r'  ÚMINr&   r'   r)  r¡   r.   r"  r*  s        r   Údescriptor_atomic_minz$TritonSemantic.descriptor_atomic_minf  ó…   € Ø× Ñ   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  r6  r9   r  r   r'  ÚMAXr&   r'   r)  r¡   r.   r"  r*  s        r   Údescriptor_atomic_maxz$TritonSemantic.descriptor_atomic_maxm  r:  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.   rh   r/   rj   ri   r  r   r'  ÚANDr&   r'   r)  r¡   r"  r*  s        r   Údescriptor_atomic_andz$TritonSemantic.descriptor_atomic_andt  ó¤   € Ø× Ñ   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.   rh   r/   rj   ri   r  r   r'  ÚORr&   r'   r)  r¡   r"  r*  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.   rh   r/   rj   ri   r  r   r'  ÚXORr&   r'   r)  r¡   r"  r*  s        r   Údescriptor_atomic_xorz$TritonSemantic.descriptor_atomic_xor‚  rA  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 r4  zdescriptor gather of ú must have at least ú columns, but got Fr  )rc   r.   r  rL  r÷  r<  r9   r”  r5  r  r'   Úcreate_descriptor_gatherr¡   r£   r&   )
r   r  Ú	x_offsetsÚy_offsetr¾  rÍ  r9   Úmin_colsrp   rq   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   rJ  r   r   rK  rL  rM  z6descriptor scatter must have at least 8 rows, but got r4  zdescriptor scatter of rN  rO  Fr  )rc   r.   r  rL  r÷  r<  Úshapaer9   r”  r  r'   Úcreate_descriptor_scatterr¡   r&   r"  )r   r  rH   rQ  rR  r9   rS  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,   rp   rþ  rx  rK  r~  r.   rf   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           	     ó(  — |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                  «       rV| j                  ||j                   j                  «       «      }|�*| j                  ||j                   j                  «       «      }|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  ú"Mask must have boolean scalar type)rp   r~   rw   r,   r   rK  r~  rx  rþ  r.   rf   r  r	  r
  r‚   r&   r'   Úcreate_storer¡   r"  r¥  Úcreate_masked_store)	r   r   r[  r  rö  r¿  rÎ  r  r  s	            r   Ú_store_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×ÑÔØ×+Ñ+¨C°·±×1JÑ1JÓ1LÓMˆCØÐØ×0Ñ0°°s·x±x×7PÑ7PÓ7RÓS�à—‘—‘ˆØ×"Ñ"ˆð ”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Ï  rp   Úis_constr~   r,   rw   rþ  rK  r\  ra  )	r   r   r[  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)é   r4  r’   z9atomic_cas only supports elements with width {16, 32, 64})ræ  rñ  rp   r~   rþ  r”  r,   r&   r'   Úcreate_atomic_casr¡   )r   r   Úcmpr[  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 rc  r§   Úatomic_z does not support fp16z does not support bf16rg  z does not support T)rp   r~   rw   r,   r   rd  rþ  r.   rI   rJ   Úint16Úuint16r”  r‚  rK  r~  rx  r‚   r'   re   rf   r¢   rM  r£   r&   )r   r   r[  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‚   rf   )r   rq   r´  ÚidtypeÚixÚsignbits         r   Ú_signbitzTritonSemantic._signbit6  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   )rr  ræ  rñ  rp   r~   rR   rÀ   r&   r'   Úcreate_atomic_rmwr   Ú	ATOMIC_OPr<  r¡   ÚUMAXr.   rK   rM   r<   r/   ri   r÷   r	  rh   rj   ry  rþ   r•   ÚUMINÚwhere©r   r   r[  r  rå  rð  Úsca_tyÚi_typeÚi_valÚi_ptrÚui_typeÚui_valÚui_ptrÚnegÚposÚpos_retÚneg_retrš   s                     r   Ú
atomic_maxzTritonSemantic.atomic_max=  ó²  € Ø×8Ñ8¸¸cÀ4ÈÓO‰ˆˆ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   )rr  ræ  rñ  rp   r~   rR   rÀ   r&   r'   r{  r   r|  r8  r¡   r~  r.   rK   rM   r<   r/   ri   r÷   r	  rh   rj   ry  rþ   r•   r}  r  r€  s                     r   Ú
atomic_minzTritonSemantic.atomic_minc  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§   )rr  ræ  rñ  rp   r~   rx   r   r|  ÚFADDr(  r&   r'   r{  r¡   )r   r   r[  r  rå  rð  r�  ro  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)rr  ræ  rñ  r&   r'   r{  r   r|  r?  r¡   rp   ©r   r   r[  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)rr  ræ  rñ  r&   r'   r{  r   r|  rC  r¡   rp   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)rr  ræ  rñ  r&   r'   r{  r   r|  rF  r¡   rp   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)rr  ræ  rñ  r&   r'   r{  r   r|  ÚXCHGr¡   rp   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                ó  — |j                  «       | j                  j                  j                  v s+J d| j                  j                  j                  › d|› �«       ‚|j	                  «       }|dk(  rd}t        t        j                  |«      S )Nzinput_precision must be one of ú. Got ÚTF32X3ÚTF32x3)Úlowerr'   r“   Úallowed_dot_input_precisionsÚupperrA  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Ü”r×)Ñ)¨?Ó;Ð;r   c           
     ó‚  — |j                   j                  «       r|j                   j                  «       sJ ‚|j                  j                  «       r|j                  j                  «       r�n|j                  t        j
                  t        j                  t        j                  t        j                  t        j                  fv sJ d|j                  › �«       ‚|j                  t        j
                  t        j                  t        j                  t        j                  t        j                  fv sJ d|j                  › �«       ‚|j                  |j                  k(  s!J d|j                  › d|j                  › �«       ‚|j                  j                  «       s|j                  j                  «       rwd| j                  j                  j                  v rt        j                  d«       | j!                  |t        j                  «      }| j!                  |t        j                  «      }|€ | j                  j                  j"                  }| 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(                  d   j*                  |j(                  d   j*                  k(  sVJ d|j(                  › d|j(                  › d|j(                  d   j*                  › d|j(                  d   j*                  › d�	«       ‚| j                  j,                  j/                  d«      	 €J d«       ‚ | j                  j,                  d   |j                   |j                   «      }	|j(                  d   j*                  |	d   k\  r>|j(                  d   j*                  |	d   k\  r|j(                  d   j*                  |	d   k\  sJ d|	d   › d|	d   › d|	d   › �«       ‚|j                   j0                  j3                  «       rZ|j                   j0                  t        j
                  k(  sJ d«       ‚| j                  j5                  d«      }
t        j6                  }n×|j9                  «       rt;        d«      ‚|j                   j0                  j=                  «       s$|j                   j0                  j9                  «       r,| j                  j?                  d«      }
t        j                  }nH|jA                  «       r| j                  jC                  d«      n| j                  j?                  d«      }
|}|j                   j(                  d   }|j                   j(                  d   }|j                   j(                  d   }|dk(  r|j                   j(                  d   nd }t	        jD                  ||r|||gn||g«      }|€6| j                  jG                  |jI                  | j                  «      |
«      }n|jJ                  }|j                   |k(  sJ ‚|€X|j                  j                  «       r;|j                  j                  «       r!| j                  j                  jL                  }nNd}nK|j                  j                  «       r1|j                  j                  «       r||kD  rt;        d|› d|› d�«      ‚| jO                  | j                  jQ                  |jJ                  |jJ                  |||«      |«      S )NzUnsupported lhs dtype zUnsupported rhs dtype z&Both operands must be same dtype. Got r   Úfp8e4b15z–the use of fp8e4b15 is deprecated on Hopper and later architectures and can cause significant slow down. It will be removed in a future triton releaser   é   ú+Both inputs must be either 2D or 3D; (lhs: ú	 vs rhs: r  r^  éþÿÿÿzFirst input shape (z) and second input shape z= are not compatible for matmul (second index of first shape (z0) must be equal to first index of second shape (Úmin_dot_sizez2target doesn't provide lower shape bounds for dot.r   r   zInput shapes should have M >= z, N >= z
 and K >= zonly int8 supported!zhout_dtype=bfloat16 is unsupported. Please use out_dtype=float32/float16 and cast with `.to(tl.bfloat16)`zmax_num_imprecise_acc (z) must be <= K ())rp   rK  r9   rQ   r.   r  Úuint8rI   rJ   rK   rŸ  r'   r“   Ú!deprecated_fp8_dot_operand_dtypesÚwarningsÚwarnr‚   Údefault_dot_input_precisionr«  rL  r<  rH   r   r¡  r~   rR   Ú	get_int32r/   rP   r,   rN   Úget_fp32rO   Úget_fp16r5  rM  r£   r¡   Úmax_num_imprecise_acc_defaultr&   Ú
create_dot)r   r„   r…   Úaccrª  Úmax_num_imprecise_accÚ	out_dtypeÚlhs_rankÚrhs_rankr²  r  Úret_scalar_tyÚMÚNÚKÚBrÃ   Ú
acc_handles                     r   ÚdotzTritonSemantic.dot»  sh  € à�x‰x× Ñ Ô" s§x¡x×'8Ñ'8Ô':Ð:Ð:à�9‰9×ÑÔ #§)¡)×"2Ñ"2Ô"4áà—9‘9¤§¡¬"¯(©(´B·J±JÄÇÁÜ!#§¡ð!-ñ -ð SØ0FÀsÇyÁyÀkÐ.RóSð -à—9‘9¤§¡¬"¯(©(´B·J±JÄÇÁÜ!#§¡ð!-ñ -ð 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àÐ"Ø"Ÿ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à-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Ò%Ð%Ð%ð !Ð(Ø�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: r^   )rA  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³  rn  r÷   )r   r[  rË  Ú	triton_tyÚunsigned_tys        r   Ú_bitcast_to_fp_typez"TritonSemantic._bitcast_to_fp_type  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                ó~  — |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  }|
s||z  }t        j                  ||r|||gn||g«      }| j                  j                  d«      }|€6| j                  j                  |j                  | j                  «      |«      }n|j                  }|j                   |k(  sJ ‚|rd n|j                  }|rd n|j                  } | 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   )rp   rK  rL  r<  rH   rÍ  rc   r.   rn   rØ  r5  r'   r¹  rM  r£   r¡   r&   Úcreate_dot_scaled)!r   r„   Ú	lhs_scaleÚ
lhs_formatr…   Ú	rhs_scaleÚ
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Ç  Úrhs_scale_handleÚlhs_scale_handles!                                    r   Ú
dot_scaledzTritonSemantic.dot_scaled!  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Ü—‘˜y±q¨1¨a°©)¸qÀ!¸fÓEˆØ�\‰\×"Ñ" 1Ó%ˆØˆ;ØŸ™×2Ñ2°6·<±<ÀÇÁÓ3MÈrÓR‰JàŸ™ˆJØ—8‘8˜vÒ%Ð%Ð%Ù#4™4¸)×:JÑ:JÐÙ#4™4¸)×:JÑ:JÐØ�{‰{Ø�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.   rf   rµ  r¶  r‚   r�   rp   rK  rƒ   r&   r'   Úcreate_selectr¡   )r   Ú	conditionrq   rØ   r‡  rÃ   s         r   r  zTritonSemantic.whereU  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.   r5  r&   )r   rq   rV   rˆ  Úres_tys        r   Úwrap_tensorzTritonSemantic.wrap_tensori  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)Tr_  N)rS  rE  rH   )rn  Útr   s     €r   rp  z+TritonSemantic.reduction.<locals>.<genexpr>s  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)   )rp   r<  )rn  rø  r<  s     €r   rp  z+TritonSemantic.reduction.<locals>.<genexpr>z  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_resultrp   r~   )rn  r|  ÚinputsÚ	reduce_oprˆ  r   s     €€€€r   rp  z+TritonSemantic.reduction.<locals>.<genexpr>€  s@   øè ø€ ò uØ\]ˆD×Ñ˜Y×1Ñ1°!Ó4°f¸Q±i·n±n×6KÑ6KÈY×Wñuùó   ƒAA)Útuplerp   r<  rL  ry  Úallr'   Úcreate_reducer¡   Úverifyr;  )r   rý  r1   Úregion_builder_fnÚrankr|  rR  rø  rþ  rˆ  r<  s   ``      @@@r   Ú	reductionzTritonSemantic.reductionq  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û  )rn  r|  rý  Úscan_opr   r<  s     €€€€r   rp  z2TritonSemantic.associative_scan.<locals>.<genexpr>˜  s;   øè ø€ ÒwÐ_`�T×%Ñ% g×&8Ñ&8¸Ó&;¸VÀA¹Y¿^¹^×=RÑ=RÐTY×ZÑwùrÿ  )	rp   r<  rL  r'   Úcreate_scanr¡   r  r   r;  )	r   rý  r1   r  Úreverser  rø  r	  r<  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   rL  rp   r<  r;  r'   Úcreate_gatherr¡   rõ  r~   )r   ÚsrcÚindexr1   r  ro  Ú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                óè  — 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 inputr^  )rL  r<  r9   rR   r~  rp   r~   r¥  r,   r¡   r&   r'   Úcreate_histogramr.   r5  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å   rL  r<  r,   r¡   Úset_attrr   Ú	make_attrÚget_context©r   rq   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©rL  r<  r,   r¡   r  r   r  r  r  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.constancyr  r  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ú rw  r^  r   )ÚendswithrL  Ú
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
                  |«      t        j                  «      S r)   )r'   r“   Údebugr&   Úcreate_assertr¡   r.   r"  )r   r�   rž   s      r   r˜   zTritonSemantic.device_assertã  sB   € Ø�|‰|×#Ñ#×)Ò)ØØ�{‰{˜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.assumeè  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&                  k7  r	|sJ d«       ‚|j                   S J dt)        |«      › �«       ‚)Nr\   r]   z@Block pointers only support 64 bit `shape/strides`, got a value z which is out of the rangerZ   r[   zFBlock pointers only support 32 bit `offsets/block_shape`, got a value r   z*Expected a scalar in shape/strides/offsetsz8Expected an integer scalar type in shape/strides/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: )rc   rg   r.   rn   rH   rd   r'   re   r¬  r¸  r&   rE  r9   rR   ri   r¤   r¡   Úget_int64_tyrÀ   r/   rp   )r   rø  r  s      r   Ú_convert_elem_to_ir_valuez(TritonSemantic._convert_elem_to_ir_valueë  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ô  r8  )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)rZ   r[   N)rc   rg   )rn  rø  s     r   rp  z0TritonSemantic.make_block_ptr.<locals>.<genexpr>  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)   )rL  )rn  r:  r÷  s     €r   rp  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  rp   rw   rþ  rK  r,   r.   rf   r‚   r	  r  r
  rô  rc   rn   rH   r  rq  rr  r;  rL  r'   Úcreate_make_block_ptrr¡   r&   r5  )	r   Úbaser<  Ústridesr  r÷  Úorderrø  r¡   s	        `   r   Úmake_block_ptrzTritonSemantic.make_block_ptr
  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¡   rp   )r   r?  r  s      r   ÚadvancezTritonSemantic.advance0  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
<   |d
   dk7  rt        d|d
   › �«      ‚|D �cg c]"  }| j                  |t        j                  «      ‘Œ$ }}|D �cg c]"  }| j                  |t        j                  «      ‘Œ$ }}t	        j                  |«      }t        |j                  t        j
                  «      sJ ‚t	        j                  |j                  j                  |«      }	|j                  }
|j                  j                  j!                  «       }| j"                  j%                  |
|D �cg c]  }|j                  ‘Œ c}|D �cg c]  }|j                  ‘Œ c}||«      }t	        j&                  ||||	«      S c c}w c c}w c c}w c c}w )Nr   é   z Expected 1 <= ndim <= 5 but got z dimensionsz	Expected z strides but got zExpected block_shape to have z dimensions but got rM  r^  rg  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 )rL  r,   rc   r9   r.   r	  rþ  r”  rU  rF  r/   ri   Ú_unwrap_shaperp   r5  r¡   rÀ   r'   Úcreate_make_tensor_descriptorÚtensor_descriptor)r   r?  r<  r@  r÷  r  Ú	elem_sizeÚcontig_dim_sizerq   rp   Úbase_handleÚis_signed_intrR  r¡   s                 r   Úmake_tensor_descriptorz%TritonSemantic.make_tensor_descriptor7  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Ó:ˆ�‰Ø�2‰;˜!ÒÜÐLÈWÐUWÉ[ÈMÐZÓ[Ð[à8=Ö>°1�×!Ñ! !¤R§X¡XÕ.Ð>ˆÐ>Ø:AÖB°Q�4×#Ñ# A¤r§x¡xÕ0ÐBˆÐBô ×&Ñ& {Ó3ˆä˜$Ÿ)™)¤R§_¡_Ô5Ð5Ð5Ü�}‰}˜TŸY™Y×1Ñ1°;Ó?ˆØ—k‘kˆØŸ	™	×,Ñ,×:Ñ:Ó<ˆà—‘×;Ñ;¸KÐ\aÖIbÐWXÈ!Ï(Ë(ÒIbØOVÖ<WÈ!¸Q¿X»XÒ<WÐYdÐfsóuˆä×#Ñ# F¨E°7¸DÓAÐAùò ?ùÚBùò JcùÚ<Ws   Ä1'JÅ'JÈ6J
ÉJ)r1   rg   Úreturnr   )r=   útl.dtyper>   rQ  rP  rQ  )r=   rQ  rS   rd   r>   rQ  rT   rd   rU   rd   rP  rQ  )T)rr   rd   )r   rQ  r   rQ  ry   rd   rP  ÚNone)FFTF)r„   úTensorTy | numbers.Numberr…   rS  rP  úTuple[TensorTy, TensorTy])r„   r   r…   r   r™   Úcallable)r©   rS  rª   rS  r”   rd   rP  r   )r©   rS  rª   rS  rP  r   )r©   rS  rª   rS  rÆ   rd   rP  r   )rq   r   rØ   r   rÙ   útl.PropagateNan)rq   r   rä   r   rå   r   rÙ   rV  )r©   r   rª   r   rP  rT  )r©   r   rª   r   rP  r   )r©   r   )r©   r   rP  r   )r  r   rP  útl.block_type)r7  rg   r8  rg   rÃ   rW  rP  r   )r9   rQ  rP  r   )r<  ú	List[int]r9   rQ  rP  r   )rH   r   r<  rX  rP  r   )r©   r   rP  rX  rQ  rd   rP  r   )r©   r   r1   rg   rP  r   )r„   r   r…   r   rQ  rd   rP  r   )ra  r   rb  r   rP  r   )ra  r   rP  rT  )r©   r   rt  ú
Tuple[int]rP  r   )r©   r   r<  rY  rP  r   )r„   r   r…   r   rP  r   )r‘  úOptional[str])r©   r   r–  rQ  rP  r   r)   )r©   r   r–  rQ  r®  rZ  rP  r   )r   r   r  úOptional[TensorTy]rª   r[  rö  r   rÖ  r‚  r¾  r‚  rÍ  r‚  r  rd   rP  r   )r  útl.tensor_descriptor_baser¾  r‚  rÍ  r‚  rP  r   )r  r\  rH   r   rP  rR  )r  r\  rH   r   rP  r   )r¾  r‚  rÍ  r‚  rP  r   )rH   r   rP  r   )r   r   r[  r   r  r[  r¾  r‚  rÍ  r‚  rP  r   )r   r   ri  r   r[  r   rå  r‚  rð  r‚  rP  r   )
r   r   r[  r   r  r   ro  r‚  rP  z#Tuple[TensorTy, TensorTy, TensorTy])rq   r   rP  r   )r   r   r[  r   r  r   rå  r‚  rð  r‚  rP  r   )r„   r   r…   r   r½  r   rª  rZ  r¾  rg   r¿  rQ  rP  r   )rË  r‚  )r[  r   rË  r‚  )r„   r   rÛ  r   rÜ  r‚  r…   r   rÝ  r[  rÞ  r‚  r½  zTensorTy | Nonerß  rd   rà  rd   rá  rd   r¿  rQ  rP  r   )rò  r   rq   r   rØ   r   rP  r   )rý  úSequence[TensorTy]r1   rg   rP  úTuple[TensorTy, ...])rý  r]  r1   rg   r  rd   rP  r^  )r  r   r  r   r1   rg   rP  r   )r©   r   r  rg   r  r[  rP  r   )rq   r   r€  rX  rP  r   )rP  r   )r)  r‚  r*  úList[TensorTy]r+  rd   rP  r   )r�   r   rž   r‚  rP  r   )r?  r   rP  r   )
r?  r   r<  r_  r@  r_  r÷  zList[tl.constexpr]rP  ztl.tensor_descriptor)yr    r!   r"   r.   r&   Ú__annotations__Úlangr   r2   r5   rC   rX   ro   rz   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  r>  rk   rF  rI  rH  rS  rX  r\  rf  rk  ru  r~  rƒ   r’  r÷   r‚   rÀ  rÅ  rÏ  rØ  ræ  rñ  rú  r  r  r  r  r  r#  r+  r4  r6  r9  r=  r@  rD  rG  rT  rX  r\  ra  re  rj  rr  ry  rŒ  r�  r’  r–  rš  r�  r¡  r«  rÈ  rÍ  rØ  rï  r  rõ  r  r  r  r  r  r  r   r#  r/  r˜   r5  r8  r  rB  rE  rO  r
  r   r   r%   r%      sþ  … ØŸY™Y€FˆNÓ&Ø€DàÓòóOó
QóDð05Ø*.ð05Ø3;ó05ôd#óR	@ð 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ó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*$ðXWØ"ðWØ'/óWó(uðØ'*ðØ/Róó8+ó$)óL$)óL%óyóxóyóò<ðLyØ#&ðLyØ3;ðLyØ@HóLyó\ó0ð$.tØ0ð.tØ>Að.tØHWð.tØdhð.tà#ð.tà15ð.tàBJð.tàOWó.tóheò(&óuð,xØ"&ðxØ+?óxó.Kó.	@óóóóCóaóRó
MòYó4Hó
$fóLYð'Bàð'Bð ð'Bð  ð	'Bð
 (ð'Bð 
ô'Br   r%   )Ú
__future__r   rµ  Útypingr   r   r   r   r   r	   r
   r|   Útriton.runtimer   Ú_C.libtritonr   rI  r   r.   r   r   Ú	Exceptionr   r%   r
  r   r   ú<module>rg     sV   ðÝ "Û ç J× JÑ JÛ å !å Ý áˆCƒL€Ù�:Ó€ôF 	ô FôEB�W˜XÑ&õ EBr   